CUDA编程指南阅读笔记(六)
4. CUDA C语言编程接口
接上文。
4.3 CUDA C Runtime
4.3.3 共享内存(Shared Memory)
- // Matrices are stored in row-major order:
- // M(row, col) = *(M.elements + row * M.width + col)
- typedef struct {
- int width;
- int height;
- float *elements;
- } Matrix;
- // Thread block size
- #define BLOCK_SIZE 16
- // Forward declaration of the matrix multiplication kernel
- __global__ void MatMulKernel(const Matrix, const Matrix, Matrix);
- // Matrix multiplication - Host code
- // Matrix dimensions are assumed to be multiples of BLOCK_SIZE
- void MatMul(const Matrix A, const Matrix B, Matrix C) {
- // Load A and B to device memory
- Matrix d_A;
- d_A.width = A.width; d_A.height = A.height;
- size_t size = A.width * A.height * sizeof(float);
- cudaMalloc(&d_A.elements, size);
- cudaMemcpy(d_A.elements, A.elements, size, cudaMemcpyHostToDevice);
- Matrix d_B;
- d_B.width = B.width; d_B.height = B.height;
- size = B.width * B.height * sizeof(float);
- cudaMalloc(&d_B.elements, size);
- cudaMemcpy(d_B.elements, B.elements, size, cudaMemcpyHostToDevice);
- // Allocate C in device memory
- Matrix d_C;
- d_C.width = C.width; d_C.height = C.height;
- size = C.width * C.height * sizeof(float);
- cudaMalloc(&d_C.elements, size);
- // Invoke kernel
- dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
- dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
- MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);
- // Read C from device memory
- cudaMemcpy(C.elements, d_c.elements, size, cudaMemcpyDeviceToHost);
- // Free device memory
- cudaFree(d_A.elements);
- cudaFree(d_B.elements);
- cudaFree(d_C.elements);
- }
- // Matrix multiplication kernel called by MatMul()
- __global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) {
- // Each thread computes one element of C
- // by accumulating results into Cvalue
- float Cvalue = 0;
- int row = blockIdx.y * blockDim.y + threadIdx.y;
- int col = blockIdx.x * blockDim.x + threadIdx.xl
- for (int e = 0; e < A.width; ++e)
- Cvalue += A.elements[row * A.width + e] * B.elements[e * B.width + col];
- C.elements[row * C.width + col] = Cvalue;
- }
可以看出,为了计算矩阵C的任何一个元素,程序都需要从全局内存(global memory)中获得矩阵A的一行和矩阵B的一列。因此,完成这一计算矩阵A被读取了B.width次,矩阵B被读取了A.height次。
现在我们来使用共享内存(shared memory)实现矩阵乘法。假设矩阵C可以被划分为若干个较小的子方阵Csub,我们使用一个线程块(thread block)来负责某一子方阵的计算,线程块中的每一个线程(thread)正好负责子方阵Csub中一个元素的计算。这样划分后,任何一个结果子方阵Csub'(尺寸为block_size * block_size)都是与该方阵具有相同行索引的尺寸为A.width * block_size的A的子矩阵Asub和与该方阵具有相同列索引的尺寸为block_size * B.height的B的子矩阵Bsub相乘所得到。
- // Matrices are stored in row-major order;
- // M(row, col) = *(M.elements + row * M.stride + col)
- typedef struct {
- int width;
- int height;
- int stride;
- float* elements;
- } Matrix;
- // Get a matrix element
- __device__ float GetElement(const Matrix A, int row, int col) {
- return A.elements[row * A.stride + col];
- }
- // Set a matrix element
- __device__ void SetElement(Matrix A, int row, int col, float value) {
- A.elements[row * A.stride + col] = value;
- }
- // Get the BLOCK_SIZExBLOCK_SIZE sub-matrix Asub of A that is
- // located col sub-matrices to the right and row sub-matrices down
- // from the upper-left corner of A
- __device__ Matrix GetSubMatrix(Matrix A, int row, int col) {
- Matrix Asub;
- Asub.width = BLOCK_SIZE;
- Asub.height = BLOCK_SIZE;
- Asub.stride = A.stride;
- Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row + BLOCK_SIZE * col];
- return Asub;
- }
- // Thread block size
- #define BLOCK_SIZE 16
- // Forward declaration of the matrix multiplication kernel
- __global__ void MatMulKernel(const Matrix, const Matrix, Matrix);
- // Matrix multiplication - Host code
- // Matrix dimensions are assumed to be multiples of BLOCK_SIZE
- void MatMul(const Matrix A, const Matrix B, Matrix C) {
- // Load A and B to device memory
- Matrix d_A;
- d_A.width = d_A.stride = A.width;
- d_A.height = A.height;
- size_t size = A.width * A.height * sizeof(float);
- cudaMalloc(&d_A.elements, size);
- cudaMemcpy(d_A.elements, A.elements, size, cudaMemcpyHostToDevice);
- Matrix d_B;
- d_B.width = d_B.stride = B.width;
- d_B.height = B.height;
- size = B.width * B.height * sizeof(float);
- cudaMalloc(&d_B.elements, size);
- cudaMemcpy(d_B.elements, B.elements, size, cudaMemcpyHostToDevice);
- // Allocate C in device memory
- Matrix d_C;
- d_C.width = d_C.stride = C.width;
- d_C.height = C.height;
- size = C.width * C.height * sizeof(float);
- cudaMalloc(&d_C.elements, size);
- // Invoke kernel
- dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
- dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
- MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);
- // Read C from device memory
- cudaMemcpy(C.elements, d_C.elements, size, cudaMemcpyDeviceToHost);
- // Free device memory
- cudaFree(d_A.elements);
- cudaFree(d_B.elements);
- cudaFree(d_C.elements);
- }
- // Matrix multiplication kernel called by MatMul()
- __global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) {
- // Block row and column
- int blockRow = blockIdx.y;
- int blockCol = blockIdx.x;
- // Each thread block computes one sub-matrix Csub of C
- Matrix Csub = GetSubMatrix(C, blockRow, blockCol);
- // Each thread computes one element of Csub
- // by accumulating results into Cvalue
- float Cvalue = 0;
- // Thread row and column within Csub
- int row = threadIdx.y;
- int col = threadIdx.x;
- // Look over all the sub-matrices of A and B that are required to compute Csub
- // Multiply each pair of sub-matrices together and accumulate the results
- for (int m = 0; m < (A.width / BLOCK_SIZE); ++m) {
- // Get sub-matrix Asub of A
- Matrix Asub = GetSubMatrix(A, blockRow, m);
- // Get sub-matrix Bsub of B
- Matrix Bsub = GetSubMatrix(B, m, blockCol);
- // Shared memory used to store Asub and Bsub respectively
- __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
- __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
- // Load Asub and Bsub from device memory to shared memory
- // Each thread loads one element of each sub-matrix
- As[row][col] = GetElement(Asub, row, col);
- Bs[row][col] = GetElement(Bsub, row, col);
- // Synchronize to make sure the sub-matrices are loaded
- // before starting the computation
- __syncthreads();
- // Multiply Asub and Bsub together
- for (int e = 0; e < BLOCK_SIZE; ++e)
- Cvalue += As[row][e] * Bs[e][col];
- // Synchronize to make sure that the preceding computation is done before
- // loading two new sub-matrices of A and B in the next iteration
- __syncthreads();
- }
- // Write Csub to device memory
- // Each thread writes one element
- SetElement(Csub, row, col, Cvalue);
- }
同一个block中的线程在95行的for循环中获取到的Asub,Bsub,Csub是一样的,每个线程就负责Csub内元素的计算
http://blog.csdn.net/csgxy123/article/details/10018531
CUDA编程指南阅读笔记(六)相关推荐
- CUDA编程指南阅读笔记
随着多核CPU和众核GPU的到来,并行编程已经得到了业界越来越多的重视,CPU-GPU异构程序能够极大提高现有计算机系统的运算性能,对于科学计算等运算密集型程序有着非常重要的意义.这一系列文章是根据& ...
- Java并发编程艺术阅读笔记(一)
Java并发编程艺术阅读笔记(一) 1.什么是上下文切换 CPU通过时间片分配算法循环执行任务,执行完一个任务的时间片后就会切换到下一个任务.但是在切换任务之前会保存上一个任务的状态,在切换回该任务, ...
- CUDA PTX ISA阅读笔记(一)
不知道这是个啥的看这里:Parallel Thread Execution ISA Version 5.0. 简要来说,PTX就是.cu代码编译出来的一种东西,然后再由PTX编译生成执行代码.如果不想 ...
- iPhone编程指南学习笔记
为什么80%的码农都做不了架构师?>>> UIWindow 在创建应用程序窗口时,您应该总是将其初始的边框尺寸设置为整个屏幕的大小.如果您的窗口是从nib文件装载得到,Inte ...
- wrappers.php,PHP源码阅读笔记六:stream_get_wrappers函数
PHP源码阅读笔记stream_get_wrappers函数 stream_get_wrappers (PHP 5) stream_get_wrappers - 返回注册的数据流列表 Descript ...
- Google C++编程风格指南阅读笔记之命名、注释和格式
文章目录 前言 命名约定 类型命名 变量命名 枚举命名 宏的命名 注释 注释风格 文件注释 类注释 函数注释 变量注释 类的数据成员 全局变量 实现注释 TODO注释 格式 行长度 空格还是制表符 函 ...
- [转载]高质量c/c++编程指南读书笔记
一个strcpy函数的代码 能考查三个方面 (1) 编程风格 (2) 出错处理 (3) 算法复杂度分析(用于提供性能) 定义编程老手和编程高手 定义1:能长期稳定地编写出高质量程序的程序员称为编程老手 ...
- 程序员健康指南阅读笔记
编程需要程序员全神贯注,这就常常导致我们忽略生活的其他方面,其中最被忽略的就是,健康状况. 职业不应该让你犯职业病,只要用对了方法,它就不会造成伤害. 做出改变 努力健康起来吧! 永别了椅子 灵活的饮 ...
- 高质量c/c++编程指南读书笔记1
一个strcpy函数的代码 能考查三个方面 (1) 编程风格 (2) 出错处理 (3) 算法复杂度分析(用于提供性能) 定义编程老手和编程高手 定义1:能长期稳定地编写出高质量程序的程序员称为编程老手 ...
最新文章
- [转] Zend studio中开发常用
- iPhone游戏编程教程一步步教你游戏开发
- 学习笔记Spark(四)—— Spark编程基础(创建RDD、RDD算子、文件读取与存储)
- 密码学-hash加密
- 《笨办法学python》(《learn python the hard way 3thrd》)习题48(ex48)的代码实现
- 计算机网络相关的知识,计算机网络知识整理
- superset可视化-word cloud
- Java中父类的静态变量和静态方法的继承问题
- volley全然解析
- gtest的介绍和使用
- JavaScript 优先队列
- 项目三大文档: 项目章程, 初步的项目范围说明书 ,详细范围说明书,项目计划书
- rmvb 转 avi 方法
- java获取指定日期当月和下个月的第一天
- hotmail邮箱收件服务器主机名,hotmail邮箱foxmail(pop3服务器)设置方法
- 基于Hi3516DV300rtmp交叉编译移植
- Linux系统编程 50 -stat和stat函数 穿透和非穿透
- Elasticsearch(十)【NEST高级客户端--搜索查询】
- 2022年国家高新技术企业申报秘笈来了
- 详解Python中的File(文件)操作
热门文章
- 两个主机mtu不相同_案例详解:MTU不一致导致主机和RAC不断重启
- 八十八、CSS两列三列的布局方式
- 二十七、深入==与equals的区别(下篇)
- 六十六,完成SpringBoot项目中的员工增删查改功能
- 我跑了ERNIE和BERT两个模型,结果出乎意料......
- 2018 ACM-ICPC亚洲区域赛 北京赛区
- POJ 2826 An Easy Problem?! 叉积求多边形面积 【计算几何】
- php实现上传文件功能,简单实现php上传文件功能
- 中的挂起是什么意思_书房装饰挂什么画好 书法字画给你想要的诗意生活
- Springboot项目中配置tomcta监控日志