• 源码安装 HIPIFY 和应用示例,将cuda生态源码转化成HIP生态源码

    1，源码下载

    GitHub - ROCm/HIPIFY: HIPIFY: Convert CUDA to Portable C++ CodeHIPIFY: Convert CUDA to Portable C++ Code. Contribute to ROCm/HIPIFY development by creating an account on GitHub.icon-default.png?t=N7T8https://github.com/ROCm/HIPIFY.git

    git clone --recursive https://github.com/ROCm/HIPIFY.git
    sudo apt install clang-dev

     2，编译并安装

    2.1 通常方式

    hipify-clang 文档：

    https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/hipify-clang.md

    编译命令：

    cmake -DCMAKE_INSTALL_PREFIX=../dist -DCMAKE_BUILD_TYPE=Release  ..
    make -j install

    此时 hipify-clang 会被安装到 HIPIFY/dist/bin 中，

    测试：

    1. cd  ../dist/bin
    2. hipify --help

    如果系统中存在多个llvm版本，在执行翻译命令时，比如hipify-clang ./vectorAdd.cu  --cuda-path=/usr/local/cuda-12.1可能会发生错误，如下提示：

    CommandLine Error: Option 'static-func-full-module-prefix' registered more than once!
    LLVM ERROR: inconsistency in registered CommandLine option

    这时需要使用自制的LLVM，如下2.2节所示。

    2.2 自制LLVM的方式

    2.2.1  下载llvm源码

    wget https://github.com/llvm/llvm-project/archive/refs/tags/llvmorg-17.0.6.tar.gz

    解压，tar zxf llvmorg.....

    2.2.2  配置编译LLVM

    1. cd llvmorg.....
    2.  
    3. mkdir -p build ../dist/local
    4.  
    5. cd build
    6.  
    7. cmake -G "Unix Makefiles" ../llvm      \
    8. -DLLVM_ENABLE_PROJECTS="clang;clang-tools-extra;compiler-rt"           \
    9. -DLLVM_BUILD_EXAMPLES=ON           -DLLVM_TARGETS_TO_BUILD="host"      \
    10. -DCMAKE_BUILD_TYPE=Release           -DLLVM_ENABLE_ASSERTIONS=ON       \
    11. -DLLVM_ENABLE_RUNTIMES=all             -DLLVM_BUILD_LLVM_DYLIB=ON      \
    12. -DCMAKE_INSTALL_PREFIX=../../dist/local

    make -j

    make -j install

    测试时，llvm 被install在如下文件夹：

    /home/hipper/ex_dock_hipify/dist/local

    ls /home/hipper/ex_dock_hipify/dist/local 如图：

    2.2.3 配置编译HIPIFY

    指定 LLVM 安装目录的配置方法：

    -DCMAKE_PREFIX_PATH=/home/hipper/ex_dock_hipify/dist/local

    1.  
    2. cmake \
    3.  -DCMAKE_BUILD_TYPE=Release \
    4.  -DCMAKE_INSTALL_PREFIX=../dist \
    5.  -DCMAKE_PREFIX_PATH=/home/hipper/ex_dock_hipify/dist/local \
    6.  -DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-12.1  ..

    make -j install

    3. 示例

    3.1翻译 .cu 文件到 .hip 文件

    命令：

    /home/hipper/ex_dock_hipify/HIPIFY/dist/bin/hipify-clang ./vectorAdd.cu  --cuda-path=/usr/local/cuda-12.1

    会在 ./ 目录中生成 vectoreAdd.cu.hip 的文件。

    其中，hipify-clang 并不检查输入文件的扩展名，比如这里的.cu，它只检查文件内部的内容，将cuda生态的关键字有机地翻译成 hip生态的关键字，输出文件会在原文件名的基础上加上 .hip 后缀；

    源代码分别如下。

    使用 cuda samples中的vectoradd.cu为例，源码如下：

    vectorAdd.cu

    1. #include 
    2. #include 
    3.  
    4. __global__ void vectorAdd(const float *A, const float *B, float *C,
    5.                           int numElements) {
    6.   int i = blockDim.x * blockIdx.x + threadIdx.x;
    7.  
    8.   if (i < numElements) {
    9.     C[i] = A[i] + B[i] + 0.0f;
    10.   }
    11. }
    12.  
    13. int main(void) {
    14.   cudaError_t err = cudaSuccess;
    15.   int numElements = 50000;
    16.   size_t size = numElements * sizeof(float);
    17.   printf("[Vector addition of %d elements]\n", numElements);
    18.  
    19.   float *h_A = (float *)malloc(size);
    20.   float *h_B = (float *)malloc(size);
    21.   float *h_C = (float *)malloc(size);
    22.   if (h_A == NULL || h_B == NULL || h_C == NULL) {
    23.     fprintf(stderr, "Failed to allocate host vectors!\n");
    24.     exit(EXIT_FAILURE);
    25.   }
    26.  
    27.   for (int i = 0; i < numElements; ++i) {
    28.     h_A[i] = rand() / (float)RAND_MAX;
    29.     h_B[i] = rand() / (float)RAND_MAX;
    30.   }
    31.  
    32.   float *d_A = NULL;
    33.   err = cudaMalloc((void **)&d_A, size);
    34.  
    35.   if (err != cudaSuccess) {
    36.     fprintf(stderr, "Failed to allocate device vector A (error code %s)!\n",
    37.             cudaGetErrorString(err));
    38.     exit(EXIT_FAILURE);
    39.   }
    40.  
    41.   float *d_B = NULL;
    42.   err = cudaMalloc((void **)&d_B, size);
    43.  
    44.   if (err != cudaSuccess) {
    45.     fprintf(stderr, "Failed to allocate device vector B (error code %s)!\n",
    46.             cudaGetErrorString(err));
    47.     exit(EXIT_FAILURE);
    48.   }
    49.  
    50.   float *d_C = NULL;
    51.   err = cudaMalloc((void **)&d_C, size);
    52.  
    53.   if (err != cudaSuccess) {
    54.     fprintf(stderr, "Failed to allocate device vector C (error code %s)!\n",
    55.             cudaGetErrorString(err));
    56.     exit(EXIT_FAILURE);
    57.   }
    58.  
    59.   printf("Copy input data from the host memory to the CUDA device\n");
    60.   err = cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    61.  
    62.   if (err != cudaSuccess) {
    63.     fprintf(stderr,
    64.             "Failed to copy vector A from host to device (error code %s)!\n",
    65.             cudaGetErrorString(err));
    66.     exit(EXIT_FAILURE);
    67.   }
    68.  
    69.   err = cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
    70.  
    71.   if (err != cudaSuccess) {
    72.     fprintf(stderr,
    73.             "Failed to copy vector B from host to device (error code %s)!\n",
    74.             cudaGetErrorString(err));
    75.     exit(EXIT_FAILURE);
    76.   }
    77.  
    78.   int threadsPerBlock = 256;
    79.   int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock;
    80.   printf("CUDA kernel launch with %d blocks of %d threads\n", blocksPerGrid,
    81.          threadsPerBlock);
    82.   vectorAdd<<>>(d_A, d_B, d_C, numElements);
    83.   err = cudaGetLastError();
    84.  
    85.   if (err != cudaSuccess) {
    86.     fprintf(stderr, "Failed to launch vectorAdd kernel (error code %s)!\n",
    87.             cudaGetErrorString(err));
    88.     exit(EXIT_FAILURE);
    89.   }
    90.  
    91.   printf("Copy output data from the CUDA device to the host memory\n");
    92.   err = cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
    93.  
    94.   if (err != cudaSuccess) {
    95.     fprintf(stderr,
    96.             "Failed to copy vector C from device to host (error code %s)!\n",
    97.             cudaGetErrorString(err));
    98.     exit(EXIT_FAILURE);
    99.   }
    100.  
    101.   for (int i = 0; i < numElements; ++i) {
    102.     if (fabs(h_A[i] + h_B[i] - h_C[i]) > 1e-5) {
    103.       fprintf(stderr, "Result verification failed at element %d!\n", i);
    104.       exit(EXIT_FAILURE);
    105.     }
    106.   }
    107.  
    108.   printf("Test PASSED\n");
    109.  
    110.   err = cudaFree(d_A);
    111.  
    112.   if (err != cudaSuccess) {
    113.     fprintf(stderr, "Failed to free device vector A (error code %s)!\n",
    114.             cudaGetErrorString(err));
    115.     exit(EXIT_FAILURE);
    116.   }
    117.  
    118.   err = cudaFree(d_B);
    119.  
    120.   if (err != cudaSuccess) {
    121.     fprintf(stderr, "Failed to free device vector B (error code %s)!\n",
    122.             cudaGetErrorString(err));
    123.     exit(EXIT_FAILURE);
    124.   }
    125.  
    126.   err = cudaFree(d_C);
    127.  
    128.   if (err != cudaSuccess) {
    129.     fprintf(stderr, "Failed to free device vector C (error code %s)!\n",
    130.             cudaGetErrorString(err));
    131.     exit(EXIT_FAILURE);
    132.   }
    133.  
    134.   free(h_A);
    135.   free(h_B);
    136.   free(h_C);
    137.  
    138.   printf("Done\n");
    139.   return 0;
    140. }

    生成的 vectorAdd.cu.hip :

    1. #include 
    2. #include 
    3.  
    4. __global__ void vectorAdd(const float *A, const float *B, float *C,
    5.                           int numElements) {
    6.   int i = blockDim.x * blockIdx.x + threadIdx.x;
    7.  
    8.   if (i < numElements) {
    9.     C[i] = A[i] + B[i] + 0.0f;
    10.   }
    11. }
    12.  
    13. int main(void) {
    14.   hipError_t err = hipSuccess;
    15.   int numElements = 50000;
    16.   size_t size = numElements * sizeof(float);
    17.   printf("[Vector addition of %d elements]\n", numElements);
    18.  
    19.   float *h_A = (float *)malloc(size);
    20.   float *h_B = (float *)malloc(size);
    21.   float *h_C = (float *)malloc(size);
    22.   if (h_A == NULL || h_B == NULL || h_C == NULL) {
    23.     fprintf(stderr, "Failed to allocate host vectors!\n");
    24.     exit(EXIT_FAILURE);
    25.   }
    26.  
    27.   for (int i = 0; i < numElements; ++i) {
    28.     h_A[i] = rand() / (float)RAND_MAX;
    29.     h_B[i] = rand() / (float)RAND_MAX;
    30.   }
    31.  
    32.   float *d_A = NULL;
    33.   err = hipMalloc((void **)&d_A, size);
    34.  
    35.   if (err != hipSuccess) {
    36.     fprintf(stderr, "Failed to allocate device vector A (error code %s)!\n",
    37.             hipGetErrorString(err));
    38.     exit(EXIT_FAILURE);
    39.   }
    40.  
    41.   float *d_B = NULL;
    42.   err = hipMalloc((void **)&d_B, size);
    43.  
    44.   if (err != hipSuccess) {
    45.     fprintf(stderr, "Failed to allocate device vector B (error code %s)!\n",
    46.             hipGetErrorString(err));
    47.     exit(EXIT_FAILURE);
    48.   }
    49.  
    50.   float *d_C = NULL;
    51.   err = hipMalloc((void **)&d_C, size);
    52.  
    53.   if (err != hipSuccess) {
    54.     fprintf(stderr, "Failed to allocate device vector C (error code %s)!\n",
    55.             hipGetErrorString(err));
    56.     exit(EXIT_FAILURE);
    57.   }
    58.  
    59.   printf("Copy input data from the host memory to the CUDA device\n");
    60.   err = hipMemcpy(d_A, h_A, size, hipMemcpyHostToDevice);
    61.  
    62.   if (err != hipSuccess) {
    63.     fprintf(stderr,
    64.             "Failed to copy vector A from host to device (error code %s)!\n",
    65.             hipGetErrorString(err));
    66.     exit(EXIT_FAILURE);
    67.   }
    68.  
    69.   err = hipMemcpy(d_B, h_B, size, hipMemcpyHostToDevice);
    70.  
    71.   if (err != hipSuccess) {
    72.     fprintf(stderr,
    73.             "Failed to copy vector B from host to device (error code %s)!\n",
    74.             hipGetErrorString(err));
    75.     exit(EXIT_FAILURE);
    76.   }
    77.  
    78.   int threadsPerBlock = 256;
    79.   int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock;
    80.   printf("CUDA kernel launch with %d blocks of %d threads\n", blocksPerGrid,
    81.          threadsPerBlock);
    82.   vectorAdd<<>>(d_A, d_B, d_C, numElements);
    83.   err = hipGetLastError();
    84.  
    85.   if (err != hipSuccess) {
    86.     fprintf(stderr, "Failed to launch vectorAdd kernel (error code %s)!\n",
    87.             hipGetErrorString(err));
    88.     exit(EXIT_FAILURE);
    89.   }
    90.  
    91.   printf("Copy output data from the CUDA device to the host memory\n");
    92.   err = hipMemcpy(h_C, d_C, size, hipMemcpyDeviceToHost);
    93.  
    94.   if (err != hipSuccess) {
    95.     fprintf(stderr,
    96.             "Failed to copy vector C from device to host (error code %s)!\n",
    97.             hipGetErrorString(err));
    98.     exit(EXIT_FAILURE);
    99.   }
    100.  
    101.   for (int i = 0; i < numElements; ++i) {
    102.     if (fabs(h_A[i] + h_B[i] - h_C[i]) > 1e-5) {
    103.       fprintf(stderr, "Result verification failed at element %d!\n", i);
    104.       exit(EXIT_FAILURE);
    105.     }
    106.   }
    107.  
    108.   printf("Test PASSED\n");
    109.  
    110.   err = hipFree(d_A);
    111.  
    112.   if (err != hipSuccess) {
    113.     fprintf(stderr, "Failed to free device vector A (error code %s)!\n",
    114.             hipGetErrorString(err));
    115.     exit(EXIT_FAILURE);
    116.   }
    117.  
    118.   err = hipFree(d_B);
    119.  
    120.   if (err != hipSuccess) {
    121.     fprintf(stderr, "Failed to free device vector B (error code %s)!\n",
    122.             hipGetErrorString(err));
    123.     exit(EXIT_FAILURE);
    124.   }
    125.  
    126.   err = hipFree(d_C);
    127.  
    128.   if (err != hipSuccess) {
    129.     fprintf(stderr, "Failed to free device vector C (error code %s)!\n",
    130.             hipGetErrorString(err));
    131.     exit(EXIT_FAILURE);
    132.   }
    133.  
    134.   free(h_A);
    135.   free(h_B);
    136.   free(h_C);
    137.  
    138.   printf("Done\n");
    139.   return 0;
    140. }
    141.

    3.2 编译运行 vectorAdd.cu.hip

    编译：

    $ /opt/rocm/bin/hipcc ./vectorAdd.cu.hip -o vectorAdd

    运行效果如下图：

  • 相关阅读:
    FlaskUser type object ‘User‘ has no attribute ‘get_user_by_token‘
    揭秘Spring事务失效场景分析与解决方案
    国际化配置（ant-design-vue设置成中文）
    面试官：说说Vue 3.0中Treeshaking特性？
    磁盘被未知资源耗尽lsof -n|grep deleted
    Java中如何执行多条shell/bat命令 2023年验证通过
    Mac电脑版鼠标连点工具 RapidClick for Mac
    Apache Seata -- 一款开源的分布式事务解决方案
    一、ubuntu-django+nginx+uwsgi：ubuntu系统部署django项目，前后端不分离项目
    【小白学机器学习13】一文理解假设检验的反证法，H0如何设计的，什么时候用左侧检验和右侧检验，等各种关于假设检验的基础知识
  • 原文地址：https://blog.csdn.net/eloudy/article/details/136364439