⚡

HIP-ROCm跨平台GPU开发 hip-rocm

HIP-ROCm跨平台GPU开发技能是专门用于AMD GPU高性能计算的工具集，支持将CUDA代码转换为HIP兼容代码，实现NVIDIA和AMD GPU的跨平台开发。该技能涵盖hipify代码转换、ROCm工具链配置、GPU内核开发、性能分析优化等核心功能，适用于机器学习、深度学习、科学计算等GPU加速应用开发。关键词：AMD HIP ROCm GPU开发 跨平台 CUDA转换 高性能计算 并行计算 GPU加速 机器学习加速 深度学习框架

嵌入式软件 0 次安装 0 次浏览 更新于 2/25/2026

name: hip-rocm description: 用于跨平台GPU开发的AMD HIP和ROCm生态系统。执行hipify转换工具，生成HIP兼容的内核代码，处理CUDA/HIP API差异，配置ROCm工具链，并使用rocprof进行性能分析。 allowed-tools: Bash(*) 读取 写入 编辑 全局搜索 Grep 网络获取 metadata: author: babysitter-sdk version: “1.0.0” category: cross-platform backlog-id: SK-009

hip-rocm

您是hip-rocm - 一个专门用于AMD HIP和ROCm生态系统开发的技能。该技能为面向AMD GPU的跨平台GPU编程提供专家级能力。

概述

此技能支持AI驱动的AMD GPU开发，包括：

  • 执行hipify转换工具（hipify-perl、hipify-clang）
  • 生成HIP兼容的内核代码
  • 处理CUDA/HIP API差异
  • 配置ROCm工具链编译
  • 使用rocprof和omniperf进行性能分析
  • 支持MI100/MI200/MI300架构
  • 维护单一源代码的NVIDIA/AMD代码
  • 跨平台性能基准测试

前提条件

  • ROCm 5.0+
  • HIP运行时
  • hipify工具
  • AMD GPU（或支持HIP的NVIDIA GPU）

能力

1. CUDA到HIP转换

将CUDA代码转换为HIP：

# 使用hipify-perl（快速转换）
hipify-perl cuda_file.cu > hip_file.cpp

# 使用hipify-clang（更准确）
hipify-clang cuda_file.cu -o hip_file.cpp

# 批量转换
hipify-perl -inplace *.cu
hipconvertinplace.sh .

# 生成转换统计信息
hipify-perl --print-stats cuda_file.cu

# 排除特定模式
hipify-perl --skip-includes cuda_file.cu > hip_file.cpp

2. HIP内核开发

编写HIP兼容的内核：

#include <hip/hip_runtime.h>

// HIP内核（可移植到CUDA和AMD）
__global__ void vectorAdd(const float* a, const float* b, float* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

// 启动语法（与CUDA相同）
int main() {
    // 分配内存
    float *d_a, *d_b, *d_c;
    hipMalloc(&d_a, size);
    hipMalloc(&d_b, size);
    hipMalloc(&d_c, size);

    // 复制到设备
    hipMemcpy(d_a, h_a, size, hipMemcpyHostToDevice);
    hipMemcpy(d_b, h_b, size, hipMemcpyHostToDevice);

    // 启动内核
    int blockSize = 256;
    int numBlocks = (n + blockSize - 1) / blockSize;
    hipLaunchKernelGGL(vectorAdd, dim3(numBlocks), dim3(blockSize),
        0, 0, d_a, d_b, d_c, n);

    // 替代启动语法
    vectorAdd<<<numBlocks, blockSize>>>(d_a, d_b, d_c, n);

    // 同步并复制回主机
    hipDeviceSynchronize();
    hipMemcpy(h_c, d_c, size, hipMemcpyDeviceToHost);

    // 清理
    hipFree(d_a);
    hipFree(d_b);
    hipFree(d_c);
}

3. API兼容性宏

处理CUDA/HIP差异：

// 平台检测
#ifdef __HIP_PLATFORM_AMD__
    // AMD特定代码
#elif defined(__HIP_PLATFORM_NVIDIA__)
    // NVIDIA HIP代码
#elif defined(__CUDA_ARCH__)
    // CUDA特定代码
#endif

// 通用兼容性头文件
#if defined(__HIPCC__) || defined(__HIP__)
    #include <hip/hip_runtime.h>
    #define DEVICE_SYNC hipDeviceSynchronize
    #define MALLOC hipMalloc
    #define FREE hipFree
    #define MEMCPY hipMemcpy
#else
    #include <cuda_runtime.h>
    #define DEVICE_SYNC cudaDeviceSynchronize
    #define MALLOC cudaMalloc
    #define FREE cudaFree
    #define MEMCPY cudaMemcpy
#endif

// 线程束大小处理
#ifdef __HIP_PLATFORM_AMD__
    #define WARP_SIZE 64  // AMD波前
#else
    #define WARP_SIZE 32  // NVIDIA线程束
#endif

4. ROCm编译

编译HIP代码：

# 为AMD GPU编译
hipcc -o program program.cpp

# 指定目标架构
hipcc --offload-arch=gfx90a -o program program.cpp  # MI200
hipcc --offload-arch=gfx942 -o program program.cpp  # MI300

# 多目标
hipcc --offload-arch=gfx908 --offload-arch=gfx90a -o program program.cpp

# 带优化
hipcc -O3 -o program program.cpp

# 生成汇编
hipcc -S --offload-arch=gfx90a program.cpp

# 详细编译
hipcc -v -o program program.cpp

# CMake配置
set(CMAKE_CXX_COMPILER hipcc)
set(GPU_TARGETS "gfx90a" CACHE STRING "GPU架构")

5. 使用rocprof进行性能分析

分析AMD GPU应用程序：

# 基本性能分析
rocprof ./program

# 收集特定指标
rocprof -i metrics.txt ./program

# 生成跟踪
rocprof --hip-trace ./program
rocprof --hsa-trace ./program

# 系统跟踪
rocprof --sys-trace ./program

# 导出到JSON
rocprof --stats --json ./program

# 指标文件示例（metrics.txt）
# pmc: SQ_WAVES, SQ_INSTS_VALU, SQ_INSTS_SMEM
# pmc: TCC_HIT_sum, TCC_MISS_sum

6. Omniperf分析

深度性能分析：

# 分析应用程序
omniperf profile -n workload_name ./program

# 分析性能数据
omniperf analyze -p workload_name

# 基于Web的GUI
omniperf analyze -p workload_name --gui

# 比较性能数据
omniperf analyze -p baseline -p optimized --compare

# 特定分析部分
omniperf analyze -p workload_name --metric-set memory
omniperf analyze -p workload_name --metric-set compute

7. 架构特定优化

针对AMD架构优化：

// 波前感知编程（64线程波前）
__device__ int waveReduceSum(int val) {
    #pragma unroll
    for (int offset = 32; offset > 0; offset >>= 1) {
        val += __shfl_down(val, offset);
    }
    return val;
}

// 高效使用LDS（本地数据共享）
__shared__ __align__(16) float lds[256];

// AMD内存合并（256字节粒度）
__global__ void coalescedKernel(float4* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float4 val = data[idx];  // 16字节对齐加载
        // 处理...
        data[idx] = val;
    }
}

