术语
SM : streaming Multiprocessor
基础模块,有4block ( 16 fp32, 8 fp64, 16 int32, 128K L1 cache, 64k register)
SIMT : single instruction multiple threads
一组32 cores执行相同指令不同数据
gridDim.x
blockIdx.x
blockDim.x
threadIdx.x
nvcc --ptxas-options=-v
: 显示寄存器和 shared mem 信息
概念
- Page-Locked Memory : Unified Virtual Addressing (UVA),CPU 和 GPU 共享 VA,实现 zero copy,所有对 CPU 内存访问走 PCIe
在 CPU 使用 page-locked host memory
使用
cudaHostAlloc
和 cudaFreeHost
分配新内存
使用 cudaHostRegister
将 malloc 的内存变为 page-locked
- Unified Memory : GPU 自动管理内存,page fault 时从 CPU 拷贝到 GPU
usage
golbal
: host调用device执行函数
__syncthreads()
同步一个block内的thread
shared int a[]
: 每个block一个共享内存中的a
T __ldg(const T address)
从address读取一个数据T (从read-only data cache中读取)
thrust:
cuda 中类似 C++ STL
thrust::device_ptr<int> d_arr = thrust::device_pointer_cast<int>(d_bin_counter);
// host指针转device_ptr
thrust::exclusive_scan(d_arr, d_arr + SEGBIN_NUM, d_arr); // 每个位置变为自己之前不包括自己的前缀和
atomicAdd(int* address, int val)
将 val 原子加到 address
cudaDeviceSynchronize()
block直到所有设备同步
GPU中所有 active thread 被分配了单独的寄存器,当切换线程时不需要交换寄存器
cuda 中所有 kernel launch 都为异步
在开始和结束 CPU timer 之前都需要调用 cudaDeviceSynchronize
cuda 中 default stream (stream 0) 可以保证所有之前任何 stream 的调用均在前执行,且所有之后任何 stream 的调用均在后执行 (serializing)
事件计时:
cudaEvent_t start, stop;
float time;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord( start, 0 );
kernel<<<grid,threads>>> ( d_odata, d_idata, size_x, size_y,
NUM_REPS);
cudaEventRecord( stop, 0 );
cudaEventSynchronize( stop );
cudaEventElapsedTime( &time, start, stop );
cudaEventDestroy( start );
cudaEventDestroy( stop );
只能在 stream 0 中用此方法计时,其他 stream 无法保证顺序
多GPU
int deviceCount;
cudaGetDeviceCount(&deviceCount); // device 数量
cudaGetDeviceProperties(&deviceProp, 1); // 获取 device 1 的 properties
cudaSetDevice(0); // 设置当前 device 为 0,之后的内存分配和 kernel launch 等操作均在 device 0 上进行
特性
在 CUDA 中用户可以创建多个 stream,但同时不会有超过 32 个 kernel 执行
在一些 device 上, L1 cache 和 shared memory 共享硬件资源
optimization
#pragma unroll [factor]
: 强制循环展开,factor为展开参数,factor=1为不展开
restrict
: 指针重名优化(Pointer Aliasing),显示指定没有重名指针
A100
tensor core 理论 315TFLOPS
A100 dram 1500GB/s
A100 L2 40MB 所有SM共享 V100 L2 6MB
compute capability :
A100 : 8.0
V100 : 7.0
P100 : 6.0
额,请在编译时指定正确的 compute capability
hip
hipLaunchKernelGGL(MyKernel, dim3(gridDim), dim3(groupDim), 0/*dynamicShared*/, 0/*stream*/, a, b, c, n /*kernel函数参数*/);
// Replace MyKernel<<<dim3(gridDim), dim3(gridDim), 0, 0>>> (a,b,c,n);