线程网络的分配

最新推荐文章于 2022-11-01 20:45:29 发布

__Sunny__

最新推荐文章于 2022-11-01 20:45:29 发布

阅读量2k

点赞数

分类专栏： CUDA 文章标签： gpu cuda

CUDA 专栏收录该内容

14 篇文章 1 订阅

订阅专栏

转载自从0开始学习《GPU高性能运算之CUDA》——2

5 GPU也不允许偏心

并行的事情多了，我们作为 GPU 的指令分配者，不能偏心了——给甲做的事情多，而乙没事做，个么甲肯定不爽的来。所以，在 GPU 中，叫做线程网络的分配。

我们将具体点的，在主机函数中如果我们分配的是这样的一个东西：

dim3 blocks(32,32);
dim3 threads(16,16);

dim3 是神马？dim3 是一个内置的结构体，和 linux 下定义的线程结构体是个类似的意义的东西，dim3 结构变量有 x，y，z，表示 3 维的维度。不理解没关系，慢慢看。

kernelfun<<<blocks, threads>>>();

我们调用 kernelfun 这个内核函数，将 blocks 和 threads 传到 <<<,>>> 里去，这句话可牛逼大了——相当于发号施令，命令那些线程去干活。这里使用了 32*32 * 16*16 个线程来干活。你看明白了吗？blocks 表示用了二维的 32*32 个 block 组，而每个 block 中又用了 16*16 的二维的 thread 组。好吧，我们这个施令动用了 262144 个线程！我们先不管 GPU 内部是如何调度这些线程的，反正我们这一句话就是用了这么多线程。

那我们的内核函数 kernelfun() 如何知道自己执行的是哪个线程？这就是线程网络的特点啦，为什么叫网络，是有讲究的，网络就可以定格到网点：

比如 int tid = threadIdx.x + blockIdx.x * 16

这里有一个讲究，block 是有维度的，一维、二维、三维。

对于一维的 block : tid = threadIdx.x
对于（Dx，Dy）二维的 block : tid = threadIdx.x + Dx*threadIdx.y
对于（Dx，Dy，Dz）三维的 block : tid = threadIdx.x + Dx*threadIdx.y + Dz*Dy*threadIdx.z

我习惯的用这样的模式去分配，比较通用：

dim3 dimGrid();
dim3 dimBlock();
kerneladd<<<dimGrid, dimBlock>>>();

这可是万金油啊，你需要做的事情是填充 dimGrid 和 dimBlock 的结构体构造函数变量，比如，dimGrid(16, 16) 表示用了 16*16 的二维的 block 线程块。

（0,0）（0,1）（0,2）……（0,15）

（1,0）（1,1）（1,2）……（1,15）

（2,0）（2,1）（2,2）……（2,15）

　……

（15,0）（15,1）（15,2）……（15,15）

（，） 是 （dimGrid.x， dimGrid.y）的网格编号。

我们这么理解吧，现在又一群人，我们分成 16*16 个小组（block），排列好，比如第 3 行第 4 列就指的是（2,3）这个小组。

而 dimBlock（16,16）表示每个小组有 16*16个成员，如果你想点名第 3 行第 4 列这个小组的里面的第 3 行第 4 列那个同学，那么，你就是在（2,3）这个 block 中选择了（2,3）这个线程。这样应该有那么一点可以理解进去的意思了吧？不理解透彻么什么关系，这个东西本来就是 cuda 中最让我纠结的事情。我们且不管如何分配线程，能达到最优化，我们的目标是先让 GPU 正确地跑起来，计算出结果即可，管他高效不高效，管他环保不环保。

唠叨了这么多，下面我们用一个最能说明问题的例子来进一步理解线程网络分配机制来了解线程网络的使用。

一维网络线程

eg：int arr[1000]，对每个数组元素进行加 1 操作。

idea：我们最直接的想法，是调度 1000 个线程去干这件事情。

first pro：我想用一个小组的 1000 个人员去干活。这里会存在这样一个问题 —— 一个小组是不是有这么多人员呢？是的，这个事情你必须了解，连自己组内多少人都不知道，你也不配作指挥官呀。对的，这个参数叫做 maxThreadsPerBlock，如何取得呢？

好吧，cuda 定义了一个结构体 cudaDeviceProp，里面存入了一系列的结构体变量作为 GPU 的参数，除了 maxThreadsPerBlock，还有很多信息哦，我们用到了再说。

