企业官网建设流程全解析

B站：2.2 CUDA线程模型_哔哩哔哩_bilibili

3.3. The CUDA Driver API — CUDA Programming Guide

CUDA Runtime API :: CUDA Toolkit Documentation

CUDA编程模型

线程模型

不同block数据不共享

内存架构

nvcc编译流程

https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#nvcc-command-options

内核函数定义

HelloWorld

HelloWorldKernel.cu

#include <stdio.h> __global__ void helloFromGPU() { printf("Hello World From GPU\n"); } int main(int argc, char **argv) { printf("Hello World from CPU\n"); helloFromGPU<<<1, 10>>>(); cudaDeviceReset(); return 0; }

(py310) root@a658a13300af:/workspace# nvcc HelloWorldKernel.cu --output-file HelloWorldKernel (py310) root@a658a13300af:/workspace# ls -alh total 490M drwxr-xr-x 1 root root 4.0K Jul 25 14:59 . drwxr-xr-x 1 root root 4.0K Jul 25 12:57 .. drwxr-xr-x 2 root root 4.0K Jul 25 14:56 .ipynb_checkpoints -rwxr-xr-x 1 root root 975K Jul 25 14:59 HelloWorldKernel -rw-r--r-- 1 root root 234 Jul 25 14:59 HelloWorldKernel.cu -rw-r--r-- 1 root root 245M Jul 25 13:15 flash_attn-2.8.3.post1+cu12torch2.4cxx11abiFALSE-cp310-cp310-linux_x86_64.whl -rw-r--r-- 1 root root 245M Jul 25 13:08 flash_attn-2.8.3.post1+cu12torch2.4cxx11abiTRUE-cp310-cp310-linux_x86_64.whl drwxr-xr-x 5 root root 4.0K Jul 25 13:25 nano-vllm -rw-r--r-- 1 root root 411K Jul 25 13:01 nano-vllm-main.zip (py310) root@a658a13300af:/workspace# ./HelloWorldKernel Hello World from CPU Hello World From GPU Hello World From GPU Hello World From GPU Hello World From GPU Hello World From GPU Hello World From GPU Hello World From GPU Hello World From GPU Hello World From GPU Hello World From GPU (py310) root@a658a13300af:/workspace#

GetThreadIdxKernel.cu

#include <stdio.h> __global__ void helloFromGPU() { printf("blockDim:x=%d, y=%d, z=%d, gridDim:x=%d, y=%d, z=%d Current ThreadIdx: x=%d, y=%d, z=%d\n", blockDim.x, blockDim.y, blockDim.z, gridDim.x, gridDim.y, gridDim.z, threadIdx.x, threadIdx.y, threadIdx.z); } int main(int argc, char **argv) { printf("Hello World from CPU\n"); dim3 grid; grid.x = 2; grid.y = 2; dim3 block; block.x = 2; block.y = 2; helloFromGPU<<<grid, block>>>(); cudaDeviceReset(); return 0; }

