一、课程目标 00:04
●
● 理论目标: 进一步认识GPU并行原理
● 技能目标:
○ 实现基本并行的矩阵乘法
○ 实现进阶版矩阵乘法
二、内容回顾 00:15
- GPU并行基础概念
●
● 核函数: 设备端运行的函数，使用global关键字修饰
● 线程组织:
○ 主机端通过<<<…>>>语法传递线程组织方式
○ 线程块(block)和网格(grid)的索引可以是三维的 - CUDA编程三步骤
●
● Part1:
○ 主机端申请显存
○ 内存内容拷贝到显存
● Part2:
○ 构建核函数在设备端运行
● Part3:
○ 设备端内容回传到主机端 - 线程配置优化
●
● 线程块大小: 使用256个线程构成每个线程块
● 网格大小: 计算ceil(n/256)作为网格维度
● 三维索引:
○ dim3 DimGrid(n/256,1,1)
○ dim3 DimBlock(256,1,1)
三、矩阵乘法基础 - 矩阵乘法原理
●
● 维度要求:
○ 矩阵A维度为m×km×km×k
○ 矩阵B维度为k×nk×nk×n
○ 结果矩阵C维度为m×nm×nm×n
● 计算规则:
○ Cij=∑k=1nAik×BkjC_{ij} = \sum_{k=1}^{n} A_{ik} × B_{kj}Cij=k=1∑nAik×Bkj
○ 即A的第i行与B的第j列对应元素相乘后累加 - GPU并行优势
● 适合并行原因:
○ 矩阵乘法中各元素计算相互独立
○ 可充分利用GPU的并行计算能力
○ 相比向量加法能更显著体现GPU优势
四、矩阵乘法 01:39 - 矩阵乘法的CPU实现 02:42
●
● 算法实现：
○ 定义三个矩阵AAA、BBB、CCC在内存中的存储空间
○ 使用三重嵌套循环：
■ 外层循环iii遍历AAA的每一行（000到M−1M-1M−1）
■ 中层循环jjj遍历BBB的每一列（000到N−1N-1N−1）
■ 内层循环kkk进行元素级运算（000到K−1K-1K−1）
● 核心计算：C(i,j)=C(i,j)+A(i,k)×B(k,j)C(i,j) = C(i,j) + A(i,k) \times B(k,j)C(i,j)=C(i,j)+A(i,k)×B(k,j)
● 时间复杂度：O(M×N×K)O(M \times N \times K)O(M×N×K)
○ 当MMM、NNN、KKK相近时，可简化为O(n3)O(n^3)O(n3)
○ 相比向量加法O(n)O(n)O(n)，矩阵乘法更能体现GPU优势 - 矩阵乘法的GPU实现
●
● 实现步骤：
○ 在CPU内存定义AcpuA_{cpu}Acpu、BcpuB_{cpu}Bcpu、CcpuC_{cpu}Ccpu
○ 在GPU内存定义AgpuA_{gpu}Agpu、BgpuB_{gpu}Bgpu、CgpuC_{gpu}Cgpu
○ 内存拷贝：Acpu→AgpuA_{cpu} \rightarrow A_{gpu}Acpu→Agpu，Bcpu→BgpuB_{cpu} \rightarrow B_{gpu}Bcpu→Bgpu
○ 设置线程块维度：dim3 dimBlock(16,16)
○ 设置网格维度：dim3 dimGrid(N/dimBlock.x, M/dimBlock.y)
○ 调用核函数：matrixMul<<<dimGrid,dimBlock>>>(A_gpu,B_gpu,C_gpu,K)
○ 结果回传：Cgpu→CcpuC_{gpu} \rightarrow C_{cpu}Cgpu→Ccpu
1）GPU适用性分析
● 缓存特性：
○ GPU缓存小，适合缓存访问简单的计算
○ 矩阵乘法内存访问模式简单，满足要求
● 控制复杂度：
○ GPU缺乏复杂分支预测单元
○ 矩阵乘法无复杂分支判断，控制简单
● 计算特性：
○ GPU有大量简单计算单元
○ 矩阵乘法仅需乘加操作，计算密集
● 并行性：
○ GPU采用SIMD架构（单指令多数据）
○ 矩阵乘法中CCC的各元素计算相互独立，并行度高 - 矩阵乘法的GPU实现 03:31
1）GPU实现三步骤 05:12
●
● 步骤1：内存定义与拷贝
○ 在主机端定义矩阵A、B、C的内存地址
○ 在设备端定义对应的显存地址A_gpu、B_gpu、C_gpu
○ 使用memcpy函数实现主机内存到显存的互相拷贝
● 步骤2：核函数书写
○ 核心计算部分，采用并行计算方式
○ 使用dim3 dimBlock(16,16)定义二维线程块
○ 使用dim3 dimGrid(N/dimBlock.x,M/dimBlock.y)定义线程块组合体维度
● 步骤3：结果回传与释放
○ 将计算结果从设备端显存拷贝回主机端内存
○ 释放主机端和设备端的内存资源
2）矩阵乘法GPU核心核函数 06:08
● 线程块与线程索引 06:20
○
○ 线程组织方式
■ 每个线程处理矩阵C中的一个元素
■ 采用二维索引的线程块和线程块组合体
■ 使用16×16的二维线程块组织方式
○ 索引计算
■ 行索引：i = blockIdx.y * blockDim.y + threadIdx.y
■ 列索引：j = blockIdx.x * blockDim.x + threadIdx.x
● 核函数参数与计算过程 06:41
○ 参数说明
■ 输入矩阵A_gpu(m×k)、B_gpu(k×n)
■ 输出矩阵C_gpu(m×n)
■ 参数K表示矩阵A的列数和矩阵B的行数
○ 计算过程
■ 初始化累加器accu=0
■ 循环k=0到K-1：
● accu += A_gpu(i,k) * B_gpu(k,j)
■ 将结果存入C_gpu(i,j) = accu
● 内存访问与效率考量 07:23
○
○ 访问复杂度分析
■ 矩阵A的每个元素被读取n次（C的列数）
■ 矩阵B的每个元素被读取m次（C的行数）
■ 总访问次数：2mnk次（每次乘法需读取A和B各一次）
○ 性能瓶颈
■ 全局内存访问时间占比大
■ 计算复杂度主要来自内存访问而非计算本身
● 共享内存的概念与优势 09:03
○
○ 基本概念
■ 线程块内的共享存储空间
■ 可由线程块内所有线程共同访问
■ 读写速度远快于全局内存（约100倍）
○ 优化原理
■ 将频繁访问的数据放入共享内存
■ 减少全局内存的重复读取次数
■ 利用共享内存低延迟的优势
● 共享内存的申请方式 11:11
○
○ 静态申请
■ 编译时确定共享内存大小
■ 使用__shared__关键字声明
■ 示例：shared int s[64]
○ 动态申请
■ 运行时确定共享内存大小
■ 使用extern shared__关键字声明
■ 调用核函数时需传入共享内存大小参数
○ 使用限制
■ 共享内存大小仅几十KB
■ 过度使用会降低程序并行性
● 共享内存的使用注意事项 12:51
○ 使用步骤
■ 从全局内存拷贝数据到共享内存
■ 注意数据边界处理
■ 使用同步函数确保数据一致性
○ 边界问题
■ 需明确哪些数据需要拷贝
■ 避免不同线程访问冲突
■ 特别注意边界条件的处理
● 线程同步函数_syncthreads() 13:36
○
○ 功能作用
■ CUDA内置的块内线程同步函数
■ 确保共享内存数据的一致性
■ 防止读写时序错乱
○ 使用规范
■ 尽量使用简单的同步结构
■ 避免复杂分支条件下的同步
■ 确保所有需要同步的线程都能到达同步点
○ 错误示例
■ 不同条件分支下分别同步
■ 可能导致部分线程读取未同步数据
■ 计算结果可能出现错误
● 静态申请共享内存示例 16:17
○
○ 实现代码
○ 调用方式
■ 与普通核函数调用相同
■ 示例：staticReverse<<<1,n>>>(d_d, n)
● 动态申请共享内存示例 16:50
○
○ 实现代码
○ 调用方式
■ 需要额外指定共享内存大小
■ 示例：dynamicReverse<<<1,n,n*sizeof(int)>>>(d_d, n)
3）使用共享内存复用全局内存数据 18:40
● 基础版矩阵乘法效率不高原因分析 18:58
○
○ 重复读取问题：在计算C=ABC=ABC=AB时，矩阵A的每行需要被读取nnn次，矩阵B的每列需要被读取mmm次，导致全局内存访问次数过多。
○ 单元素计算瓶颈：每次只计算CCC中的一个元素cijc{ij}cij，需要从全局内存读取A的第iii行和B的第jjj列各kkk次，共2k2k2k次全局内存访问。
○ 改进方向：
■ 将频繁读取的数据放入共享内存
■ 每次计算多个CCC的元素（矩阵块）
● 平铺矩阵乘法 20:27
○
○ 核心思想：将矩阵乘法分解为多个阶段，每个阶段集中处理数据的一个子集（tile）。
○ 分块计算原理：
■ 计算CCC中一个block_size×block_sizeblock_size \times block_sizeblock_size×block_size的块需要对应AAA的行块和BBB的列块
■ 将大矩阵乘法分解为多个小矩阵块相乘再累加：Cblock=Ablue×Bblue+Aorange×Borange+…C{block} = A_{blue} \times B_{blue} + A_{orange} \times B_{orange} + …Cblock=Ablue×Bblue+Aorange×Borange+…
○ 内存访问优化：
■ 原始方法需要block_size2×2kblock_size^2 \times 2kblock_size2×2k次全局内存访问
■ 平铺方法只需2k×block_size2k \times block_size2k×block_size次全局内存访问
■ 理论加速比为block_sizeblock_sizeblock_size倍（实际略低）
○ 平铺矩阵乘法核函数 25:59
■
■ 共享内存声明：
■ 关键步骤：
● 将全局内存数据加载到共享内存
● 使用__syncthreads()确保所有线程完成加载
● 从共享内存读取数据进行计算
● 再次同步确保计算完成
■ 索引计算：
● 行索引：i = blockIdx.y * blockDim.y + threadIdx.y
● 列索引：j = blockIdx.x * blockDim.x + threadIdx.x
■ 循环结构：
● 外层循环遍历所有tile（K/block_sizeK/block_sizeK/block_size次）
● 内层循环计算单个tile的乘积累加
■ 同步要求：每次共享内存加载后必须同步
● 性能分析
○ 全局内存访问次数：
■ 原始矩阵乘法：2mnk2mnk2mnk次
■ 平铺矩阵算法：2mnk/block_size2mnk/block_size2mnk/block_size次
○ 实际加速比：
■ 理论最大加速比：block_sizeblock_sizeblock_size倍
■ 实际加速比略低（考虑同步和共享内存访问开销）
○ 优化关键：
■ 合理选择block_sizeblock_sizeblock_size值
■ 平衡共享内存使用和计算并行度
五、矩阵乘法文档实现 30:16 - CPU矩阵乘法实现 30:27
●
● 矩阵维度:
○ A矩阵：M×KM×KM×K
○ B矩阵：K×NK×NK×N
○ C矩阵：M×NM×NM×N
● 实现原理:
○ 使用三重循环结构
○ 外层两重循环遍历C矩阵每个元素
○ 内层循环进行KKK次乘加运算
○ 时间复杂度为O(M×N×K)O(M×N×K)O(M×N×K)
○
● 函数参数:
○ array_A, array_B, array_C：三个矩阵的内存地址
○ M_p, K_p, N_p：对应矩阵维度
● 核心算法: - GPU基础版矩阵乘法 31:02
●
● 实现特点:
○ 使用二维线程结构
○ 每个线程计算C矩阵一个元素
○ 线程索引计算：
● 存在问题:
○ 重复加载相同数据到全局内存
○ 单线程仅计算一个元素，效率较低 - GPU共享内存矩阵乘法优化 31:43
●
● 优化措施:
○ 使用共享内存减少全局内存访问
○ 采用平铺矩阵算法
○ 定义共享内存：
● 实现流程:
○ 计算线程对应的行列索引
○ 循环处理每个矩阵块
○ 将数据从全局内存拷贝到共享内存
○ 执行同步操作__syncthreads()
○ 计算矩阵块乘积
○ 再次同步后写入结果
○
● 边界处理:
○ 超出矩阵范围的元素置为0
○ 条件判断： - 主函数与编译 32:51
●
● 主要流程:
○ 定义矩阵维度：M=512M=512M=512,K=512K=512K=512,N=512N=512N=512
○ 内存分配与初始化
○ CPU矩阵乘法计时
○ GPU内存分配与数据传输
○ GPU核函数调用与计时
○
● 编译指令: - 运行时间与性能比较 33:53
●
● 实测性能:
○ CPU版本：0.78秒
○ GPU基础版：0.08秒
○ GPU共享内存版：0.006秒
● 性能分析:
○ 内存拷贝和同步操作影响实际加速比
○ 建议单独测量核函数执行时间
六、作业 34:40
●
● 作业1:
○ 测试不同矩阵尺寸(64×64, 512×512)下三种实现的运行时间
○ 填写性能对比表格
● 作业2:
○ 研究BLOCK_SIZE取值对性能的影响
○ 寻找共享内存使用与并行度的最佳平衡点
七、知识小结
| 知识点 | 核心内容 | 考试重点/易混淆点 | 难度系数 |
|---|---|---|---|
| GPU与CPU基础结构对比 | GPU缓存小、计算单元多，适合并行计算；CPU适合复杂分支判断 | GPU适用场景判断标准（内存访问频率、控制复杂度等） | ⭐⭐ |
| CUDA编程模型 | 核函数（__global__修饰）、线程块与网格组织（三维索引） | 核函数参数传递（线程块/网格维度定义） | ⭐⭐⭐ |
| 矩阵乘法基础算法 | 时间复杂度O(mnk)，元素计算独立性强，适合GPU并行 | 行列匹配限制（a列数=b行数） | ⭐⭐ |
| 基础版GPU矩阵乘法 | 每个线程计算一个c[i][j]，全局内存重复读取（2mnk次） | 性能瓶颈（全局内存访问延迟500周期） | ⭐⭐⭐ |
| 共享内存优化 | 静态/动态申请（__shared__关键字）、平铺矩阵分块计算 | 边界同步问题（__syncthreads()使用时机） | ⭐⭐⭐⭐ |
| 平铺矩阵乘法 | 分块计算（block_size²），全局内存访问降至2mnk/block_size | block_size选择平衡（共享内存大小限制 vs 并行效率） | ⭐⭐⭐⭐ |
| 实践对比 | CPU耗时0.78s vs GPU基础版0.08s vs 共享内存版0.006s | 显存拷贝时间影响（需排除cudaMemcpy干扰） | ⭐⭐⭐ |
扫一扫
817
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
