增加元素 CUDA 内核的算术强度的技术

Posted 2023-03-23

【中文标题】增加元素 CUDA 内核的算术强度的技术

【英文标题】：Techniques for increasing arithmetic intensity of elementwise CUDA kernel

【发布时间】：2021-05-16 07:03:56

【问题描述】：

我编写了一个 CUDA 内核，用于计算一组源粒子和目标粒子之间的成对相互作用。

我的 M 目标粒子看起来像，

[[x_1,y_1,z_1],...,[x_m, y_m, z_m]]

而我的 N 个源粒子看起来像

[[x_1,y_1,z_1],...,[x_n, y_n, z_n]]

与 M

我首先将所有源/目标数据传输到 GPU，然后循环处理成批的源粒子并评估与所有目标的成对交互。

像这样（在cupy语法中）

for i in range(n_blocks):
    left_idx = i*width
    right_idx = (i+1)*width
    gpu_func(
       grid_dimensions, block_dimensions, targets, 
       sources[left_idx:right_idx,:], width
    )

其中sources 和targets 是GPU 上包含源数据和目标数据的数组。

我的问题是我必须做些什么来避免这个循环？我是 CUDA 的新手。我的想法是检查来自每个线程的全局线程索引是否满足我在主机设备循环中的“left_idx”和“right_index”条件，这是正确的吗？有没有更好的方法来做到这一点？我觉得我目前对 GPU 的利用严重不足，因为源/目标的数量明显少于我机器上的 CUDA 核心数量。

【问题讨论】：

建议如何增加核的算术强度需要核的知识。对内核的一句话描述不足以提供这样的建议。代码很重要

【参考方案1】：

我的问题是我必须做些什么来避免这个循环？我是 CUDA 的新手。

并行编程（CUDA 或其他）的一个基本概念是，您让不同的处理元素/线程/自动机/任何东西并行处理所有“循环迭代”，而不是按时间顺序发生的循环：N 个处理元素每个人做 1 件工作，而不是 1 个处理元件使用循环做 N 件工作。

在 CUDA 中，在 M 和 N 的情况下，这可能意味着有一个 M x N 的 2D 网格，每个 CUDA 线程处理一对源和目标粒子集。或者更小的网格，每个 CUDA 线程在几对上运行一个循环，但对更少。

您可能还记得 CUDA vectorAdd 示例，它计算两个向量的元素相加。这是一个可能的内核：

__global__ void vecAdd(int *A, int *B, int *C, int N)

   int i = blockIdx.x * blockDim.x + threadIdx.x;
   C[i] = A[i] + B[i]; 

你看到了吗？没有任何循环。正是网格中的大量线程确保了每个可能的 i（介于 0 和 N-1 之间）都计算了 C 的相应元素。