【发布时间】：2021-10-27 18:34:13
【问题描述】：
您好，我正在尝试了解 cuda 内核的一些行为。这是我拥有的两个 cuda 内核。我发现gpuReduce 需要的持续时间是gpuReduceOpt 的两倍。是发散引起的吗？
#include <cuda_runtime.h>
#include <stdio.h>
#include <chrono>
void initData_int(int *p, int size){
for (int t=0; t<size; t++){
p[t] = (int)(rand()&0xff);
}
}
__global__ void gpuReduce(int *in, int *out, int size)
{
int tid = threadIdx.x;
int* data = in + blockIdx.x*blockDim.x;
if (tid >= size)
return;
for (int stride = 1; stride < blockDim.x; stride*=2)
{
if((tid%(2*stride)) == 0){
data[tid] += data[tid+stride];
}
__syncthreads();
}
if (tid == 0){
out[blockIdx.x] = data[0];
}
}
__global__ void gpuReduceOpt(int *in, int *out, int size)
{
int tid = threadIdx.x;
int* data = in + blockIdx.x*blockDim.x;
if (tid >= size)
return;
for (int stride = 1; stride < blockDim.x; stride*=2)
{
int index = 2*stride*tid;
if(index < blockDim.x){
data[index] += data[index+stride];
}
__syncthreads();
}
if (tid == 0){
out[blockIdx.x] = data[0];
}
}
int main(int agrc, char **argv)
{
int size = 1<<24;
int blocksize = 1024;
dim3 block(blocksize, 1);
dim3 grid((size-1)/block.x+1, 1);
int nBytes = sizeof(int)*size;
int *a_h = (int*)malloc(nBytes);
int *tmp = (int*)malloc(sizeof(int)*grid.x);
int *tmp1 = (int*)malloc(sizeof(int)*grid.x);
initData_int(a_h, size);
int *a_d, *output;
cudaMalloc((int**)&a_d, nBytes);
cudaMalloc((int**)&output, grid.x*sizeof(int));
int *a_d1, *output1;
cudaMalloc((int**)&a_d1, nBytes);
cudaMalloc((int**)&output1, grid.x*sizeof(int));
cudaMemcpy(a_d1, a_h, nBytes, cudaMemcpyHostToDevice);
cudaMemcpy(a_d, a_h, nBytes, cudaMemcpyHostToDevice);
auto start2 = std::chrono::system_clock::now();
gpuReduce<<<grid, block>>>(a_d, output, size);
cudaMemcpy(tmp, output, grid.x*sizeof(int), cudaMemcpyDeviceToHost);
int gpu_result;
for (int i =0; i < grid.x; i++)
{
gpu_result += tmp[i];
}
cudaDeviceSynchronize();
auto end2 = std::chrono::system_clock::now();
std::chrono::duration<double>diff2 = end2 - start2;
printf("Gpu reduce take:%2f s\n", diff2.count());
auto start3 = std::chrono::system_clock::now();
gpuReduceOpt<<<grid, block>>>(a_d1, output1, size);
cudaMemcpy(tmp1, output1, grid.x*sizeof(int), cudaMemcpyDeviceToHost);
int gpu_result1;
for (int i =0; i < grid.x; i++)
{
gpu_result1 += tmp1[i];
}
cudaDeviceSynchronize();
auto end3 = std::chrono::system_clock::now();
std::chrono::duration<double>diff3 = end3 - start3;
printf("Gpu reduce opt take:%2f s\n", diff3.count());
printf("Result from gpuReduce and gpuReduceOpt are %6d and %6d\n", gpu_result, gpu_result1);
cudaFree(a_d);
cudaFree(output);
free(a_h);
free(tmp);
cudaDeviceReset();
return 0;
}
这是我得到的性能数据：
Gpu reduce take:0.004238 s
Gpu reduce opt take:0.002606 s
Result from gpuReduce and gpuReduceOpt are 2139353471 and 2139353471
【问题讨论】：
-
这两个内核都包含未定义的行为，并且您没有解释任何关于您如何编译、运行和基准测试的内容，因此无法说出原因。猜测将是整数模运算符。但这是一个猜测
-
您的内核时序包括内存传输和主机操作。您还没有测量 内核 以了解它们的性能。并且永远不要通过运行一次来进行基准测试。运行多次（理想情况下执行几秒钟的挂钟），丢弃第一次调用和平均，因为这可能包括第一次在 GPU 上运行代码时发生的运行时开销，并且诸如帧刷新显示之类的事情可能会影响单个计时如果您使用的是显示 GPU
-
然后删除您不感兴趣的内容，这是创建minimal reproducible example 的一部分。此外，您发布的代码不会编译。此外，您的
gpuReduce和gpuReduceOpt不会产生相同的结果，那么为什么还要比较它们呢？gpuReduceOpt已损坏，如果您希望它产生与其他 2 次缩减相同的结果。如果您对快速并行缩减感兴趣，则不推荐使用这些扫描方案。使用canonical training中提供的那个。 -
嗨罗伯特。我删除了不相关的代码。我确实验证了
gpuReduce和gpuReduceOpt的结果是一样的。是的，我还实现了另一个交错减少，它要快得多。我只是好奇上面两个内核的性能改进源。 -
第一个使用慢模，第二个不使用。内存访问也很重要（我不仅仅是它们相似）。请注意，跨步访问通常很慢。
标签： cuda