CUDA C编程10：内存访问模式之全局内存读取

系列文章目录

---

文章目录

• 系列文章目录
• 前言
• 
  

  一、内存访问模式之全局内存读取

• • 1. 内存访问模式基础知识
  • • 1.1 对齐和合并访问
• 
  

  二、全局内存读取示例

• • 1.代码实现
  • 2.NVIDIA Visual Profiler运行结果分析
• 总结
• 参考资料

---

前言

之前在复习现代C++的新特性，没有继续CUDA C编程的学习，今天开始继续之前的学习，这里跟大家分享内存访问模式中全局内存读取的知识。

---

一、内存访问模式之全局内存读取 1. 内存访问模式基础知识

我们所编写的GPU程序容易受到内存带宽的限制，因此，最大限度利用全局内存带宽是调控核函数性能的基本。

为了达到读写数据时的最佳性能，内存访问 *** 作必须满足一定条件。

CUDA程序的显著特征之一就是指令必须以线程束为单位进行发布和执行。

1.1 对齐和合并访问

全局内存通过缓存来加载或者存储，注意，全局内存是逻辑内存空间，可通过核函数（ _ _ g l o b a l _ _ \_\_global\_\_ __global__关键字）访问。

应用程序数据最初存于DRAM上（物理设备内存），核函数的内存请求通常在DRAM设备和片上内存以128字节或32字节内存事务来实现。

所有对全局内存访问都通过二级缓存，许多访问会通过一级缓存，这取决于访问类型和GPU架构。

如果两级缓存都被使用，则内存访问是由128字节内存事务实现；如果仅使用二级缓存，那么内存访问由32字节的内存事务实现。

优化应用程序时，关注两个特性：
（1）对齐内存访问
当设备内存事务的第一个地址是缓存粒度的偶数倍时（32字节的二级缓存或128字节的一级缓存），就会出现对齐内存访问，运行非对齐的加载会造成带宽的浪费。

（2）合并内存访问
当一个线程束的全部32个线程访问一个连续的内存块时，就会出现合并内存访问。

对齐合并内存访问的最理想状态是线程束从对齐内存地址开始访问一个连续的内存块，这种情况下，只需要一个128字节内存事务从设备内存中读取数据，如下图所示：
对非对齐和未合并的内存访问，可能需要3个128字节的内存事务从设备内存中读取数据，这将造成带宽浪费，如下图所示：
### 1.2 全局内存读取
在SM中，有三种缓存路径进行传输：
（1）一级和二级缓存（默认路径）
（2）常量缓存
（3）只读缓存

可通过如下编译指令禁用一级缓存：

-xptxas-dlcm=cg

如果一级缓存被禁用，所有对全局内存的加载请求将直接进入二级缓存；如果二级缓存缺失，则由DRAM完成请求；

可通过如下编译指令开启一级缓存：

-xptxas-dlcm=ca

二、全局内存读取示例

这里使用偏移量来读取内存数据，并验证对齐与非对齐对全局内存读取效率的影响。

1.代码实现

#include 
#include 
#include 
#include 
#include 
#include 
#include "CudaUtils.h"

void InitData(float* data, size_t nElem)
{
	for (size_t i = 0; i < nElem; i++)
	{
		data[i] = i % 255;
	}
}

void CheckResults(float* hostRef, float* gpuRef, size_t nElem)
{
	bool bSame = true;
	for (size_t i = 0; i < nElem; i++)
	{
		if (abs(gpuRef[i] - hostRef[i]) > 1e-5)
		{
			bSame = false;
		}
	}

	if (bSame)
	{
		printf("Result is correct!\n");
	}
	else
	{
		printf("Result is error!\n");
	}
}

__global__ void readOffset(float* A, float* B, float* C,
	const int n, const int offset)
{
	unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
	unsigned int k = i + offset;
	if (k < n)
		C[i] = A[k] + B[k];
}

//check the result of kernel function
void sumArrayOnHost(float* A, float* B, float* C,
	const int n, const int offset)
{
	for (int idx = offset, k = 0; idx < n; idx++, k++)
	{
		C[k] = A[idx] + B[idx];
	}
}

int main(int argc, char** argv)
{
	int dev = 0;
	cudaDeviceProp deviceProp;
	cudaGetDeviceProperties(&deviceProp, dev);
	printf("%s starting reduction at ", argv[0]);
	printf("device %d： %s ", dev, deviceProp.name);
	cudaSetDevice(dev);

	int nElem = 1 << 20;
	printf(" with array size %d\n", nElem);
	size_t nBytes = nElem * sizeof(float);

	int blockSize = 512;
	int offset = 0;
	if (argc > 1)
		offset = atoi(argv[1]);
	if (argc > 2)
		blockSize = atoi(argv[2]);

	dim3 block(blockSize, 1);
	dim3 grid((nElem + block.x - 1) / block.x, 1);

	//alloc host memory
	float* h_A = (float*)malloc(nBytes);
	float* h_B = (float*)malloc(nBytes);
	float* hostRef = (float*)malloc(nBytes);
	float* gpuRef = (float*)malloc(nBytes);

	InitData(h_A, nElem);
	memcpy(h_B, h_A, nBytes);
	
	//summary at host side
	sumArrayOnHost(h_A, h_B, hostRef, nElem, offset);

	//alloc device memory
	float* d_A, *d_B, *d_C;
	cudaMalloc(&d_A, nBytes);
	cudaMalloc(&d_B, nBytes);
	cudaMalloc(&d_C, nBytes);

	//copy data from cpu to device
	cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);
	cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);

	//kernel 1: warm up
	CudaUtils::Time::Start();
	readOffset<<<grid, block>>>(d_A, d_B, d_C, nElem, offset);
	cudaDeviceSynchronize();
	CudaUtils::Time::End();
	double iElaps = CudaUtils::Time::Duration<CudaUtils::Time::TIME_UNIT::MS>();
	printf("warmup <<<%4d, %4d>>> offset %4d elapsed %.3f ms\n", grid.x, block.x, offset, iElaps);

	CudaUtils::Time::Start();
	readOffset << <grid, block >> > (d_A, d_B, d_C, nElem, offset);
	cudaDeviceSynchronize();
	CudaUtils::Time::End();
	iElaps = CudaUtils::Time::Duration<CudaUtils::Time::TIME_UNIT::MS>();
	printf("readOffset <<<%4d, %4d>>> offset %4d elapsed %.3f ms\n", grid.x, block.x, offset, iElaps);

	cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);
	CheckResults(hostRef, gpuRef, nElem - offset);

	cudaFree(d_A);
	cudaFree(d_B);
	cudaFree(d_C);
	free(h_A);
	free(h_B);
	free(hostRef);
	free(gpuRef);

	cudaDeviceReset();



	system("pause");
	return 0;
}

2.NVIDIA Visual Profiler运行结果分析

offset 0:

offset 11:

offset 128:
从上述Global Memory Load Efficiency结果可知，偏移量11（非对齐内存访问）的内存访问效率比偏移量0或128（对齐内存访问）低。

---

总结

提高CUDA程序的运行效率，提高带宽利用率很重要，提高内存访问效率可以提高带宽利用率。

参考资料

《CUDA C编程权威指南》