[转]浅析GPU计算—

在《浅析GPU计算——CPU和GPU的选择》一文中，我们分析了在遇到什么瓶颈时需要考虑使用GPU去进行计算。本文将结合cuda编程来讲解实际应用例子。（转载请指明出于breaksoftware的csdn博客）

之前我们讲解过，CPU是整个计算机的核心，它的主要工作是负责调度各种资源，包括其自身的计算资源以及GPU的计算计算资源。比如一个浮点数相乘逻辑，理论上我们可以让其在CPU上执行，也可以在GPU上执行。那这段逻辑到底是在哪个器件上执行的呢？cuda将决定权交给了程序员，我们可以在函数前增加修饰词来指定。

关键字	执行位置
__host__	CPU
__global__	GPU
__device__	GPU

一般来说，我们只需要2个修饰词就够了，但是cuda却提供了3个——2个执行位置为GPU。这儿要引入一个“调用位置”的概念。父函数调用子函数时，父函数可能运行于CPU或者GPU，相应的子函数也可能运行于CPU或者GPU，但是这绝不是一个2*2的组合关系。因为GPU作为CPU的计算组件，不可以调度CPU去做事，所以不存在父函数运行于GPU，而子函数运行于CPU的情况。

关键字	调用位置
__host__	CPU
__global__	CPU
__device__	GPU

__global__描述的函数就是“被CPU调用，在GPU上运行的代码”，同时它也打通了__host__和__device__修饰的函数。

如果一段代码既需要运行于CPU，也要运行于GPU，怎么办？难道要写两次？当然不用，我们可以同时使用__host__和__device__修饰。这样编译器就会帮我们生成两份代码逻辑。

#include "cuda.h" #include "cuda_runtime.h"#include "device_launch_parameters.h"#include <stdio.h> __host__ __device__ int run_on_cpu_or_gpu() { return 1;} __global__ void run_on_gpu() {   printf("run_on_cpu_or_gpu GPU: %d\n", run_on_cpu_or_gpu());} int main() { printf("run_on_cpu_or_gpu CPU: %d\n", run_on_cpu_or_gpu());   run_on_gpu<<<1, 1>>>();   return 0;}

上例中，第11行调用的是run_on_cpu_or_gpu的GPU实现，而第15行的调用的是的CPU实现。

可以看到，为了实现上面的例子，我引入了一个__global__函数——run_on_gpu用于衔接CPU和GPU，那么有人可能会问：如果__global__和__host__、__device__一起修饰函数不就行了？实际上cuda禁止了这种方案，而且__global__不能和它们中任何一个修饰符一起使用。

__global__这个修饰符的使命使得它足够的特殊。比如它修饰的函数是异步执行的。我们稍微修改一下上面main函数

int main() { printf("run_on_cpu_or_gpu CPU: %d\n", run_on_cpu_or_gpu());   run_on_gpu<<<1, 1>>>();   printf("will end\n"); return 0;}

如果上述代码都是同步执行的，那么will end将是最后一条输出，然而输出却是

可见__global__修饰的函数是被异步执行的。

__global__修饰的函数只能是void类型。我们假设其可以有返回值，则其调用者是可以捕获这个值的。但是__global__函数是异步调用的，当函数返回时，接受返回值的变量可能已经被销毁了。所以设计其有返回值也没太多意义。

