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

Posted 2023-03-23

【中文标题】CUDA中的一个简单的缩减程序

【英文标题】：A simple reduction program in CUDA

【发布时间】：2015-09-16 02:44:57

【问题描述】：

在下面的代码中，我试图实现一个简单的并行缩减，块大小和每个块的线程数为 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]);

【参考方案1】：

这段代码有很多问题。

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

任何时候您在使用 CUDA 代码时遇到问题，您都应该使用proper cuda error checking 并使用cuda-memcheck 运行您的代码在寻求帮助之前。如果您这样做了，您的注意力就会集中在下面的第 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);

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

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

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

__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
$

【讨论】：

对于使用 global 和 shared 的部分，这是一个复制粘贴问题，并且在我的本地文件中。感谢我犯了不以大小形式传递字节的错误部分。至于第 5 条评论，我为函数提供了一个额外的参数，用于存储输出，并将输入设为 const。它的工作！谢谢

你好，是否可以仅将值减少技术用于增加变量？

我不知道仅用于增加变量的值减少技术是什么。