// 架构特定内核
#if __gfx90a__ || __gfx942__
    // MI200/MI300优化
    // 使用矩阵核心（MFMA指令）
#elif __gfx908__
    // MI100优化
#endif

8. hipBLAS和rocBLAS

GPU数学库：

#include <hipblas/hipblas.h>
// 或ROCm原生
#include <rocblas/rocblas.h>

hipblasHandle_t handle;
hipblasCreate(&handle);

// GEMM操作
float alpha = 1.0f, beta = 0.0f;
hipblasSgemm(handle,
    HIPBLAS_OP_N, HIPBLAS_OP_N,
    M, N, K,
    &alpha,
    d_A, M,
    d_B, K,
    &beta,
    d_C, M);

// 带显式流的rocBLAS
rocblas_handle roc_handle;
rocblas_create_handle(&roc_handle);
rocblas_set_stream(roc_handle, stream);

rocblas_sgemm(roc_handle,
    rocblas_operation_none, rocblas_operation_none,
    M, N, K,
    &alpha, d_A, M, d_B, K, &beta, d_C, M);

9. RCCL集合操作

AMD的NCCL等效库：

#include <rccl/rccl.h>

// 初始化RCCL（与NCCL相同API）
rcclComm_t comm;
rcclUniqueId id;
rcclGetUniqueId(&id);
rcclCommInitRank(&comm, worldSize, id, rank);

// 全归约
rcclAllReduce(sendbuff, recvbuff, count, rcclFloat, rcclSum, comm, stream);

// 清理
rcclCommDestroy(comm);

流程集成

此技能与以下流程集成：

  • hip-porting-cross-platform.js - 跨平台移植
  • multi-gpu-programming.js - 多GPU开发

输出格式

{
  "operation": "hipify",
  "status": "success",
  "input_files": ["kernel.cu", "main.cu"],
  "output_files": ["kernel.cpp", "main.cpp"],
  "conversion_stats": {
    "cuda_calls_converted": 45,
    "manual_review_needed": 3,
    "warnings": ["__shfl_sync无法直接移植到HIP"]
  },
  "target_architectures": ["gfx90a", "gfx942"],
  "recommendations": [
    "检查归约内核中的波前大小（64 vs 32）",
    "考虑使用rocBLAS进行BLAS操作"
  ]
}

依赖项

  • ROCm 5.0+
  • HIP运行时
  • hipify-perl或hipify-clang
  • rocprof/omniperf（用于性能分析）

限制

  • 线程束/波前大小不同（32 vs 64）
  • 某些CUDA内置函数需要手动移植
  • 纹理内存API不同
  • CUDA特定功能可能无法移植