CUDA 中不同块和线程的性能优化

Posted 2023-03-23

【中文标题】CUDA 中不同块和线程的性能优化

【英文标题】：Performance optimization with different blocks and threads in CUDA

【发布时间】：2015-02-12 21:26:11

【问题描述】：

我编写了一个程序来计算直方图，其中每个 char 字节的 256 个值都被计算在内：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "..\..\common\book.h"
#include <stdio.h> 
#include <cuda.h>
#include <conio.h>

#define SIZE (100*1024*1024)

__global__ void histo_kernel(unsigned char *buffer, long size, unsigned int *histo)

__shared__ unsigned int temp[256];
temp[threadIdx.x] = 0;
__syncthreads();

int i = threadIdx.x + blockIdx.x * blockDim.x;
int offset = blockDim.x * gridDim.x;

while (i < size) 
    atomicAdd(&temp[buffer[i]], 1);
    i += offset;

__syncthreads();
atomicAdd(&(histo[threadIdx.x]), temp[threadIdx.x]);


int main()



unsigned char *buffer = (unsigned char*)big_random_block(SIZE);

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);

unsigned char *dev_buffer;
unsigned int *dev_histo;
cudaMalloc((void**)&dev_buffer, SIZE);
cudaMemcpy(dev_buffer, buffer, SIZE, cudaMemcpyHostToDevice);

cudaMalloc((void**)&dev_histo, 256 * sizeof(long));
cudaMemset(dev_histo, 0, 256 * sizeof(int));

cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
int blocks = prop.multiProcessorCount; 

histo_kernel << <blocks * 256 , 256>> >(dev_buffer, SIZE, dev_histo);

unsigned int histo[256];
cudaMemcpy(&histo, dev_histo, 256 * sizeof(int), cudaMemcpyDeviceToHost);

cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapsed_time;
cudaEventElapsedTime(&elapsed_time, start, stop);
printf("Time to generate: %f ms\n", elapsed_time);

long sum = 0;
for (int i = 0; i < 256; i++)
    sum += histo[i];

printf("The sum is %ld", sum);

cudaFree(dev_buffer);
cudaFree(dev_histo);
free(buffer);

getch();
return 0;

我在书中读到，以 CUDA 为例，根据经验发现以多处理器数量两倍的块数启动内核是最佳解决方案。然而，当我以 8 倍的块数启动它时，运行时间被缩短了。

我已经运行内核：1.Blocks 与多处理器数量相同，2.Blocks 是多处理器数量的两倍，3.Blocks 4 倍，等等。

使用（1），我得到的运行时间是 112ms 使用（2）我得到的运行时间是 73ms 使用（3）我得到的运行时间是 52ms 有趣的是，在块数是多处理器数的 8 倍之后，运行时间并没有显着变化。就像块是多处理器数量的 8 倍、256 倍和 1024 倍一样。

如何解释？

【问题讨论】：

您是否使用blocks * 256 块启动内核？？

多个块可以驻留在多处理器上。对于 cc 3.0 及更高版本，每个 SM 的最大线程数为 2048，即 8 个块，每个 256 个线程。这与您的结果密切相关。

是的，在这个例子中，我使用 blocks * 256 blocks 来启动它。

void_ptr，当 SM 获得的线程数超过其运行能力时会发生什么？它会排队吗？

@user2580446 - 是的，这是在块级别完成的。块被安排在任何有足够可用资源的 SM 上执行。

【参考方案1】：

这种行为是典型的。 GPU 是一种隐藏延迟的机器。为了隐藏延迟，当它遇到停顿时，它需要额外的新工作可用。通过为 GPU 提供大量块和线程，您可以最大限度地增加可用的额外新工作量。

一旦你给它足够的工作来尽可能地隐藏延迟，再给它额外的工作也无济于事。机器饱和。然而，拥有额外的工作通常/通常也不会造成太大的损害。与块和线程相关的开销很小。

您在 CUDA by Example 中读到的任何内容对于特定情况可能都是正确的，但要启动的正确块数等于多处理器数量的两倍肯定不是普遍正确的。更好的目标（通常）是每个多处理器 4-8 个块。

对于块和线程，通常越多越好，而且在极少情况下，拥有任意数量的块和线程实际上会导致性能显着下降。这与典型的 CPU 线程编程相反，在典型的 CPU 线程编程中，当您超过核心数量时，拥有大量 OMP 线程可能会导致性能显着下降。

当您为最后 10% 的性能调整代码时，您会看到人们将他们启动的块数量限制在某个数字（通常是 SM 数量的 4-8 倍），并将他们的线程块构建为循环遍历数据集。但是，在大多数情况下，这通常只会带来百分之几的性能提升。作为一个合理的 CUDA 编程起点，至少要针对数万个线程和数百个块。仔细调整的代码可能能够用更少的块和线程使机器饱和，但此时它将变得依赖于 GPU。正如我已经说过的那样，拥有数百万个线程和数千个块几乎不会对性能造成太大影响。