cuda中threadIdxblockIdxblockDim和gridDim的使用
Posted FILWY_M
tags:
篇首语:本文由小常识网(cha138.com)小编为大家整理,主要介绍了cuda中threadIdxblockIdxblockDim和gridDim的使用相关的知识,希望对你有一定的参考价值。
一、直观的感觉线程、线程块、线程格
为了直观的感觉线程、线程块、线程格,画了下面一个示意图。分为了两部分,一部分为线程格,另一部分为线程块,在图中线程格和线程块都画成了3维的,实际也可以是一维或者二维的。其中线程格里面最小的单元为线程块,而一个线程块里面最小的单元为线程。
二、threadIdx、blockIdx、blockDim和gridDim
可以把线程格和线程块都看作一个三维的矩阵。这里假设线程格是一个3*4*5
的三维矩阵, 线程块是一个4*5*6
的三维矩阵。
-
gridDim
gridDim.x
、gridDim.y
、gridDim.z
分别表示线程格各个维度的大小,所以有gridDim.x=3 gridDim.y=4 gridDim.z=5
-
blockDim
blockDim.x
、blockDim.y
、blockDim.z
分别表示线程块中各个维度的大小,所以有blockDim.x=4 blockDim.y=5 blockDim.z=6
-
blockIdx
blockIdx.x
、blockIdx.y
、blockIdx.z
分别表示当前线程块所处的线程格的坐标位置 -
threadIdx
threadIdx.x
、threadIdx.y
、threadIdx.z
分别表示当前线程所处的线程块的坐标位置
通过 blockIdx.x
、blockIdx.y
、blockIdx.z
、threadIdx.x
、threadIdx.y
、threadIdx.z
就可以完全定位一个线程的坐标位置了。
线程格里面总的线程个数N即可通过下面的公式算出
N = gridDim.x*gridDim.y*gridDim.z*blockDim.x*blockDim.y*blockDim.z
三、举例
将所有的线程排成一个序列,序列号为 0 , 1 , 2 , … , N 0,1,2,\\dots,N 0,1,2,…,N,如何找到当前的序列号?
-
先找到当前线程位于线程格中的哪一个线程块
blockId
blockId = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;
-
找到当前线程位于线程块中的哪一个线程
threadId
threadId = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;
-
计算一个线程块中一共有多少个线程M
M = blockDim.x*blockDim.y*blockDim.z
-
求得当前的线程序列号
idx
idx = threadId + M*blockId;
下面是通过GPU并行计算实现的两个向量相减的例子
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdlib.h>
#include <iostream>
using namespace std;
//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);
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;
以上是关于cuda中threadIdxblockIdxblockDim和gridDim的使用的主要内容,如果未能解决你的问题,请参考以下文章