maxThreadsPerBlock 这个参数值是随着 GPU 级别有递增的，早起的显卡可能 512 个线程，我的 GT520 可以跑 1024 个线程，办公室的 GTX650ti2G 可以跑 1536 个，无可非议，当然多多益善。一开始，我在想，是不是程序将每个 block 开的线程开满是最好的呢？这个问题留在以后在说，一口吃不成胖子啦。

好吧，我们的数组元素 1000 个，是可以在一个 block 中干完的。

内核函数：

#define N 1000
__gloabl__ void kerneladd(int *dev_arr)
{
    int tid = threadIdx.x;
    if (tid < 1000)
        dev_arr[tid] ++;
}

int main()
{
    int *arr, *dev_arr;// 习惯的我喜欢在内核函数参数变量前加个dev_作为标示

    // 开辟主机内存
    arr = (int*)malloc(N*sizeof(int));

    // 开辟设备内存

    // 主机拷贝到设备

    kerneladd<<<1, N>>>(dev_arr);

    // 设备拷贝到主机

    // 打印

    // 释放设备内存

    // 释放主机内存

    return 0;

}

呀，原来这么简单，个么 CUDA 也忒简单了哇！这中想法是好的，给自己提高信心，但是这种想法多了是不好的，因为后面的问题多了去了。

盆友说，1000 个元素，还不如 CPU 来的快，对的，很多情况下，数据量并行度不是特别大的情况下，可能 CPU 来的更快一些，比较设备与主机之间互相调度操作，是会有额外开销的。
有人就问了，一个 10000 个元素的数组是不是上面提供的 idea 就解决不了啦？
对，一个 block 人都没怎么多，如何完成！
这个情况下有两条路可以选择：

第一，我就用一个组的 1000 人来干活话，每个人让他干 10 个元素好了。

这个解决方案，我们需要修改的是内核函数：

__global__ void kernelarr(int *dev_arr)
{
    int tid = threadIdx.x;
    if(tid < 1000) // 只用0~999号线程
    { 
    //每个线程处理10个元素，比如0号线程处理0、1001、2001、……9001
        for(int i = tid; i<N; i=i+1000)
        {
            dev_arr[tid] ++;
        }
    }

}

第二，我多用几个组来干这件事情，比如我用 10 个组，每个组用 1000 人。
这个解决方案就稍微复杂了一点，注意只是一点点哦~因为，组内部怎么干活和最原始的做法是一样的，不同之处是，我们调遣了 10 个组去干这件事情。

首先我们来修改我们的主机函数：

int main()
{
……
    kerneladd<<<10, 1000>>>(dev_arr); //我们调遣了10个组，每个组用了1000人
……

}

盆友要问了，10 个组每个组 1000 人，你怎么点兵呢？很简单啊，第 1 组第 3 个线程出列，第 9 组第 9 个线程出列。每个人用组号和组内的编号定了位置。在线程网络中，blockIdx.x 和 threadIdx.x 就是对应的组号和组内编号啦，我必须要这里开始形象点表示这个对应关系，如果这个对应关系是这样子的[blockIdx.x，threadIdx.x]，那么我们的数组 arr[10000] 可以这样分配给这 10 个组去干活：

(0,0)->arr[0]，　　　(0,1)->arr[1]，　　　　……(0,999)->arr[999]

(1,0)->arr[0+1*1000]，(1,1)->arr[1+1*1000]，…… (1,999)->arr[999+1*1000]

……

(9,0)->arr[0+9*1000]，(9,1)->arr[1+9*1000]，……(9,999)->arr[999+9*1000]

是不是很有规律呢？对的，用 blockIdx.x 和 threadIdx.x 可以很好的知道哪个线程干哪个元素，这个元素的下表就是 threadIdx.x + 1000*blockIdx.x。

这里我想说的是，如果我们哪天糊涂了，画一画这个对应关系的表，也许，就更加清楚的知道我们分配的线程对应的处理那些东西啦。

一维线程网络，就先学这么多了。

二维网络线程

eg2：int arr[32][16]二维的数组自增 1。

第一个念头，开个 32*16 个线程好了哇，万事大吉！好吧。但是，朕现在想用二维线程网络来解决，因为朕觉得一个二维的网络去映射一个二维的数组，朕看的更加明了，看不清楚自己的士兵，如何带兵打仗！

我还是画个映射关系：

一个 block 中，现在是一个二维的 thread 网络，如果我用了 16*16 个线程。

(0,0)，(0,1)，……(0,15)

(1,0)，(1,1)，……(1,15)

