Cuda中Global memory中coalescing例程解释

Global memory是cuda中最常见的存储类型，又叫做Device memory，位于Host主机区域上，它的生命周期是在整个Grid里面，大约具有500个cycle latency。在cuda并行程序中，尽量用Coalesing accessing的策略来最大化带宽bandwidth。什么是Coalesing accessing呢？如图所示：

当半个Warp的16个threads在一次memory transaction中coalesced时，Global memory中的带宽得到了最大的利用。其中，需要注意的是，Device在一次transaction中，从global memory中可以一次读取32-bit，64-bit，128-bit，例如

64 bytes - each thread reads a word: int, float, …

128 bytes - each thread reads a double-word: int2, float2, …

32 bytes (compute capability 1.2+) - each thread reads a short int.

下面有两个实例来说明Global memory中的coalescing问题：

1）float3型Uncoalesced

__global__ void accessFloat3(float3 *d_in,

float3* d_out)

{

int index = blockIdx.x * blockDim.x + threadIdx.x;

float3 a = d_in[index];

a.x += 2;

a.y += 2;

a.z += 2;

d_out[index] = a;

}

在这段代码中，float3有12个bytes，不等于要求的4,8,16 bytes，半个warp读取3个64bytes中非连续区域，如图：

有三种方法可以解决这个问题

1:使用shared memory，也叫做3-step approach

假如每个block中使用256个threads，这样一个thread block需要 sizeof(float3)*256 bytes的share memory空间，每个thread读取3个单独的float型，这实质上是指讲输入定义为float型，在核函数里面讲读取在share memory中的float变量转换为float3型并进行操作，最后再转换成float型输出，如图;

代码如下：

如果不好理解的话，假设我们的blockDim=4，取4个float3型变量，我们会发现，每一个thread中输入操作（输出操作一样）为：

Thread 0：

S_data[0]=g_in[0]; S_data[4]=g_in[4]; S_data[8]=g_in[8];

Thread 1：

S_data[1]=g_in[1]; S_data[5]=g_in[5]; S_data[9]=g_in[9];

Thread 2：

S_data[2]=g_in[2]; S_data[6]=g_in[6]; S_data[10]=g_in[10];

Thread 3：

S_data[3]=g_in[3]; S_data[7]=g_in[7]; S_data[11]=g_in[11];

可以看出，对于每个thread同一时刻（similar step）的数据读入，地址均是连续，这样就达到了coalescing。

2）使用数组的结构体(SOA)来取代结构体的数组(AOS)

3）使用alignment specifiers

__align__(X), where X = 4, 8, or 16

struct __align__(16) {float x; float y; float z; };

尽管这损失了比较多的空间：

2）第二个实例：矩阵转置 Matrix Transpose.

一般做法：Uncoalesced Transpose，GMEM为Global memory

我们发现一般的做法，在写output时，地址是不连续的，即uncoalesced，因此我们利用shared memory存储输入数据，根据转置的关系，来实现coalescing，SMEM为shared memory，如下图：

代码如下：

__global__ void transpose(float *odata, float *idata, int width, int height)

