CUDA笔记002：CUDA基础入门
前言
本篇为了巩固基础的cuda语法，学习部分cuda的相关函数
使用资源
谭升的CUDA博客：2.0-2.4：https://face2ai.com/CUDA-F-2-0-CUDA编程模型概述1/
file:///F:/new_interview/%E7%9F%A5%E8%AF%86%E5%BA%93/cuda/cuda%E6%95%99%E7%A8%8B/Professional%20CUDA%20C%20Programming.pdf
谭升的CUDA博客
2.0
- 每个block内部有自己独享的shared memory
- grid内部有全局的global memory
一个经常会发生的错误就是混用设备和主机的内存地址！！
所以，为了避免混淆，提出了一个方案：给内存变量加后缀，_d表示device，_h表示host
不同block内线程不能相互影响！他们是物理隔离的！
计算thread index：依赖blockIdx和threadIdx
blockIdx.x
blockIdx.y
blockIdx.z
threadIdx.x
threadIdx.y
threadIdx.z
以上两个index的范围，由dim3类型的blockDim和gridDim指定
<<<grid,block>>>
- grid: grid中的block num
- block: block中的thread num
2.1
一个很重要的思想是：把串行的for展开，然后可以并行化。
这个过程通常需要借助blockIdx和threadIdx，计算自己的thread_index，例如：
__global__ void sumArraysOnGPU(float *A, float *B, float *C) {
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
Q: 一个编写良好（至少正确）的核函数，grid_num和block_num不会影响其正确性，<<<32, 32>>> 和 <<<1, 1>>>结果相同，只是执行时间有区别。这种说法正确吗？
A: 不完全正确，编程实践中是会依赖gridDim 和 blockDim的。
Q: 对于线程数量不足的情况，可以添加for使得该线程处理更多任务，而当线程够多时则每个线程处理自己的。这样的编程思想正确吗？
A: 正确的，这种思想被称为 “循环展开（loop unrolling）” 或 “任务批处理（work batching）”，是 CUDA 编程中常见的一种优化策略。尤其是在支持的最大线程数不足时使用。
所以，不必强求gridDim 和 blockDim即使变化也要兼容，只需要确定预先设置正确即可
注意，大部分cuda函数的返回类型都是cudaError_t ，看起来类似Linux中的errno
一个处理函数的小技巧：
#define CHECK(call)\
{\
const cudaError_t error=call;\
if(error!=cudaSuccess)\
{\
printf("ERROR: %s:%d,",__FILE__,__LINE__);\
printf("code:%d,reason:%s\n",error,cudaGetErrorString(error));\
exit(1);\
}\
}
CHECK(cudaMemcpy(res_from_gpu_h,res_d,nByte,cudaMemcpyDeviceToHost));
2.2
并行程序中的函数计时：
clock()函数计算的是cpu time，在并行场景下计算的是所有thread的cpu time。所以并行场景下需要使用gettimeofday
否则你可能在计算某个函数的性能占比时，算出来超过100%，这个情况在我写trigger benchmark的时候真的遇到过。
在性能测试中发现：当数据不能被完整切块的时候性能滑铁卢了，这个我们可以使用一点小技巧，比如只传输可完整切割数据块，然后剩下的1，2个使用cpu计算
尽量使得线程数量和数据数量对齐，没必要折磨自己也折磨机器
nvprof是cuda的性能分析工具，做性能优化使用。
nvprof ./test 尝试运行vectorAdd，结果如下：
==54506== NVPROF is profiling process 54506, command: ./test
==54506== Warning: Unified Memory Profiling is not supported on the current configuration because a pair of devices without peer-to-peer support is detected on this multi-GPU setup. When peer mappings are not available, system falls back to using zero-copy memory. It can cause kernels, which access unified memory, to run slower. More details can be found at: http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#um-managed-memory
最大误差: 0
==54506== Profiling application: ./test
==54506== Warning: 26 API trace records have same start and end timestamps.
This can happen because of short execution duration of CUDA APIs and low timer resolution on the underlying operating system.
==54506== Profiling result:
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 60.46% 2.4588ms 1 2.4588ms 2.4588ms 2.4588ms [CUDA memcpy DtoH]
38.46% 1.5640ms 2 782.01us 681.73us 882.28us [CUDA memcpy HtoD]
1.09% 44.192us 1 44.192us 44.192us 44.192us add(float*, float*, float*, int)
API calls: 92.08% 253.26ms 3 84.419ms 412.90us 249.43ms cudaMalloc
3.78% 10.404ms 1 10.404ms 10.404ms 10.404ms cudaLaunchKernel
2.06% 5.6755ms 3 1.8918ms 831.60us 3.9078ms cudaMemcpy
1.79% 4.9199ms 114 43.157us 0ns 4.9073ms cuDeviceGetAttribute
0.28% 759.50us 3 253.17us 181.90us 377.10us cudaFree
0.00% 8.1000us 1 8.1000us 8.1000us 8.1000us cuDeviceGetPCIBusId
0.00% 1.1000us 3 366ns 100ns 700ns cuDeviceGetCount
0.00% 1.0000us 2 500ns 300ns 700ns cuDeviceGet
0.00% 1.0000us 1 1.0000us 1.0000us 1.0000us cuDeviceGetName
0.00% 200ns 1 200ns 200ns 200ns cuDeviceTotalMem
0.00% 200ns 1 200ns 200ns 200ns cuDeviceGetUuid
0.00% 100ns 1 100ns 100ns 100ns cuModuleGetLoadingMode
观察到，实际的add只有1.09%，时间占比绝大多数是内存拷贝
以上是当前的数据性能，我们需要学会评估当前设备的理论计算极限，尽可能优化代码接近这个极限。
各个设备的理论极限可以通过其芯片说明计算得到，以下是谭升给出的一个计算实例:
Tesla K10 单精度峰值浮点数计算次数：745MHz核心频率 x 2GPU/芯片 x（8个多处理器 x 192个浮点计算单元 x 32 核心/多处理器） x 2 OPS/周期 =4.58 TFLOPS
Tesla K10 内存带宽峰值： 2GPU/芯片 x 256 位 x 2500 MHz内存时钟 x 2 DDR/8位/字节 = 320 GB/s
指令比：字节 4.58 TFLOPS/320 GB/s =13.6 个指令： 1个字节
2.3
这里我们回顾了之前thread_index的计算方式，比如说，对于二维矩阵，理论上：
ix=threadIdx.x+blockIdx.x×blockDim.x
iy=threadIdx.y+blockIdx.y×blockDim.y
matrix[ix][iy] // 当前矩阵的位置
当然，由于我们一般使用一位矩阵来模拟，所以这里要接着上面做一次映射：
int idx=ix+iy*ny;
matrix[idx] // 当前矩阵
事实上，对于二维甚至三维矩阵，一般情况下优秀的工程实践是使用
float *matrix来模拟二维/三维，而不是使用float ***matrix
多次指针dereference，导致非连续内存访问，影响性能。
且cudaMalloc也不支持float ***matrix格式的内存申请
事实上，我们一般使用cudaMallocPitch() 和 cudaMalloc3D() 来优化2D和3D内存的分配，对于内存对齐有好处。
配合cudaMemcpy2D()和cudaMemcpy3D()一起使用。
感谢gpt老师给出了一个例子，这里的row代表了rowHead，指的是每行首的指针地址，要学好c才能看懂哦:->
__global__ void printThreadIndex(float *A, int pitch, int nx, int ny) {
int ix = threadIdx.x + blockIdx.x * blockDim.x;
int iy = threadIdx.y + blockIdx.y * blockDim.y;
if (ix < nx && iy < ny) {
// 使用 pitch 计算二维索引
float *row = (float *)((char *)A + iy * pitch); // 计算行地址
float value = row[ix]; // 读取列值
}
}
谭升这里举了一个例子很有趣：
__global__ void sumMatrix(float * MatA,float * MatB,float * MatC,int nx,int ny)
{
int ix=threadIdx.x+blockDim.x*blockIdx.x;
int iy=threadIdx.y+blockDim.y*blockIdx.y;
int idx=ix+iy*ny;
if (ix<nx && iy<ny)
{
MatC[idx]=MatA[idx]+MatB[idx];
}
}
这个矩阵相加函数依赖idx作为全局所有thread的总id，在一维矩阵上进行操作，每个线程只操作一次
所以，
<<<(128,128),(32,32)>>>
<<<(524288,1),(32,1)>>>
<<<(128,4096),(32,1)>>>
都能够运行成功，只是由于内存物理排版导致的速度不同
2.4
nvidia-smi，在不会cuda的时候就应该知道如何去查看显卡配置了。
有各种API来支持你查找GPU的各种参数，此处不在多说，用到什么查什么即可。
浙公网安备 33010602011771号