【发布时间】：2019-06-03 02:09:13
【问题描述】：
以下主机代码test.c 和设备代码test0.cu 旨在提供相同的结果。
test.c
$ cat test.c
#include <stdio.h>
#include <string.h>
int main()
{
int data[32];
int dummy[32];
for (int i = 0; i < 32; i++)
data[i] = i;
memcpy(dummy, data, sizeof(data));
for (int i = 1; i < 32; i++)
data[i] += dummy[i - 1];
memcpy(dummy, data, sizeof(data));
for (int i = 2; i < 32; i++)
data[i] += dummy[i - 2];
memcpy(dummy, data, sizeof(data));
for (int i = 4; i < 32; i++)
data[i] += dummy[i - 4];
memcpy(dummy, data, sizeof(data));
for (int i = 8; i < 32; i++)
data[i] += dummy[i - 8];
memcpy(dummy, data, sizeof(data));
for (int i = 16; i < 32; i++)
data[i] += dummy[i - 16];
printf("kernel : ");
for (int i = 0; i < 32; i++)
printf("%4i ", data[i]);
printf("\n");
}
$
test0.cu
$ cat test0.cu
#include <stdio.h>
__global__ void kernel0(int *data)
{
size_t t_id = threadIdx.x;
if (1 <= t_id)
data[t_id] += data[t_id - 1];
if (2 <= t_id)
data[t_id] += data[t_id - 2];
if (4 <= t_id)
data[t_id] += data[t_id - 4];
if (8 <= t_id)
data[t_id] += data[t_id - 8];
if (16 <= t_id)
data[t_id] += data[t_id - 16];
}
int main()
{
int data[32];
int result[32];
int *data_d;
cudaMalloc(&data_d, sizeof(data));
for (int i = 0; i < 32; i++)
data[i] = i;
dim3 gridDim(1);
dim3 blockDim(32);
cudaMemcpy(data_d, data, sizeof(data), cudaMemcpyHostToDevice);
kernel0<<<gridDim, blockDim>>>(data_d);
cudaMemcpy(result, data_d, sizeof(data), cudaMemcpyDeviceToHost);
printf("kernel0 : ");
for (int i = 0; i < 32; i++)
printf("%4i ", result[i]);
printf("\n");
}
$
如果我编译并运行它们，它们会给出与我预期相同的结果。
$ gcc -o test test.c
$ ./test
kernel : 0 1 3 6 10 15 21 28 36 45 55 66 78 91 105 120 136 153 171 190 210 231 253 276 300 325 351 378 406 435 465 496
$ nvcc -o test_dev0 test0.cu
$ ./test_dev0
kernel0 : 0 1 3 6 10 15 21 28 36 45 55 66 78 91 105 120 136 153 171 190 210 231 253 276 300 325 351 378 406 435 465 496
$
但是，如果我在设备代码中使用共享内存而不是全局内存，如test1.cu，则会给出不同的结果。
test1.cu
$ cat test1.cu
#include <stdio.h>
__global__ void kernel1(int *data)
{
__shared__ int data_s[32];
size_t t_id = threadIdx.x;
data_s[t_id] = data[t_id];
if (1 <= t_id)
data_s[t_id] += data_s[t_id - 1];
if (2 <= t_id)
data_s[t_id] += data_s[t_id - 2];
if (4 <= t_id)
data_s[t_id] += data_s[t_id - 4];
if (8 <= t_id)
data_s[t_id] += data_s[t_id - 8];
if (16 <= t_id)
data_s[t_id] += data_s[t_id - 16];
data[t_id] = data_s[t_id];
}
int main()
{
int data[32];
int result[32];
int *data_d;
cudaMalloc(&data_d, sizeof(data));
for (int i = 0; i < 32; i++)
data[i] = i;
dim3 gridDim(1);
dim3 blockDim(32);
cudaMemcpy(data_d, data, sizeof(data), cudaMemcpyHostToDevice);
kernel1<<<gridDim, blockDim>>>(data_d);
cudaMemcpy(result, data_d, sizeof(data), cudaMemcpyDeviceToHost);
printf("kernel1 : ");
for (int i = 0; i < 32; i++)
printf("%4i ", result[i]);
printf("\n");
}
$
如果我编译 test1.cu 并运行它，它会给出与 test0.cu 或 test.c 不同的结果。
$ nvcc -o test_dev1 test1.cu
$ ./test_dev1
kernel1 : 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31
$
warp 同步不应该与共享内存一起使用吗？
对此问题的一些调查：
使用 CUDA8.0 时，如果我使用 -arch=sm_61 选项编译 test1.cu（我正在使用 GTX 1080 进行测试），它会得到与 test0.cu 和 test.c 相同的结果。
$ nvcc -o test_dev1_arch -arch=sm_61 test1.cu
$ ./test_dev1_arch
kernel1 : 0 1 3 6 10 15 21 28 36 45 55 66 78 91 105 120 136 153 171 190 210 231 253 276 300 325 351 378 406 435 465 496
$
但这不适用于较新版本的 CUDA。如果我使用任何比 8.0 更新的版本，即使我给出了-arch=sm_61 选项，测试结果也会有所不同。
【问题讨论】：
标签： cuda gpu-shared-memory gpu-warp