深度学习CUDA编程⼲货-kernel的编写和调⽤
title: 深度学习CUDA编程⼲货-kernel的编写和调⽤
五年级语文试卷date: 2020-06-09 12:08:54
category: 默认分类
本⽂介绍 深度学习CUDA编程⼲货-kernel的编写和调⽤
深度学习CUDA编程⼲货-kernel的编写和调⽤
腌牛肉的腌制方法本⽂由林⼤佬原创，转载请注明出处，来⾃腾讯、阿⾥等⼀线AI算法⼯程师组成的QQ交流群欢迎你的加⼊: 1037662480
上⼀篇给⼤家分享了⼀些CUDA编程的⼲货，这⼀篇来夯实⼀下，我们主要看⼀些基础的cuda概念。
三个层级
cuda编程主要是三个层级，分别是 thread，block和grid。
多个thread组成block；
多个block组成grid；
这样的话，thread实际上就是cuda⾥⾯最⼩的运⾏单位，实际上也是如此，你在kernel⾥⾯定义的threadIdx是唯⼀的标志。所以当我们要进⾏⼀个并⾏计算的时候，我们第⼀步要做的，实际上是如何设计这个grid，block以及thread配套使⽤。但是在这之前，我们还是需要了解⼀下CUDA kernel的调⽤⽅式以及它的意义。
CUDA的kernel
我们这⾥有⼀个kernel，很简单的：
__global__ void kernel(param list){ }
kernel<<<Dg,Db, Ns, S>>>(param list);
我们主要关⼼⼀下这⾥的 <<<Dg, Db, Ns, S>>>是什么意思？
事实上总结起来很简单：
三个尖括号没什么⾼级的，就是告诉编译器，接下来我要调⽤cuda的kernel了；
⾥⾯的kernel参数（不是kernel函数的参数），是告诉编译器⽤多少个线程来执⾏它；
⾥⾯有4个参数，其中都是可选的；
Dg，第⼀个参数指的是定义Grid的维度和尺⼨，为dim3的类型，也就是说⼀个grid有多少个block；
Db，同理指的是block的维度和尺⼨，是dim3的类型，也就是⼀个block有多少个thread；
Ns，是可选的，它告诉编译器，它设置的是每个block除了静态分配的sharedmemory以外，最多能分配的sharedmemory⼤⼩，单位是byte；
S是cudastream的可选参数。
因此，实际上我们需要⽤到只有两个参数，⼤多数情况下，这两个参数，分别是⼀个grid⾥⾯多少个block，⼀个是⼀个block⾥⾯多少个thread。
上⾯说道，这⾥的值是⼀个dim。这个dim实际上是⼀个xyz的维度，它的构成是：
dim3 dim = {x, y, z};
dim3 dim = {23, 1, 1};
这⾥的z，实际上是grid的数⽬，xy是block的size。同理定义thread的⽅式也⼀样，因此我们有
总共block有 Dg.x * Dg.y * Dg.z 个blocks，然后thread的总数是 Db.x * Db.y * Db.z个，其中Db.x 和 Db.y的最⼤值为512，计算能⼒⽐较decent的显卡上最⼤值通常是有限制的。
kernel中gridIdx和blockIdx⽤法
上⾯只是理论，但是实际上我们不是传⼊⼀个dim3的⽅式给kernel的参数，⽽通常是这样的：
sum<<< 512, 512 >>>()
那么我们的问题来了：这⾥⾯的Idx循环到底是怎么循环的呢？。这⾥的int变量和dim3其实是⼀个意思，只不过这⾥指的是⼀维的排布。简单来说，就是⼀个grid⾥⾯有512个block，⼀个block⾥⾯有512个thread。
⽤⼀张图来表⽰grid和block以及thread的排布⽅式可以如图：
那么，如果我们有这样的定义：奶茶店装修设计
dim3 grid(3, 2, 1), block(5, 3, 1)
这个意思就是，grid⾥⾯有3x2排布的block，block⾥⾯是5x3的排布。
假如我们传⼊的kernel参数就上上⾯这个，那么在kernel⾥⾯，我们这样定义Idx：
int idx = blockDim.x * blockIdx.x + threadIdx.x;
这⾥⾯跑的循环是多少呢？
实际上我们循环了3x2x5x3=90次。实际上当你定义了dim都是两个三维的排布⽅式，那么相应的idx取值的时候也需要使⽤三维的排布⽅式。
三维的其实不是很好理解，我们⽤两维的来展⽰⼀下：
⽐如⼀个kernel这样调⽤：
kernel<<<4, 8>>>()
然后它的Idx排布是这样的：
我们有4个线程块，每个线程块有8个线程，那么我们如何获取每⼀个线程的idx？
int idx = blockIdx.x * blockDim.x + threadIdx.x
助学贷款申请
最后不得不说，如果是3维的情况，会⽐较复杂。这⾥给⼤家提供⼀段3维情况下的打印threadID的例⼦：星期六晚上
#include <stdio.h>
#include <cuda.h>
__global__ void printThreadIndex() {
int ix = threadIdx.x + blockIdx.x * blockDim.x;
int iy = threadIdx.y + blockIdx.y * blockDim.y;
身份证多长时间能办下来unsigned int idx = iy*blockDim.x * gridDim.x + ix;
printf("thread_id (%d,%d) block_id (%d,%d) coordinate (%d, %d), global index %2d \n",
threadIdx.x, threadIdx.y, blockIdx.x, blockIdx.y, ix, iy, idx);
}
int main(void) {
婚假有工资吗
dim3 grid(2, 3, 1), block(4, 8, 1);
printThreadIndex<<<grid, block>>>();
cudaRetDevice();
return 0;
}
CUDA blocksize设计的讲究
对于我们炼丹的来说，我们不关⼼太复杂的kernel设计，最简单的，对图⽚的w和h进⾏并⾏化的操作，那么讲道理，我只需要这样调⽤kernel就可以了：
kernel<<<244, 244>>>
元宵节在什么时候这样的话，w和h就并⾏了，⼤⼤的提⾼了效率，但是很多时候你会看到⼈们都是这么⽤的 ：
BLOCK=512;
dim3 cudaGridSize(uint n)
{
uint k = (n - 1) /BLOCK + 1;
uint x = k ;
uint y = 1 ;
if (x > 65535 )
{
x = ceil(sqrt(x));
y = (n - 1 )/(x*BLOCK) + 1;
}
dim3 d = {x,y,1} ;
return d;
}
int num = w*h;
FCOSforward_kernel << < cudaGridSize(num), BLOCK >> > ()
所以说，这⾥的BLOCK的设计是有什么讲究吗 ？如果你去⾕歌你会发现都是⼀些说合理设置BLOCK使得sharedmemory占满的说法，笔者试验下来，好像对于我们这么⼀些简单的使⽤场合，差别并不是很⼤。
重点是看起来那些优化⽐较复杂并且难以控制。等到后⾯我们更加深⼊的学习之后再来分享⼀波把。