int main() { printf("run_on_cpu_or_gpu CPU: %d\n", run_on_cpu_or_gpu());   {       int ret = run_on_gpu<<<1, 1>>>(); // error!!！even if run_on_gpu return int!!  }   printf("will end\n"); return 0;}

还有人会问，上面main函数怎么没有用修饰符修饰?cuda编程规定如果没有使用修饰符修饰的默认就是__host__类型。这种设计让大家熟悉的规则成为默认的规则，可以让更多第三方代码不用修改就直接被cuda编译器编译使用。

cuda是一个GPU编程环境，所以它对__device__修饰的函数进行了比较多的优化。比如它会根据它的规则，让某个__device__修饰函数成为内联函数（inline）。这些规则是程序员不可控，但是如果我们的确对是否内联有需求，cuda也提供了方式：使用__noinline__修饰函数不进行内联优化；使用 __forceinline__修饰函数强制进行内联优化。当然这两种修饰符不能同时使用。

也许你已经发现，__global__函数调用方式非常特别——使用“<<<>>>”标志。这儿就需要引入cuda的并行执行的线程模型来解释了。在同一时刻，一个cuda核只能运行一个线程，而线程作为逻辑的运行载体有其自己的ID。这个ID和我们在linux或windows系统上CPU相关的线程ID有着不同的表达方式。比如在Linux系统上可以使用gettid方法获取一个pid_t值，比如3075。但是cuda的表达方式是一个三维空间，表达这个空间的是一个叫block的概念。比如单个block定义其有(Dx, Dy, 0)个线程，则每个线程ID为x+yDx；再比如有(Dx, Dy, Dz)个线程，则每个线程ID为x+yDx+zDxDy。

__global__ void run_on_gpu() {   printf("GPU thread info X:%d Y:%d Z:%d\t block info X:%d Y:%d Z:%d\n",        threadIdx.x, threadIdx.y, threadIdx.z, blockIdx.x, blockIdx.y, blockIdx.z);} int main() {   dim3 threadsPerBlock(2, 3, 4);  int blocksPerGrid = 1; run_on_gpu<<<blocksPerGrid, threadsPerBlock>>>(); return 0;}

其运行结果如下

承载block的是一个叫做grid的概念。block在grid中的位置也是通过一个三维结构来表达，比如下面代码标识的是一个一个Grid包含(3,3,3)结构的Block，一个Block包含(3,3,3)结构的Thread

 dim3 blocksPerGrid(3, 3, 3);    dim3 threadsPerBlock(3, 3, 3);  run_on_gpu<<<blocksPerGrid, threadsPerBlock>>>();

对于上例中的各个线程的ID算法就更加复杂了，详细的计算规则可以见《CUDA（10）之深入理解threadIdx》。

为什么cuda的线程要设计的这么复杂？我想其可能和GPU设计的初始目的有关——图像运算。而我们肉眼的感官就是三维的，所以GPU有大量三维计算的需求。

个人觉得大家不要拘泥于threadID的计算，而要学会如何利用blockIdx、threadIdx的三维坐标来进行并行计算。比如我们把问题的维度降低，看看一个Grid只有一个Block，一个Block中N*N个线程，计算两个N*N矩阵相加的例子

// Kernel definition __global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {   int i = threadIdx.x;   int j = threadIdx.y;   C[i][j] = A[i][j] + B[i][j]; } int main() {   ...     // Kernel invocation with one block of N * N * 1 threads    int numBlocks = 1;     dim3 threadsPerBlock(N, N);     MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);  ... }

矩阵相加是一个非常好说明GPU并行计算的例子。结合上面的代码，我们假设GPU中有大于N*N个空闲的cuda核，且假设调度器同时让这N*N个线程运行，则整个计算的周期可以认为是一个元的计算周期。而CPU却要串行处理每个元的计算（不考虑CPU中向量计算单元）。

那矩阵相加的什么特性让其成为一个经典的案例呢？那就是“可并行性”！因为每个元的计算都不依赖于其他元的计算结果，所以这种计算是适合并行进行的。如果一个逻辑的“可并行计算单元”越多越连续，其就越适合使用GPU并行计算来优化性能。

可能有人提出一个疑问：假如N非常大，比如是128，那么就需要100*100=10000个线程来执行。而目前最高配的GPU只有5120个cuda核，那这些线程是如何在cuda核上调度的呢？这儿要引入一个叫做warp的概念，它是一个线程集合。假如GPU的warp包含32个线程，则我们的任务被打包成10000/32 =313个warp（除不尽。其中312个warp包含32个我们的计算线程，而剩下的16将和16个空线程被打包成一个warp）。这313个warp分批次在GPU上调度执行，这样设计有几个好处：可以成批有序的执行线程逻辑，且发生等待操作时，可以切换其他warp去工作，从而提供cuda核心的利用率。

我们再看下计算的数据源。一般情况下，数据源是由CPU发射到GPU上去的，于是连接GPU和主板的PCIe接口带宽至关重要。目前设计中的PCIe5的理论最大带宽可以达到64GB/s（https://zh.wikipedia.org/wiki/PCI_Express），而CPU和内存之间的访问带宽也不到100GB/s。通过这两者比较，我们发现内存和PCIe之间，明显PCIe带宽是瓶颈。那这个瓶颈重要么？我们再看下目前NV的显存带宽是多少——超过700GB/s。可以说PCIe和显存带宽差一个数量级。下图可以让我们更直观的看到问题

如果我们把数据从DRAM一次性保存到GDRAM，然后GPU直接访问GDRAM上的数据，这和频繁通过PCIe读写数据相比，无疑大幅提高了效率。等到数据都处理完毕，我们再将GDRAM的数据发送回CPU。cuda为我们提供了cudaMemcpy系列方法，让我们可以在DRAM和GDRAM之间传递数据。我们把上面矩阵相加的例子加以改造

// Device code __global__ void VecAdd(float* A, float* B, float* C, int N) {     int i = blockDim.x * blockIdx.x + threadIdx.x;    if (i < N)       C[i] = A[i] + B[i]; }  // Host code int main() {  int N = ...;   size_t size = N * sizeof(float);   // Allocate input vectors h_A and h_B in host memory    float* h_A = (float*)malloc(size);     float* h_B = (float*)malloc(size);     // Initialize input vectors     ...     // Allocate vectors in device memory    float* d_A; cudaMalloc(&d_A, size);     float* d_B; cudaMalloc(&d_B, size);     float* d_C; cudaMalloc(&d_C, size);     // Copy vectors from host memory to device memory   cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);     cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);     // Invoke kernel    int threadsPerBlock = 256;     int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;  VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);     // Copy result from device memory to host memory    // h_C contains the result in host memory   cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);     // Free device memory   cudaFree(d_A);  cudaFree(d_B);  cudaFree(d_C);  // Free host memory     ... }

通过这两章的浅析，我们可以大致了解GPU并行计算的相关概念，以及使用cuda实现并行计算的基本操作。如果大家还想有更加深入的了解，可以参见《cuda c programming guide》。还可以参考周斌老师《NVIDIA CUDA初级教程视频》。

[转]浅析GPU计算——cuda编程相关推荐

浅析GPU计算——cuda编程
在<浅析GPU计算--CPU和GPU的选择>一文中,我们分析了在遇到什么瓶颈时需要考虑使用GPU去进行计算.本文将结合cuda编程来讲解实际应用例子.(转载请指明出于breaksoftwa ...
浅析GPU计算——CPU和GPU的选择
目前市面上介绍GPU编程的博文很多,其中很多都是照章宣科,让人只能感受到冷冷的技术,而缺乏知识的温度.所以我希望能写出一篇可以体现技术脉络感的文章,让读者可以比较容易理解该技术,并可以感悟到cuda编 ...
linux运行并行计算cuda,并行化计算与CUDA编程
原标题:并行化计算与CUDA编程近年来,显卡的更新换代也比较快,NVIDIA今年的发布会也宣布了RTX3080TI即将到来.显卡的运算能力也越来越强. 很多人对显卡的了解可能源于游戏,但是你的显卡不 ...
2023年的深度学习入门指南(10) - CUDA编程基础
2023年的深度学习入门指南(10) - CUDA编程基础上一篇我们走马观花地看了下SIMD和GPGPU的编程.不过线条太粗了,在开发大模型时遇到问题了肯定还会晕. 所以我们还是需要深入到CUDA中 ...
计算引擎-GPU计算的优势
基于GPU的大规模计算早有耳闻,直到中午吃饭范W聊到他的多层神经网络程算法在GPU和CPU上的运行效率是40:1的时候,我才恍然大悟. 中央处理CPU 用来解释机器指令以及处理计算数据.首先是控制部件 ...
CUDA 编程简单入门 Advance CUDA 编程基础 (C++ programming)
Advance CUDA编程基础 (C++ programming) GPU 架构 CUDA 编程基础基本代码框架 CUDA Execution Model Case Study : Vector ...
CUDA编程基础与实践学习笔记（一）
开个坑- 用pytorch太不灵活了,学习一下CUDA编程. 研究生都第二年了,坚持不下去就有点丢人了奥. 书名见标题,樊哲勇老师的著作,清华大学出版社出版. 重点看前面十二章. 初次阅读,如有理解错 ...
CUDA编程（一）：GPU计算与CUDA编程简介
CUDA编程(一):GPU计算与CUDA编程简介 GPU计算 GPU硬件资源 GPU软件资源 GPU存储资源 CUDA编程 GPU计算 NVIDIA公司发布的CUDA是建立在GPU上的一个通用并行计算 ...
CUDA刷新：GPU计算生态系统
CUDA刷新:GPU计算生态系统 CUDA Refresher: The GPU Computing Ecosystem 这是CUDA Refresher系列的第三篇文章,其目标是刷新CUDA中的关键 ...

[转]浅析GPU计算——cuda编程

[转]浅析GPU计算——cuda编程相关推荐

最新文章

热门文章