CUDA ---- Constant Memory

CONSTANT  MEMORY

constant Memory对于device来说只读但是对于host是可读可写。constant Memory和global Memory一样都位于DRAM,并且有一个独立的on-chip cache,比直接从constant Memory读取要快得多。每个SM上constant Memory大小限制为64KB。

constant Memory的获取方式不同于其它的GPU内存,对于constant Memory来说,最佳获取方式是warp中的32个thread获取constant Memory中的同一个地址。如果获取的地址不同的话,只能串行的服务这些获取请求了。

constant Memory使用__constant__限定符修饰变量。

constantMemory的生命周期伴随整个应用程序,并且可以被同一个grid中的thread和host中调用的API获取。因为constant Memory对device来说是可读的,所以只能在host初始化,使用下面的API:

cudaError_t cudaMemcpyToSymbol(const void *symbol, const void * src, size_t count, size_t offset, cudaMemcpyKind kind)

Implementing a 1D Stencil with Constant Memory

实现一个1维Stencil(数值分析领域的东,卷积神经网络处理图像的时候那个stencil),简单说就是计算一个多项式,系数放到constant Memory中,即y=f(x)这种东西,输入是九个点,如下:

{x − 4h, x − 3h, x − 2h, x − h, x, x + h, x + 2h, x + 3h, x + 4h}

在内存中的过程如下:

公式如下:

那么要放到constant Memory中的便是其中的c0、c1、c2 ……

因为每个thread使用九个点来计算一个点,所以可以使用shared memory来降低延迟。

__shared__ float smem[BDIM + 2 * RADIUS];

RADIUS定义了x两边点的个数,对于本例,RADIUS就是4。如下图所示,每个block需要RADIUS=4个halo(晕)左右边界:

#pragma unroll用来告诉编译器,自动展开循环。

__global__ void stencil_1d(float *in, float *out) {
// shared memory
__shared__ float smem[BDIM + 2*RADIUS];
// index to global memory
int idx = threadIdx.x + blockIdx.x * blockDim.x;
// index to shared memory for stencil calculatioin
int sidx = threadIdx.x + RADIUS;
// Read data from global memory into shared memory
smem[sidx] = in[idx];
// read halo part to shared memory
if (threadIdx.x < RADIUS) {
smem[sidx - RADIUS] = in[idx - RADIUS];
smem[sidx + BDIM] = in[idx + BDIM];
}
// Synchronize (ensure all the data is available)
__syncthreads();
// Apply the stencil
float tmp = 0.0f;
#pragma unroll
for (int i = 1; i <= RADIUS; i++) {
tmp += coef[i] * (smem[sidx+i] - smem[sidx-i]);
}
// Store the result
out[idx] = tmp;
}

Comparing with the Read-only Cache

Kepler系列的GPU允许使用texture pipeline作为一个global Memory只读缓存。因为这是一个独立的使用单独的带宽的只读缓存,所以对带宽限制的kernel性能有很大的提升。

Kepler的每个SM有48KB大小的只读缓存,一般来说,在读地址比较分散的情况下,这个只读缓存比L1表现要好,但是在读同一个地址的时候,一般不适用这个只读缓存,只读缓存的读取粒度为32比特。

有两种方式来使用只读缓存:

  • 使用__ldg限定
  • 指定特定global Memory称为只读缓存

下面代码片段对于第一种情况:

__global__ void kernel(float* output, float* input) {
    ...
    output[idx] += __ldg(&input[idx]);
    ...
}

下面代码对应第二种情况,使用__restrict__来指定该数据的要从只读缓存中获取:

void kernel(float* output, const float* __restrict__ input) {
    ...
    output[idx] += input[idx];
}

一般使用__ldg是更好的选择。通过constant缓存存储的数据必须相对较小而且必须获取同一个地址以便获取最佳性能,相反,只读缓存则可以存放较大的数据,且不必地址一致。

下面的代码是之前stencil的翻版,使用过了只读缓存来存储系数,二者唯一的不同就是函数的声明:

