CUDA 共享内存编程不起作用

Question

全部：

我正在学习共享内存如何加速 GPU 编程过程。我正在使用下面的代码来计算每个元素的平方值加上其左右邻居的平均值的平方值。 代码运行了，结果却不如预期。

打印出来的前 10 个结果是 0,1,2,3,4,5,6,7,8,9，而我期望的结果是 25,2,8, 18,32,50,72 ,98,128,162;

代码如下，参考here；

你能告诉我哪一部分出了问题吗？非常感谢您的帮助。

#include <stdio.h>
#include <stdlib.h>
#include <iostream>
#include <cuda.h>

const int N=1024;

 __global__ void compute_it(float *data)
 {
 int tid = threadIdx.x;
 __shared__ float myblock[N];
 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:(N-1)] + myblock[tid<(N-1)?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] = myblock[tid];
 __syncthreads();
  }

int main (){

char key;

float *a;
float *dev_a;

a = (float*)malloc(N*sizeof(float));
cudaMalloc((void**)&dev_a,N*sizeof(float));

for (int i=0; i<N; i++){
a [i] = i;
}


cudaMemcpy(dev_a, a, N*sizeof(float), cudaMemcpyHostToDevice);

compute_it<<<N,1>>>(dev_a);

cudaMemcpy(a, dev_a, N*sizeof(float), cudaMemcpyDeviceToHost);


for (int i=0; i<10; i++){
std::cout<<a [i]<<",";
}

std::cin>>key;

free (a);
free (dev_a);

Answer 1

内核代码中最直接的问题之一是：

data[tid] = myblock[tid];

我想你可能是这个意思：

data[tid] = tmp;

此外，您将启动 1024 个块，每个块包含一个线程。这不是使用 GPU 的一种特别有效的方法，这意味着每个线程块中的 tid 变量为 0（并且只有 0，因为每个线程块只有一个线程。）

这种方法有很多问题，但这里会遇到一个直接的问题：

tmp = (myblock[tid>0?tid-1:(N-1)] + myblock[tid<31?tid+1:0]) * 0.5f;

由于tid 始终为零，因此共享内存数组 (myblock) 中没有其他值被填充，因此该行中的逻辑不合理。当tid 为零时，您选择myblock[N-1] 作为分配给tmp 的第一项，但myblock[1023] 永远不会填充任何内容。

看来你对各种CUDA层次结构不太了解：

• 网格是与内核启动相关的所有线程
• 网格由线程块组成
• 每个线程块是一组在单个 SM 上一起工作的线程
• 共享内存资源是每个 SM 资源，而不是设备范围的资源
• __synchthreads() 也在线程块基础上运行（不是设备范围的）
• threadIdx.x 是一个内置变量，它为线程块内的所有线程提供唯一的线程 ID，但不是在整个网格中全局提供。

相反，您应该将问题分解为大小合理的线程块组（即多个线程）。然后，每个线程块将能够以大致如您所概述的方式运行。然后，您需要对每个线程块的起点和终点（在您的数据中）的行为进行特殊处理。

你也没有做正确的cuda error checking，这是推荐的，尤其是当你遇到 CUDA 代码问题时。

如果您在内核代码中进行了我首先指出的更改，并颠倒您的块和网格内核启动参数的顺序：

compute_it<<<1,N>>>(dev_a);

我认为，正如 Kristof 所说，您会得到接近您想要的东西。但是，如果不对代码进行其他更改，您将无法方便地将其扩展到 N=1024 以上。

这行代码也不正确：

free (dev_a);

由于dev_a 是使用cudaMalloc 在设备上分配的，您应该像这样释放它：

cudaFree (dev_a);

Answer 2

由于每个块只有一个线程，因此您的 tid 将始终为 0。

尝试以这种方式启动内核： 计算_it>>(dev_a);

而不是 计算_it>(dev_a);