c - CUDA的__shared__内存什么时候有用?

标签 c cuda gpu

有人可以帮我举一个关于如何使用共享内存的非常简单的例子吗? Cuda C 编程指南中包含的示例似乎被不相关的细节弄得一团糟。

例如,如果我将一个大数组复制到设备全局内存并想对每个元素进行平方,如何使用共享内存来加快速度?还是在这种情况下没有用?

最佳答案

在您提到的特定情况下,共享内存没有用,原因如下:每个数据元素仅使用一次。要使共享内存有用,您必须多次使用传输到共享内存的数据,使用良好的访问模式,才能发挥作用。原因很简单:仅从全局内存读取需要 1 次全局内存读取和零次共享内存读取;首先将其读入共享内存需要 1 次全局内存读取和 1 次共享内存读取,这需要更长的时间。

这是一个简单的例子,其中 block 中的每个线程计算相应的值,平方,加上它的左右邻居的平均值,平方:

  __global__ void compute_it(float *data)
  {
     int tid = threadIdx.x;
     __shared__ float myblock[1024];
     float tmp;

     // load the thread's data element into shared memory
     myblock[tid] = data[tid];

     // ensure that all threads have loaded their values into
     // shared memory; otherwise, one thread might be computing
     // on unitialized data.
     __syncthreads();

     // compute the average of this thread's left and right neighbors
     tmp = (myblock[tid > 0 ? tid - 1 : 1023] + myblock[tid < 1023 ? tid + 1 : 0]) * 0.5f;
     // square the previousr result and add my value, squared
     tmp = tmp*tmp + myblock[tid] * myblock[tid];

     // write the result back to global memory
     data[tid] = tmp;
  }

请注意,这预计仅使用一个 block 即可工作。扩展到更多 block 应该很简单。假设 block 维度 (1024, 1, 1) 和网格维度 (1, 1, 1)。

关于c - CUDA的__shared__内存什么时候有用?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/8011376/

相关文章:

c - 将空格替换为 %20,但不打印最后一个单词

c - LabView 与 WM_Copydata 的通信

c++ - 如何将其他 .wo、.so、.a 文件链接到当前 cuda 编译

c++ - cuda header 包含失败

cuda - 将 cuda 计算模式切换为默认模式

c++ - Tensorflow C++ 不使用 GPU

c - 仅在将数组设为私有(private)后才进行多线程加速

java - 是否有用于加速 vector 计算的 Java 库?

c++ - 基于 cuda 推力的方法对 TCP 流中的数据包进行分组

多个文件错误 : dereferencing pointer to incomplete type 中的 c 结构