……

(15,0)，(15,1)，……(15,15)

呀，现在一个组内的人称呼变了嘛，一维网络中，你走到一个小组里，叫 3 号出列，就出来一个，你现在只是叫 3 号，没人会出来！这个场景是这样的，现在你班上有两个人同名的人，你只叫名，他们不知道叫谁，你必须叫完整点，把他们的姓也叫出来。所以，二维网络中的 (0,3) 就是原来一维网络中的 3，二维中的 (i,j) 就是一维中的 (j+i*16)。不管怎么样，一个 block 里面能处理的线程数量总和还是不变的。

一个 grid 中，block 也可以是二维的，一个 block 中已经用了 16*16 的 thread 了，那我们一共就 32*16 个元素，我们用 2 个 block 就行了。

先给出一个代码清单吧，程序员都喜欢看代码，这段代码是我抄袭的。第一次这么完整的放上代码，因为我觉得这个代码可以让我说明我想说的几个问题：

第一，二维数组和二维指针的联系。
第二，二维线程网络。
第三，cuda 的一些内存操作，和返回值的判断。

#include <stdio.h> 
#include <stdlib.h> 
#include <cuda_runtime.h> 

#define ROWS 32 
#define COLS 16 
#define CHECK(res) if(res!=cudaSuccess){exit(-1);} 

__global__ void Kerneltest(int **da, unsigned int rows, unsigned int cols) 
{ 
    unsigned int row = blockDim.y*blockIdx.y + threadIdx.y; 
    unsigned int col = blockDim.x*blockIdx.x + threadIdx.x; 
    if (row < rows && col < cols) 
    { 
        da[row][col] = row*cols + col; 
    } 
} 
int main(int argc, char **argv) 
{ 
    int **da = NULL; 
    int **ha = NULL; 
    int *dc = NULL; 
    int *hc = NULL; 

    cudaError_t res; 
    int r, c; 
    bool is_right=true; 

    res = cudaMalloc((void**)(&da), ROWS*sizeof(int*));
    CHECK(res)

    res = cudaMalloc((void**)(&dc), ROWS*COLS*sizeof(int));
    CHECK(res)

    ha = (int**)malloc(ROWS*sizeof(int*)); 
    hc = (int*)malloc(ROWS*COLS*sizeof(int)); 


    for (r = 0; r < ROWS; r++) 
    { 
        ha[r] = dc + r*COLS; 
    }

    res = cudaMemcpy((void*)(da), (void*)(ha), ROWS*sizeof(int*), cudaMemcpyHostToDevice);
    CHECK(res)

    dim3 dimBlock(16,16); 
    dim3 dimGrid((COLS+dimBlock.x-1)/(dimBlock.x), (ROWS+dimBlock.y-1)/(dimBlock.y));

    Kerneltest<<<dimGrid, dimBlock>>>(da, ROWS, COLS); 


    res = cudaMemcpy((void*)(hc), (void*)(dc), ROWS*COLS*sizeof(int), cudaMemcpyDeviceToHost);
    CHECK(res)

    for (r = 0; r < ROWS; r++) 
    { 
        for (c = 0; c < COLS; c++) 
        { 
            printf("%4d ", hc[r*COLS+c]); 
            if (hc[r*COLS+c] != (r*COLS+c)) 
            { 
                is_right = false; 
            } 
        } 
        printf("\n"); 
    } 

    printf("the result is %s!\n", is_right? "right":"false"); 

    cudaFree((void*)da); 
    cudaFree((void*)dc); 

    free(ha); 
    free(hc); 

    getchar(); 
    return 0; 
}

简要的来学习一下二维网络这个知识点，

dim3 dimBlock(16,16); //定义block内的thread二维网络为16*16

dim3 dimGrid((COLS+dimBlock.x-1)/(dimBlock.x), (ROWS+dimBlock.y-1)/(dimBlock.y)); //定义grid内的block二维网络为1*2

unsigned int row = blockDim.y*blockIdx.y + threadIdx.y; //二维数组中的行号

unsigned int col = blockDim.x*blockIdx.x + threadIdx.x; //二维线程中的列号

三维网络线程

dim3 定义了三维的结构，但是，貌似二维之内就能处理很多事情啦，所以，我放弃学习三维。网上看到的不支持三维网络是什么意思呢？先放一放。

给自己充充电

同一块显卡，不管你是二维和三维或一维，其计算能力是固定的。比如一个 block 能处理 1024 个线程，那么，一维和二维线程网络是不是处理的线程数一样呢？

