cuda中的blockIdx，threadIdx，blockDim以及gridDim的使用

转载于：DSP Tian的博客

threadIdx是一个uint3类型，表示一个线程的索引。

blockIdx是一个uint3类型，表示一个线程块的索引，一个线程块中通常有多个线程。

blockDim是一个dim3类型，表示线程块的大小。

gridDim是一个dim3类型，表示网格的大小，一个网格中通常有多个线程块。

下面这张图比较清晰的表示的几个概念的关系：

cuda 通过<<< >>>符号来分配索引线程的方式，我知道的一共有15种索引方式。

下面程序展示了这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.yblockDim.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.yblockDim.x + threadIdx.zblockDim.xblockDim.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.ygridDim.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.ygridDim.x + blockIdx.zgridDim.xgridDim.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.xblockIdx.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.yblockDim.x;
int i = threadId_2D+ (blockDim.xblockDim.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.yblockDim.x + threadIdx.zblockDim.xblockDim.y;
int i = threadId_3D + (blockDim.xblockDim.yblockDim.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.ygridDim.x;
int i = threadIdx.x + blockDim.xblockId_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.ygridDim.x + blockIdx.zgridDim.xgridDim.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.yblockDim.x;
int blockId_2D = blockIdx.x + blockIdx.ygridDim.x;
int i = threadId_2D + (blockDim.xblockDim.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.yblockDim.x + threadIdx.zblockDim.xblockDim.y;
int blockId_2D = blockIdx.x + blockIdx.y*gridDim.x;
int i = threadId_3D + (blockDim.xblockDim.yblockDim.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.yblockDim.x;
int blockId_3D = blockIdx.x + blockIdx.ygridDim.x + blockIdx.zgridDim.x*gridDim.y;
int i = threadId_2D + (blockDim.xblockDim.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.yblockDim.x + threadIdx.zblockDim.xblockDim.y;
int blockId_3D = blockIdx.x + blockIdx.ygridDim.x + blockIdx.zgridDim.x*gridDim.y;
int i = threadId_3D + (blockDim.xblockDim.yblockDim.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(</span><span style="color: rgba(128, 0, 128, 1)">0</span><span style="color: rgba(0, 0, 0, 1)">);cudaMalloc((</span><span style="color: rgba(0, 0, 255, 1)">void</span>**)&amp;dev_c, size * <span style="color: rgba(0, 0, 255, 1)">sizeof</span>(<span style="color: rgba(0, 0, 255, 1)">int</span><span style="color: rgba(0, 0, 0, 1)">));
cudaMalloc((</span><span style="color: rgba(0, 0, 255, 1)">void</span>**)&amp;dev_a, size * <span style="color: rgba(0, 0, 255, 1)">sizeof</span>(<span style="color: rgba(0, 0, 255, 1)">int</span><span style="color: rgba(0, 0, 0, 1)">));
cudaMalloc((</span><span style="color: rgba(0, 0, 255, 1)">void</span>**)&amp;dev_b, size * <span style="color: rgba(0, 0, 255, 1)">sizeof</span>(<span style="color: rgba(0, 0, 255, 1)">int</span><span style="color: rgba(0, 0, 0, 1)">));cudaMemcpy(dev_a, a, size </span>* <span style="color: rgba(0, 0, 255, 1)">sizeof</span>(<span style="color: rgba(0, 0, 255, 1)">int</span><span style="color: rgba(0, 0, 0, 1)">), cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, size </span>* <span style="color: rgba(0, 0, 255, 1)">sizeof</span>(<span style="color: rgba(0, 0, 255, 1)">int</span><span style="color: rgba(0, 0, 0, 1)">), cudaMemcpyHostToDevice);</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">testThread1&lt;&lt;&lt;1, size&gt;&gt;&gt;(dev_c, dev_a, dev_b);</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s;s.x = size/5;s.y = 5;s.z = 1;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">testThread2 &lt;&lt;&lt;1,s&gt;&gt;&gt;(dev_c, dev_a, dev_b);</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s; s.x = size / 10; s.y = 5; s.z = 2;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">testThread3&lt;&lt;&lt;1, s &gt;&gt;&gt;(dev_c, dev_a, dev_b);</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">testBlock1&lt;&lt;&lt;size,1 &gt;&gt;&gt;(dev_c, dev_a, dev_b);</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s; s.x = size / 5; s.y = 5; s.z = 1;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">testBlock2&lt;&lt;&lt;s, 1 &gt;&gt;&gt;(dev_c, dev_a, dev_b);</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s; s.x = size / 10; s.y = 5; s.z = 2;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">testBlock3&lt;&lt;&lt;s, 1 &gt;&gt;&gt;(dev_c, dev_a, dev_b);</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">testBlockThread1&lt;&lt;&lt;size/10, 10&gt;&gt;&gt;(dev_c, dev_a, dev_b);</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s1; s1.x = size / 100; s1.y = 1; s1.z = 1;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s2; s2.x = 10; s2.y = 10; s2.z = 1;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">testBlockThread2 &lt;&lt; &lt;s1, s2 &gt;&gt; &gt;(dev_c, dev_a, dev_b);</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s1; s1.x = size / 100; s1.y = 1; s1.z = 1;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s2; s2.x = 10; s2.y = 5; s2.z = 2;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">testBlockThread3 &lt;&lt; &lt;s1, s2 &gt;&gt; &gt;(dev_c, dev_a, dev_b);</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s1; s1.x = 10; s1.y = 10; s1.z = 1;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s2; s2.x = size / 100; s2.y = 1; s2.z = 1;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">testBlockThread4 &lt;&lt; &lt;s1, s2 &gt;&gt; &gt;(dev_c, dev_a, dev_b);</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s1; s1.x = 10; s1.y = 5; s1.z = 2;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s2; s2.x = size / 100; s2.y = 1; s2.z = 1;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">testBlockThread5 &lt;&lt; &lt;s1, s2 &gt;&gt; &gt;(dev_c, dev_a, dev_b);</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s1; s1.x = size / 100; s1.y = 10; s1.z = 1;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s2; s2.x = 5; s2.y = 2; s2.z = 1;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">testBlockThread6 &lt;&lt; &lt;s1, s2 &gt;&gt; &gt;(dev_c, dev_a, dev_b);</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s1; s1.x = size / 100; s1.y = 5; s1.z = 1;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s2; s2.x = 5; s2.y = 2; s2.z = 2;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">testBlockThread7 &lt;&lt; &lt;s1, s2 &gt;&gt; &gt;(dev_c, dev_a, dev_b);</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s1; s1.x = 5; s1.y = 2; s1.z = 2;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">uint3 s2; s2.x = size / 100; s2.y = 5; s2.z = 1;
</span><span style="color: rgba(0, 128, 0, 1)">//</span><span style="color: rgba(0, 128, 0, 1)">testBlockThread8 &lt;&lt;&lt;s1, s2 &gt;&gt;&gt;(dev_c, dev_a, dev_b);</span>

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</span>*<span style="color: rgba(0, 0, 255, 1)">sizeof</span>(<span style="color: rgba(0, 0, 255, 1)">int</span><span style="color: rgba(0, 0, 0, 1)">), cudaMemcpyDeviceToHost);cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);cudaGetLastError();

}

int main()
{
const int n = 1000;

</span><span style="color: rgba(0, 0, 255, 1)">int</span> *a = <span style="color: rgba(0, 0, 255, 1)">new</span> <span style="color: rgba(0, 0, 255, 1)">int</span><span style="color: rgba(0, 0, 0, 1)">[n];
</span><span style="color: rgba(0, 0, 255, 1)">int</span> *b = <span style="color: rgba(0, 0, 255, 1)">new</span> <span style="color: rgba(0, 0, 255, 1)">int</span><span style="color: rgba(0, 0, 0, 1)">[n];
</span><span style="color: rgba(0, 0, 255, 1)">int</span> *c = <span style="color: rgba(0, 0, 255, 1)">new</span> <span style="color: rgba(0, 0, 255, 1)">int</span><span style="color: rgba(0, 0, 0, 1)">[n];
</span><span style="color: rgba(0, 0, 255, 1)">int</span> *cc = <span style="color: rgba(0, 0, 255, 1)">new</span> <span style="color: rgba(0, 0, 255, 1)">int</span><span style="color: rgba(0, 0, 0, 1)">[n];</span><span style="color: rgba(0, 0, 255, 1)">for</span> (<span style="color: rgba(0, 0, 255, 1)">int</span> i = <span style="color: rgba(128, 0, 128, 1)">0</span>; i &lt; n; i++<span style="color: rgba(0, 0, 0, 1)">)
{a[i] </span>= rand() % <span style="color: rgba(128, 0, 128, 1)">100</span><span style="color: rgba(0, 0, 0, 1)">;b[i] </span>= rand() % <span style="color: rgba(128, 0, 128, 1)">100</span><span style="color: rgba(0, 0, 0, 1)">;c[i] </span>= b[i] -<span style="color: rgba(0, 0, 0, 1)"> a[i];
}addWithCuda(cc, a, b, n);FILE </span>*fp = fopen(<span style="color: rgba(128, 0, 0, 1)">"</span><span style="color: rgba(128, 0, 0, 1)">out.txt</span><span style="color: rgba(128, 0, 0, 1)">"</span>, <span style="color: rgba(128, 0, 0, 1)">"</span><span style="color: rgba(128, 0, 0, 1)">w</span><span style="color: rgba(128, 0, 0, 1)">"</span><span style="color: rgba(0, 0, 0, 1)">);
</span><span style="color: rgba(0, 0, 255, 1)">for</span> (<span style="color: rgba(0, 0, 255, 1)">int</span> i = <span style="color: rgba(128, 0, 128, 1)">0</span>; i &lt; n; i++<span style="color: rgba(0, 0, 0, 1)">)fprintf(fp, </span><span style="color: rgba(128, 0, 0, 1)">"</span><span style="color: rgba(128, 0, 0, 1)">%d %d\n</span><span style="color: rgba(128, 0, 0, 1)">"</span><span style="color: rgba(0, 0, 0, 1)">, c[i], cc[i]);
fclose(fp);</span><span style="color: rgba(0, 0, 255, 1)">bool</span> flag = <span style="color: rgba(0, 0, 255, 1)">true</span><span style="color: rgba(0, 0, 0, 1)">;
</span><span style="color: rgba(0, 0, 255, 1)">for</span> (<span style="color: rgba(0, 0, 255, 1)">int</span> i = <span style="color: rgba(128, 0, 128, 1)">0</span>; i &lt; n; i++<span style="color: rgba(0, 0, 0, 1)">)
{</span><span style="color: rgba(0, 0, 255, 1)">if</span> (c[i] !=<span style="color: rgba(0, 0, 0, 1)"> cc[i]){flag </span>= <span style="color: rgba(0, 0, 255, 1)">false</span><span style="color: rgba(0, 0, 0, 1)">;</span><span style="color: rgba(0, 0, 255, 1)">break</span><span style="color: rgba(0, 0, 0, 1)">;}
}</span><span style="color: rgba(0, 0, 255, 1)">if</span> (flag == <span style="color: rgba(0, 0, 255, 1)">false</span><span style="color: rgba(0, 0, 0, 1)">)printf(</span><span style="color: rgba(128, 0, 0, 1)">"</span><span style="color: rgba(128, 0, 0, 1)">no pass</span><span style="color: rgba(128, 0, 0, 1)">"</span><span style="color: rgba(0, 0, 0, 1)">);
</span><span style="color: rgba(0, 0, 255, 1)">else</span><span style="color: rgba(0, 0, 0, 1)">printf(</span><span style="color: rgba(128, 0, 0, 1)">"</span><span style="color: rgba(128, 0, 0, 1)">pass</span><span style="color: rgba(128, 0, 0, 1)">"</span><span style="color: rgba(0, 0, 0, 1)">);cudaDeviceReset();</span><span style="color: rgba(0, 0, 255, 1)">delete</span><span style="color: rgba(0, 0, 0, 1)">[] a;
</span><span style="color: rgba(0, 0, 255, 1)">delete</span><span style="color: rgba(0, 0, 0, 1)">[] b;
</span><span style="color: rgba(0, 0, 255, 1)">delete</span><span style="color: rgba(0, 0, 0, 1)">[] c;
</span><span style="color: rgba(0, 0, 255, 1)">delete</span><span style="color: rgba(0, 0, 0, 1)">[] cc;getchar();
</span><span style="color: rgba(0, 0, 255, 1)">return</span> <span style="color: rgba(128, 0, 128, 1)">0</span><span style="color: rgba(0, 0, 0, 1)">;

}

这里只保留了3D-3D方式，注释了其余14种方式，所有索引方式均测试通过。

还是能看出一些规律的：)