(py310) root@a658a13300af:/workspace# nvcc GetThreadIdxKernel.cu -o GetThreadIdxKernel (py310) root@a658a13300af:/workspace# ./GetThreadIdxKernel Hello World from CPU blockDim:x=2, y=2, z=1, gridDim:x=2, y=2, z=1 Current ThreadIdx: x=0, y=0, z=0 blockDim:x=2, y=2, z=1, gridDim:x=2, y=2, z=1 Current ThreadIdx: x=1, y=0, z=0 blockDim:x=2, y=2, z=1, gridDim:x=2, y=2, z=1 Current ThreadIdx: x=0, y=1, z=0 blockDim:x=2, y=2, z=1, gridDim:x=2, y=2, z=1 Current ThreadIdx: x=1, y=1, z=0 blockDim:x=2, y=2, z=1, gridDim:x=2, y=2, z=1 Current ThreadIdx: x=0, y=0, z=0 blockDim:x=2, y=2, z=1, gridDim:x=2, y=2, z=1 Current ThreadIdx: x=1, y=0, z=0 blockDim:x=2, y=2, z=1, gridDim:x=2, y=2, z=1 Current ThreadIdx: x=0, y=1, z=0 blockDim:x=2, y=2, z=1, gridDim:x=2, y=2, z=1 Current ThreadIdx: x=1, y=1, z=0 blockDim:x=2, y=2, z=1, gridDim:x=2, y=2, z=1 Current ThreadIdx: x=0, y=0, z=0 blockDim:x=2, y=2, z=1, gridDim:x=2, y=2, z=1 Current ThreadIdx: x=1, y=0, z=0 blockDim:x=2, y=2, z=1, gridDim:x=2, y=2, z=1 Current ThreadIdx: x=0, y=1, z=0 blockDim:x=2, y=2, z=1, gridDim:x=2, y=2, z=1 Current ThreadIdx: x=1, y=1, z=0 blockDim:x=2, y=2, z=1, gridDim:x=2, y=2, z=1 Current ThreadIdx: x=0, y=0, z=0 blockDim:x=2, y=2, z=1, gridDim:x=2, y=2, z=1 Current ThreadIdx: x=1, y=0, z=0 blockDim:x=2, y=2, z=1, gridDim:x=2, y=2, z=1 Current ThreadIdx: x=0, y=1, z=0 blockDim:x=2, y=2, z=1, gridDim:x=2, y=2, z=1 Current ThreadIdx: x=1, y=1, z=0

错误代码

common.h

#include <sys/time.h> #include <cuda_runtime.h> #include <stdio.h> cudaError_t ErrorCheck(cudaError_t status, const char *filename, int lineNumber) { if(status != cudaSuccess) { printf("CUDA API error:\r\ncode=%d, name=%s, description=%s\r\nfile=%s, line=%d\r\n", status, cudaGetErrorName(status), cudaGetErrorString(status), filename, lineNumber); return status; } return status; } #define ERROR_CHECK(status) \ ErrorCheck(status, __FILE__, __LINE__);

ErrorHandlerKernel.cu

#include "common.h" #include <stdio.h> int main(int argc, char *argv[]) { float *gpuMemory = NULL; ERROR_CHECK(cudaMalloc(&gpuMemory, sizeof(float))); ERROR_CHECK(cudaFree(gpuMemory)); ERROR_CHECK(cudaFree(gpuMemory)); ERROR_CHECK(cudaDeviceReset()); }

(py310) root@a658a13300af:/workspace# nvcc ErrorHandlerKernel.cu -o ErrorHandlerKernel (py310) root@a658a13300af:/workspace# ./ErrorHandlerKernel CUDA API error: code=1, name=cudaErrorInvalidValue, description=invalid argument file=ErrorHandlerKernel.cu, line=8

API获取GPU信息

https://docs.nvidia.com/cuda/cuda-runtime-api/structcudaDeviceProp.html#structcudaDeviceProp

GetRuntimeGPUInfoKernel.cu

