0

全部：

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

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

代码如下，参考这里；

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

#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);
4

2 回答 2

4

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

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 错误检查，尤其是在您遇到 CUDA 代码问题时。

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

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

正如克里斯托夫所指出的，我认为你会得到接近你想要的东西。但是，如果不对代码进行其他更改，您将无法方便地将其扩展到 N=1024 之外。

这行代码也不正确：

free (dev_a);

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

cudaFree (dev_a);
于 2013-11-08T17:37:35.510 回答
1

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

尝试以这种方式启动内核：compute_it<<<1,N>>>(dev_a);

而不是 compute_it<<>>(dev_a);

于 2013-11-08T18:14:30.997 回答