做网站在哪里找客户,成都锦江规划建设局网站,郑州市网站建设,如何建设网站服务器转载至https://www.cnblogs.com/1024incn/tag/CUDA/ CONSTANT MEMORY
constant Memory对于device来说只读但是对于host是可读可写。constant Memory和global Memory一样都位于DRAM#xff0c;并且有一个独立的on-chip cache#xff0c;比直接从constant Memory读取要快得多…转载至https://www.cnblogs.com/1024incn/tag/CUDA/ CONSTANT MEMORY
constant Memory对于device来说只读但是对于host是可读可写。constant Memory和global Memory一样都位于DRAM并且有一个独立的on-chip cache比直接从constant Memory读取要快得多。每个SM上constant Memory cache大小限制为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中即yf(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需要RADIUS4个halo晕左右边界 #pragma unroll用来告诉编译器自动展开循环。 View Code
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的翻版使用过了只读缓存来存储系数二者唯一的不同就是函数的声明 View Code
由于系数原本是存放在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 widthwarpSize);
该函数的作用是将var的值返回给同一个warp中lane索引为srcLane的thread。可选参数width可以设置为2的n次幂n属于[1,5]。
eg如果shuffle指令如下
int y shfl(x, 3, 16);
则thread0到thread15会获取thread3的数据xthread16到thread31会从thread19获取数据x。
当传送到shfl的lane索引相同时该指令会执行一次广播操作如下所示 另一种使用shuffle的形式如下
int __shfl_up(int var, unsigned int delta, int widthwarpSize)
该函数通过使用调用方的thread的lane索引减去delta来计算源thread的lane索引。这样源thread的相应数据就会返回给调用方这样warp中最开始delta个的thread不会改变如下所示 第三种shuffle指令形式如下
int __shfl_down(int var, unsigned int delta, int widthwarpSize)
该格式是相对__shfl_down来说的具体形式如下图所示 最后一种shuffle指令格式如下
int __shfl_xor(int var, int laneMask, int widthwarpSize)
这次不是加减操作而是同laneMask做抑或操作具体形式如下图所示 所有这些提及的shuffle函数也都支持单精度浮点值只需要将int换成float就行除此外和整型的使用方法完全一样。
