CUDA编程之快速入门-----运行机制

最新推荐文章于 2024-09-12 15:07:24 发布

sru_alo

最新推荐文章于 2024-09-12 15:07:24 发布

阅读量1k

点赞数 1

分类专栏： gpu相关

gpu相关专栏收录该内容

4 篇文章 2 订阅

订阅专栏

参考：原文：https://blog.csdn.net/AuntieLee/article/details/82492369

优达学城：https://classroom.udacity.com/courses/cs344/lessons/55120467/concepts/670611900923

CUDA Zone：https://www.nvidia.cn/object/cuda_education_cn_old.html

应用程序利用GPU实现加速的总体分工是：密集计算代码（约占5%的代码量）由GPU负责完成，剩余串行代码由CPU执行。

一、CPU与GPU的关系

1.计算机有两个不同的处理器，CPU（host）和GPU（device）

2.单纯的C语言编程（plain C program）只能让CPU工作

二、CUDA作用

1.CUDA：可以让代码在CPU和GPU上运行

2.CUDA:将GPU看作是CPU的协同处理器，并假设它们分别有各自的内存。

3.把你的程序分成小块，分别在CPU和GPU上运行，并分别为他们生成代码。

三、程序运行时CPU与GPU的关系

1.CPU和GPU的关系中,CPU占主导，它运行主程序，并向GPU发送指令，告诉GPU做什么。

*也就是说，无论是CPU向GPU发送数据还是GPU向CPU发送数据，都要CPU来控制/initiate。

2.运行时系统要做如下工作：

①将CPU中的数据送到GPU中

②将GPU的数据返回CPU

在C程序中，将数据从一个地方copy到另一个地方的命令叫Memcy；不难理解在gpu运算中，将数据传输来传输去叫做cudaMemcy。

③在GPU上分配存储

C程序中，分配存储的命令叫做Malloc；GPU运算中，就叫做cudaMalloc。

④调用kernel函数，即在GPU上并行运算的程序（行话中：The host launches kernels on the device.）

四、典型GPU程序构成

一个典型GPU程序有如下几个部分：

①CPU在GPU上分配内存

②CPU将CPU中的数据copy到GPU中

③调用内核函数来处理数据

④CPU将GPU中的数据copy到CPU中

*可以看出，四个步骤中有两个是数据的copy，因此如果你的程序需要不断地进行copy，那么运行效率会比较低，不适合利用GPU运算。一般情况下，最好的方式是，让GPU进行大量运算，同时保证计算量与通信的比值很高。

五、运行程序时GPU是如何运作的？

前面提到，我们将运算看作是一个由一系列或多个内核组成的结构。GPU有很多并行的计算单元，那么当我们写kernels的时候，应该充分利用这些硬件并行。

★Big Idea（one of the very core concept of CUDA）：

当你在写程序的时候，你把它写的像是一个串行的程序，就好像是在一个单独的线程里运行一样。当你从CPU调用核的时候，就给核规定好要启动几个线程，每一个线程都会运行这个程序（的一部分？，存疑）。

小结：

GPU擅长：①启动很多的线程 ②并行地运行很多线程因此，如果你的程序没有利用它进行多线程计算，那么就没有很好的利用GPU。

六、一个典型的程序

程序要求：对一个数组当中的所有元素进行平方计算。

先看一下C语言版本的程序：

for (int i = 0; i < 64; i++){
    out[i] = in[i] * in[i];
}

C语言的程序有以下两个特点：

①只有一个线程在执行计算，该线程遍历输入数组内所有元素。（将该线程定义为“执行代码的一个独立路径”，这一定义也将适用于GPU代码中。）

②代码中没有显式并行性，是串行代码。这段代码是一个循环64次的线程，每一个循环进行一次平方计算。

对于GPU来说，程序要分为两个部分：①在CPU上运行 ②在GPU上运行的

图 GPU程序构成

