CUDA内核启动参数解释正确吗？

Posted 2023-03-23

【中文标题】CUDA内核启动参数解释正确吗？

【英文标题】：CUDA kernel launch parameters explained right?

【发布时间】：2013-10-14 23:02:44

【问题描述】：

这里我尝试使用一些伪代码自我解释CUDA启动参数模型（或执行配置模型），但我不知道是否有一些大错误，所以希望有人帮助复习一下，给我一些建议。谢谢高级。

这里是：

/*
  normally, we write kernel function like this.
  note, __global__ means this function will be called from host codes,
  and executed on device. and a __global__ function could only return void.
  if there's any parameter passed into __global__ function, it should be stored
  in shared memory on device. so, kernel function is so different from the *normal*
  C/C++ functions. if I was the CUDA authore, I should make the kernel function more
  different  from a normal C function.
*/

__global__ void
kernel(float *arr_on_device, int n) 
        int idx = blockIdx.x * blockDIm.x + threadIdx.x;
        if (idx < n) 
                arr_on_device[idx] = arr_on_device[idx] * arr_on_device[idx];
        


/*
  after this definition, we could call this kernel function in our normal C/C++ codes !!
  do you feel something wired ? un-consistant ?
  normally, when I write C codes, I will think a lot about the execution process down to
  the metal in my mind, and this one...it's like some fragile codes. break the sequential
  thinking process in my mind.
  in order to make things normal, I found a way to explain: I expand the *__global__ * function
  to some pseudo codes:
*/

#define __foreach(var, start, end) for (var = start, var < end; ++var)

__device__ int
__indexing() 
        const int blockId = blockIdx.x * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;

        return 
                blockId * (blockDim.x * blockDim.y * blockDim.z) +
                threadIdx.z * (blockDim.x * blockDim.y) +
                threadIdx.x;


global_config =:
        
                /*
                  global configuration.
                  note the default values are all 1, so in the kernel codes,
                  we could just ignore those dimensions.
                 */ 
                gridDim.x = gridDim.y = gridDim.z = 1;
                blockDim.x = blockDim.y = blockDim.z = 1;
        ;

kernel =:
        
                /*
                  I thought CUDA did some bad evil-detail-covering things here.
                  it's said that CUDA C is an extension of C, but in my mind,
                  CUDA C is more like C++, and the *<<<>>>* part is too tricky.
                  for example:
                  kernel<<<10, 32>>>(); means kernel will execute in 10 blocks each have 32 threads.

                  dim3 dimG(10, 1, 1);
                  dim3 dimB(32, 1, 1);
                  kernel<<<dimG, dimB>>>(); this is exactly the same thing with above.

                  it's not C style, and C++ style ? at first, I thought this could be done by
                  C++'s constructor stuff, but I checked structure *dim3*, there's no proper
                  constructor for this. this just brroke the semantics of both C and C++. I thought
                  force user to use *kernel<<<dim3, dim3>>>* would be better. So I'd like to keep
                  this rule in my future codes.
                */

                gridDim  = dimG;
                blockDim = dimB;

                __foreach(blockIdx.z,  0, gridDim.z)
                __foreach(blockIdx.y,  0, gridDim.y)
                __foreach(blockIdx.x,  0, gridDim.x)
                __foreach(threadIdx.z, 0, blockDim.z)
                __foreach(threadIdx.y, 0, blockDim.y)
                __foreach(threadIdx.x, 0, blockDim.x)
                
                        const int idx = __indexing();        
                        if (idx < n) 
                                arr_on_device[idx] = arr_on_device[idx] * arr_on_device[idx];
                        
                
        ;

/*
  so, for me, gridDim & blockDim is like some boundaries.
  e.g. gridDim.x is the upper bound of blockIdx.x, this is not that obvious for people like me.
 */

/* the declaration of dim3 from vector_types.h of CUDA/include */
struct __device_builtin__ dim3

        unsigned int x, y, z;
#if defined(__cplusplus)
        __host__ __device__ dim3(unsigned int vx = 1, unsigned int vy = 1, unsigned int vz = 1) : x(vx), y(vy), z(vz) 
        __host__ __device__ dim3(uint3 v) : x(v.x), y(v.y), z(v.z) 
        __host__ __device__ operator uint3(void)  uint3 t; t.x = x; t.y = y; t.z = z; return t; 
#endif /* __cplusplus */
;

typedef __device_builtin__ struct dim3 dim3;

【问题讨论】：

有很多课程可供选择，包括介绍性课程，您可以随时观看，只需 1 小时，here。试试GPU Computing using CUDA C 系列。

你读过我下面的解释了吗？

您似乎对网格和块感到非常困惑，更不用说非常重要的底层架构了。我在 udacity 注册了免费的在线 CUDA 课程，并在一周内获得了有效（如果不是很高级）的编码。检查一下，因为 GPU 编程似乎需要坚实的基础。

@Boyko 我查过你建议的课程，特别是解释内核启动配置的部分udacity.com/course/viewer#!/c-cs344/l-55120467/m-67074291（从这个剪辑开始，然后是3个剪辑），老师仍然没有说清楚GPU如何将任务分派给线程。看看那个内核函数，编译器是如何扩展代码的？很多人只是被告知要记住规则，但是规则是如何在代码中实施的呢？所以现在，我正在阅读驱动程序 API 和 OpenCL 文档，只是想了解幕后发生的事情。

制作了更高级别的语言，因此您不必担心代码如何将其转化为 CPU 上的机器指令。而且每一个新一代的 CPU 都会带来一些新的东西，即使指令集是一样的，但在不同的架构上实际发生的事情并不一定是一样的。

【参考方案1】：

CUDA 驱动程序 API

CUDA Driver API v4.0 及更高版本使用以下函数来控制内核启动：

