CUDA Thread Indexing
内容提要
本文介绍CUDA线程索引与内存操作，通过向量相加和Sobel边缘检测示例，讲解cudaMalloc分配显存、cudaMemcpy传输数据，以及用blockIdx、blockDim、threadIdx计算全局索引。强调线程数应设为32的倍数以避免Warp浪费，数据过大时用stride循环处理，过小时需判断越界。
延伸解读
线程索引计算与边界判断
在CUDA中，全局线程索引通过blockIdx、blockDim和threadIdx计算，如index = blockIdx.x * blockDim.x + threadIdx.x。当数据量小于总线程数时，必须添加if (index < count)判断，否则越界访问可能导致不可预知错误，尤其在二维矩阵中可能将y维数据误读到x维。
Warp效率与线程数设置
Warp是block中连续的32个线程，若线程数不是32的倍数，最后一个Warp中部分线程会闲置，造成浪费。因此，设置block大小为32的倍数（如128）可提高执行效率。同时，block内线程数受硬件限制，x或y方向不超过1024，z方向不超过64，且资源有限，过多线程可能导致寄存器溢出。
数据规模不匹配的应对策略
当数据量超过总线程数时，可使用stride循环，每个线程处理多个元素，stride等于总线程数（blockDim.x * gridDim.x），循环中index每次增加stride。当数据量小于线程数时，需在核函数中判断越界。这两种方法在向量加法和Sobel边缘检测中均有体现。
Sobel边缘检测的GPU实现要点
Sobel核函数使用二维线程块（如32x32），通过x和y计算像素索引。由于卷积需要访问周围像素，必须判断边界（x>0且x<imgWidth-1等），避免越界读取。核函数中计算Gx和Gy的绝对值之和除以2得到梯度幅值，与CPU实现逻辑一致，但需注意内存访问模式。
Q&A
CUDA中如何计算一个线程的全局索引？
全局索引计算公式为：int index = threadIdx.x + blockIdx.x * blockDim.x; 其中threadIdx.x是线程在block内的索引，blockIdx.x是block在grid内的索引，blockDim.x是每个block的线程数。
为什么CUDA中建议将线程数设置为32的倍数？
因为Warp是block中连续的32个线程，如果线程数不是32的倍数，最后一个Warp中会有线程闲置，导致处理器资源浪费。因此应尽量将线程数设为32的倍数。
当数据量远大于线程数时，CUDA核函数中如何处理？
使用stride循环，每个线程处理多个数据。stride等于总线程数（blockDim.x * gridDim.x），线程从自己的索引开始，每次增加stride直到处理完所有数据。例如：for(; index < n; index += stride) z[index] = x[index] + y[index];
CUDA中cudaMalloc和cudaMemcpy的作用是什么？
cudaMalloc用于在GPU上分配显存，参数为指向设备指针的指针和字节大小。cudaMemcpy用于在CPU和GPU之间传输数据，参数包括目标地址、源地址、字节数和传输方向（如cudaMemcpyHostToDevice）。
在CUDA中，如果线程数超过数据量，如何防止越界访问？
在核函数中添加边界判断，例如：if (blockDim.x * blockIdx.x + threadIdx.x < count) { ... }，确保线程只处理有效数据，避免越界导致错误。
如何设置CUDA的grid和block大小以处理任意长度的向量？
通常设置block_size为128或256，grid_size = (N + block_size - 1) / block_size，确保有足够的线程覆盖所有数据。例如：block_size = 128; grid_size = (N + block_size - 1) / block_size;