GPU上运行的程序很简单，可以看作是一个普通的串行运算，在这个例子里就是 OUT = IN * IN，即计算平方。那么这个核函数就可以看作跟“并行”无关。

CPU上的程序则要负责分配内存、从GPU复制数值、将数值传递给GPU、启动Kernel函数。

*CPU启动kernels的过程：用64个线程，启动一个执行平方计算的kernel，即square_kernel。每一个内核的实例（instance）将运算一个平方计算。

*线程索引：分辨每一个独立的内核做的工作。

CPU code：

square_kernel<<<64>>>(input_array, output_array)

完整的代码如下：

#include <stdio.h>
 
//内核函数，就跟一个串行程序一样
//__global__ : 定义说明符，用来表明这是一个内核，而非普通函数
//kernel的参数都必须是分配在GPU上的，不然会程序崩溃
__global__ void square(float * d_out, float * d_in){
    int idx = threadIdx.x; //用来得到线程的索引值,threadIdx是一个CUDA内置变量，其本质是一个C结构，有三个成员，x、y、z。
    float f = d_in[idx];
    d_out[idx] = f * f;
}
 
//下面是CPU代码
int main(int argc, char ** argv) {
    const int ARRAY_SIZE = 64;
    const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);
 
//generate input array on the host
    float h_in[ARRAY_SIZE];
    for (int i = 0; i < ARRAY_SIZE; i++){
        h_in[i] = float(i);
    }
    float h_out[ARRAY_SIZE];
 
//declare GPU memory pointers,看起来跟c program一样，是一个普通的数组
    float * d_in;
    float * d_out;
 
//allocate GPU memory，在这里给它分配GPU内存
    cudaMalloc((void **) &d_in, ARRAY_BYTES);
    cudaMalloc((void **) &d_out, ARRAY_BYTES);
 
//transfer the array to the GPU，前三个参数跟C Memcpy一样，目标地址、源地址、字节数，最后一个参数表明转移方向cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost,cudaMemcpyDeviceToDevice
    cudaMencpy(d_in, h_in, ARRAY_BYTES,cudaMemcpyHostToDevice);
 
//launch the kernel, cuda启动运算符，在一个具有64个元素的block上启动一个核，两个参数d_in,d_out
//这句话告诉CPU，在GPU上启动64个线程，运行64个内核副本
//注意，我们只能在调用内核处理GPU数据，而不能处理CPU数据
    square<<<1,ARRAY_SIZE>>>(d_out, d_in);
 
//copy back the result array to the CPU，完成内核运行以后，结果在d_out中
    cudaMemcpy(h_out, d_out, ATTAY_BYTES, cudaMemcpyDeviceToHost);
 
//print out the resulting array
    for (int i = 0; i < ARRAY_SIZE; i++){
        printf("%f",h_out[i]);
        printf(((i % 4) != 3) ? "\t" : "\n" );
    }
 
//free GPU memory allocation
    cudaFree(d_in);
    cudaFree(d_out);
 
    return 0;
}

需要注意的是，在CUDA程序中，我们以h_开头表示host上的数据，d_开头表示GPU上的数据。

针对代码，有值得注意的几点：

(1)square<<<1,64>>>(d_out, d_in)

参数说明： 1：number of blocks 64：Threads Per Block

从硬件角度说: ①硬件可以同时运行多个块（blocks），②每个块有线程数限制，较新的GPU一个块可以支持1024个线程，旧的512.

(2)另外，这是一个一维diagram，只在一个维度发展。如果想要处理2/3等维度问题，则需要如下代码：

Kernel<<<GRID OF BLOCKS, BLOCK OF THREADS，shemem>>>(...)

说明：

GRID OF BLOCKS：线程块网格维数，取值1-3

BLOCK OF THREADS：每个块中线程维数，取值1-3

dim3(x,y,z)：线程/blocks维数 = x*y*z，xyz分别代表三个方向的值

shemem：每一个块共享的内存，字节数