线程档次的概念：
简略说，就是一个grid有多个block,一个block有多个thread.grid有多大，用gridDim示意它有多少个block，具体分为gridDim.x, gridDim.y，gridDim.z。block有多大，用blockDim示意它有多少个thread，具体分为blockDim.x，blockDim.y，blockDim.z。怎么示意thread在block中的绝对地位呢？用 threadIdx.x,threadIdx.y,threadIdx.z示意。怎么示意block在grid中的绝对地位呢？用blockIdx.x,blockIdx.y,blockIdx.z示意。顺便解释下 https://bbs.huaweicloud.com/f… 中hello_from_gpu<<<x,y>>>(); 中的x和y是什么意思？它们别离示意 gridDim和blockDim。对于上面这个函数：
示意gridDim是1，示意grid有1个block，blockDim是4。示意block有4个thread。所以对于下面的核函数，相当于有4个thread别离执行了 c[n]=a[n]+b[n]的操作，n=threadIdx.x在调用的时候，所有的CUDA核都是执行同一个函数。这与CPU多线程可能会执行不同的工作不同。
如上图所示，Thread在CUDA core中执行，Block在 SM中执行，Grid在Device中执行。那么，CUDA是如何执行的呢？看上面这张图：
如果没有block的概念，要同时进行同步、通信、合作时，整体的外围都要产生期待的行为，如要进行扩大时，扩大的越多期待也越多。所以性能会受影响。然而有block的概念后，能够实现可扩展性。用block或warp就能够很容易实现扩大了。如何找到线程该解决的数据在哪里呢？这就要提到线程索引的概念。
以上：假设每8个thread时一个block。具体的公式如下：具体的索引地位 index = blockDim.x * blockIdx.x + threadIdx.x
那么一个CUDA程序到底应该怎么写呢？以将一个CPU实现的代码转换为GPU为例：
CPU的实现过程大抵如下：（1）主程序main：先调配 源地址空间a,b，目标地址空间c，并生成a，b的随机数。而后调用 一维矩阵加的CPU函数。（2）一维矩阵加的CPU函数：遍历a,b地址空间，别离将 a[i] 与 b[i]相加，写入 c[i]地址。这个时候，请留神是要显式地进行for循环遍历。
那么，GPU该如何实现呢？（1）主程序main：因为GPU存在Host和Device内存，所以先申请host内存h_a,h_b，寄存a，b的一维矩阵的内容（也能够生成随机数），并申请host内存h_c寄存c的计算结果。而后申请device内存，这个时候，须要申请 d_a,d_b两个源device内存(cudaMalloc)，以及d_c这个目标device内存(cudaMalloc)。将h_a和h_b的内容拷贝到d_a和d_b (显然须要应用 cudaMemcpyHostToDevice)；而后调用核函数实现GPU的并行计算，后果写入h_c；最初将d_c的device内存写回到h_c(cudaMemcpyDeviceToHost)，并开释所有的host内存(应用free)和device内存（应用cudaFree)。（2）核函数这里就是重点了。核函数只须要去掉最外层的循环，并且依据后面 的index写法，将i替换成index的写法即可。如何设置Gridsize和blocksize呢？对于一维的状况：block_size=128;grid_size = (N+ block_size-1)/block_size;（没有设成什么值是最好的）每个block能够申请多少个线程呢？总数也是1024。如（1024，1，1）或者（512，2，1）grid大小没有限度。底层是以warp为单位申请。 如果blockDim为160，则正好申请5个warp。如果blockDim为161，则不得不申请6个warp。如果数据过大，线程不够用怎么办？这样子，每个线程须要解决多个数据。
比方对于上图，线程0，须要解决 0,8,16,24 四个数据。核函数须要将每一个大块都跑一遍。代码如下：
这里引入了一个stride的概念，它的大小为blockDim.x X gridDim.x 。核函数须要实现每个满足 index = index + stride * count对应的相干地址的计算。范例1：体验indexIndex_of_thread.cu#include <stdio.h>
global void hello_from_gpu()
{
//仅仅是在原先代码的根底上打印 blockIdx.x 和 threadIdx.x
const int bid = blockIdx.x;
const int tid = threadIdx.x;
printf("Hello World from block %d and thread %d!\n", bid, tid);}
int main(void)
{
hello_from_gpu<<<5, 5>>>();
//记得加上同步，不然后果会出不来。
cudaDeviceSynchronize();
return 0;}
Makefile：TEST_SOURCE = Index_of_thread.cu
TARGETBIN := ./Index_of_thread
CC = /usr/local/cuda/bin/nvcc
$(TARGETBIN):$(TEST_SOURCE)
$(CC) $(TEST_SOURCE) -o $(TARGETBIN)
.PHONY:clean
clean:
-rm -rf $(TARGETBIN)编译并执行：
范例2：实现一维向量计算：addvectorAdd.cu#include <math.h>
include <stdio.h>
void global add(const double x, const double y, double *z, int count)
{
const int n = blockDim.x * blockIdx.x + threadIdx.x;
//这里判断是避免溢出
if( n < count)
{
z[n] = x[n] + y[n];
}
}
void check(const double *z, const int N)
{
bool error = false;
for (int n = 0; n < N; ++n)
{
//查看两个值是否相等，如不等则error=true.
if (fabs(z[n] - 3) > (1.0e-10))
{
error = true;
}
}
printf("%s\n", error ? "Errors" : "Pass");}
int main(void)
{
const int N = 1000;
const int M = sizeof(double) * N;
//调配host内存
double *h_x = (double*) malloc(M);
double *h_y = (double*) malloc(M);
double *h_z = (double*) malloc(M);
//初始化一维向量的值
for (int n = 0; n < N; ++n)
{
h_x[n] = 1;
h_y[n] = 2;
}
double *d_x, *d_y, *d_z;
//调配device内存
cudaMalloc((void **)&d_x, M);
cudaMalloc((void **)&d_y, M);
cudaMalloc((void **)&d_z, M);
//host->device
cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);
//这个是公式。记住就能够了。
const int block_size = 128;
const int grid_size = (N + block_size - 1) / block_size;
//核函数计算
add<<<grid_size, block_size>>>(d_x, d_y, d_z, N);
//device->host
cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
//查看后果
check(h_z, N);
//开释host内存
free(h_x);
free(h_y);
free(h_z);
//开释device内存
cudaFree(d_x);
cudaFree(d_y);
cudaFree(d_z);
return 0;}Makefile-addTEST_SOURCE = vectorAdd.cu
TARGETBIN := ./vectorAdd
CC = /usr/local/cuda/bin/nvcc
$(TARGETBIN):$(TEST_SOURCE)
$(CC) $(TEST_SOURCE) -o $(TARGETBIN)
.PHONY:clean
clean:
-rm -rf $(TARGETBIN)编译后执行：
发表回复