RTX A5000的CUDA优化

Question

我的数据自动化系统的技能有点生疏。我试图生成一个完全由cuRand产生的随机噪声组成的5120x5120图像。我使用的是一个单一的NVIDIA A5000，其计算兼容性为8.6。我的问题是，什么应该是理想的网格尺寸和每个块的线程，以挤出最高数量的效率从这种噪音产生。对于上下文，下面是A5000的一些硬件规范：

-64 SM

-8192个数据自动化系统核心

-每SM 128个核心

-每SM 16个区块

-每毫升48根有效经纱(1536个螺纹)

Answer 1

对于基于最大化占用优化性能的一般考虑以及CURAND生成器初始化的考虑，您都希望每个块选择多个块和线程，以便产品最多等于每SM 64条SMs x 1536个线程。理想情况下，你应该瞄准这个数字。

首先编写接受已经初始化的生成器状态的内核，然后运行一个网格步长环来使用随机数生成来写入图像点。

如果占用率分析(可能基于每个线程寄存器)表明每个SM的最大线程负载是可能的(对于您刚刚编写的内核)，那么您可以这样调整网格的大小。

如果占用率分析表明内核不能拥有每个SM完整的1536个线程，那么您将相应地减少网格的启动(大小)。

因为cc8.6SM最多有1536个线程，所以不要在1024个线程上调整线程块的大小。选择512，或者其他类似256的数字。

对于CURAND，在我看来，最佳实践是在映像更新内核之前启动生成器状态初始化内核(首先是单独的)。不要尝试在执行映像生成的同一个内核中进行生成器状态初始化。

一旦您确定了最大占用率，那么您将调整生成器状态数组的大小以匹配它，然后启动一个网格步长内核来填充该数组作为CURAND init内核。

然后启动您的图像创建内核。

下面是一个简单的例子：

$ cat t2115.cu
#include <curand_kernel.h>
#include <curand.h>

const int imageW = 5120;
const int imageH = 5120;
using mt = float;

__global__ void setup_kernel(curandState *state, size_t N, const unsigned long long seed = 1, const unsigned long long offset = 0){

  for (size_t id = blockIdx.x * blockDim.x + threadIdx.x; id < N; id += gridDim.x*blockDim.x)
    curand_init(seed, id, offset, state+id);
}

template <typename T>
__global__ void image_gen(T *img, curandState *state, const size_t img_size){

  size_t id = blockIdx.x * blockDim.x + threadIdx.x;
  curandState s = state[id];
  for (;id < img_size; id += gridDim.x*blockDim.x)
    img[id] = curand_uniform(&s);
}

int main(){

  size_t is = ((size_t)imageW)*imageH;
  int grid_dim = 64 * 1536;
  int bs = 512;
  int gs = grid_dim/bs;
  mt *img;
  curandState *s;
  cudaMalloc(&s, sizeof(curandState)*grid_dim);
  cudaMallocManaged(&img, sizeof(mt)*is);
  setup_kernel<<<gs, bs>>>(s, grid_dim);
  image_gen<<<gs, bs>>>(img, s, is);
  cudaDeviceSynchronize();
}

$ nvcc -Xptxas=-v -arch=sm_86 -o t2115 t2115.cu
ptxas info    : 218048 bytes gmem, 72 bytes cmem[3]
ptxas info    : Compiling entry function '_Z9image_genIfEvPT_P17curandStateXORWOWm' for 'sm_86'
ptxas info    : Function properties for _Z9image_genIfEvPT_P17curandStateXORWOWm
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 18 registers, 376 bytes cmem[0]
ptxas info    : Compiling entry function '_Z12setup_kernelP17curandStateXORWOWmyy' for 'sm_86'
ptxas info    : Function properties for _Z12setup_kernelP17curandStateXORWOWmyy
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 39 registers, 384 bytes cmem[0]
$

我们使用-Xptxas=-v开关进行编译，这会导致编译器指示寄存器利用率。我们看到image_gen函数每个线程使用18个寄存器。这是远远低于限制什么将允许完全占用，所以我们应该与指定的发射大小，并应期望完全入住率。sm_86 SM 有支持高达65536寄存器，因此当跨1536个线程考虑时，这意味着每个线程的寄存器限制约为65536/1536 =42个。如果我们有一个比这个更大的数字，我们将相应地减少grid_dim变量。

这在很大程度上可以通过占用API和启动边界功能实现“自动化”，但首先应该理解这些基本思想。