cuda程序demo
#include <stdio.h> // C头文件
__global__ void hello_from_gpu() // __global__ 关键词表示gpu核函数
{
printf("Hello World from the the GPU\n");
}
int main(void)
{
hello_from_gpu<<<1, 1>>>(); // 1*1 个线程
cudaDeviceSynchronize(); // 同步代码，cpu、gpu协同
return 0;
}
编译命令，nvcc等价于g++
nvcc test.cu -o test
核函数
kernel函数GPU进行并行执行(自发并行执行，硬件支持) 核函数与C++函数区别：
（1）限定词 global
（2）返回值void
注意事项：
(1)核函数只能访问GPU内存（即设备显存）
(2)核函数不支持变长参数
(3)核函数不支持静态变量
(4)核函数不能使用函数指针
(5)核函数具有异步性(解释：CPU与GPU异构，CPU发起核函数调用，必须显式调用同步函数)
CUDA线程模型
-
01 线程模型结构 grid 网格(核函数启动产生的所有线程)、block线程块（一组线程）。线程是GPU编程最小单位。 配置线程 <<<grid_size, block_size>>>
-
02 线程组织管理 一维线程模型
线程在核函数中有一个唯一标识。该标识由参数<<<grid_size, block_size>>>确定。对于一维线程模型：gridDim.x等于grid_size，blockDim.x等于block_size。
线程索引(一维) blockIdx.x 指定一个线程在网格中线程块的索引值； threadIdx.x 指定一个线程在线程块中的线程索引值； 对于一维线程模型，线程唯一标识：
Idx = threadIdx.x + blockIdx.x * blockDim.x
线程索引(多维) dim3类型变量：gridDim、blockDim
定义网格、线程块：
dim3 gridDim(3,2,1) dim3 blockDim(3,2,1)
线程全局索引计算
grid 可以是一维、二维、三维，block可以是一维、二维、三维，共有9种组合。
一维网格 一维线程块
dim3 gridDim(4)； dim3 blockDim(8)；
计算方式： int id = blockIdx.x * blockDim.x + threadIdx.x;
二维网格 二维线程块
dim3 gridDim(2, 2)； dim3 blockDim(4, 4)；
计算方式：
int bid = blockId.x + blockId.y * gridDim.x; // 计算线程块索引 int tid = threadIdx.y * blockDim.x + threadIdx.x; // 计算线程索引 int id = bid * (blockDim.x * blockDim.y) + tid; // 线程块索引+线程索引 -> 线程在grid中唯一索引
三维网格 三维线程块
dim3 gridDim(2, 2, 2)； dim3 blockDim(4, 4, 2)；
int bid = blockId.x + blockId.y * gridDim.x + blockIdx.z * gridDim.x * gridDim.y int tid = threadIdx.z * blockDim.x * blockDim.y + threadIdx.z * blockDim.x + threadIdx.x int id = bid * (blockDim.x * blockDim.y * blockDim.z) + tid
GPU架构与计算能力
计算能力：标识计算能力版本号；X.Y
CUDA12.2 文档：
| 主计算能力 | 架构 |
|---|---|
| sm_80,sm_86,sm_87 | Ampere GPU arch |
| sm_90,sm_90a | Hopper support |
注意：
GPU架构Tesla 与 型号Tesla表示的含义不同，型号Tesla设备的架构可以是MaxWell、Tesla，这里容易迷惑。
GPU性能：计算能力与性能没有关系，性能（浮点运算峰值、显存容量、带宽）
| GPU型号 | 计算能力 | 显存容量 | 带宽 | 浮点运算峰值 |
|---|---|---|---|---|
| Tesla V100 | 7.0 | 32GB | 900GB/s | 7(14) TFLOPS |
| GeForce RTX 2070 | 7.5 | 8GB | 448GB/s | 0.2(6.5) TFLOPS |
CUDA矩阵加法程序
CUDA程序基本框架
#include <头文件>
__global__ void 函数名（参数...）
{
核函数实现
}
int main(void)
{
设置GPU设备
分配Host、Device内存
初始化Host中数据
数据从Host复制到Device
调用Kernel，在Device中执行计算
计算结果数据从Device传给Host
释放Host、Device内存
}
获取Device数量
int iDeviceCount = 0; cudaGetDeviceCount(&iDeviceCount);
设置Device
int iDev = 0; cudaSetDevice(iDev);
__host____device__cudaError_t cudaGetDeviceCount ( int* count ) // 在host、device都可执行
__host__cudaError_t cudaSetDevice ( int device ) // 只能在host执行
内存管理 内存分配、数据拷贝、内存初始化、内存释放
| 标准C内存管理函数 | CUDA内存管理函数 |
|---|---|
| malloc | cudaMalloc |
| memcpy | cudaMemcpy |
| memset | cudaMemset |
| free | cudaFree |
// 内存分配
float *fpHost_A;
cudaMalloc((float **)&fpHost_A, nBytes); // 传入双重指针
// copy
cudaMemcpy(Device_A, Host_A, nBytes, cudaMemcpyHostToDevice )
// init
cudaMemset(fpDevice_A, 0, nBytes);
// free
cudaFree(pDevice_A)
device、kernel、host function
- device function(设备函数)
只能在GPU设备上执行，只能被核函数或其他设备函数调用，__device__修饰
- kernel function(设备函数)
由主机调用，在GPU设备执行；global__修饰，不能使用__device、__host__修饰符
- host function(设备函数)
主机端C++函数__host__可省略，可使用__device__、__host__同时修饰减少冗余，编译器分别编译函数用于主机、设备
事件event计时
cuda 事件可为Host代码、Device代码计时。
cudaEvent start, stop;
cudaEventCreate(&start, __FILE__, __LINE__);
cudaEventCreate(&stop, __FILE__, __LINE__);
cudaEventQuery(start);
多维数组线程模型
数据存储方式：先存储第一行，再第二行...C/C++ 以行为主，Matlab是以列为主。
二维网格、二维线程块
线程与二维矩阵映射关系:
ix 表示全局内存中x维度索引值，iy表示y维度索引值
ix = blockIdx.x * blockDim.x + threadIdx.x; iy = blockIdx.y * blockDim.y + threadIdx.y;
idx = iy * nx + ix //线程全局索引
二维网格、一维线程块
ix = blockIdx.x * blockDim.x + threadIdx.x; iy = threadIdx.y;
idx = iy * nx + ix //线程全局索引
