HIP编程

这篇文章主要讲共享内存

国内互联网博客中，讲CUDA的很多，但讲AMD显卡的HIP编程语言的极少，刚好前段时间使用过国内超算平台，这个系列就来讲讲HIP语言。

1. 什么是HIP语言及相关有用查询网址；

2.怎么将CUDA转换为HIP语言；

3.共享内存

 

使用参考网址：

http://rocm-developer-tools.github.io/HIP/group__API.html

https://rocmdocs.amd.com/en/latest/Installation_Guide/HIP.html

https://github.com/ROCm-Developer-Tools/HIP

书籍推荐：

GPU高性能编程CUDA实战——有中文版的

---

1.核函数的调用

HIP语言中，调用核函数的函数为hipLaunchKernelGGL，例子如下：

hipLaunchKernelGGL(kernel_phase2, dim3(round, 2), dim3(blockSize, blockSize), blockSize * blockSize * sizeof(int) * 2, 0, r, num_node, pathDistanceBuffer, pathBuffer);

参数1：核函数的名称；

参数2：gird的大小，和CUDA一样是dim3类型；

参数3：block的大小；

参数4：所需分配的共享内存的大小，默认可以写为0；

参数5：

参数5以后：核函数的所带的参数；

1.1 关于第三个参数--共享内存动态申请

当需要动态申请共享内存大小时，首先在共享内存前加extern关键字

1	extern __shared__ float a[];

然后，在调用内核函数的时候，需要使用第三个参数来指明所需分配的共享内存的字节大小。

const size_t smemSize=(TPB+ 2*RAD)*sizeof(float);
hipLaunchKernelGGL(ddkernel, Grids, Blocks, smemSize, 0, paramenter);

注意：

1. 这里的paramenter里包含多个核函数对应的参数；

2.既然叫共享内存，就需要明白，一个block中的各个线程是共用这个内存的，赋值的时候需要注意多线程的逻辑，防止出错；

3.一般是按照线程号，将共享内存进行划分；

4.注意使用__syncthreads();函数，如果一个线程中需要使用其他线程中赋值的共享内存，需要等待线程同步；这里__syncthreads();函数一定要确保所有线程都能运行到，不能耍小聪明，只让需要同步的线程执行这个函数，否则会出现所有线程无限等待下去的情况。（其实也是NVIDIA的现实硬件考虑）

关于共享内存使用的文章链接如下：

https://www.cnblogs.com/cofludy/p/7622254.html

https://blog.csdn.net/q583956932/article/details/78872146

共享内存（Shared memory）是位于每个流处理器组（SM）中的高速内存空间，主要作用是存放一个线程块（Block）中所有线程都会频繁访问的数据。流处理器（SP）访问它的速度仅比寄存器（Register）慢，它的速度远比全局显存快。但是他也是相当宝贵的资源，一般只有几十KByte。

1.2 核函数的参数存放地址

kernel的参数是存放在常量存储器中的。或者精确的说，存放在显存中，并被常量缓存缓冲。

2.核函数中定义的变量

局部变量：存在寄存器中 
静态数组：存在本地内存（栈）中 
动态变量：存在全局内存中 
__shared__（静态）：存在共享内存中

2.1 局部变量

如果在global核函数中定义了局部变量，其是存在寄存器中的（注意：寄存器的大小是有限的，所以不能定义占用空间太多的局部变量）。

int x = 45;

 

2.2 全局变量

 

kernel内部使用new[], delete[]分配存储空间。

int* array = new int[num];
delete[] array;

 

2.3 共享内存

 

main函数外：
const int blockSize = 32;

核函数中：
__shared__ float shared_dist[blockSize * blockSize];

 

 

6.遇到的问题

6.1 需要打印的数据太多，希望打印到txt中

可以在运行时加入重定向符号，如通过编译生成了tep执行体。

./tep > output.txt

6.2 使用如下的hipMemcpy函数导致内存超出

这个问题还没找到原因，我才定义了6个4G大小的矩阵，运行起来就提示内存超了或者段错误

hipMemcpy(din, arc, sizeof(float) * num_node * num_node, hipMemcpyHostToDevice);