HIP编程

Dear 丶 2023-02-28 01:42 17阅读 0赞

这篇文章主要讲共享内存

国内互联网博客中,讲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,例子如下:

  1. 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以后:核函数的所带的参数;

watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L2tpc3Nnb29kYnllMjAxMg_size_16_color_FFFFFF_t_70

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

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








1

extern shared float a[];

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

  1. const size_t smemSize=(TPB+ 2*RAD)*sizeof(float);
  2. 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核函数中定义了局部变量,其是存在寄存器中的(注意:寄存器的大小是有限的,所以不能定义占用空间太多的局部变量)。

  1. int x = 45;

2.2 全局变量

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

  1. int* array = new int[num];
  2. delete[] array;

2.3 共享内存

  1. main函数外:
  2. const int blockSize = 32;
  3. 核函数中:
  4. __shared__ float shared_dist[blockSize * blockSize];

6.遇到的问题

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

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

  1. ./tep > output.txt

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

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

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

发表评论

表情:
评论列表 (有 0 条评论,17人围观)

还没有评论,来说两句吧...

相关阅读

    相关 HIP编程

    > 这篇文章主要讲共享内存 国内互联网博客中,讲CUDA的很多,但讲AMD显卡的HIP编程语言的极少,刚好前段时间使用过国内超算平台,这个系列就来讲讲HIP语言。 1. 什

    相关 Socket编程——UDP编程

    UDP协议(用户数据协议) 是无连接、不可靠的、无序的 特点在于速度比较快 UDP协议以数据报作为数据传输的载体 进行数据传输时,首先需要将要传输的数据定义成数据报(

    相关 编程编程思想

    \[TIP\]什么是编程思想?答案可能很会复杂,但也可以很简单。一句话来讲就是,用计算机来解决人们实际问题的思维方式,即编程思想。    ![2011121411044659

    相关 编程编程语言

    转自计蒜客 编程就是人类让计算机为解决某个问题而使用某种程序设计语言编写程序代码,并最终得到结果的过程。为了使计算机能够理解人的意图,人类就必须要将解决的问题的