CUDA 计算线程索引的一般公式

最新推荐文章于 2022-10-26 22:37:20 发布

lumping

最新推荐文章于 2022-10-26 22:37:20 发布

阅读量868

点赞数

分类专栏： CUDA 文章标签：高性能计算

原文链接：https://www.cnblogs.com/kangtian0/p/7296198.html

版权

CUDA 专栏收录该内容

7 篇文章 1 订阅

订阅专栏

第一种方法：

CUDA thread index:

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

int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)
+ threadIdx.z * (blockDim.x * blockDim.y)
+ threadIdx.y * blockDim.x
+ threadIdx.x;

从上到下分别是：

block的3,2,1维；

thread的3,2,1维。

如果没有对应维度，删除对应的计算部分即可。

维度的值最小是1，但是索引的值最小是0.

这里定义的格式为:

dim3 blocks(blockDim.x,blockDim.y,blockDim.z);

dim3 grids(gridDim.x,gridDim.y,gridDim.z); // 这里需要说明的是如果的你的数组满足 M[NX][NY][NZ] 的形似，在设置blocks 和 grids时候应该满足 blockDim.x * gridDim.x=NX, blockDim.y * gridDim.y =NY, blockDim.z * gridDim.z=NZ;

ps: 最少的情况是都只有1维，最多的情况是都有3维。

所以最简单的都只有1维时：blockIdx.x * blockDim.x + threadIdx.x

第二种方法：

分别计算 idx,idy,idz

idx = threadIdx.x + blockDim.x * blockIdx.x;

idy = threadIdx.y + blockDim.y * blockIdx.y;

idz = threadIdx.z + blockDim.z * blockIdx.z;

总的index (一维情况下)

id = idx + NX * idy + NX * NY * idz;

上面两种方法计算的结果是一样的，但是需要注意的是，第一种方法只能在index在一维情况下使用，而如果 index是二维的情况下的话，建议使用第二种计算方法，此时在设备端数组可以写成

M_dev[idx][idy][idz];

但是个人建议还是使用第一种方法，因为第二种方法涉及到二级指针问题，处理起来稍微会麻烦一点。

参考：https://www.cnblogs.com/tiandsp/p/9458734.html 有15种检索方式

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <stdlib.h>
#include <iostream>

using namespace std;

//thread 1D
__global__ void testThread1(int *c, const int *a, const int *b)
{
    int i = threadIdx.x;
    c[i] = b[i] - a[i];
}

//thread 2D
__global__ void testThread2(int *c, const int *a, const int *b)
{
    int i = threadIdx.x + threadIdx.y*blockDim.x;
    c[i] = b[i] - a[i];
}

//thread 3D
__global__ void testThread3(int *c, const int *a, const int *b)
{
    int i = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;
    c[i] = b[i] - a[i];
}

//block 1D
__global__ void testBlock1(int *c, const int *a, const int *b)
{
    int i = blockIdx.x;
    c[i] = b[i] - a[i];
}

//block 2D
__global__ void testBlock2(int *c, const int *a, const int *b)
{
    int i = blockIdx.x + blockIdx.y*gridDim.x;
    c[i] = b[i] - a[i];
}

//block 3D
__global__ void testBlock3(int *c, const int *a, const int *b)
{
    int i = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;
    c[i] = b[i] - a[i];
}

//block-thread 1D-1D
__global__ void testBlockThread1(int *c, const int *a, const int *b)
{
    int i = threadIdx.x + blockDim.x*blockIdx.x;
    c[i] = b[i] - a[i];
}

//block-thread 1D-2D
__global__ void testBlockThread2(int *c, const int *a, const int *b)
{
    int threadId_2D = threadIdx.x + threadIdx.y*blockDim.x;
    int i = threadId_2D+ (blockDim.x*blockDim.y)*blockIdx.x;
    c[i] = b[i] - a[i];
}

//block-thread 1D-3D
__global__ void testBlockThread3(int *c, const int *a, const int *b)
{
    int threadId_3D = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;
    int i = threadId_3D + (blockDim.x*blockDim.y*blockDim.z)*blockIdx.x;
    c[i] = b[i] - a[i];
}

//block-thread 2D-1D
__global__ void testBlockThread4(int *c, const int *a, const int *b)
{
    int blockId_2D = blockIdx.x + blockIdx.y*gridDim.x;
    int i = threadIdx.x + blockDim.x*blockId_2D;
    c[i] = b[i] - a[i];
}