回答此问题，先给出网络配置的参数形式——<<<Dg,Db,Ns,S>>>，各个参数含义如下：

Dg：定义整个 grid 的维度，类型 Dim3，但是实际上目前显卡支持两个维度，所以，dim3<<Dg.x, Dg.y, 1>>> 第 z 维度默认只能为 1，上面显示出这个最大有 65536*65536*1，每行有 65536 个 block，每列有 65536 个 block，整个 grid 中一共有 65536*65536*1 个 block。
Db：定义了每个 block 的维度，类型 Dim3，比如 512*512*64，这个可以定义 3 维尺寸，但是，这个地方是有讲究了，三个维度的积是有上限的，对于计算能力 1.0、1.1 的 GPU，这个值不能大于 768，对于 1.2、1.3 的不能大于1024，对于我们试一试的这块级别高点的，不能大于 1536。这个值可以获取哦 —— maxThreadsPerBlock
Ns：这个是可选参数，设定最多能动态分配的共享内存大小，比如 16k，单不需要是，这个值可以省略或写 0。
S：也是可选参数，表示流号，默认为 0。流这个概念我们这里不说。

接着，我想解决几个你肯定想问的两个问题，因为我看很多人想我这样的问这个问题：

block 内的 thread 我们是都饱和使用吗？
答：不要，一般来说，我们开 128 或 256 个线程，二维的话就是 16*16。
grid 内一般用几个 block 呢？
答：牛人告诉我，一般来说是你的流处理器的 4 倍以上，这样效率最高。

回答这两个问题的解释，我想抄袭牛人的一段解释，解释的好的东西就要推广呀：

GPU 的计算核心是以一定数量的 Streaming Processor(SP) 组成的处理器阵列，NV 称之为 Texture Processing Clusters(TPC)，每个 TPC 中又包含一定数量的 Streaming Multi-Processor(SM)，每个 SM 包含 8 个 SP。SP 的主要结构为一个 ALU（逻辑运算单元），一个 FPU （浮点运算单元）以及一个 Register File(寄存器堆)。SM 内包含有一个 Instruction Unit、一个 Constant Memory、一个 Texture Memory，8192 个 Register、一个 16KB 的 Share Memory、8 个 Stream Processor(SP) 和两个 Special Function Units（SFU）。（ GeForce9300M GS 只拥有 1 个 SM ） Thread 是 CUDA 模型中最基本的运行单元，执行最基本的程序指令。Block 是一组协作 Thread，Block 内部允许共享存储，每个Block 最多包含 512 个 Thread。Grid 是一组 Block，共享全局内存。Kernel 是在 GPU 上执行的核心程序，每一个 Grid 对应一个 Kernel 任务。在程序运行的时候，实际上每 32 个T hread 组成一个 Warp，每个 warp 块都包含连续的线程，递增线程 ID 。Warp 是 MP 的基本调度单位，每次运行的时候，由于 MP 数量不同，所以一个 Block 内的所有 Thread 不一定全部同时运行，但是每个 Warp 内的所有 Thread 一定同时运行。因此，我们在定义 Block Size 的时候应使其为 Warp Size 的整数倍，也就是 Block Size 应为 32 的整数倍。理论上 Thread 越多，就越能弥补单个 Thread 读取数据的 latency ，但是当 Thread 越多，每个 Thread 可用的寄存器也就越少，严重的时候甚至能造成 Kernel 无法启动。因此每个 Block 最少应包含 64 个 Thread，一般选择 128 或者 256，具体视 MP 数目而定。一个 MP 最多可以同时运行 768 个 Thread，但每个 MP 最多包含 8 个 Block，因此要保持 100% 利用率，Block 数目与其 Size 有如下几种设定方式： Ø 2 blocks x 384 threads Ø 3 blocks x 256 threads Ø 4 blocks x 192 threads Ø 6 blocks x 128 threads Ø 8 blocks x 96 threads

这些电很重要啊，必须要充！不然，我就很难理解为什么网络线程如何分配的。

__Sunny__

关注

0
点赞
踩
2

收藏

觉得还不错? 一键收藏
0
评论
线程网络的分配

转载自从0开始学习《GPU高性能运算之CUDA》——25 GPU也不允许偏心并行的事情多了，我们作为 GPU 的指令分配者，不能偏心了——给甲做的事情多，而乙没事做，个么甲肯定不爽的来。所以，在 GPU 中，叫做线程网络的分配。
复制链接

扫一扫