__global__ void stencil_1d_read_only (float* in, float* out, const float *__restrict__ dcoef) {
// shared memory
__shared__ float smem[BDIM + 2*RADIUS];
// index to global memory
int idx = threadIdx.x + blockIdx.x * blockDim.x;
// index to shared memory for stencil calculatioin
int sidx = threadIdx.x + RADIUS;
// Read data from global memory into shared memory
smem[sidx] = in[idx];
// read halo part to shared memory
if (threadIdx.x < RADIUS) {
smem[sidx - RADIUS] = in[idx - RADIUS];
smem[sidx + BDIM] = in[idx + BDIM];
}
// Synchronize (ensure all the data is available)
__syncthreads();
// Apply the stencil
float tmp = 0.0f;
#pragma unroll
for (int i=1; i<=RADIUS; i++) {
tmp += dcoef[i]*(smem[sidx+i]-smem[sidx-i]);
}
// Store the result
out[idx] = tmp;
}

由于系数原本是存放在global Memory中的,然后读进缓存,所以在调用kernel之前,我们必须分配和初始化global Memory来存储系数,代码如下:

const float h_coef[] = {a0, a1, a2, a3, a4};
cudaMalloc((float**)&d_coef, (RADIUS + 1) * sizeof(float));
cudaMemcpy(d_coef, h_coef, (RADIUS + 1) * sizeof(float), cudaMemcpyHostToDevice);

下面是运行在TeslaK40上的结果,从中可知,使用只读缓存性能较差。

Tesla K40c array size: 16777216 (grid, block) 524288,32
3.4517ms stencil_1d(float*, float*)
3.6816ms stencil_1d_read_only(float*, float*, float const *)

总的来说,constant缓存和只读缓存对于device来说,都是只读的。二者都有大小限制,前者每个SM只能有64KB,后者则是48KB。对于读同一个地址,constant缓存表现好,只读缓存则对地址较分散的情况表现好。

The Warp Shuffle Instruction

之前我们有介绍shared Memory对于提高性能的好处,在CC3.0以上,支持了shuffle指令,允许thread直接读其他thread的寄存器值,只要两个thread在 同一个warp中,这种比通过shared Memory进行thread间的通讯效果更好,latency更低,同时也不消耗额外的内存资源来执行数据交换。

这里介绍warp中的一个概念lane,一个lane就是一个warp中的一个thread,每个lane在同一个warp中由lane索引唯一确定,因此其范围为[0,31]。在一个一维的block中,可以通过下面两个公式计算索引:

laneID = threadIdx.x % 32

warpID = threadIdx.x / 32

例如,在同一个block中的thread1和33拥有相同的lane索引1。

Variants of the Warp Shuffle Instruction

有两种设置shuffle的指令:一种针对整型变量,另一种针对浮点型变量。每种设置都包含四种shuffle指令变量。为了交换整型变量,使用过如下函数:

int __shfl(int var, int srcLane, int width=warpSize);

该函数的作用是将var的值返回给同一个warp中lane索引为srcLane的thread。可选参数width可以设置为2的n次幂,n属于[1,5]。

eg:如果shuffle指令如下:

int y = shfl(x, 3, 16);

则,thread0到thread15会获取thread3的数据x,thread16到thread31会从thread19获取数据x。

当传送到shfl的lane索引相同时,该指令会执行一次广播操作,如下所示:

另一种使用shuffle的形式如下:

int __shfl_up(int var, unsigned int delta, int width=warpSize)

该函数通过使用调用方的thread的lane索引减去delta来计算源thread的lane索引。这样源thread的相应数据就会返回给调用方,这样,warp中最开始delta个的thread不会改变,如下所示:

第三种shuffle指令形式如下:

int __shfl_down(int var, unsigned int delta, int width=warpSize)

该格式是相对__shfl_down来说的,具体形式如下图所示:

最后一种shuffle指令格式如下:

int __shfl_xor(int var, int laneMask, int width=warpSize)

这次不是加减操作,而是同laneMask做抑或操作,具体形式如下图所示:

所有这些提及的shuffle函数也都支持单精度浮点值,只需要将int换成float就行,除此外,和整型的使用方法完全一样。

参考书:《professional cuda c programming》

NVIDIA CUDA板块:https://developer.nvidia.com/cuda-zone

CUDA在线文档:http://docs.nvidia.com/cuda/index.html#

时间: 2024-10-13 18:58:53

CUDA ---- Constant Memory的相关文章

CUDA ---- Shared Memory

CUDA SHARED MEMORY shared memory在之前的博文有些介绍,这部分会专门讲解其内容.在global Memory部分,数据对齐和连续是很重要的话题,当使用L1的时候,对齐问题可以忽略,但是非连续的获取内存依然会降低性能.依赖于算法本质,某些情况下,非连续访问是不可避免的.使用shared memory是另一种提高性能的方式. GPU上的memory有两种: · On-board memory · On-chip memory global memory就是一块很大的on

