网站设计 工作,龙岗网络科技有限公司,网站制作公司网站,一般做网站费用Nvidia CUDA初级教程6 CUDA编程一 视频#xff1a;https://www.bilibili.com/video/BV1kx411m7Fk?p7 讲师#xff1a;周斌 GPU架构概览
GPU特别使用于#xff1a; 密集计算#xff0c;高度可并行计算图形学 晶体管主要被用于#xff1a; 执行计算而不是 缓存数据控制指令…Nvidia CUDA初级教程6 CUDA编程一 视频https://www.bilibili.com/video/BV1kx411m7Fk?p7 讲师周斌 GPU架构概览
GPU特别使用于 密集计算高度可并行计算图形学 晶体管主要被用于 执行计算而不是 缓存数据控制指令流 图中分别是CPU、GPU各个部件所占的芯片面积。可以看到CPU芯片中大量部分是缓存和控制逻辑而GPU中则绝大部分都是计算单元。
CUDA编程相关简介
CUDA的一些信息
层次化线程集合共享存储同步
CUDA术语
主机端和设备端
HOST - 主机端通常指CPU 采用ANSI标准C语言编程 Device - 设备端通常指GPU数据可并行 采用ANSI标准C的扩展语言编程 CUDA C HOST 和 Device 拥有各自的存储器CUDA编程 包括主机端和设备端两部分代码
核
Kernel 数据并行处理函数 通过调用 Kernel 函数在设备端创建轻量级的线程线程由硬件负责创建并调度
类似于 OpenCL 的 shader 核函数会在 N 个不同的 CUDA 线程上并行执行 // 定义核函数
__global__ void VecAdd(float* a, float* B, float* C) {int i threadIdx.x;C[i] A[i] B[i];
}int main() {// ...// 在N个线程上调用核函数VecAdd1, N(A, B, C);
}CUDA程序的执行
CUDA程序执行的流程大体上是这样的当我们在CPU端的代码是串行执行的这里简单地认为指令在CPU上串行执行当遇到需要并行大量处理数据时会调用核函数在GPU上进行计算计算完成后将结果返回给CPU。 线程层次
Grid - 一维或多维线程块block 一维或二维 Block - 一维线程 一维二维或三维一个 Grid 中的每个 Block 的线程数是一样的Block 内部的每个线程可以 同步 Synchronize访问共享存储器 shared memory
一个线程可以类比为一个员工一个 block 是一个科室grid 是整个公司。 线程ID
每一个线程都有一个索引threadIdx
一维 Block Thread ID Thread Index二维 Block (Dx, Dy) 索引为 (x, y) 的 Thread ID x yDy 三维 Block (Dx, Dy, Dz) 索引为 (x, y) 的 Thread ID x yDy zDxDy
代码实例
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {int i threadIdx.x;int j threadIdx.y;C[i][j] A[i][j] B[i][j];
}int main() {int numBlocks 1;dim3 threadsPerBlock(N, N);MatAddnumBlocks, threadsPerBlock(A, B, C);
}每个 Block 中的线程的索引是二维的这在我们处理二维数据矩阵时可以很方便地进行对应。
线程数
Thread Block 线程块
线程的的集合 G80 和 GT200多达512个线程Fermi多达1024个线程 位于相同的处理器核相同的SM共享所在核的存储器 块索引
块索引blockIdx维度blockDim 一维二维或三维
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {int i blockIdx.x * blockDim.x threadIdx.x;int j blockIdx.y * blockDim.y threadIdx.y;if (i N j N)C[i][j] A[i][j] B[i][j];
}int main() {// ...dim3 threadsPerBlock(16, 16);dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);MatAddnumBlocks, threadsPerBlock(A, B, C);
}例如 N 32
每个块有16x16个线程跟N无关 threadIdx([0, 15], [0, 15]) Grid 里面有 2x2 个线程块 block blockIdx([0, 1], [0, 1])blockDim 16
int i blockIdx.x * blockDim.x threadIdx.x;
int j blockIdx.y * blockDim.y threadIdx.y;i [0, 1] * 16 [0, 15]
线程块之间
线程块彼此之间独立执行
任意顺序并行或串行被任意数量的处理器以任意顺序调度 处理器的数量具有可扩展性
一个块内部的线程
一个块内部的线程有一些重要的特性
共享容量有限的低延迟存储器 (shared memory)同步执行 合并访存__syncThreads() barrier - 块内线程一起等待所有的线程都轻量级线程
CUDA内存传输
主机端与设备端
CUDA内存传输 device 端代码可以 读写该线程的 registers读写该线程的local memory读写该线程所属的块的 shared memory读写grid的 global memory只读grid的 constant memory host 端代码可以 读写grid的 global memory 和 constant memory host 可以从 device 往返传输数据 global memory 全局存储器constant memory 常量存储器
CUDA内存传输函数
在设备端分配 global memorycudaMalloc()释放存储空间 cudaFree()
float* Md;
int size Width * Width * sizof(float);
cudaMalloc((void**)Md, size);
//...
cudaFree(Md);注意这里的指针 Md 是指向 deviceGPU上的存储空间。
内存传输cudaMemcpy() host to hosthost to devicedevice to hostdevice to device
cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);
cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);示例矩阵相乘 Matrix Multiply
矩阵相乘简介
向量点乘行优先或列优先每次点乘结果输出一个元素 矩阵相乘CPU实现
void MatrixMulOnHost(float* M, float* N, float* P, int width) {for (int i0; iwidth; i) {for (int j0; jwidth; j) {float sum 0;for (int k0; kwidth; k) {float a M[i * width k];float b N[k * width j]:sum a * b;}P[i * width j] sum;}}
}CUDA算法框架
三步走
int main(void) {// 1 分配device空间// 2 在GPU上并行计算MatrixMulOnDevice(M, N, P, width);// 3 将结果拷贝回CPU并释放device空间return 0;
}伪代码如下
void MatrixMulOnDevice(float* M, float* N, float* P, int Width) {int size Width * Width * sizeof(float);// 1cudaMalloc(Md, size);cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);cudaMalloc(Nd, size);cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);cudaMalloc(Pd, size);// 2 调用cuda核函数并行计算// 3cudaMemcpy(P. Pd, size, cudaMemcpyDeivceToHost)cudaFree(Md); cudaFree(Nd); cudaFree(Pd);
}CUDA C 实现
矩阵相乘样例 目前版本矩阵相乘的问题
在上述算法实现中最主要的性能问题是什么 访存 主要的限制是什么 访存带宽