cuFuncSetCacheConfig
cuFuncSetSharedMemConfig
cuLaunchKernel

在 v4.0 中引入 cuLaunchKernel 之前，使用了以下 CUDA 驱动程序 API 函数。

cuFuncSetBlockShape()
cuFuncSetSharedSize()
cuParamSetSize,i,fv()
cuLaunch
cuLaunchGrid

关于这些函数的更多信息可以在 cuda.h 中找到。

CUresult CUDAAPI cuLaunchKernel(CUfunction f,
    unsigned int gridDimX,
    unsigned int gridDimY,
    unsigned int gridDimZ,
    unsigned int blockDimX,
    unsigned int blockDimY,
    unsigned int blockDimZ,
    unsigned int sharedMemBytes,
    CUstream hStream,
    void **kernelParams,
    void **extra);

cuLaunchKernel 将整个启动配置作为参数。

有关详细信息，请参阅 NVIDIA 驱动程序 API[执行控制]1。

CUDA 内核启动

cuLaunchKernel 将 1.验证启动参数 2.更改共享内存配置 3.改变本地内存分配 4. 将流同步令牌推入命令缓冲区，确保流中的两个命令不重叠 4. 将启动参数推送到命令缓冲区 5.将启动命令推入命令缓冲区 6. 向设备提交命令缓冲区（在 wddm 驱动程序上，此步骤可能会延迟） 7. 在 wddm 上，内核驱动程序将分页设备内存中所需的所有内存

GPU 将 1.验证命令 2. 将命令发送给计算工作分配器 3. 将启动配置和线程块分派给 SM

当所有线程块都完成后，工作分配器将刷新缓存以遵循 CUDA 内存模型，并将内核标记为已完成，以便流中的下一项可以向前推进。

线程块的调度顺序因架构而异。

计算能力 1.x 设备将内核参数存储在共享内存中。 计算能力 2.0-3.5 的设备将 kenrel 参数存储在常量内存中。

CUDA 运行时 API

CUDA 运行时是一个 C++ 软件库和构建在 CUDA 驱动程序 API 之上的工具链。 CUDA Runtime 使用以下函数来控制内核启动：

cudaConfigureCall cudaFuncSetCacheConfig cudaFuncSetSharedMemConfig cuda启动 cudaSetupArgument

见 NVIDIA Runtime API[执行控制]2

>> CUDA 语言扩展是最常用的启动内核的方法。

在编译期间，nvcc 将为每个使用 >> 调用的内核函数创建一个新的 CPU 存根函数，并将用对存根函数的调用替换 >>。

例如

__global__ void kernel(float* buf, int j)

    // ...


kernel<<<blocks,threads,0,myStream>>>(d_buf,j);

生成

void __device_stub__Z6kernelPfi(float *__par0, int __par1)__cudaSetupArgSimple(__par0, 0U);__cudaSetupArgSimple(__par1, 4U);__cudaLaunch(((char *)((void ( *)(float *, int))kernel)));

您可以通过在 nvcc 命令行中添加 --keep 来检查生成的文件。

cudaLaunch 调用 cuLaunchKernel。

CUDA 动态并行

CUDA CDP 的工作方式类似于上述的 CUDA 运行时 API。

【讨论】：

谢谢！这对我很有用……几乎是这里唯一有用的答案。

【参考方案2】：

通过使用<<<...>>>，您将在 GPU 中启动多个线程。这些线程被分组为块并形成一个大网格。所有线程都会执行调用的内核函数代码。

在内核函数中，threadIdx 和blockIdx 等内置变量使代码能够知道它运行的是哪个线程并执行计划的部分工作。

编辑

基本上，<<<...>>> 简化了启动内核的配置过程。如果不使用它，一个内核启动可能需要调用 4~5 个 API，就像 OpenCL 方式一样，它只使用 C99 语法。

事实上，您可以检查 CUDA 驱动程序 API。它可能会提供所有这些 API，因此您无需使用 <<<>>>。

【讨论】：

我很久以前就读过那个文档，以及其他一些相关的文档。我认为CUDA在kernel函数部分破坏了C和C++的语义，这使得它难以理解，容易混淆。所以，你说通过使用>>，你正在启动...，是的，你是对的，就像所有那些文档所说的那样，但是像这样解释CUDA执行配置太糟糕了。你只是告诉了如何，而不是WHY，哦，不，你只是告诉了what，甚至没有告诉how。

我已经编辑了答案。您可能希望通过添加 WHY 和 how 来修改您的问题，以避免所有观众误解您的问题。

【参考方案3】：

基本上，GPU 分为单独的“设备”GPU（例如 GeForce 690 有 2 个）-> 多个 SM（流式多处理器）-> 多个 CUDA 内核。据我所知，块或网格的维度只是与硬件无关的逻辑分配，但块的总大小（x*y*z）非常重要。

块中的线程必须位于同一个 SM 上，才能使用其共享内存和同步功能。因此，您不能拥有比 SM 中包含的 CUDA 内核更多线程的块。

如果我们有一个简单的场景，我们有 16 个 SM，每个 32 个 CUDA 内核，并且我们有 31x1x1 的块大小和 20x1x1 的网格大小，我们将丧失至少 1/32 的卡处理能力。每次运行一个块时，一个 SM 的 32 个内核中只有 31 个处于忙碌状态。块将加载以填满 SM，我们将大致同时完成 16 个块，并且随着前 4 个 SM 释放，它们将开始处理最后 4 个块（不一定是块 #17-20）。

欢迎评论和指正。

【讨论】：

我想我想弄清楚的部分是：CUDA编译器如何解析内核代码，并将任务分派给设备中的线程。在我原来的帖子中，我想象一种自我解释机制的方法。