ICode9

精准搜索请尝试: 精确搜索
首页 > 编程语言> 文章详细

MindSpore:CUDA编程(三)线程层次

2022-08-11 18:31:34  阅读:185  来源: 互联网

标签:MindSpore const thread blockDim int double 线程 CUDA block


线程层次的概念:

简单说,就是一个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/forum/thread-194449-1-1.html 中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:体验index

Index_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:完成一维向量计算:add

vectorAdd.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-add
TEST_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)

 

编译后执行:

标签:MindSpore,const,thread,blockDim,int,double,线程,CUDA,block
来源: https://www.cnblogs.com/skytier/p/16577107.html

本站声明: 1. iCode9 技术分享网(下文简称本站)提供的所有内容,仅供技术学习、探讨和分享;
2. 关于本站的所有留言、评论、转载及引用,纯属内容发起人的个人观点,与本站观点和立场无关;
3. 关于本站的所有言论和文字,纯属内容发起人的个人观点,与本站观点和立场无关;
4. 本站文章均是网友提供,不完全保证技术分享内容的完整性、准确性、时效性、风险性和版权归属;如您发现该文章侵犯了您的权益,可联系我们第一时间进行删除;
5. 本站为非盈利性的个人网站,所有内容不会用来进行牟利,也不会利用任何形式的广告来间接获益,纯粹是为了广大技术爱好者提供技术内容和技术思想的分享性交流网站。

专注分享技术,共同学习,共同进步。侵权联系[81616952@qq.com]

Copyright (C)ICode9.com, All Rights Reserved.

ICode9版权所有