#include "common.h" #include <cuda_runtime.h> #include <stdio.h> int main(int argc, char **argv) { printf("%s Starting...\n", argv[0]); int deviceCount = 0; cudaGetDeviceCount(&deviceCount); if(deviceCount == 0) { printf("There are no available device(s) that support CUDA\n"); return 1; } else { printf("Detected %d CUDA Capable device(s)\n", deviceCount); } int dev = 0, driverVersion = 0, runtimeVersion = 0; cudaSetDevice(dev); cudaDeviceProp deviceProp; cudaGetDeviceProperties(&deviceProp, dev); printf("Device %d: \"%s\"\n", dev, deviceProp.name); cudaDriverGetVersion(&driverVersion); cudaRuntimeGetVersion(&runtimeVersion); printf(" CUDA Driver Version / Runtime Version %d.%d / %d.%d\n", driverVersion / 1000, (driverVersion % 100)/10, runtimeVersion / 1000, (runtimeVersion % 100) / 10); printf(" CUDA Capability Major/Minor version number: %d.%d\n", deviceProp.major, deviceProp.minor); printf(" Total amount of global memory: %.2f GBytes (%llu bytes)\n", (float)deviceProp.totalGlobalMem / pow(1024.0,3), (unsigned long long)deviceProp.totalGlobalMem); printf(" GPU Clock rate: %.0f MHZ (%0.2f GHz)\n", deviceProp.clockRate * 1e-3f, deviceProp.clockRate * 1e-6f); printf(" Memory Clock rate: %.0f MHz\n", deviceProp.memoryClockRate *1e-3f); printf(" Memory Bus Width: %d-bit\n", deviceProp.memoryBusWidth); if(deviceProp.l2CacheSize) { printf(" L2 Cache Size: %d bytes\n", deviceProp.l2CacheSize); } printf(" Max Texture Dimension Size (x,y,z) 1D=(%d), 2D=(%d,%d), 3D=(%d,%d,%d)\n", deviceProp.maxTexture1D, deviceProp.maxTexture2D[0], deviceProp.maxTexture2D[1], deviceProp.maxTexture3D[0], deviceProp.maxTexture3D[1], deviceProp.maxTexture3D[2]); printf(" Total amount of constant memory: %lu bytes\n", deviceProp.totalConstMem); printf(" Total amount of shared memory per block: %lu bytes\n", deviceProp.sharedMemPerBlock); printf(" Total number of registers available per block:%d\n", deviceProp.regsPerBlock); printf(" Warp size: %d\n", deviceProp.warpSize); printf(" Maximum number of MultiProcessor: %d\n", deviceProp.multiProcessorCount); printf(" Maximum number of threads per multiprocessor: %d\n", deviceProp.maxThreadsPerMultiProcessor); printf(" Maximum number of threads per block: %d\n", deviceProp.maxThreadsPerBlock); printf(" Maximum sizes of each dimension of a block: %d x %d x %d\n", deviceProp.maxThreadsDim[0], deviceProp.maxThreadsDim[1], deviceProp.maxThreadsDim[2]); printf(" Maximum sizes of each dimension of a grid: %d x %d x %d\n", deviceProp.maxGridSize[0], deviceProp.maxGridSize[1], deviceProp.maxGridSize[2]); return 0; }

(py310) root@a658a13300af:/workspace# nvcc GetRuntimeGPUInfoKernel.cu -o GetRuntimeGPUInfoKernel (py310) root@a658a13300af:/workspace# ./GetRuntimeGPUInfoKernel ./GetRuntimeGPUInfoKernel Starting... Detected 1 CUDA Capable device(s) Device 0: "NVIDIA GeForce RTX 3080 Ti" CUDA Driver Version / Runtime Version 13.2 / 12.4 CUDA Capability Major/Minor version number: 8.6 Total amount of global memory: 11.63 GBytes (12491292672 bytes) GPU Clock rate: 1665 MHZ (1.66 GHz) Memory Clock rate: 9501 MHz Memory Bus Width: 384-bit L2 Cache Size: 6291456 bytes Max Texture Dimension Size (x,y,z) 1D=(131072), 2D=(131072,65536), 3D=(16384,16384,16384) Total amount of constant memory: 65536 bytes Total amount of shared memory per block: 49152 bytes Total number of registers available per block:65536 Warp size: 32 Maximum number of MultiProcessor: 80 Maximum number of threads per multiprocessor: 1536 Maximum number of threads per block: 1024 Maximum sizes of each dimension of a block: 1024 x 1024 x 64 Maximum sizes of each dimension of a grid: 2147483647 x 65535 x 65535

GPU架构

流处理器&Fermi架构

Fermi架构

线程束

不能整除32，浪费

ptx兼容性

设置ptx架构版本

二进制兼容性

设置不兼容

CUDA运行时库 cuda编程核心API