HibisciDai 主页关于我归档标签分类相册
CUDA基础
|CUDA
字数总计: 2.4k|阅读时长: 8 分钟

CUDA基础

[TOC]

CUDA基础

CUDA简介

CUDA(Compute Unified Device Architecture)，NVIDIA推出的的通用并行计算架构，该架构使GPU能够解决复杂的计算问题。

将GPU视为数据并行计算设备，在其上进行计算的分配和管理，而无需将其映射到图形API。

CUDA的安装与编译

下载并安装最新版本的CUDA SDK包。

  • CUDA Installation Guide for Microsoft Windows
  • NVIDIA CUDA Installation Guide for Linux

主机与设备

GPU可以看做并行执行非常多个线程的计算设备（compute device）。
CPU作为主机（host），一个系统可以存在一个主机和若干个设备。
CPU、GPU各自拥有相互独立的存储地址空间：主机端的内存和设备端的显存。

并行线程的组织结构

Thread

并行的基本单位，具有IDs。

Thread block

互相合作的线程组
Cooperative Thread Array
允许彼此同步
通过快速共享内存交换数据
以1维、2维或3维组织

Grid：一组thread block

以1维、2维或3维组织
共享全局内存

Kernel：在GPU上执行的核心程序

一个kernel函数中只有一个grid

GPU计算单元

SM代表流多处理器

每个SM包含8个或更多标量流处理器SP，以及少量的其它计算单元

每个SM具有单指令多数据（SIMD）架构

每个多处理器具有下列四种类型的片上存储器

  • 每个处理器有一组本地32位寄存器。
  • 并行数据高速缓存（shared memory），由所有处理器共享并实现共享内存空间。
  • 只读constant cache由所有处理器共享并加速从常量内存空间的读取。
  • 只读texture cache由所有处理器共享并加速从纹理内存空间的读取。

Warp

若干线程捆绑在一起执行成为warp。

每个block根据thread ID将thread分为多个warp。

warp是SM内的基本调度单位。在任何时刻，硬件只选择一个warp执行。

划分warp的方式始终相同，每个warp包含thread ID连续递增的线程，其中第一个warp从0开始。

SIMD

SIMD (Single Instruction Multiple Data)

指一条指令作用在多个数据上面（Intel X86的SSE/AVX指令）

执行模型

  • CUDA采用了SIMT（Single Instruction Multiple Thread，单指令多线程）执行模型，是SIMD的一种改进。
  • 重要的不同在于SIMD组织方法会向应用暴露SIMD宽度，而SIMT指定单线程的执行和分支行为。
  • 与SIMD向量机相反，SIMT允许程序员为独立标量线程编写线程级并行代码，也为协作线程编写数据并行代码。
  • 为了正确性，程序员可忽略SIMT行为；然而只要维护束内线程很少分支的代码就可显著提升性能。

多个设备

  • 使用多个GPU作为CUDA设备的情况下，仅当这些GPU的类型相同时，应用程序才能保证工作。

异构计算

运行在GPU上的CUDA并行计算函数称为kernel（内核函数）。

一个完整的CUDA程序是由一系列的设备端kernel函数并行步骤和主机端的串行处理步骤共同组成。

GPU线程映射关系

线程由SP（Scalar Processor）执行。

Thread blocks在SM上执行

Thread blocks do not migrate

一个内核函数作为一个grid启动

同一时刻GPU上只能运行一个内核函数

CUDA语法基础

Grid、Block及Kernel函数

  • Grid
1
dim3 Grid(2,2);//2*2个block
  • Block
1
dim3 Block(4,2,2);//每个Block有4*2*2线程
  • Kernel
1
kernelFunction<<<Grid,Block>>>(参数1,参数2...)
  • Kernel函数调用例子
1
2
3
4
5
__global__ void KernelFunc(...);
dim3 DimGrid(4,8);
dim3 DimBlock(4,4);
size_t SharedMemBytes=32;//64 bytes of shared memory
KernelFunc<<<DimGrid,DimBlock,SharedMemBytes>>>();//4*8x4*4

Kernel函数的定义与调用

运行在GPU上的程序成为kernel，必须通过__global__函数类型限定定义。

只能在host代码中调用。

代码演示对长度为N的两个向量A和B求和，结果存在向量C中。

1
2
3
4
5
6
7
8
9
//kernel定义
__global__ void VecAdd(float* A,float* B,float* C){
    int i=threadIdx.x;
    C[i]=A[i]+B[i];
}
int main(){
    //kernel调用
    VecAdd<<<1,N>>>(A,B,C);
}

线程层次

将长度为N * N的两个矩阵A和B相加，存入矩阵C。

由于块内的所有线程必须存在于同一个处理器核心中且共享该核心有限的存储器资源，因此，一个块内的线程数目是有限的。在目前的GPU上，一个线程块可以包含多达1024个线程。

一个内核可被多个同样大小的线程块执行，所以总的线程数等于每个块内的线程数乘以线程块数。 

1
2
3
4
5
6
7
8
9
10
11
12
// 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 调用每个blockN*N*1个线程
   int numBlocks = 1;   
   dim3 threadsPerBlock(N, N);  
   MatAdd<<<numBlocks,   threadsPerBlock>>>(A, B, C);  
}

线程结构

线程被组织成1、2、3维线程网格。

threadId是系统自定义，直接使用即可

