中科曙光HPC培训教程汇总：D31-并行编程—CUDA程序设计简介.ppt

CUDA编程举例vector add 在GPU下运行的kernel函数的文件名后缀为.cu 核函数（__global__） 线程索引号的计算 线程私有变量的申请 shared memory的申请等等 设备端函数（__device__），只可以在device端调用，host端不可调用 在CPU运行的程序文件名可以为.c .cpp .cu 文件的读写 内存的申请 显存（global memory）的申请 CPU和GPU之间数据的交互 GPU端线程的申请 CPU端函数以及kernel函数的调用 内存的释放 显存的释放 CUDA编程举例vector add Step1: Step2: CUDA编程举例vector add Step3: Step4: Step5: Step6: CUDA编程举例vector add Step7: 函数执行 位置： 可以调用的 设备： __device__ float DeviceFunc ( ) GPU GPU __global__ void KernelFunc ( ) GPU CPU __host__ float HostFunc ( ) CPU CPU 对比: 减少Warps分支 什么是Warps 一个block中的每32个线程组成一个warps 这是一种实现的方式，并不是CUDA编程模型中的一部分（在分配线程是可以对线程分配方式进行调整，起到程序优化的作用） Warps的执行单元是SM 在一个Warps中的threads的执行方式是SIMD 什么是Warps分支 在同一个Warps中的线程执行了不同的操作（比如if判断导致的分支） 产生分支之后再一个Warps中的分支会串行进行 过多的Warps分支会导致性能的下降 Warps分支举例 If(threadIdx.x 2) { } 在一个Warps中产生了两个不同的操作，thread0, 1, 2执行相同的操作，而剩下的threads执行另外的操作。 shared memory的运用 合理的运用shared memory在进行程序优化时有着非常重要的作用，shared memory相比于global memory有着非常明显的优势 访存延时非常小，基本可以忽略 可以随机的进行访问 一个block中的threads可以共享，其他的threads不能进行访问，保证了数据的安全性。 一个block中的threads可以进行同步，在很多需要线程同步的程序中非常有用。 shared memory也有这自己的其他的特点 存储的空间不大 只能一个block内的threads共享，在保证数据安全性的同时，在其他的thread需要用到数据时，需要更多的操作。 shared memory的运用 图像在进行卷积操作是可以运用shared memory进行优化 假如卷积核较大，在进行整个图像卷积时，相同的像素会访问多次。 图像可能较大，可以进行分快处理。 可减少线程访问global memory的次数 shared memory的运用 shared memory的运用 注意上图中if括号内表示的是把数据从global memory中读入少shared memory中，在读的过程中不同的block会读取不同的数据，分三个通道分别读取，在图像的边缘部分，进行卷积的时候回超出图像，所以需要在其他的部分置为0，也就是上图中的else部分，每个shared memory中存有计算小块图像卷积的所有信息。 shared memory的运用 CUDA流水线技术 CUDA stream 声明：cudaStream_t stream0; 创建：cudaStreamCreate(stream0); 销毁：cudaStreamDestroy(stream0); 优化策略 Kernel函数和内存拷贝通过并行流水线的方式进行 cudaMemcpyAsync(deviceInput, hostInput, sizef, cudaMemcpyHostToDevice, stream0); Kernel函数之间通过并行流水线的方式进行 KernelDimGrid, DimBlock, 0, stream0(deviceInput, deviceOutput, Size); 利用Nvidia的函数库 利用Nvidia函数库对程序进行加速 Nvidia公司提供了很多有关线性代数，快速傅里叶变换和矩阵求解等函数库，并且进行了深层次的优化，可以在CUDA编程时直接调用。 矩阵运算的函数库：cublas 快速傅里叶变换的函数库：cufft 有关深度学习的函数库：cudnn 稀疏矩阵库：cuSPARSE 举例（cublas） cublas中数据的存储方式为列存储，所以在编程时