CUDA ---- Memory Model

Memory kernel性能高低是不能单纯的从warp的执行上来解释的.比如之前博文涉及到的,将block的维度设置为warp大小的一半会导致load efficiency降低,这个问题无法用warp的调度或者并行性来解释.根本原因是获取global memory的方式很差劲. 众所周知,memory的操作在讲求效率的语言中占有极重的地位.low-latency和high-bandwidth是高性能的理想情况.但是购买拥有大容量,高性能的memory是不现实的,或者不经济的.因此,我们就要尽量

并行程序设计---cuda memory

CUDA存储器模型: GPU片内:register,shared memory: host 内存: host memory, pinned memory. 板载显存:local memory,constant memory, texture memory, texture memory,global memory; register: 訪问延迟极低: 基本单元:register file (32bit/each) 计算能力1.0/1.1版本号硬件:8192/SM. 计算能力1.2/1.3版本号硬

CUDA memory

原文链接 CUDA存储器类型: 每个线程拥有自己的register and loacal memory; 每个线程块拥有一块shared memory; 所有线程都可以访问global memory; 还有,可以被所有线程访问的只读存储器:constant memory and texture memory 1.  寄存器Register 寄存器是GPU上的高速缓存器,其基本单元是寄存器文件,每个寄存器文件大小为32bit. Kernel中的局部(简单类型)变量第一选择是被分配到Register

【Cuda并行编程之二】Cuda Memory Hierarchy_Cuda内存层次结构

要想编写高效的程序,那么一定要对内存结构有比较深刻的认识,就像C/C++里面的堆内存,栈内存,全局存储区,静态存储区,常量区等.Cuda是并行计算框架,而GPU的内存有限,那么如果想编写高效的Cuda程序,首先要对其内存结构有一个简单的认识. 首先我们先上一张图,然后通过解释一些名词和代码来进行解释. 各种存储器比较: 存储器  位置 拥有缓存 访问权限 变量生存周期 register GPU片内 N/A device可读/写 与thread相同 local memory 板载显存 无 devi

CUDA系列学习(二)CUDA memory &amp; variables

本文来介绍CUDA的memory和变量存放,分为以下章节: (一).CPU Memory 结构 (二).GPU Memory结构 (三).CUDA Context (四).kernel设计 (五).变量 & Memory 5.1 global arrays 5.2 global variables 5.3 Constant variables 5.4 Register 5.5 Local Array 5.6 Shared Memory 5.7 Texture Memory 5.8 总结 (一).

CUDA内存类型memory

分享一下我老师大神的人工智能教程吧.零基础!通俗易懂!风趣幽默!还带黄段子!希望你也加入到我们人工智能的队伍中来!http://www.captainbed.net http://www.cnblogs.com/traceorigin/archive/2013/04/11/3015482.html CUDA存储器类型: 每个线程拥有自己的register and loacal memory; 每个线程块拥有一块shared memory; 所有线程都可以访问global memory; 还有,可

CUDA从入门到精通

CUDA从入门到精通(零):写在前面 在老板的要求下,本博主从2012年上高性能计算课程开始接触CUDA编程,随后将该技术应用到了实际项目中,使处理程序加速超过1K,可见基于图形显示器的并行计算对于追求速度的应用来说无疑是一个理想的选择.还有不到一年毕业,怕是毕业后这些技术也就随毕业而去,准备这个暑假开辟一个CUDA专栏,从入门到精通,步步为营,顺便分享设计的一些经验教训,希望能给学习CUDA的童鞋提供一定指导.个人能力所及,错误难免,欢迎讨论. PS:申请专栏好像需要先发原创帖超过15篇...

Caffe + Ubuntu 14.04 64bit + CUDA 6.5 配置说明2

1. 安装build-essentials 安装开发所需要的一些基本包 sudo apt-get install build-essential 2. 安装NVIDIA驱动 (3.4.0) 2.1 准备工作(2014-12-03更新) 在关闭桌面管理 lightdm 的情况下安装驱动似乎可以实现Intel 核芯显卡 来显示 + NVIDIA 显卡来计算.具体步骤如下: 1. 首先在BIOS设置里选择用Intel显卡来显示或作为主要显示设备 2. 进入Ubuntu, 按 ctrl+alt+F1 进