//block-thread 3D-1D
__global__ void testBlockThread5(int *c, const int *a, const int *b)
{
    int blockId_3D = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;
    int i = threadIdx.x + blockDim.x*blockId_3D;
    c[i] = b[i] - a[i];
}

//block-thread 2D-2D
__global__ void testBlockThread6(int *c, const int *a, const int *b)
{
    int threadId_2D = threadIdx.x + threadIdx.y*blockDim.x;
    int blockId_2D = blockIdx.x + blockIdx.y*gridDim.x;
    int i = threadId_2D + (blockDim.x*blockDim.y)*blockId_2D;
    c[i] = b[i] - a[i];
}

//block-thread 2D-3D
__global__ void testBlockThread7(int *c, const int *a, const int *b)
{
    int threadId_3D = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;
    int blockId_2D = blockIdx.x + blockIdx.y*gridDim.x;
    int i = threadId_3D + (blockDim.x*blockDim.y*blockDim.z)*blockId_2D;
    c[i] = b[i] - a[i];
}

//block-thread 3D-2D
__global__ void testBlockThread8(int *c, const int *a, const int *b)
{
    int threadId_2D = threadIdx.x + threadIdx.y*blockDim.x;
    int blockId_3D = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;
    int i = threadId_2D + (blockDim.x*blockDim.y)*blockId_3D;
    c[i] = b[i] - a[i];
}

//block-thread 3D-3D
__global__ void testBlockThread9(int *c, const int *a, const int *b)
{
    int threadId_3D = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;
    int blockId_3D = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;
    int i = threadId_3D + (blockDim.x*blockDim.y*blockDim.z)*blockId_3D;
    c[i] = b[i] - a[i];
}


void addWithCuda(int *c, const int *a, const int *b, unsigned int size)
{
    int *dev_a = 0;
    int *dev_b = 0;
    int *dev_c = 0;

    cudaSetDevice(0);

    cudaMalloc((void**)&dev_c, size * sizeof(int));
    cudaMalloc((void**)&dev_a, size * sizeof(int));
    cudaMalloc((void**)&dev_b, size * sizeof(int));

    cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);

    //testThread1<<<1, size>>>(dev_c, dev_a, dev_b);

    //uint3 s;s.x = size/5;s.y = 5;s.z = 1;
    //testThread2 <<<1,s>>>(dev_c, dev_a, dev_b);

    //uint3 s; s.x = size / 10; s.y = 5; s.z = 2;
    //testThread3<<<1, s >>>(dev_c, dev_a, dev_b);

    //testBlock1<<<size,1 >>>(dev_c, dev_a, dev_b);

    //uint3 s; s.x = size / 5; s.y = 5; s.z = 1;
    //testBlock2<<<s, 1 >>>(dev_c, dev_a, dev_b);

    //uint3 s; s.x = size / 10; s.y = 5; s.z = 2;
    //testBlock3<<<s, 1 >>>(dev_c, dev_a, dev_b);

    //testBlockThread1<<<size/10, 10>>>(dev_c, dev_a, dev_b);

    //uint3 s1; s1.x = size / 100; s1.y = 1; s1.z = 1;
    //uint3 s2; s2.x = 10; s2.y = 10; s2.z = 1;
    //testBlockThread2 << <s1, s2 >> >(dev_c, dev_a, dev_b);

    //uint3 s1; s1.x = size / 100; s1.y = 1; s1.z = 1;
    //uint3 s2; s2.x = 10; s2.y = 5; s2.z = 2;
    //testBlockThread3 << <s1, s2 >> >(dev_c, dev_a, dev_b);

    //uint3 s1; s1.x = 10; s1.y = 10; s1.z = 1;
    //uint3 s2; s2.x = size / 100; s2.y = 1; s2.z = 1;
    //testBlockThread4 << <s1, s2 >> >(dev_c, dev_a, dev_b);

    //uint3 s1; s1.x = 10; s1.y = 5; s1.z = 2;
    //uint3 s2; s2.x = size / 100; s2.y = 1; s2.z = 1;
    //testBlockThread5 << <s1, s2 >> >(dev_c, dev_a, dev_b);

    //uint3 s1; s1.x = size / 100; s1.y = 10; s1.z = 1;
    //uint3 s2; s2.x = 5; s2.y = 2; s2.z = 1;
    //testBlockThread6 << <s1, s2 >> >(dev_c, dev_a, dev_b);

    //uint3 s1; s1.x = size / 100; s1.y = 5; s1.z = 1;
    //uint3 s2; s2.x = 5; s2.y = 2; s2.z = 2;
    //testBlockThread7 << <s1, s2 >> >(dev_c, dev_a, dev_b);

    //uint3 s1; s1.x = 5; s1.y = 2; s1.z = 2;
    //uint3 s2; s2.x = size / 100; s2.y = 5; s2.z = 1;
    //testBlockThread8 <<<s1, s2 >>>(dev_c, dev_a, dev_b);

    uint3 s1; s1.x = 5; s1.y = 2; s1.z = 2;
    uint3 s2; s2.x = size / 200; s2.y = 5; s2.z = 2;
    testBlockThread9<<<s1, s2 >>>(dev_c, dev_a, dev_b);

    cudaMemcpy(c, dev_c, size*sizeof(int), cudaMemcpyDeviceToHost);

    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);

    cudaGetLastError();
}


