乐趣区

关于机器学习:MindSporeCUDA编程三线程层次

线程档次的概念:

简略说,就是一个 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)编译后执行:
退出移动版