对于一维的block，线程的threadID就是threadIdx.x；

对于大小为（Dx,Dy）的二维block，线程的threadID是threadIdx.x+threadIdx.y*Dx；

对于大小为（Dx,Dy,Dz）的三维block，线程的threadID是threadIdx.x+threadIdx.y*Dx+threadIdx.z*Dx*Dy；

CUDA关键字

device

储存于GPU上的 global memory 空间

和应用程序具有相同的生命期（lifetime）

可被grid中所有线程存取

constant

储存于GPU上的常量 constant memory 空间

和应用程序具有相同的生命期（lifetime）

可被grid中所有线程存取，CPU代码通过runtime函数存取

shared

存储于GPU上的thread block内的共享存储器

和thread block 具有相同的生命期

只能被thread block内的线程存取

Local变量

存储于SM内的寄存器和local memory

和thread具有相同的生命期

thread私有

CUDA函数及变量定义

函数定义方式 执行 调用
__decice__ float DeviceFunc() GPU GPU
__global__ void KernelFunc() GPU CPU
__host__ float HostFunc() CPU CPU
  • __global__定义kernel函数

必须返回void

  • device和host可以组合使用

则被定义的很熟在CPU和GPU上都被编译

  • __device__函数

不能被&运算符取地址
不支持递归调用
不支持静态变量
不支持可变长度参数函数调用

CUDA编程框架

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
//全局变量声明
__host__,...,__device__,...,__global__,__constant__,__texture__//函数原型声明

__global__ void kernelOne(...)//内核函数

float handyFunc(...)//CPU函数

main(){
    cudaMalloc(&d_GlblVarPtr,bytes);//在设备上分配空间
    
    cudaMemcpy(d_GlblVarPtr,h_Gl...);//从主机端传输数据到设备端
    
    kernelOne<<<arg1,arg2>>>(arg...);//arg1：grid参数；arg2：block参数
    
    cudaMemcpy(h_GlblVarPtr...);//从设备端传输数据到主机端
}

__global__void kernelOne(type args,...){//内核函数
	//局部变量声明
    __local__,__shared__//自动变量被默认分配到register或local memory
}

CUDA软件体系

CUDA的核心是CUDA C语言，它包含对C语言的最小扩展集和一个运行时库使用这些扩展和运行时库的源文件必须通过nvcc编译器进行编译。

在一个程序中只能使用CUDA运行时API与CUDA驱动API中的一种，不能混合使用。

nvcc编译器

nvcc编译器根据配置编译CUDA C代码，可以生成三中不同的输出：PTX、CUDA二进制序列和标准C。

PTX (Parallel Thread eXecution) 类假于汇编语言，是为动态编译器 JIT ( JIT包含在标准的 NVIDIA驱功中)设计的输入指令序列。

内核函数可以通过PTX编写，但通常还是通过CUDA C一类的高级语言进行编写。

1
2
3
4
5
6
7
8
9
10
11
graph TD
   subgraph virtual
   A[C/C++ CUDA Application] --> B[NVCC]
   B[NVCC] --> C[PTX Code]
   end
   C[PTX Code] --> D[PTX to Target Complier]
   subgraph physcical
   D[PTX to Target Complier] --> E1[G80]
   D[PTX to Target Complier] --> E2[...]
   D[PTX to Target Complier] --> E3[GPU]
   end

CUDA函数库

目前CUDA中有CUFFT、CUBLAS和CUDPP三个函数库，提供了简单高效的常用函数。

CUFFT库是一个利用GPU进行傅里叶变换的函数库。

CUBLAS库是一个基本的矩阵与向量的运算库。

CUDPP库提供了很多基本的常用的并行操作函数，如排序、搜索等。

线程同步

void __syncthreads()使得block内所有线程同步。

只有当所有线程都达到同步语句处，后续指令才能继续执行。

用于避免访问共享或全局内存时出现的RAW/WAR/WAW冲突。

在最理想的情况下，调用一次__syncthreads()需要至少四个时钟周期，但

一般调用都需要更多的时钟周期，因此尽量避免或节约使用__syncthreads()。

kernel间通信

不同block之间数据相对独立，只能通过Global Memory实现。

  • GPU内部的block之间都缺乏有效的通信手段，多个GPU间的低延迟通信就更加无从谈起。

  • 由于每个设备拥有各自独立的device端存储器，因而GPU间的通信需要通过host端内存进行。

  • 在CUDA2.2中加入的mapped memory允许多个设备从内核程序中直接访问同一块pinned memory，可以提供在某些应用中的效率。
    pinned固定内存，不会被替换，提高效率。

GPU与CPU线程同步

在CUDA主机端代码中使用cudaThreadSynchronize()，可以实现GPU与CPU线程的同步。

kernel启动后控制权将异步返回，利用该函数可以确定所有设备端线程均已结束。

类似函数：cudaStreamSynchronize()和cudaEventSynchronize()，它们阻塞所有的流或事件，直到此前的所有CUDA调用均已完成。

文章作者: HibisciDai
文章链接: http://hibiscidai.com/2020/05/07/CUDA基础/
版权声明: 本博客所有文章除特别声明外，均采用 CC BY-NC-SA 4.0 许可协议。转载请注明来自 HibisciDai！
CUDA高性能计算并行计算
好用、实惠、稳定的梯子,点击这里