int main()
{
    const int n = 1000;

    int *a = new int[n];
    int *b = new int[n];
    int *c = new int[n];
    int *cc = new int[n];

    for (int i = 0; i < n; i++)
    {
        a[i] = rand() % 100;
        b[i] = rand() % 100;
        c[i] = b[i] - a[i];
    }

    addWithCuda(cc, a, b, n);

    FILE *fp = fopen("out.txt", "w");
    for (int i = 0; i < n; i++)
        fprintf(fp, "%d %d\n", c[i], cc[i]);
    fclose(fp);

    bool flag = true;
    for (int i = 0; i < n; i++)
    {
        if (c[i] != cc[i])
        {
            flag = false;
            break;
        }
    }

    if (flag == false)
        printf("no pass");
    else
        printf("pass");

    cudaDeviceReset();

    delete[] a;
    delete[] b;
    delete[] c;
    delete[] cc;

    getchar();
    return 0;
}

其他补充内容，二维数组分块计算。

ix=threadIdx.x+blockDim.x*blockIdx.x
iy=threadIdx.y+blockDim.y*blockIdx.y

k=ix+iy*nx

blockId  =blockIdx.x+gridDim.x*blockIdx.y+gridDim.x*gridDim.y*blockIdx.z
threadId=blockId*blockDim.x*blockIdDim.y*blockDim.z+
               threadIdx.z*blockDim.x*blockDim.y
               threadIdx.y*blockDim.x
               threadIdx.x

dim3  thread(tx   ,ty   )
dim3  block (NX/tx,NY/ty)

gridDim.x=NX/tx
gridDim.y=NY/ty

blockDim.x=tx
blockDim.y=ty
               

a[ix][iy]

索引计算方法1：
k=iy*NX+ix
  =(threadIdx.y+ty*blockIdx.y)*NX+tx*blockIdx.x     +threadIdx.x
  =tx*blockIdx.x    +NX*ty*blockIdx.y+NX*threadIdx.y+threadIdx.x

索引计算方法2：
k=(blockIdx.x+NX/tx*blockIdx.y)*tx*ty+threadIdx.y*tx +threadIdx.x
  =tx*ty*blockIdx.x+NX*ty*blockIdx.y +tx*threadIdx.y +threadIdx.x

虽然上面两种计算方法的表达式不同，但是计算的结果都是一样的，索引值均是从0到NX*NY-1，总共NX*NY个坐标，证明如下，

dim3  thread(tx   ,ty   )
dim3  block (NX/tx,NY/ty)
 
max{blockIdx.x} =NX/tx-1
max{blockIdx.y} =NY/ty-1
max{threadIdx.x}=tx-1
max{threadIdx.y}=ty-1           
 
将上式带入到两种索引计算公式中，求得最大索引数
 
k = tx*blockIdx.x + NX*ty*blockIdx.y + NX*threadIdx.y + threadIdx.x
  = tx*(NX/tx -1) + NX*ty*(NY/ty-1)  + NX*(ty-1)      + tx-1
  = NX*NY-1

k = tx*ty*blockIdx.x + NX*ty*blockIdx.y + tx*threadIdx.y + threadIdx.x
  = tx*ty*(NX/tx-1)  + NX*ty*(NY/ty-1)  + tx*(ty-1)      + tx-1
  = NX*NY-1

lumping

关注

0
点赞
踩
4

收藏

觉得还不错? 一键收藏
1
评论
CUDA 计算线程索引的一般公式

CUDA thread index:int blockId = blockIdx.z * (gridDim.x*gridDim.y) + blockIdx.y * gridDim.x + blockIdx.x;int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z) + threadIdx.z * (bloc...
复制链接

扫一扫