我开始学习 CUDA,我想写一个简单的程序,将一些数据复制到 GPU,修改它,然后将它传回。我已经用谷歌搜索并试图找出我的错误。我很确定问题出在我的内核中,但我不完全确定哪里出了问题。
这是我的内核:
__global__ void doStuff(float* data, float* result)
{
if (threadIdx.x < 9) // take the first 9 threads
{
int index = threadIdx.x;
result[index] = (float) index;
}
}
下面是我的 main
的相关部分:
#include <stdlib.h>
#include <stdio.h>
int main(void)
{
/*
Setup
*/
float simple[] = {-1.0, -2.0, -3.0, -4.0, -5.0, -6.0, -7.0, -8.0, -9.0};
float* data_array;
float* result_array;
size_t data_array_pitch, result_array_pitch;
int width_in_bytes = 3 * sizeof(float);
int height = 3;
/*
Initialize GPU arrays
*/
cudaMallocPitch(&data_array, &data_array_pitch, width_in_bytes, height);
cudaMallocPitch(&result_array, &result_array_pitch, width_in_bytes, height);
/*
Copy data to GPU
*/
cudaMemcpy2D(data_array, data_array_pitch, simple, width_in_bytes, width_in_bytes, height, cudaMemcpyHostToDevice);
dim3 threads_per_block(16, 16);
dim3 num_blocks(1,1);
/*
Do stuff
*/
doStuff<<<num_blocks, threads_per_blocks>>>(data_array, result_array);
/*
Get the results
*/
cudaMemcpy2D(simple, width_in_bytes, result_array, result_array_pitch, width_in_bytes, height, cudaMemcpyDeviceToHost);
for (int i = 1; i <= 9; ++i)
{
printf("%f ", simple[i-1]);
if(!(i%3))
printf("\n");
}
return 0;
}
当我运行它时,我得到第一行的 0.000000 1.000000 2.00000
和其他两行的垃圾。
最佳答案
如果您刚开始学习 cuda,我不确定我是否会专注于二维数组。
如果您在问题中手动输入代码,也很好奇,因为您定义了一个 threads_per_block
变量,但随后您在内核调用中使用了 threads_per_blocks
。
无论如何,您的代码有几个问题:
- 使用二维数组时,几乎总是需要通过 pitch
参数(以某种方式)到内核。
cudaMallocPitch
在每一行的末尾分配带有额外填充的数组,以便 下一行从一个很好对齐的边界开始。这通常会 导致分配粒度为 128 或 256 字节。所以你的第一个 行有 3 个有效数据实体,后跟足够的空白空间来填充 向上,比如说 256 字节(等于你的 pitch 变量是什么)。因此,我们必须更改内核调用和内核本身来解决这个问题。 - 您的内核本质上是一维内核(例如,它不理解或使用
threadIdx.y
)。因此,启动 2D 网格没有意义。虽然在这种情况下它不会造成任何伤害,但它会产生冗余,这在其他代码中可能会造成混淆和麻烦。
根据上述评论,这是一个更新的代码,显示了一些可以为您带来预期结果的更改:
#include <stdio.h>
__global__ void doStuff(float* data, float* result, size_t dpitch, size_t rpitch, int width)
{
if (threadIdx.x < 9) // take the first 9 threads
{
int index = threadIdx.x;
result[((index/width)*(rpitch/sizeof(float)))+ (index%width)] = (float) index;
}
}
int main(void)
{
/*
Setup
*/
float simple[] = {-1.0, -2.0, -3.0, -4.0, -5.0, -6.0, -7.0, -8.0, -9.0};
float* data_array;
float* result_array;
size_t data_array_pitch, result_array_pitch;
int height = 3;
int width = 3;
int width_in_bytes = width * sizeof(float);
/*
Initialize GPU arrays
*/
cudaMallocPitch(&data_array, &data_array_pitch, width_in_bytes, height);
cudaMallocPitch(&result_array, &result_array_pitch, width_in_bytes, height);
/*
Copy data to GPU
*/
cudaMemcpy2D(data_array, data_array_pitch, simple, width_in_bytes, width_in_bytes, height, cudaMemcpyHostToDevice);
dim3 threads_per_block(16);
dim3 num_blocks(1,1);
/*
Do stuff
*/
doStuff<<<num_blocks, threads_per_block>>>(data_array, result_array, data_array_pitch, result_array_pitch, width);
/*
Get the results
*/
cudaMemcpy2D(simple, width_in_bytes, result_array, result_array_pitch, width_in_bytes, height, cudaMemcpyDeviceToHost);
for (int i = 1; i <= 9; ++i)
{
printf("%f ", simple[i-1]);
if(!(i%3))
printf("\n");
}
return 0;
}
您可能还会找到 this question有趣的阅读。
编辑:回答评论中的问题:
result[((index/width)*(rpitch/sizeof(float)))+ (index%width)] = (float) index;
1 2 3
要计算倾斜数组中正确的元素索引,我们必须:
- 根据线程索引计算(虚拟)行索引。为此,我们将线程索引除以每个(非间距)行的宽度(以元素为单位,而不是字节)。
- 将行索引乘以每个间距 行的宽度。每个 pitched 行的宽度由 pitched 参数给出,以字节为单位。要将这个倾斜的 byte 参数转换为倾斜的 element 参数,我们除以每个元素的大小。然后通过将数量乘以在步骤 1 中计算的行索引,我们现在已经索引到正确的行。
- 通过线程索引除以宽度(以元素为单位)的余数(模除法),从线程索引计算(虚拟)列索引。一旦我们有了列索引(在元素中),我们就将它添加到在步骤 2 中计算的正确行开始索引中,以确定该线程将负责的元素。
以上是相对简单操作的相当多的努力,这就是为什么我建议首先关注基本 cuda 概念而不是倾斜数组的一个例子。例如,在处理倾斜阵列之前,我会想出如何处理 1 维和 2 维线程 block 以及 1 维和 2 维网格。在某些情况下,倾斜数组是访问 2D 数组(或 3D 数组)的有用的性能增强器,但它们绝不是处理 CUDA 中的多维数组所必需的。
关于c++ - CUDA:二维数组索引给出意想不到的结果,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/16905899/