【参加CUDA线上训练营】第1天课程知识梳理和实验记录（上午知识梳理）

Brain Compiler

已于 2023-02-07 21:49:57 修改

阅读量248

点赞数

文章标签： ubuntu linux

于 2023-02-06 22:15:39 首次发布

本文链接：https://blog.csdn.net/qq_41663081/article/details/128909169

版权

课程说明

在这里插入图片描述

实验平台

Jetson Nano
Jetson Nano’s Compute Capability: 5.3（所以架构是SM53）
CUDA GPUs - Compute Capability | NVIDIA Developer

在这里插入图片描述

优点

统一存储，减少了耗时的PCIE传输数据操作
对于个人用户来说，一个低成本的小jetson设备可能是比较适合入门的。价格便宜，能学习GPU开发，还能学习Linux上的CPU开发，和熟悉ARM CPU等等。

下午实验前置知识

Linux

比较简单，所以不细致记录

Ubuntu的文件管理：目录结构

在这里插入图片描述

这里的d就是指这是一个文件夹，如果不是d是-，则说明是一个文件而非文件夹。

权限控制

对于权限修改，有字母法和数字法
r (read) -> 4
w (write) -> 2
x (excute) -> 1

字母法个人感觉比较繁琐，形如：
chmod (u g o a) (+ - =) (r w x) (文件名)
例如：
chmod ugo+r file.txt
chmod ug+w, o-w file1.txt file2.txt

其他收获

内容比较多的文本文件建议用more来代替cat看，好处是分页，滚动查看（要不然行会很多，占用屏幕空间）

device, global和host修饰符

__device__ float Func()执行位置device，调用位置device

__global__ void KernelFunc()执行位置是device，调用位置是host或device(arch>3.2)。

__global__定义一个kernel函数，CPU上调用，GPU上执行。
必须返回的是void类型
2.x: Fermi，3.0/3.2: Kepler (3.2老式嵌入式的TK1）

__host__ float HostFunc()执行位置是host，调用位置是host。

问：device和host可以同时修饰一个函数吗？可以！同时有__device__和__host__等于分别生成了两份代码，一份在GPU上，一份在CPU上。方便你将一些通用的小例程，同时得到CPU+GPU的版本。
另外，host修饰的函数不能直接调用device修饰的函数。

变量一般的起名规则

d_xxx 的名字，常用来表示GPU上的变量(d = device, GPU)
h_xxx 的名字，常用来表来CPU上的变量(h = host, CPU)

编译

Makefile的用法

就是实现自动化编译，把编译命令都写在一个Makefile文件里，编译的时候只需要运行make（或make -f makefile的文件名）
make时文件路径下如果有一个文件叫clean怎么办：在makefile里写一行.PHONY: clean，声明clean是个操作。./PHONY表示声明，有声明的会按照声明执行，没有声明的会寻找makefile同级的文件。这里.PHONY 意思表示 clean 是一个“伪目标”，在rm命令前面加了一个小减号的意思就是，也许某些文件出现问题，但不要管，继续做后面的事。注意 clean 的规则不要放在文件的开头，不然，这就会变成 make 的默认目标，一般我们都会放在文件结尾处。

关于Makefile高阶教程，参考Makefile详细教程_扫地的小何尚的博客-CSDN博客

单文件编译

.cu的编译器是nvcc
/usr/local/cuda/bin/nvcc -arch=compute_53 -code=sm_53 hello_cuda.cu -o hello_cuda -run
code需要大于等于arch。不同的架构，可以理解成不同代的卡（或者嵌入式的带有GPU的Jetson设备）。通过这里的arch和code设定，可以告诉编译器，在哪种显卡上生成代码。我们实验课采用的Jetson Nano算力是5.3，所以架构是SM53。

多文件编译

主文件main.c，引用了add.h，sub.h，div.h，mul.h

OBJ = main.o add.o sub.o mul.o div.o
CC = gcc
app: $(OBJ)
	$(CC) -o app $(OBJ)
main.o: main.c
	$(CC) -c main.c
add.o: add.c
	$(CC) -c add.c
sub.o:sub.c
	$(CC) -c sub.c
mul.o: mul.c
	$(CC) -c mul.c
div.o: div.c
	$(CC) -c div.c
.PHONY : clean //防止有个文件叫clean
clean :
	-rm $(OBJ) app

头文件

（1）.h里放声明
（2）.cu里面放实现代码
（3）用nvcc进行编译和链接操作，得到可执行文件（注意命令行）。

.cuh和.h等头文件
回答：头文件扩展名无所谓。可以用.h, .hpp, .cuh, 或者.txt也可以。（因为头文件是被#include 后，在包含它的那个文件中.c或者.cpp或者.cu中参与编译的, 最终的文件类型看宿主文件。所以它自身无所谓）

CUDA线程

Thread

线程层次

threadIdx为三维向量
一个thread block由大量thread组成
二维block:
Block size: (Dx,Dy), thread(x,y) 的ID=x+yDx
三维block:
Block size: (Dx,Dy,Dz), thread(x,y,z) 的
ID=x+yDx+zDxDy

Block

Block也叫线程协作组（CTA）
线程以block为单位分配到GPU处理核心（SM）中
SM内部计算和存储资源有限
在目前的NVIDIA GPU中，单程序多数据一个block中最多包含1024个thread
同一kernel中的所有block大小都一致
一个kernel可以被大量block执行
总线程数=block数量*一个block中的线程数量
通常block数量大于GPU SM数量
Block组成一维、二维、三位的grid

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	int j = blockIdx.y * blockDim.y + threadIdx.y;
	if (i < N && j < N){
		C[i][j] = A[i][j] + B[i][j];
	}
	
int main()
{
	...
	// Kernel invocation
	dim3 threadsPerBlock(16, 16);
	dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
	MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
	...
}

在硬件中，单程序多数据每个block独立执行
不同block可以并行或串行执行
不同block可以按照任意顺序被调度和执行
保证了程序的可扩展性（scalability）

同一个block中的thread可以通过shared memory共享数据
- Shared memory是由程序员控制的高速访问存储空间，单程序多数据访问速度类似L1 cache
同一个block中的thread可以进行同步
- 通过在kernel中调用__syncthreads()函数设置同步点
- 同步表示必须等待所有thread均到达该路障后才可以继续执行之后的代码

threadIdx.x 是执行当前kernel函数的线程在block中的x方向的序号  

blockIdx.x 是执行当前kernel函数的线程所在block，在grid中的x方向的序号

threadIdx整体是3个uint分量的向量类型（有threadIdx.x, threadIdx.y, threadIdx.z三个分量），类型是unsigned int, 或者说，uint32_t，32位的无符号的整数。

dim3类型等于是一个特殊的uint3(没有用到的元素是1，而不是0）。

Grid

一个Grid里有多个Block。
1个block里面能容纳的线程数量较少，x/y/z方向上安排的线程数量分别有限制（1024个，1024个，和64个），但是总大小（它们的乘积，例如我32行，每行32个线程）有限制，最多不能超过1024个线程（三个方向的乘积），不过grid里的blocks数量就比较大了。

我们谈论warp的时候，AMD上叫Wavefront，大小有两种，64（老A卡）和32（新A卡）。