※ [本文转录自 VideoCard 看板] 作者: a5000ml (咖啡里的海洋蓝) 看板: VideoCard 标题: [分享] CUDA 程式设计(6) -- 记忆体 时间: Wed Oct 22 19:17:12 2008 因为学弟们觉得需要练习, 所以从这单元开始加入习题, 并把一些原来在本文中的内容 移到习题中 ~ ^_^= 第六章 记忆体种类 ◆ 简介 ◆ CUDA 的记忆体种类很繁杂, 依重要顺序排列, 可分为下六种 (1) 暂存器 register (2) 全域记忆体 global memory (3) 共享记忆体 shared memory (5) 常数记忆体 constant memory (4) 材质快取 texture cache (6) 区域记忆体 local memory 其中最重要的前 4 种, 基本的功能如下 (1) 暂存器 : 执行绪使用到的一般变数 (快速, 预设, 执行绪内部) (2) 共享记忆体 : 同一区块内的执行绪共用 (快速, 区块内交换资料用) (3) 常数记忆体 : 存放整个程式共用的常数 (快速, 有快取, 全域共用) (4) 全域记忆体 : 显示卡中的 DRAM (很慢, 无快取, 全域共用) 材质快取是独立於其它记忆体的存取机制, GPU 使用上要透过 API, 但是它有快取, 即使没有做什麽复杂的最佳化, 都可以达到不错的效能, 我们将另辟一个单元讨论它. 最後的区域记忆体是被动产生的, 我们无法在 CUDA 的 C 程式语法中主动控制它 (除非嵌入组合语言 ptx), 而且它对效能的影响是负面的, 要避免它的出现. ◆ 延伸标签 ◆ 共享记忆体、全域记忆体、常数记忆体三者具有延伸标签, 我们可以使用这些标签直接 宣告变数, 并在 GPU 核心中使用传统 C/C++ 的方式进行一般操作, 取值或取址. (1) 全域记忆体 __device__ 在档案范围宣告 (2) 常数记忆体 __constant__ 在档案范围宣告 (3) 共享记忆体 __shared__ 在函式范围宣告 在主机的存取上, 用这些标签宣告的变数不能直接存取, 必需先经过 API 进行位址的 解析, 或直接使用 API 进行存取。(详见操作特性 7) 范例 6.1: 以下程式码会将阵列 p[] 周期性地填成 0,1,2,0,1,2,0,1,2,0,1,2,.... __device__ int a[3]={0,1,2}; //在 file scope 宣告 __global__ void set_array(int* p, int n){ int j=blockIdx.x*blockDim.x+threadIdx.x; int m=gridDim.x*blockDim.x; for(int k=j; k<n; k+=m){ p[k]=a[k%3]; //周期性地填入 } } 范例 6.2: 同样的程式使用常数快取来加速 __constant__ int a[3]={0,1,2}; //在 file scope 宣告 __global__ void set_array(int* p, int n){ int j=blockIdx.x*blockDim.x+threadIdx.x; int m=gridDim.x*blockDim.x; for(int k=j; k<n; k+=m){ p[k]=a[k%3]; } } 在两个范例中的指位器 p 指向的是全域记忆体, 位在显示卡的 DRAM 上. ◆ 共享记忆体 ◆ 共享记忆体特殊之处在於它的共用范围是整个区块, 会产生这种设计的原因在於效能　 和硬体成本上的考量, 其用途很广 (1) 执行绪之间快速资料交换. (基本功能) (2) 存放会多次使用资料, 当做可控制的快取. (3) 存放区块中的共用变数, 减少暂存器的使用量. (4) 做为全域记忆体的合并读取缓冲 (coalesced read) (5) 做为资料顺序动态重整单元, 避免区域记忆体的介入. 使用时必需配合 __syncthreads() 对执行绪做同步化, 以确保资料已存妥, 避免後续 其它执行绪因不同步而读取到尚未写入的记忆体位址, 造成读取上的错误。 范例 6.3: 一维阵列的区域平均 (单一区块), 未使用共享记忆体 #define BLOCK_DIM 10 __global__ void local_average_1(float *r, float *a){ int j=threadIdx.x; //j 为执行绪索引 if(j==0){ r[j]=(2*a[j]+a[j+1])/4; //左边界=0 } else if(j==BLOCK_DIM-1){ r[j]=(a[j-1]+2*a[j])/4; //右边界=0 } else{ r[j]=(a[j-1]+2*a[j]+a[j+1])/4; //输出加权平均 } } local_average_1<<<1,BLOCK_DIM>>> (r, a); Note: 它的缺点主要在於要计算一个 r[j], 全域记忆体要读取 3 次, a[j-1], a[j], a[j+1], 在边界点上要做条件分支, 且分支里面 带的程式码延迟很重. 范例 6.4: 一维阵列的区域平均 (单一区块), 使用共享记忆体 #define BLOCK_DIM 10 __global__ void local_average_2(float *r, float *a){ int j=threadIdx.x; //j 为执行绪索引 __shared__ float s[BLOCK_DIM+2]; //宣告共享记忆体 s[j+1]=a[j]; //多执行绪一起将资料载入共享记忆体 //使用 +1 的偏移, 0 和 BLOCK_DIM+1 //两点做为阵列边界 if(j==0){ //只用一个执行绪设定边界值 s[0]=s[BLOCK_DIM+1]=0; } __syncthreads(); //同步化, 确保资料己存好 r[j]=(s[j]+2*s[j+1]+s[j+2])/4; //输出加权平均 } local_average_2<<<1,BLOCK_DIM>>> (r, a); 使用乱数测试的结果为 +------+------+------+------+------+------+------+------+------+------+ | 7 | 8 | 6 | 3 | 3 | 9 | 8 | 5 | 9 | 7 | +------+------+------+------+------+------+------+------+------+------+ | 5.5 | 7.25 | 5.75 | 3.75 | 4.5 | 7.25 | 7.5 | 6.75 | 7.5 | 5.75 | +------+------+------+------+------+------+------+------+------+------+ Note: 其中变数重用度为 3, 亦即要计算 1 个 r[j], 全域记忆体只要读取 1 次, 因为计算很少, 效能主要限制於记忆体, 所以加速差不多是 3x. ◆ 操作特性 ◆ (1) 暂存器是所有记忆体中最快的, 执行绪中大部份的局部变数都预设使用暂存器, 包括阵列(array)也是, 但有些情况下, 它会被编译器以较慢的区域记忆体取代, 这些情况包括在执行绪中「同时占用」过多变数, 以致於使用的暂存器数目 超过编译器的限制 (可使用 --maxrregcount N 选项来限制, 预设 N=128 ), 或是使用动态变数做为索引存取阵列 (因为此时需要引入阵列的顺序结构). Note:「同时占用」是指在某一时间点上, 在後续还会被用到的变数才称为占用, 因为编译器会自动最佳化, 暂存器的使用是自动调配的, 程式码中宣告的 变数个数并不等於使用的暂存器数 (单一变数并非直接对应於单一暂存器), 可使用 --ptxas-options=-v 选项观察, 会发现有时候明明宣告很多变数, 结果暂存器使用到的却很少, 或是稍微调整一下变数使用的次序就会改变 暂存器使用数目. (2) 共享记忆体的应用范围是区块, 只有同一区块里的执行绪才可以共用它, 使用上要 注意它的大小限制, 以及存取前後要对执行绪同步化, 避免因资料读写的先後顺序 错误而导致不可预期的资料错乱, 其效能仅次於暂存器, 是最佳化的一个重点项目. 另外, 我们无法在设计时期对它进行初始化, 必需等到核心执行时期才能设定它, 而且也无法在主机中对它进行存取, 现阶段 CUDA 只提供 API 指定其大小. (3) 区域记忆体是暂存器不够用的时候, 编译器自动将资料置换到全域记忆体的产物, 有点像作业系统的页置换 (page swap), 它对效能的影响是负面的, 而且非常地 难以捉摸, 所以在最佳化程式的时候, 时常要用编译器选项 --ptxas-options=-v 来追踪它, 深怕一不小心它就蹦出来, 但有时候为了加大区块中的执行绪数目 (blockDim), 必需使用编译器选项 --maxrregcount N 来限制执行绪最大暂存器 使用量, 没有它又无法达成这种限制 (同时占用的变数就这麽多, 一定要置换出去), 两者之间往往必需进行妥协, 或使用共享记忆体进行手动置换. (4) 全域记忆体除了使用 __device__ 标签宣告, 另外直接透过 cudaMalloc() 等 API 直接配置的记忆体也算,「全域」顾名思义就是所有执行单元都可以对它进行操作, 所以凡是放在显示卡的 DRAM 中, 能够被所有执行绪操作 (包括在不同区块中的 执行绪) 皆称为全域记忆体. 它读取写入是没有经过快取的, 跟材质快取在硬体上 属於不同的 port, 必需配合「合并读取」(coalesced read, 也就是半个 warp 的 执行绪同时读取记忆体中的连续区块, 使记忆体控制器做一次性的合并发出) 如此 才能增进其效能 (约 5~10x), 这部份将在後面单元中详细介绍. (5) 常数记忆体虽然和材质快取属於同一个层次, 但因其大小受到限制, 使快取失误率 非常低, 所以在执行时期除了第一次使用需要载入时间外, 之後使用和共享记忆体 一样快, 它在核心的存取是唯读的, 只能在档案中使用初始值的方法设定, 或是在 主机中透过 API 进行存取, 使用范围是全域性的. (6) 材质快取因为有快取做为缓冲, 所以读取上不需要做合并, 但也因此比直接存取 全域记忆体稍微慢一点 (数个到数十个周期) , 但仍比「未合并读取」全域记忆体 快上甚多, 所以如果在「合并读取」很复杂的情况下, 使用材质快取是不错的选择. 材质快取是唯读的, 在快取的区域性上, 除了传统微处理机的 1D 快取模式外, 因绘图需求的缘故, CUDA 亦提供 2D 和 3D 的材质快取, 使用范围是全域性的, 这部份也在後面单元中详细介绍. (7) 所有的标签 __shared__, __device__, __constant__ 宣告的变数所对应的位址 只有在核心中能直接使用, CUDA 将这些变数称为 Symbol, 在主机中不能直接以 C/C++ 原生的方式处理 (进行取值或取址), 必需透过 API a. __device__ 必需先透过 cudaGetSymbolAddress() 取得位址, 然後才可以呼叫 其它的主机 API, 例如 cudaMemcpy() 等或丢给其它 kernel 进行操作, 取得的 位址和 cudaMalloc() 所配置的位址地位相等, 而且可在主机中进行位址偏移 (offset, 也就是 ptr+offset 这种动作), 它对映的是显示卡中实体的位址, 另外亦可透过和 __constant__ 同样的方式存取. b. __constant__ 只能透过 cudaMemcpyToSymbol() 和 cudaMemcpyFromSymbol() 进行存取. (note: cudaGetSymbolAddress() 不能用) c. __shared__ 主机无法直接存取, 只能设定其大小. (8) 这些记忆体拥有各自的位址空间, 而且载入/储存的指令不一样 (see ptx code), 所以指位器不能混用 (例如无法使用动态的条件来切换同一个指位器, 让它可以 在一些条件下指向共享记忆体, 而在其它条件下指向全域记忆体, 这会造成编译器 无法由本文关系解析指位器对应的是何种记忆体, 对後续程式进行编码). ◆ 特性表 ◆ 以下是这些记忆体的特性表, 资料 compute 1.0~1.3 通用 +-----------+--------------+--------+-------------+----------+--------------+ |种类 | 延伸标签 | 生命期 | 存取范围 | 主机存取 | 硬体配置 | +-----------+--------------+--------+-------------+----------+--------------+ |暂存器 | (无) | 区块 | 执行绪 R/W | | on chip | |区域记忆体 | (无) | 区块 | 执行绪 R/W | | DRAM | +-----------+--------------+--------+-------------+----------+--------------+ |共享记忆体 | __shared__ | 区块 | 区块 R/W | size | on chip | +-----------+--------------+--------+-------------+----------+--------------+ |材质快取 | (无) | 程式 | 全域 R(API) | R/W (API)| DRAM + cache | |常数记忆体 | __constant__ | 程式 | 全域 R | R/W (API)| DRAM + cache | |全域记忆体 | __device__ | 程式 | 全域 R/W | R/W (API)| DRAM | +-----------+--------------+--------+-------------+----------+--------------+ 表 6.1 记忆体种类 (按照存取范围划分) +-----------+------------------+--------------------------------+ |种类 | 存取时间(clocks) | 影响效能因素 | +-----------+------------------+--------------------------------+ |暂存器 | immediate | | |共享记忆体 | 4 | 记忆库冲突 (16KB/MP, 16 banks) | +-----------+------------------+--------------------------------+ |常数记忆体 | 4,(失误)400~600 | 快取失误 (cache 8KB/MP) | |材质快取 | 4,(失误)400~600 | 快取失误 (cache 6~8KB/MP) | +-----------+------------------+--------------------------------+ |区域记忆体 | 400~600 | 不易控制 (compiler auto) | |全域记忆体 | 400~600 | 记忆库冲突, 未合并读取 | +-----------+------------------+--------------------------------+ 表 6.2 记忆体存取速度 (按照存取速度划分) +--------------+------------+----------+----------+--------+-----------+ | 标签 | 大小限制 | 使用范围 | 变数种类 | 初始值 | 外部变数 | +--------------+------------+----------+----------+--------+-----------+ | __shared__ | 16KB/block | block | static | X | no extern | | __device__ | | file | | O | no extern | | __constant__ | 64KB | file | static | O | no extern | +--------------+------------+----------+----------+--------+-----------+ 表 6.3 标签特性 ◆ 练习题 ◆ 练习 6.1: 将范例 6.3 和 6.4 改为多区块的版本, 测试大阵列并比较效能。 练习 6.2: 试用共享记忆体做矩阵乘法和 transpose，并比较未使用共享记忆体的效能。 练习 6.3: 使用 cudaMemcpyFromSymbol() 和 cudaMemcpyToSymbol() 在执行时期修改 范例 6.1 和 6.2 的 __device__ 和 __constant__ 阵列的内容, 并使用 cudaGetSymbolAddress() 和 cudaMemcpy() 做同样的事情, __constant__ 都能成功吗？若否, 请使用 cudaGetErrorString() 秀出错误讯息。 练习 6.4: 使用编译器选项 --maxrregcount N 和 --ptxas-options=-v 观察和调整 练习 6.2 的暂存器与其它记忆体的使用量, 并测试对程式效能的影响. 练习 6.5: 测试使用动态的条件切换同一个指位器, 让它可以在某些条件下指向 共享记忆体, 在其它条件下指向全域记忆体, 例如以下两段单一区块 的程式码, 执行的结果何者正确？ ------------------------------------------------------- 函式功能: (使用单一区块, 部份平方) k=0~99 r[k] = (k<n) ? a[k] : a[k]*a[k]; ------------------------------------------------------- __global__ void part_square_1(int *r, int *a, int n){ int j=threadIdx.x; //j 为执行绪索引 __shared__ int s[100]; //宣告共享记忆体 s[j]=a[j]*a[j]; //先将每个数的平方存入共享记忆体中 __syncthreads(); //同步化, 确保资料己存好 int* p; //使用共同指位器 (混合指向) if(j<n){ p=a+j; //指向全域记忆体 } else{ p=s+j; //指向共享记忆体 } r[j]=*p; //指位器取值 } ------------------------------------------------------- __global__ void part_square_2(int *r, int *a, int n){ int j=threadIdx.x; __shared__ int s[100]; s[j]=a[j]*a[j]; __syncthreads(); if(j<n){ int* p; //使用不同指位器 p=a+j; //指向全域记忆体 r[j]=*p; } else{ int* p; //使用不同指位器 p=s+j; //指向共享记忆体 r[j]=*p; } } ------------------------------------------------------- 主机呼叫方式 part_square_1<<<1,100>>>(r,a,10); part_square_2<<<1,100>>>(r,a,10); ------------------------------------------------------- --

※ 发信站: 批踢踢实业坊(ptt.cc) ◆ From: 114.45.208.109 ※ 编辑: a5000ml 来自: 114.45.208.109 (10/22 19:22) --

※ 发信站: 批踢踢实业坊(ptt.cc) ◆ From: 114.45.208.109