有时候，在调用编译器的时候，我们像移除掉编译器的某些参数（例如这些参数是 cmake 等工具生成的，我们由不好修改cmake的内部机制）。 一种典型的场景是，英伟达的 nvcc 对很多 gcc/clang 的参数都不支持，当使用 nvcc 作为编译器时，需要移除掉一些参数或者进行特殊处理。
这时候，可以考虑用我御用的脚本。 下面的脚本是一个参考，支持移除编译器参数、修改或者添加编译器参数。 这里用的编译器是 hipcc（在英伟达平台上，hipcc 会调用 nvcc，然后 nvcc 调用系统的gcc），大家可以在此基础上进行修改。
最近在用 HIP 写 SpMV（稀疏矩阵向量乘），在算法实现过程中，遇到了一些 wavefront/block 内线程通信的问题，在此记录下。
我们都知道 __syncthreads() 可用于让 block 内的线程同步。
在 AMD GPU 上（ROCm），__syncthreads 会被编译成 s_barrier 指令（注：链接中的地址是 AMD GPU GFX9的内容），并加上必要的的全局访存(global memory) 和 LDS 访存 (shared memory) 的同步。
一般地，但 block 中的线程都会操作 LDS（如往其中写入数据）， 但后续执行过程中，线程又需要用到 LDS 中的数据时（如从其中取数据），通常会在用数据之前加上 __syncthreads， 以保证前面block 内所有的线程操作 LDS 的步骤都已经完成了。
我们考虑下面这个示例代码：
#include <stdio.h>
constexpr int THREADS_PER_BLOCK = 256;
constexpr int N = 8;
constexpr int VECTOR_SIZE = 4;
constexpr int REDUCE_SIZE = 8;
__global__ void test_kernel(int *x, int *y, int alpha) {
const int g_tid = threadIdx.x + blockDim.x * blockIdx.x; // global thread id
const int g_bid = blockIdx.x; // global block id
const int tid_in_block = g_tid % blockDim.x;
__shared__ int SH[THREADS_PER_BLOCK];
constexpr int VECTOR_NUM = THREADS_PER_BLOCK / VECTOR_SIZE; // vectors in block
const int g_vector_id = g_tid / VECTOR_SIZE;
const int tid_in_vector = g_tid % VECTOR_SIZE;
const int vec_id_in_block = tid_in_block / VECTOR_SIZE;
__shared__ int lds_y[VECTOR_NUM];
int K = 0;
for (int i = 0; i < N; i++) {
const int index = i * THREADS_PER_BLOCK + g_tid;
SH[tid_in_block] = x[index];
__syncthreads(); // label:sync1:
// reduce in vector
if (vec_id_in_block < THREADS_PER_BLOCK / REDUCE_SIZE) { // label1:
int sum = 0;
for (int j = 0; j < REDUCE_SIZE / VECTOR_SIZE; j++) { // label2:
const int lds_index = vec_id_in_block * REDUCE_SIZE + tid_in_vector + j * VECTOR_SIZE;
sum += SH[lds_index];
}
for (int j = VECTOR_SIZE >> 1; j > 0; j >>= 1) {
sum += __shfl_down(sum, j, VECTOR_SIZE);
}
// store sum value to y with memory coalescing
if (tid_in_vector == 0) { // label3:
lds_y[vec_id_in_block] = sum;
}
// }
__syncthreads(); // label:sync2:
// if (vec_id_in_block < THREADS_PER_BLOCK / REDUCE_SIZE) {
if (tid_in_block < THREADS_PER_BLOCK / REDUCE_SIZE) { // label4:
const int local_sum = lds_y[tid_in_block];
y[K + tid_in_block] = alpha * local_sum;
}
}
K += THREADS_PER_BLOCK / REDUCE_SIZE;
}
}
int main() {
constexpr int DATA_SIZE = THREADS_PER_BLOCK * N;
int *hx = new int[DATA_SIZE];
int *hy = new int[DATA_SIZE/REDUCE_SIZE];
for (int i = 0; i < DATA_SIZE; i++) {
hx[i] = i;
}
int *x = nullptr;
int *y = nullptr;
cudaMalloc(&x, DATA_SIZE * sizeof(int));
cudaMalloc(&y, DATA_SIZE / REDUCE_SIZE * sizeof(int));
cudaMemcpy(x, hx, DATA_SIZE * sizeof(int), cudaMemcpyHostToDevice);
test_kernel<<<1, THREADS_PER_BLOCK>>>(x, y, 1);
cudaDeviceSynchronize();
cudaMemcpy(hy, y, DATA_SIZE / REDUCE_SIZE * sizeof(int), cudaMemcpyDeviceToHost);
for (int i = 0; i < DATA_SIZE / REDUCE_SIZE; i++) {
// let R <- REDUCE_SIZE;
// hy[i] shoule be: R*(2*R*i+R-1)/2
int R = REDUCE_SIZE;
printf("%d\n", hy[i] == (R * (2 * R * i + R - 1) / 2));
}
}
HIP (Heterogeneous-Compute Interface for Portability) 是 AMD 开发的一款异构计算的接口工具。 HIP 允许只用写一套代码(hip代码), 就可以将程序同时在 NVIDIA GPU 和 AMD GPU 及 DCU 上编译运行。
HIP is a C++ Runtime API and Kernel Language that allows developers to create portable applications for AMD and NVIDIA GPUs from single source code.
HIP 的 API 和 CUDA 的API十分类似，例如 CUDA 中内存拷贝用cuMemory, 在 hip 中用hipMemcpy，且参数也十分一致。
因此，会 CUDA 的开发者可以很轻松地转移到 hip 上。
并且，hip 还提供了hipfy 工具，将 CUDA 代码转换为 hip 代码。
HIP 在不降低性能的前提下，统一了CUDA API 和AMD GPU 编程API，可谓极大地降低了各个平台的适配与移植工作，
做到了一套代码，在多个异构平台上运行。
可以说, "舍弃 CUDA，进入HIP时代"。
那么，在 NV GPU下，如何安装并使用 hip 呢？