CUDA编程 2 线程组织【参加CUDA线上训练营】

2.1 使用多个线程的核函数

在调用核函数时填入<<< , >>>的是gridDim和blockDim。分别表示一个grid有多少个block和一个block有多少个thread。

<<<1, 1>>>相当于单线程。

此外，我们可以通过使用 threadIdx.x 和blockIdx.x 来分别获得执行当前kernel函数的线程在block中的x方向的序号 和 执行当前kernel函数的线程所在block，在grid中的x方向的序号。

#include <stdio.h>

__global__ void hello_from_gpu()
{
    const int bid = blockIdx.x;
    const int tid = threadIdx.x;
    printf("Hello World from block %d and thread %d!\n", bid, tid);
}

int main(void)
{
    hello_from_gpu<<<5, 5>>>();
    cudaDeviceSynchronize();
    return 0;
}

以上是核函数中打印执行该核函数的线程编号和所在线程块编号的实例代码。

2.2 使用线程索引

Cuda编程的一个重点就是对线程索引的处理。

处理向量时，index = blockIdx.x * blockDim.x + threadIdx.x

课后作业：  

1. 如果我们设置的线程数过大，比如设置grid_size = (N + block_size - 1) / block_size+10000，会产生什么后果？如何避免这种后果？

在 GPU 上可用的线程数是有限的，如果线程数过大，超出限制后会导致引发运行时错误。

要避免这种后果，

可以使用 CUDA 库函数 cudaDeviceGetAttribute 获取 GPU 可用的线程数，并在根据该信息设置线程数时保证不超过该限制。

也可以使用 cudaOccupancyMaxPotentialBlockSize 函数来检查在 GPU 上的某个块大小下的可用线程数。

#include <cuda_runtime.h>
#include <iostream>

int main() {
  int deviceCount;
  cudaGetDeviceCount(&deviceCount);

  for (int i = 0; i < deviceCount; i++) {
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, i);

    int maxThreadsPerBlock;
    cudaDeviceGetAttribute(&maxThreadsPerBlock, cudaDevAttrMaxThreadsPerBlock, i);
    std::cout << "Device ID: " << i << std::endl;
    std::cout << "Device Name: " << deviceProp.name << std::endl;
    std::cout << "Max Threads Per Block: " << maxThreadsPerBlock << std::endl;
    std::cout << std::endl;
  }

  return 0;
}

2. 如果我们的要处理的数据太多，远远超过我们能申请的线程数怎么办？

分块处理：将要处理的数据分成多个块，分别使用 CUDA 函数进行处理，并将结果合并。

重构代码：考虑将数据处理流程重构为更高效的方式，以降低数据量。

使用多个 GPU：如果有多个 GPU 可供使用，可以使用多个 GPU 并行处理数据，以提高处理速度。

增加线程数限制：如果 GPU 支持，可以通过修改驱动程序或系统设置来增加线程数限制，从而获得更多的线程。

根据数据处理需求，可以从这几种方法中选择一种或多种方法来解决问题。