c++ - 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 个 block ，每个 block 一个线程。这不是一种使用 GPU 的特别有效的方法，它意味着每个线程 block 中的 tid 变量都是 0(而且只有 0，因为每个线程 block 只有一个线程。)

这种方法存在很多问题，但这里会遇到一个直接的问题:

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层次结构:

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

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

你也没做对 cuda error checking这是推荐的，尤其是在您遇到 CUDA 代码问题时。

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

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

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

这行代码也不正确:

free (dev_a);

因为 dev_a 是使用 cudaMalloc 在设备上分配的，所以你应该像这样释放它:

cudaFree (dev_a);