{

__shared__ float block[BLOCK_DIM*BLOCK_DIM];

unsigned int xBlock = blockDim.x * blockIdx.x;

unsigned int yBlock = blockDim.y * blockIdx.y;

unsigned int xIndex = xBlock + threadIdx.x;

unsigned int yIndex = yBlock + threadIdx.y;

unsigned int index_out, index_transpose;

if (xIndex < width && yIndex < height)

{

unsigned int index_in = width * yIndex + xIndex;

unsigned int index_block = threadIdx.y * BLOCK_DIM + threadIdx.x;

block[index_block] = idata[index_in];

index_transpose = threadIdx.x * BLOCK_DIM + threadIdx.y;

index_out = height * (xBlock + threadIdx.y) + yBlock + threadIdx.x;

}

__syncthreads();

if (xIndex < width && yIndex < height)

odata[index_out] = block[index_transpose];

程序的逻辑关系有时还挺绕的，我们以一个4*4矩阵为例，将逻辑关系展示如下：

设dim3 gridDim(4,1), dim3 blockDim(1,4)，以橙色block为例，如输入数据时，将其放入到sharememory中，代码体现在：

unsigned int index_in = width * yIndex + xIndex;

unsigned int index_block = threadIdx.y * BLOCK_DIM + threadIdx.x;

block[index_block] = idata[index_in];

接下来的代码实际上是将block的区域给换了,如左下图所示，block换成了一列四种不同颜色的，最终转置的矩阵如右下图所示，从图示可以看出，最终结果的坐标系Height、Width、blockIdx.x、blockIdx.y均对位变换了，这时我们只需要找threadIdx.x'、threadIdx.y'与threadIdx.x、threadIdx.y之间的关系，其实可以看出，一个block里面的坐标系没有发生变换，则threadIdx.x'=threadIdx.x，threadIdx.y'=threadIdx.y，所以代码如下：

index_transpose = threadIdx.x * BLOCK_DIM + threadIdx.y;

index_out = height * (xBlock + threadIdx.y) + yBlock + threadIdx.x;

odata[index_out] = block[index_transpose];

总体来说，Global memory中coalescing就是保证其在数据读取或者写入时，使用连续的地址，且地址所存储的变量尺寸为32、64、128 bit，我们常常使用share memory来解决coalescing问题。

Cuda中Global memory中coalescing例程解释相关推荐

cuda的global memory介绍
CUDA Memory Model 对于程序员来说,memory可以分为下面两类: Programmable:我们可以灵活操作的部分. Non-programmable:不能操作,由一套自动机制来达到 ...
Share memory中bank conflict问题
Share memory是片上资源,生命周期是整个block中,它的数据读写十分快,有1个cycle latency.在Share memory中,经常存在bank conflict问题,如果没有ba ...
通过设置PYTORCH_CUDA_ALLOC_CONF中的max_split_size_mb解决Pytorch的显存碎片化导致的CUDA:Out Of Memory问题
问题的出现最近在基友的带动下开始投身ai绘画的大潮,于是本地部署了stable diffusion web ui,利用手上的24G显存开始了愉快的跑高分辨率图片之旅.然而某天在用inpaint功能修 ...
PyTorch中“CUDA out of memory”的调试笔记
1 问题描述 "RuntimeError: CUDA out of memory."是PyTorch写作中常见的一种运行错误,这里我们将记录一下调试过程中发现的一些解决方案: 2 ...
在 CUDA C/C++ kernel中使用内存
在 CUDA C/C++ kernel中使用内存如何在主机和设备之间高效地移动数据.本文将讨论如何有效地从内核中访问设备存储器,特别是全局内存 . 在 CUDA 设备上有几种内存,每种内存的作用域 ...
php static 和 global,php中global static和$GLOBALS使用与区别
本文章来总结介绍关于php中global static和$GLOBALS使用与区别有需要了解他们三区别的朋友可参考参考. 1.global在整个页面起作用. 例1 代码如下复制代码 global ...
#define barrier() __asm__ __volatile__(: : :memory) 中的memory是gcc的东西
gcc内嵌汇编简介在内嵌汇编中,可以将C语言表达式指定为汇编指令的操作数,而且不用去管如何将C语言表达式的值读入哪个寄存器,以及如何将计算结果写回C 变量,你只要告诉程序中C语言表达式与汇编指令操作 ...
Php global echo,php中global和$GLOBALS[]的分析之一
PHP 的全局变量和 C 语言有一点点不同,在 C 语言中,全局变量在函数中自动生效,除非被局部变量覆盖这可能引起一些问题,有些人可能漫不经心的改变一个全局变量.PHP 中全局变量在函数中使用时必须 ...
php中global什么意思,php中global与$GLOBAL的用法及区别(转载)
php中global 与 $GLOBALS[""] 差别原本觉得global和$GLOBALS除了写法不一样觉得,其他都一样,可是在实际利用中发现2者的差别还是很大的! 先看下面 ...

Cuda中Global memory中coalescing例程解释

Cuda中Global memory中coalescing例程解释相关推荐

最新文章

热门文章