CUDA中的一个简单的缩减程序

Question

在下面的代码中，我试图实现一个简单的并行缩减，块大小和每个块的线程数为 1024。但是，在实现部分缩减之后，我希望看看我的实现是否正确过程 我让程序打印主机内存的第一个元素（在数据从设备内存复制到主机内存之后）。 我的主机内存初始化为“1”，并被复制到设备内存以减少内存。并且缩减过程之后的 printf 语句仍然在数组的第一个元素处给我 '1'。

我要打印的内容是否有问题，或者在减少的实施中是否符合逻辑？ 此外，内核中的 printf 语句不会打印任何内容。我的语法或对 printf 语句的调用有问题吗？ 我的代码如下：

    ifndef CUDACC
define CUDACC
endif
include "cuda_runtime.h"
include "device_launch_parameters.h"
include
include
ifndef THREADSPERBLOCK
define THREADSPERBLOCK 1024
endif
ifndef NUMBLOCKS
define NUMBLOCKS 1024
endif

global void reduceKernel(int *c)
{
extern shared int sh_arr[];

int index = blockDim.x*blockIdx.x + threadIdx.x;
int sh_index = threadIdx.x;

// Storing data from Global memory to shared Memory
sh_arr[sh_index] = c[index];
__syncthreads();

for(unsigned int i = blockDim.x/2; i>0 ; i>>=1)
{
    if(sh_index < i){
        sh_arr[sh_index] += sh_arr[i+sh_index];
    }
    __syncthreads();
}

if(sh_index ==0)
    c[blockIdx.x]=sh_arr[sh_index];
printf("value stored at %d is %d \n", blockIdx.x, c[blockIdx.x]);
return;

}

int main()
{
int *h_a;
int *d_a;
int share_memSize, h_memSize;
size_t d_memSize;

share_memSize = THREADSPERBLOCK*sizeof(int);
h_memSize = THREADSPERBLOCK*NUMBLOCKS;

h_a = (int*)malloc(sizeof(int)*h_memSize);

d_memSize=THREADSPERBLOCK*NUMBLOCKS;
cudaMalloc( (void**)&d_a, h_memSize*sizeof(int));

for(int i=0; i<h_memSize; i++)
{
    h_a[i]=1;    
};

//printf("last element of array %d \n", h_a[h_memSize-1]);

cudaMemcpy((void**)&d_a, (void**)&h_a, h_memSize, cudaMemcpyHostToDevice);
reduceKernel<<<NUMBLOCKS, THREADSPERBLOCK, share_memSize>>>(d_a);
cudaMemcpy((void**)&h_a, (void**)&d_a, d_memSize, cudaMemcpyDeviceToHost);

printf("sizeof host memory %d \n", d_memSize); //sizeof(h_a));
printf("sum after reduction %d \n", h_a[0]);

}

Answer 1

这段代码有很多问题。

1. 您发布的大部分内容都是无效代码。举几个例子，你的 global 和 shared 关键字应该有双下划线前后，像这样：__global__ 和 __shared__。我认为这是某种复制粘贴错误或格式错误。您的定义语句也存在问题。你应该努力发布没有这些问题的代码。

2. 任何时候您在使用 CUDA 代码时遇到问题，您都应该使用proper cuda error checking 并使用cuda-memcheck 运行您的代码在寻求帮助之前。如果您这样做了，您的注意力就会集中在下面的第 3 项上。

3. 您的cudaMemcpy 操作在几个方面被破坏：

   cudaMemcpy((void**)&d_a, (void**)&h_a, h_memSize, cudaMemcpyHostToDevice);
   

   首先，不像cudaMalloc，但像memcpy，cudaMemcpy只接受普通的指针参数。其次，传输的大小（如memcpy）是以bytes为单位的，所以你的大小需要放大sizeof(int)：

   cudaMemcpy(d_a, h_a, h_memSize*sizeof(int), cudaMemcpyHostToDevice);
   

   内核之后的那个也是类似的。

4. printf 来自大型内核中的每个线程（例如具有 1048576 个线程的内核）可能不是一个好主意。您实际上不会得到您期望的所有输出，并且在 Windows 上（似乎您正在 Windows 上运行）由于内核执行时间过长，您可能会遇到 WDDM 看门狗超时。如果您需要从大型内核中使用 printf，请谨慎选择并将您的 printf 设置为 threadIdx.x 和 blockIdx.x

5. 以上内容可能足以获得一些合理的打印输出，并且正如您指出的那样，您还没有完成：“我希望看看我的实现是否正确”。但是，精心设计的这个内核会用输出数据覆盖其输入数据：

   __global__ void reduceKernel(int *c)
   ...
       c[blockIdx.x]=sh_arr[sh_index];
   

   这将导致竞争条件。我建议您将输出数据与输入数据分开，而不是试图为您解决这个问题。更好的是，你应该研究cuda reduction sample code，它也有一个关联的presentation。

这是您的代码的修改版本，已修复上述大部分问题。 仍然不正确。它仍然存在上面的缺陷 5。 我不会完全重写代码来修复缺陷 5，而是将您定向到上面提到的 cuda 示例代码。

$ cat t820.cu
#include <stdio.h>

#ifndef THREADSPERBLOCK
#define THREADSPERBLOCK 1024
#endif
#ifndef NUMBLOCKS
#define NUMBLOCKS 1024
#endif

__global__ void reduceKernel(int *c)
{
extern __shared__ int sh_arr[];

int index = blockDim.x*blockIdx.x + threadIdx.x;
int sh_index = threadIdx.x;

// Storing data from Global memory to shared Memory
sh_arr[sh_index] = c[index];
__syncthreads();

for(unsigned int i = blockDim.x/2; i>0 ; i>>=1)
{
    if(sh_index < i){
        sh_arr[sh_index] += sh_arr[i+sh_index];
    }
    __syncthreads();
}

if(sh_index ==0)
    c[blockIdx.x]=sh_arr[sh_index];
// printf("value stored at %d is %d \n", blockIdx.x, c[blockIdx.x]);
return;

}

int main()
{
int *h_a;
int *d_a;
int share_memSize, h_memSize;
size_t d_memSize;

share_memSize = THREADSPERBLOCK*sizeof(int);
h_memSize = THREADSPERBLOCK*NUMBLOCKS;

h_a = (int*)malloc(sizeof(int)*h_memSize);

d_memSize=THREADSPERBLOCK*NUMBLOCKS;
cudaMalloc( (void**)&d_a, h_memSize*sizeof(int));

for(int i=0; i<h_memSize; i++)
{
    h_a[i]=1;
};

//printf("last element of array %d \n", h_a[h_memSize-1]);

cudaMemcpy(d_a, h_a, h_memSize*sizeof(int), cudaMemcpyHostToDevice);
reduceKernel<<<NUMBLOCKS, THREADSPERBLOCK, share_memSize>>>(d_a);
cudaMemcpy(h_a, d_a, d_memSize*sizeof(int), cudaMemcpyDeviceToHost);

printf("sizeof host memory %d \n", d_memSize); //sizeof(h_a));
printf("first block sum after reduction %d \n", h_a[0]);
}
$ nvcc -o t820 t820.cu
$ cuda-memcheck ./t820
========= CUDA-MEMCHECK
sizeof host memory 1048576
first block sum after reduction 1024
========= ERROR SUMMARY: 0 errors
$