CUDA线程全局索引：确定当前线程处理的数据位置

线程全局索引：确定当前线程处理的数据位置

 一维索引计算​：

直接计算线程的全局一维索引：const int n = blockDim.x * blockIdx.x + threadIdx.x;

全局线程ID = 当前块在网格中的偏移量 + 线程在块内的偏移量

                   = (blockIdx.x * blockDim.x) + threadIdx.x

三维索引计算​：

// 块全局ID（三维网格→一维）
unsigned int bid = blockIdx.z * gridDim.x * gridDim.y + 
                   blockIdx.y * gridDim.x + 
                   blockIdx.x;

计算逻辑​​

• ​​blockIdx.z​​：当前块在网格的z轴索引

  • 乘以gridDim.x * gridDim.y：跳过所有前层的xy平面

• ​​blockIdx.y​​：当前块在网格的y轴索引

  • 乘以gridDim.x：跳过当前层的前行

• ​​blockIdx.x​​：当前块在网格的x轴索引

// 线程局部ID（三维块→一维）
unsigned int tid = threadIdx.z * blockDim.x * blockDim.y + 
                   threadIdx.y * blockDim.x + 
                   threadIdx.x;

​​计算逻辑​​

• ​​threadIdx.z​​：当前线程在块的z轴索引

  • 乘以blockDim.x * blockDim.y：跳过所有前层的xy平面

• ​​threadIdx.y​​：当前线程在块的y轴索引

  • 乘以blockDim.x：跳过当前层的前行

• ​​threadIdx.x​​：当前线程在块的x轴索引

最终全局ID需结合两者：gid = bid * (blockDim.x*blockDim.y*blockDim.z) + tid

示例：

(1) 一维核函数（向量加法）

__global__ void add_1d(float *x, float *y, int N) {
    int n = blockDim.x * blockIdx.x + threadIdx.x;
    if (n < N) y[n] += x[n];
}
// 启动配置：<<<(N+255)/256, 256>>>

(2) 二维核函数（矩阵处理）​

__global__ void matrixMultiply(int *A, int *B, int *C, int N) {
    // 计算每个线程负责的矩阵元素位置
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
 
    if (row < N && col < N) {
        int value = 0;
        // 计算矩阵 C 的元素 C[row][col]
        for (int k = 0; k < N; ++k) {
            value += A[row * N + k] * B[k * N + col];
        }
        C[row * N + col] = value;
    }
}

注：A[row * N + k]​​：访问矩阵 A的第 row行、第 k列的元素；

  B[k * N + col]​​：访问矩阵 B的第 k行、第 col列的元素

(3) 三维核函数（矩阵处理）​

__global__ void process_3d(float *data, int width, int height, int depth) {
    // 计算全局3D坐标
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int z = blockIdx.z * blockDim.z + threadIdx.z;
    
    if (x < width && y < height && z < depth) {
        int idx = z * (width*height) + y * width + x;
        data[idx] *= 2.0f;
    }
}
// 启动配置：dim3 block(8,8,8); dim3 grid((width+7)/8, (height+7)/8, (depth+7)/8);