CUDA
参考资料
官方资料:
其他
英伟达各型号 GPU 的关键参数: https://en.wikipedia.org/wiki/List_of_Nvidia_graphics_processing_units
一些未仔细核验的资料:
基本概念
术语方面, GPU 硬件与 CUDA 软件术语有些使用了同样的词。
一个 GPU 由如下部分组成 (参考 https://medium.com/analytics-vidhya/cuda-memory-model-823f02cef0bf)
多个 SM (Streaming Multiprocessor, 有时候也被简称为 Multiprocessor)
多个 CUDA Cores
Shared Memory
L1 Cache
Read-Only DataCache
L2 Cache
DRAM (GPU主存)
机器信息:
GeForce MX250: Windows + WSL2
Tesla V100-PCIE-16GB: 虚拟化 linux 云主机
GeForce GTX 1650: Windows + WSL2
使用 cuda-samples 可以查看 GPU 的设备信心
git clone https://github.com/NVIDIA/cuda-samples.git
cd cuda-samples
git checkout v11.2 # 切换到 nvcc -V 的版本
cd /Samples/deviceQuery
make
./deviceQuery输出 (重新 format 了一下, 便于对比):
NVIDIA GeForce MX250 Tesla V100-PCIE-16GB NVIDIA GeForce GTX 1650
CUDA Driver Version / Runtime Version 12.2 / 11.2 12.2 / 11.5 12.2 / 11.2
CUDA Capability Major/Minor version number: 6.1 7.0 7.5
Total amount of global memory: 2048 MBytes (2147352576 bytes) 16151 MBytes (16935419905 bytes) 4096 MBytes (4294639616 bytes)
( 3) Multiprocessors, (128) CUDA Cores/MP: 384 CUDA Cores (080) Multiprocessors, (064) CUDA Cores/MP: 5120 CUDA Cores (14) Multiprocessors, (64) CUDA Cores/MP: 896 CUDA Cores
GPU Max Clock rate: 1038 MHz (1.04 GHz) 1380 MHz (1.38 GHz) 1710 MHz (1.71 GHz)
Memory Clock rate: 3004 Mhz 877 MHz 6001 Mhz
Memory Bus Width: 64-bit 4096-bit 128-bit
L2 Cache Size: 524288 bytes 6291456 bytes 1048576 bytes
Maximum Texture Dimension Size (x,y,z) 1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384) 1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384) 1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
Maximum Layered 1D Texture Size, (num) layers 1D=(32768), 2048 layers 1D=(32768), 2048 layers 1D=(32768), 2048 layers
Maximum Layered 2D Texture Size, (num) layers 2D=(32768, 32768), 2048 layers 2D=(32768, 32768), 2048 layers 2D=(32768, 32768), 2048 layers
Total amount of constant memory: 65536 bytes 65536 bytes 65536 bytes
Total amount of shared memory per block: 49152 bytes 49152 bytes 49152 bytes
Total shared memory per multiprocessor: 98304 bytes 98304 bytes 65536 bytes
Total number of registers available per block: 65536 65536 65536
Warp size: 32 32 32
Maximum number of threads per multiprocessor: 2048 2048 1024
Maximum number of threads per block: 1024 1024 1024
Max dimension size of a thread block (x,y,z): (1024, 1024, 64) (1024, 1024, 64) (1024, 1024, 64)
Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535) (2147483647, 65535, 65535) (2147483647, 65535, 65535)
Maximum memory pitch: 2147483647 bytes 2147483647 bytes 2147483647 bytes
Texture alignment: 512 bytes 512 bytes 512 bytes
Concurrent copy and kernel execution: Yes with 1 copy engine(s) Yes with 7 copy engine(s) Yes with 6 copy engine(s)
Run time limit on kernels: Yes No Yes
Integrated GPU sharing Host Memory: No No No
Support host page-locked memory mapping: Yes Yes Yes
Alignment requirement for Surfaces: Yes Yes Yes
Device has ECC support: Disabled Enabled Disabled
Device supports Unified Addressing (UVA): Yes Yes Yes
Device supports Managed Memory: Yes Yes Yes
Device supports Compute Preemption: Yes Yes Yes
Supports Cooperative Kernel Launch: Yes Yes Yes
Supports MultiDevice Co-op Kernel Launch: No Yes No
Device PCI Domain ID / Bus ID / location ID: 0 / 1 / 0 0 / 177 / 0 0 / 1 / 0
Compute Mode:
< Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >这里我们先关注这几个值 (以 V100 为例):
// int dev = 0;
// cudaSetDevice(dev);
// cudaDeviceProp deviceProp;
// cudaGetDeviceProperties(&deviceProp, dev);
(080) Multiprocessors, (064) CUDA Cores/MP: 5120 CUDA Cores // deviceProp.multiProcessorCount, _ConvertSMVer2Cores(deviceProp.major, deviceProp.minor)
Maximum number of threads per multiprocessor: 2048 // deviceProp.maxThreadsPerMultiProcessor
Maximum number of threads per block: 1024 // deviceProp.maxThreadsPerBlock
Max dimension size of a thread block (x,y,z): (1024, 1024, 64) // deviceProp.maxThreadsDim[0], deviceProp.maxThreadsDim[1], deviceProp.maxThreadsDim[2]
Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535) // deviceProp.maxGridSize[0], deviceProp.maxGridSize[1], deviceProp.maxGridSize[2]deviceProp.maxThreadsPerBlock=1024: CUDA 编程模型视角里 block 的三个维度之积不能超过 1024deviceProp.maxThreadsPerMultiProcessor=2048: 一个 SM 可以同时执行的最大硬件线程数为 2048, 注意从 CUDA 编程视角来看, 一个 block 里的所有线程都会被运行在同一个 SM 上 (注意可能不会是并发执行的), 而最终在运行 thread 时, thread 会映射到硬件线程上, 例如 block 的三个维度之积为 68 (这不是最佳实践, 假设硬件的 warp 的大小为 32), 那么这 68 个软件意义上的线程会被分为 32+32+4 三组, 其中每组内的软件线程总是会映射到同一个硬件线程束 (warp) 上并发执行的, 但这三组 warp 有可能不是并发执行的, 但无论如何, 这 68 个线程一定都在同一个 SM 上执行.deviceProp.multiProcessorCount=80: GPU 包含 80 个 SM_ConvertSMVer2Cores(deviceProp.major, deviceProp.minor)=64: 每个 SM 包含 64 个 CUDA CoredeviceProp.maxThreadsDim=(1024, 1024, 64). 这代表了 CUDA 编程模型视角里每个 block 的软件线程数在三个维度上的最大值, 注意还需满足三个维度之积不超过deviceProp.maxThreadsPerBlock=1024deviceProp.maxGridSize=(2147483647, 65535, 65535). 这代表了 CUDA 编程模型视角里每个 grid 的 block 数在三个维度上的最大值, 也就是说一个核函数最多只能由2147483647*65535*65535*1024个线程来完成整个任务.
关于 warp: warp 是硬件层的概念, 一个 warp 里的 32 个物理线程严格并发执行, 且并发执行的指令一模一样, 当然, 操作数可以是不一样的 (这种模式也被称作 SIMD, 即 Single Instruction, Multiple Data).
关于 block: block 是纯粹的软件视角的概念
warp 与 CUDA Core 的关系: 这篇 博客 里有个误解是 1 个 CUDA Core 就对应 1 个 warp, 但根据这个问答:
Now your Card has a total Number of 384 cores on 2 SMs with 192 cores each. The CUDA core count represents the total number of single precision floating point or integer thread instructions that can be executed per cycle. Do not consider CUDA cores in any calculation.
我们在使用这种方式调用核函数时:
int a, b, c; // 每个 block 的 thread 数为 a*b*c
int x, y, z; // 调用整个核函数完成整个功能需要 x*y*z 个 block
dim3 threads(a, b, c);
dim3 grid(x, y, z);
kernel_fun <<<grid, threads>>> ();理论上说:
// (a, b, c) <= (1024, 1024, 64), a*b*c <= 1024
// (x, y, z) <= (2147483647, 65535, 65535)
a=64, b=64, c=2 // 非法, a*b*c 超过了 deviceProp.maxThreadsPerBlock 的限制
a=1024, b=1, c=1, x=160, y=1, z=1 // 理论上可以利用完整个 GPU 的并发: 每个 SM 同时运行 2 个 block
a=32, b=4, c=4, x=320, y=1, z=1 // 理论上可以利用完整个 GPU 的并发: 每个 SM 同时运行 4 个 block
a=68, b=1, c=1, x=320, y=320, z=320 // a*b*c=68 不是 32 的倍数, 强烈不推荐
a=1024, b=1, c=1, x=320, y=1, z=1 // 理论上可以利用完整个 GPU 的并发, 只是需要 2 次完全并发才能完成整个任务threadIdx, blockDim, blockIdx, gridDim
https://erangad.medium.com/1d-2d-and-3d-thread-allocation-for-loops-in-cuda-e0f908537a52
__global__ void kernel_fun(){
threadIdx.x < blockDim.x;
threadIdx.y < blockDim.y;
threadIdx.z < blockDim.z;
(blockDim.x == 2) && (blockDim.y==16) && (blockDim.z==4)
blockIdx.x < gridDim.x;
blockIdx.y < gridDim.y;
blockIdx.z < gridDim.z;
(gridDim.x == 2) && (gridDim.y==3) && (gridDim.z==6)
}
int a=2, b=16, c=4;
int x=2, y=3, z=6;
dim3 block_dim(a, b, c);
dim3 thread_dim(x, y, z);
kernel_fun <<<grid, threads>>> ();注意: threadIdx.x 和 blockIdx.x 是变化最快的维度: 上面的例子里: threadIdx.x=0, threadIdx.y=1, threadIdx.z=2 的下一个 thread 是 threadIdx.x=1, threadIdx.y=1, threadIdx.z=2, 即:
// Thread 0 - Thread 31 会组成一个 warp.
Thread 0: (threadIdx.x, threadIdx.y, threadIdx.z) = (0, 0, 0)
Thread 1: (threadIdx.x, threadIdx.y, threadIdx.z) = (1, 0, 0)
Thread 2: (threadIdx.x, threadIdx.y, threadIdx.z) = (2, 0, 0)
Thread 3: (threadIdx.x, threadIdx.y, threadIdx.z) = (3, 0, 0)
Thread 4: (threadIdx.x, threadIdx.y, threadIdx.z) = (0, 1, 0)
Thread 5: (threadIdx.x, threadIdx.y, threadIdx.z) = (1, 1, 0)执行逻辑
参考: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capability-5-x
假设一次调用包含 8 个 block, 每个 block 中有 512 个 thread, 即使用这种方式进行调用, kernel_fun <<<8, 512>>> (p); 假设 warp 大小为 32, 那么每个 block 的前 32 个 thread 将会最终映射到一个 warp 上, 后 32 个也会映射到一个 warp 上 (warp 不是硬件概念, 而是意味着前 32 个 thread 在同一时钟周期内会执行相同的指令). 假设 GPU 有 80 个 SM, 而每个 SM 包含 4 个 warp scheduler, 那么当一个 block 被调度到一个 SM 上后 (block 一旦被调度到 SM, 那么一定会将其完成, 不会被切换到别的 SM 上), SM 会进一步分配给 4 个 warp scheduler 来处理, 在前面的例子里, 一个 block 被分为 16 组, 一个可能的情况是:
0号调度器：负责执行 0，4，7，9，10，14 组
1号调度器：负责执行 1，5，8，11 组
2号调度器：负责执行 2，3，6 组
3号调度器：负责执行 12，13，15 组而以 2 号调度器为例, 假设之前描述的每个 thread 只包含 2 个指令, 具体的执行顺序可能是:
block3-warp2-instruct0
block3-warp6-instruct0
block3-warp2-instruct1
block3-warp3-instruct0
block3-warp3-instruct1
block3-warp6-instruct1注意, 在 warp2, warp3, warp6 在执行过程中, 有可能会插入别的 block 的执行, 但是 warp scheduler 的同一时间, 只能执行一个 warp. 也就是说在这个例子里, 一个 SM 的最大并发量是 4 * 32 = 128 个线程.
FAQ & 杂录
compute-capability 与 cuda-architecture 是同一个意思: 问答
最大单精度浮点数计算次数: MX250: 797.2 GFLOPS, V100: 14028 GFLOPS, GTX1650: 2984 GFLOPS
矩阵乘法
benchmark for (1024 x 1024) x (1024 x 1024)
cuda-sample/matrixMul
MX250: 140 GFLOPS / 797.2 GFLOPS
V100: 3800 GFLOPS / 14028 GFLOPS
GTX1650: 426 GFLOPS / 2984 GFLOPS
Last updated
Was this helpful?