hip-rocm

HIP / ROCm

目的

HIP を使った AMD GPU プログラミングについて、HIP ランタイム API、hipcc によるコンパイル、HIPIFY（hipify-perl、hipify-clang）を使った CUDA コードの移植、ROCm ツールチェーンのセットアップ、rocprof によるプロファイリング、rocgdb によるデバッグ、HIP と CUDA の API 対応、MI300X 固有の最適化をエージェント向けにガイドします。

使用する状況

• 既存の CUDA コードベースを AMD GPU に移植する
• MI200／MI300 ハードウェア向けに Linux で ROCm をセットアップする
• AMD データセンター GPU 向けのネイティブ HIP カーネルを記述する
• rocprof または rocprofiler-sdk で HIP アプリケーションをプロファイリングする
• rocgdb またはコンピュートサニタイザーでデバイス障害をデバッグする
• HIP のポータビリティマクロを使って、複数ベンダー対応の GPU コードをビルドする

ワークフロー

1. ROCm のインストールと確認

# Ubuntu/Debian (check ROCm docs for your distro version)
sudo apt install rocm-dev rocm-libs hip-dev

# Verify
rocminfo | head -30
hipconfig --version
hipcc --version

# List devices
rocm-smi

コンパイル対象の GPU を設定します。

export AMDGPU_TARGETS=gfx942   # MI300X
export HIP_PLATFORM=amd

2. 最小限の HIP カーネル

// vector_add.hip
#include <hip/hip_runtime.h>
#include <stdio.h>

__global__ void vector_add(const float *a, const float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n)
        c[i] = a[i] + b[i];
}

int main(void) {
    const int n = 1 << 20;
    size_t bytes = n * sizeof(float);
    float *d_a, *d_b, *d_c;

    hipMalloc(&d_a, bytes);
    hipMalloc(&d_b, bytes);
    hipMalloc(&d_c, bytes);

    int threads = 256;
    int blocks = (n + threads - 1) / threads;
    hipLaunchKernelGGL(vector_add, dim3(blocks), dim3(threads), 0, 0,
                       d_a, d_b, d_c, n);
    hipDeviceSynchronize();

    hipFree(d_a); hipFree(d_b); hipFree(d_c);
    return 0;
}

hipcc -O3 --offload-arch=gfx942 -o vector_add vector_add.hip
./vector_add

3. HIPIFY による CUDA → HIP の移植

# Perl-based batch converter (quick port)
hipify-perl cuda_kernel.cu > cuda_kernel.hip

# Clang-based (more accurate, preserves structure)
hipify-clang cuda_project/ -o hip_project/ --cuda-path=/usr/local/cuda

# Convert single file in place
hipify-clang -inplace --cuda-path=/usr/local/cuda main.cu

一般的な API の対応関係：

CUDA	HIP
cudaMalloc	hipMalloc
cudaMemcpy	hipMemcpy
cudaMemcpyAsync	hipMemcpyAsync
cudaStream_t	hipStream_t
<<<grid, block>>>	hipLaunchKernelGGL または <<<>>>（HIP は CUDA 構文をサポート）
__syncthreads()	__syncthreads()（同じ）
threadIdx / blockIdx	同じ組み込み変数

両方のコンパイルに対応するポータビリティヘッダー：

#ifdef __HIP_PLATFORM_AMD__
#include <hip/hip_runtime.h>
#else
#include <cuda_runtime.h>
#define hipMalloc cudaMalloc
#define hipMemcpy cudaMemcpy
// ... more macros
#endif

4. hipcc のフラグ

# Target specific GPU architecture
hipcc --offload-arch=gfx942 -O3 -o app main.hip

# Multiple architectures
hipcc --offload-arch=gfx90a --offload-arch=gfx942 -o app main.hip

# Debug
hipcc -g -O0 --offload-arch=gfx942 -o app_debug main.hip

# Link with rocBLAS
hipcc -lrocblas -o app main.hip

5. rocprof によるプロファイリング

# Basic kernel trace
rocprof --stats ./app

# CSV metrics output
rocprof -i input.csv -o output.csv ./app

# input.csv example:
# pmc: SQ_INSTS_VALU_ADD_F32,SQ_INSTS_VALU_MUL_F32,GRBM_COUNT

# ROCm 6.x rocprofiler-sdk (preferred for new projects)
rocprofv3 --kernel-trace -- ./app

主なメトリクス（AMD の用語）：

• VALU 使用率 — コンピュートユニットの稼働状況
• LDS バンク競合 — 共有メモリ（LDS）によるストール
• メモリスループット — HBM 帯域幅の使用率

6. rocgdb によるデバッグ

# Build with debug symbols
hipcc -g -O0 --offload-arch=gfx942 -o app_debug main.hip

rocgdb ./app_debug

(rocgdb) break vector_add
(rocgdb) run
(rocgdb) info rocm kernels
(rocgdb) rocm thread 0 0 0
(rocgdb) print i

AMD は、ROCm の roc-obj-extract や、利用可能なメモリ検査ツールを通じて、compute-sanitizer に相当するツールもサポートしています。

7. MI300X 向けの最適化

# Enable MFMA (matrix fused multiply-add) instructions
hipcc --offload-arch=gfx942 -munsafe-fp-atomics -O3 -o app main.hip

最適化	MI300X に関する注意点
行列演算	GEMM には rocBLAS／hipBLASLt を使い、カスタム処理には MFMA intrinsic を使います
HBM 帯域幅	約 5.3 TB/s（MI300X のピーク値）— この値に近づけるにはメモリアクセスのコアレス化を最大化します
ウェーブフロントのサイズ	64 スレッド（CUDA のワープは 32）— リダクションのパターンを調整します
LDS（共有メモリ）	CU あたり 64 KB。バンク競合に注意します

ウェーブフロントを考慮したリダクション：

__device__ float warp_reduce_sum(float val) {
    // AMD wavefront = 64 lanes
    for (int offset = 32; offset > 0; offset >>= 1)
        val += __shfl_down(val, offset);
    return val;
}

8. ライブラリエコシステム

NVIDIA	AMD ROCm
cuBLAS	rocBLAS / hipBLAS
cuDNN	MIOpen
NCCL	rccl
Thrust	hipCUB（ポータブル）
cuFFT	rocFFT

hipcc -lrocblas -o gemm_test gemm.hip

よくある問題

症状	原因	対処法
hipErrorNoDevice	ROCm ドライバーが読み込まれていない	rocm-smi を確認し、ユーザーを render グループに追加します
アーキテクチャが誤ったバイナリ	gfx* ターゲットの不一致	rocminfo を実行し、--offload-arch を設定します
hipify で移植が不完全	CUDA 固有の API	協調グループやテクスチャ参照を手動で修正します
CUDA の参照実装より遅い	ウェーブフロントは 64、ワープは 32	ブロックサイズを 64 の倍数に調整します
HSA_STATUS_ERROR	GPU がビジー、またはメモリ不足	rocm-smi --showmeminfo を実行し、割り当てを減らします
rocprof の出力が空	カーネルが起動されていない	起動後に hipGetLastError() を確認します

関連スキル

• skills/gpu/cuda — 移植元となる CUDA のパターン
• skills/gpu/cuda-profiling — Nsight の概念と rocprof の対応関係
• skills/gpu/gpu-memory-model — ウェーブフロントとワープ、コアレス化のルール
• skills/gpu/triton-lang — Triton は ROCm バックエンドで AMD をサポート
• skills/compilers/llvm — HIP は Clang／LLVM ツールチェーンを使用