amd.md

# Using AMD Partition

For testing your application on the AMD partition,
you need to prepare a job script for that partition or use the interactive job:

```
salloc -N 1 -c 64 -A PROJECT-ID -p p03-amd --gres=gpu:4 --time=08:00:00
```
where: 
- -N 1 means allocating one server, 
- -c 64 means allocation 64 cores,  
- -A is your project, 
- -p p03-amd is AMD partition, 
- --gres=gpu:4 means allcating all 4 GPUs of the node,
- --time=08:00:00 means allocation for 8 hours.  

You have also an option to allocate subset of the resources only, by reducing the -c and --gres=gpu to smaller values. 

```
salloc -N 1 -c 48 -A PROJECT-ID -p p03-amd --gres=gpu:3 --time=08:00:00
salloc -N 1 -c 32 -A PROJECT-ID -p p03-amd --gres=gpu:2 --time=08:00:00
salloc -N 1 -c 16 -A PROJECT-ID -p p03-amd --gres=gpu:1 --time=08:00:00
```

### Note: 

p03-amd01 server has hyperthreading enabled therefore htop shows 128 cores.

p03-amd02 server has hyperthreading dissabled therefore htop shows 64 cores.


## Using AMD MI100 GPUs

The AMD GPUs can be programmed using the ROCm open-source platform (see: https://docs.amd.com/ for more information.)

ROCm and related libraries are installed directly in the system. You can find it here: 
```
/opt/rocm/
```
The actual version can be found here: 
```
[user@p03-amd02.cs]$ cat /opt/rocm/.info/version

5.5.1-74
```

## Basic HIP code

The first way how to program AMD GPUs is to use HIP. 

The basic vector addition code in HIP looks like this. This a full code and you can copy and paste it into a file. For this example we use `vector_add.hip.cpp` .  

```
#include <cstdio>
#include <hip/hip_runtime.h>



__global__ void add_vectors(float * x, float * y, float alpha, int count)
{
    long long idx = blockIdx.x * blockDim.x + threadIdx.x;

    if(idx < count)
        y[idx] += alpha * x[idx];
}

int main()
{
    // number of elements in the vectors
    long long count = 10;

    // allocation and initialization of data on the host (CPU memory)
    float * h_x = new float[count];
    float * h_y = new float[count];
    for(long long i = 0; i < count; i++)
    {
        h_x[i] = i;
        h_y[i] = 10 * i;
    }

    // print the input data
    printf("X:");
    for(long long i = 0; i < count; i++)
        printf(" %7.2f", h_x[i]);
    printf("\n");
    printf("Y:");
    for(long long i = 0; i < count; i++)
        printf(" %7.2f", h_y[i]);
    printf("\n");
    
    // allocation of memory on the GPU device
    float * d_x;
    float * d_y;
    hipMalloc(&d_x, count * sizeof(float));
    hipMalloc(&d_y, count * sizeof(float));

    // copy the data from host memory to the device
    hipMemcpy(d_x, h_x, count * sizeof(float), hipMemcpyHostToDevice);
    hipMemcpy(d_y, h_y, count * sizeof(float), hipMemcpyHostToDevice);

    int tpb = 256;
    int bpg = (count - 1) / tpb + 1;
    // launch the kernel on the GPU
    add_vectors<<< bpg, tpb >>>(d_x, d_y, 100, count);
    // hipLaunchKernelGGL(add_vectors, bpg, tpb, 0, 0, d_x, d_y, 100, count);

    // copy the result back to CPU memory
    hipMemcpy(h_y, d_y, count * sizeof(float), hipMemcpyDeviceToHost);

    // print the results
    printf("Y:");
    for(long long i = 0; i < count; i++)
        printf(" %7.2f", h_y[i]);
    printf("\n");

    // free the allocated memory
    hipFree(d_x);
    hipFree(d_y);
    delete[] h_x;
    delete[] h_y;

    return 0;
}
```

To compile the code we use `hipcc` compiler. The compiler information can be found like this: 

````
[user@p03-amd02.cs ~]$ hipcc --version 

HIP version: 5.5.30202-eaf00c0b
AMD clang version 16.0.0 (https://github.com/RadeonOpenCompute/llvm-project roc-5.5.1 23194 69ef12a7c3cc5b0ccf820bc007bd87e8b3ac3037)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/rocm-5.5.1/llvm/bin
````

The code is compiled a follows: 

```
hipcc vector_add.hip.cpp -o vector_add.x
```

The correct output of the code is: 
```
[user@p03-amd02.cs ~]$ ./vector_add.x 
X:    0.00    1.00    2.00    3.00    4.00    5.00    6.00    7.00    8.00    9.00
Y:    0.00   10.00   20.00   30.00   40.00   50.00   60.00   70.00   80.00   90.00
Y:    0.00  110.00  220.00  330.00  440.00  550.00  660.00  770.00  880.00  990.00
```

## HIP and ROCm libraries

The list of official AMD libraries can be found here: https://docs.amd.com/category/libraries. 



The libraries are installed in the same directory is ROCm 
```
/opt/rocm/
```

Following libraries are installed: 
```
drwxr-xr-x  4 root root   44 Jun  7 14:09 hipblas
drwxr-xr-x  3 root root   17 Jun  7 14:09 hipblas-clients
drwxr-xr-x  3 root root   29 Jun  7 14:09 hipcub
drwxr-xr-x  4 root root   44 Jun  7 14:09 hipfft
drwxr-xr-x  3 root root   25 Jun  7 14:09 hipfort
drwxr-xr-x  4 root root   32 Jun  7 14:09 hiprand
drwxr-xr-x  4 root root   44 Jun  7 14:09 hipsolver
drwxr-xr-x  4 root root   44 Jun  7 14:09 hipsparse
```

and 

```
drwxr-xr-x  4 root root   32 Jun  7 14:09 rocalution
drwxr-xr-x  4 root root   44 Jun  7 14:09 rocblas
drwxr-xr-x  4 root root   44 Jun  7 14:09 rocfft
drwxr-xr-x  4 root root   32 Jun  7 14:09 rocprim
drwxr-xr-x  4 root root   32 Jun  7 14:09 rocrand
drwxr-xr-x  4 root root   44 Jun  7 14:09 rocsolver
drwxr-xr-x  4 root root   44 Jun  7 14:09 rocsparse
drwxr-xr-x  3 root root   29 Jun  7 14:09 rocthrust
```



### Using hipBlas library

The basic code in HIP that uses hipBlas looks like this. This a full code and you can copy and paste it into a file. For this example we use `hipblas.hip.cpp` .  

```
#include <cstdio>
#include <vector>
#include <cstdlib>
#include <hip/hip_runtime.h>
#include <hipblas/hipblas.h>


int main()
{    
    srand(9600);

    int width = 10;
    int height = 7;
    int elem_count = width * height;


    // initialization of data in CPU memory

    float * h_A;
    hipHostMalloc(&h_A, elem_count * sizeof(*h_A));
    for(int i = 0; i < elem_count; i++)
        h_A[i] = (100.0f * rand()) / (float)RAND_MAX;
    printf("Matrix A:\n");
    for(int r = 0; r < height; r++)
    {
        for(int c = 0; c < width; c++)
            printf("%6.3f  ", h_A[r + height * c]);
        printf("\n");
    }

    float * h_x;
    hipHostMalloc(&h_x, width * sizeof(*h_x));
    for(int i = 0; i < width; i++)
        h_x[i] = (100.0f * rand()) / (float)RAND_MAX;
    printf("vector x:\n");
    for(int i = 0; i < width; i++)
        printf("%6.3f  ", h_x[i]);
    printf("\n");
    
    float * h_y;
    hipHostMalloc(&h_y, height * sizeof(*h_y));
    for(int i = 0; i < height; i++)
        h_x[i] = 100.0f + i;
    printf("vector y:\n");
    for(int i = 0; i < height; i++)
        printf("%6.3f  ", h_x[i]);
    printf("\n");

    
    // initialization of data in GPU memory

    float * d_A;
    size_t pitch_A;
    hipMallocPitch((void**)&d_A, &pitch_A, height * sizeof(*d_A), width);
    hipMemcpy2D(d_A, pitch_A, h_A, height * sizeof(*d_A), height * sizeof(*d_A), width, hipMemcpyHostToDevice);
    int lda = pitch_A / sizeof(float);

    float * d_x;
    hipMalloc(&d_x, width * sizeof(*d_x));
    hipMemcpy(d_x, h_x, width * sizeof(*d_x), hipMemcpyHostToDevice);
    
    float * d_y;
    hipMalloc(&d_y, height * sizeof(*d_y));
    hipMemcpy(d_y, h_y, height * sizeof(*d_y), hipMemcpyHostToDevice);


    // basic calculation of the result on the CPU

    float alpha=2.0f, beta=10.0f;

    for(int i = 0; i < height; i++)
        h_y[i] *= beta;
    for(int r = 0; r < height; r++)
        for(int c = 0; c < width; c++)
            h_y[r] += alpha * h_x[c] * h_A[r + height * c];
    printf("result y CPU:\n");
    for(int i = 0; i < height; i++)
        printf("%6.3f  ", h_y[i]);
    printf("\n");
    
    
    // calculation of the result on the GPU using the hipBLAS library

    hipblasHandle_t blas_handle;
    hipblasCreate(&blas_handle);

    hipblasSgemv(blas_handle, HIPBLAS_OP_N, height, width, &alpha, d_A, lda, d_x, 1, &beta, d_y, 1);
    hipDeviceSynchronize();

    hipblasDestroy(blas_handle);
    

    // copy the GPU result to CPU memory and print it
    hipMemcpy(h_y, d_y, height * sizeof(*d_y), hipMemcpyDeviceToHost);
    printf("result y BLAS:\n");
    for(int i = 0; i < height; i++)
        printf("%6.3f  ", h_y[i]);
    printf("\n");


    // free all the allocated memory
    hipFree(d_A);
    hipFree(d_x);
    hipFree(d_y);
    hipHostFree(h_A);
    hipHostFree(h_x);
    hipHostFree(h_y);

    return 0;
}
```

The code compilation can be done as follows: 
```
hipcc hipblas.hip.cpp -o hipblas.x -lhipblas
```

### Using hipSolver library

The basic code in HIP that uses hipSolver looks like this. This a full code and you can copy and paste it into a file. For this example we use `hipsolver.hip.cpp` .  

```
#include <cstdio>
#include <vector>
#include <cstdlib>
#include <algorithm>
#include <hipsolver/hipsolver.h>
#include <hipblas/hipblas.h>

int main()
{
    srand(63456);

    int size = 10;


    // allocation and initialization of data on host. this time we use std::vector

    int h_A_ld = size;
    int h_A_pitch = h_A_ld * sizeof(float);
    std::vector<float> h_A(size * h_A_ld);
    for(int r = 0; r < size; r++)
        for(int c = 0; c < size; c++)
            h_A[r * h_A_ld + c] = (10.0 * rand()) / RAND_MAX;
    printf("System matrix A:\n");
    for(int r = 0; r < size; r++)
    {
        for(int c = 0; c < size; c++)
            printf("%6.3f  ", h_A[r * h_A_ld + c]);
        printf("\n");
    }    
    
    std::vector<float> h_b(size);
    for(int i = 0; i < size; i++)
        h_b[i] = (10.0 * rand()) / RAND_MAX;
    printf("RHS vector b:\n");
    for(int i = 0; i < size; i++)
        printf("%6.3f  ", h_b[i]);
    printf("\n");

    std::vector<float> h_x(size);


    // memory allocation on the device and initialization

    float * d_A;
    size_t d_A_pitch;
    hipMallocPitch((void**)&d_A, &d_A_pitch, size, size);
    int d_A_ld = d_A_pitch / sizeof(float);

    float * d_b;
    hipMalloc(&d_b, size * sizeof(float));
    
    float * d_x;
    hipMalloc(&d_x, size * sizeof(float));

    int * d_piv;
    hipMalloc(&d_piv, size * sizeof(int));

    int * info;
    hipMallocManaged(&info, sizeof(int));

    hipMemcpy2D(d_A, d_A_pitch, h_A.data(), h_A_pitch, size * sizeof(float), size, hipMemcpyHostToDevice);
    hipMemcpy(d_b, h_b.data(), size * sizeof(float), hipMemcpyHostToDevice);
    

    // solving the system using hipSOLVER

    hipsolverHandle_t solverHandle;
    hipsolverCreate(&solverHandle);

    int wss_trf, wss_trs; // wss = WorkSpace Size
    hipsolverSgetrf_bufferSize(solverHandle, size, size, d_A, d_A_ld, &wss_trf);
    hipsolverSgetrs_bufferSize(solverHandle, HIPSOLVER_OP_N, size, 1, d_A, d_A_ld, d_piv, d_b, size, &wss_trs);
    float * workspace;
    int wss = std::max(wss_trf, wss_trs);
    hipMalloc(&workspace, wss * sizeof(float));
    
    hipsolverSgetrf(solverHandle, size, size, d_A, d_A_ld, workspace, wss, d_piv, info);
    hipsolverSgetrs(solverHandle, HIPSOLVER_OP_N, size, 1, d_A, d_A_ld, d_piv, d_b, size, workspace, wss, info);

    hipMemcpy(d_x, d_b, size * sizeof(float), hipMemcpyDeviceToDevice);
    hipMemcpy(h_x.data(), d_x, size * sizeof(float), hipMemcpyDeviceToHost);
    printf("Solution vector x:\n");
    for(int i = 0; i < size; i++)
        printf("%6.3f  ", h_x[i]);
    printf("\n");

    hipFree(workspace);

    hipsolverDestroy(solverHandle);


    // perform matrix-vector multiplication A*x using hipBLAS to check if the solution is correct

    hipblasHandle_t blasHandle;
    hipblasCreate(&blasHandle);

    float alpha = 1;
    float beta = 0;
    hipMemcpy2D(d_A, d_A_pitch, h_A.data(), h_A_pitch, size * sizeof(float), size, hipMemcpyHostToDevice);
    hipblasSgemv(blasHandle, HIPBLAS_OP_N, size, size, &alpha, d_A, d_A_ld, d_x, 1, &beta, d_b, 1);
    hipDeviceSynchronize();

    hipblasDestroy(blasHandle);

    for(int i = 0; i < size; i++)
        h_b[i] = 0;
    hipMemcpy(h_b.data(), d_b, size * sizeof(float), hipMemcpyDeviceToHost);
    printf("Check multiplication vector Ax:\n");
    for(int i = 0; i < size; i++)
        printf("%6.3f  ", h_b[i]);
    printf("\n");


    // free all the allocated memory

    hipFree(info);
    hipFree(d_piv);
    hipFree(d_x);
    hipFree(d_b);
    hipFree(d_A);

    return 0;
}
```

The code compilation can be done as follows: 
```
hipcc hipsolver.hip.cpp -o hipsolver.x -lhipblas -lhipsolver
```

### Other AMD libraries and frameworks 





Please see [gcc options](https://gcc.gnu.org/onlinedocs/gcc/AArch64-Options.html) for more advanced compilation settings.
No complications are expected as long as the application does not use any intrinsic for `x64` architecture.
If you want to use intrinsic,
[SVE](https://developer.arm.com/documentation/102699/0100/Optimizing-with-intrinsics) instruction set is available.