CUDA库之NPP(四):内存开辟和字节对齐

---

博主将在 编程语言|CUDA入门中不定期更新NPP库相关知识

---

---

一、前言

本文主要利用nppiMalloc 来开辟一块内存，并简单探讨npp中的字节对齐问题：

本文只是Demo演示，不考虑返回码的参数检查，也不考虑Free函数来释放内存。

二、nppiMalloc系列函数

npp的内存开辟函数定义在 nppi_support_functions.h中，

博主在这里记录过npp函数的命名的一些规则，还不清楚的，可以直接跳转。

CUDA库之NPP入门(一)：NVIDIA 2D Image and Signal Processing Performance Primitives：

这里以 nppiMalloc_8u_C1 和 nppiMalloc_32fc_C3函数为例：

nppiMalloc_8u_C1：返回一个8bit unsigned char的内存空间， 且通道数为1；
nppiMalloc_32f_C3： 返回一个32bit float的内存空间，通道数为3；

见函数定义

/**
 * 8-bit unsigned image memory allocator.
 * \param nWidthPixels Image width.
 * \param nHeightPixels Image height.
 * \param pStepBytes \ref line_step.
 * \return Pointer to new image data.
 */
Npp8u  * 
nppiMalloc_8u_C1(int nWidthPixels, int nHeightPixels, int * pStepBytes);

/**
 * 3 channel 32-bit floating point image memory allocator.
 * \param nWidthPixels Image width.
 * \param nHeightPixels Image height.
 * \param pStepBytes \ref line_step.
 * \return Pointer to new image data.
 */
Npp32f * 
nppiMalloc_32f_C3(int nWidthPixels, int nHeightPixels, int * pStepBytes);

• 参数 nWidthPixels ： 图像宽度，以像素为单位
• 参数 nHeightPixels ： 图像高度度，以像素为单位
• 参数 pStepBytes ：输出值，步长，也就是每行所占的字节数，对应opencv的step，bitMap中的stride

三、字节对齐

通常，Malloc的所开辟的内存，是全局存储器，也就是普通的显存，所有cuda网格上的操作，都能读写全局存储器中的任意位置，而这个读取时存在延迟的，很容易造成性能瓶颈。

所以，在访问显存时，读取和存储就必须字节对齐 ，如果没有正确对齐，读写将会被编译器拆分为多次操作，降低访存性能。

而我们上一节内容提到的参数 pStepBytes，这个就是字节对齐后，每行所占的字节数。

这个类似Bitmap中的stride，在bitmap图像中，一般是四字节对齐，而在npp中，则和显卡相关，楼主这里是$512$。

以如下例子：

用npp创建一个 $4\ast 2$ 大小的矩阵， 返回的step是$512$， 也就是每行所占的字节数为$512$, 这和我们所创建的 $2$有一些出入。 也就是，npp已经自动帮我们在每行补上了$510$ 字节的大小了。

总之，无论我们创建的长度所占多少字节，最后用npp来创建内存时，都会补成$512$ 的倍数，这就是NPP中的字节对齐。

	const int nRows = 2;
	const int nCols = 4;
	int nLineStep_npp = 0;
	Npp8u* pu8_npp = nppiMalloc_8u_C1(nRows, nCols, &nLineStep_npp);

	printf("Step = %d\n", nLineStep_npp);

字节对齐的优点：计算快，这个和CUDA的架构有关，每次读取数据时，能够提高性能；
缺点：拷贝数据时，比较复杂，需要循环开辟空间

四、NPP中的阈值函数来探讨字节对齐

前面提到，nppMalloc的内存，会在每行自动补上512倍数的内存，那么我们在拷贝数据时，需要考虑缺省的字节长度么 ？答案时，当然要考虑！

表面上我们只创建了 $2\ast 4\ast sizeof(float)$大小的空间，实际上创建了$2\ast step$ 的空间。
补齐的那些空间，我们自然也就不需要用到了，虽然空间有点浪费，但却提高了内存访问的效率。

本例子主要利用NPP的阈值函数，将大于3的数字，都替换成0。
以下是本例子验证的步骤：

• Step1：在CPU下创建一个 $2\ast 4$大小的矩阵
• Step2：用 nppiMalloc_32f_C1 函数在Cuda上创建内存，并打印字节对齐后的长度
• Step3：将CPU创建的数据，拷贝到CUDA中
• Step4：调用 nppiThreshold_Val_32f_C1R 阈值处理函数，获取输出结果
• Step5：将CUDA上的结果，拷贝到CPU下，打印验证

	const int nH = 2;
	const int nW = 4;

	float data[nH * nW] = { 1.0, 1.0, 2.0, 2.0, 3.0, 4.0, 5.0, 9.0 };

	//float* pSrc_dev = nullptr;
	//cudaMalloc((void**)&pSrc_dev, nH * nW * sizeof(float));
	//cudaMemcpy(pSrc_dev, data, nH * nW * sizeof(float), cudaMemcpyHostToDevice);

	int nLineStep_npp = 0;
	Npp32f* pSrc_dev = nppiMalloc_32f_C1(nH, nW, &nLineStep_npp);
	printf("Step = %d\n", nLineStep_npp);

	for (int i = 0; i < nH; ++i)
	{
		cudaMemcpy((unsigned char*)pSrc_dev + i * nLineStep_npp,
			(unsigned char*)data + i * nW * sizeof(float),  nW * sizeof(float), cudaMemcpyHostToDevice);
	}

	//unsigned char* pDst_dev = nullptr;
	//cudaMalloc((void**)&pDst_dev, nRows * nCols);
	Npp32f* pDst_dev = nppiMalloc_32f_C1(nH, nW, &nLineStep_npp);

	nppiThreshold_Val_32f_C1R(pSrc_dev, nLineStep_npp,
		pDst_dev, nLineStep_npp,
		NppiSize{ nW , nH }, 3.0, 0.0, NPP_CMP_GREATER);
	//nppiThreshold_LTValGTVal_8u_C1R(pSrc_dev, nCols,
	//	pDst_dev, nCols,
	//	NppiSize{ nCols , nRows }, 2u, 0u, 1u, 5u);
	//nppiThreshold_LTValGTVal_32f_C1IR(pSrc_dev, nCols,
	//	NppiSize{ nCols , nRows }, 2u, 0u, 1u, 5u);

	float* pDst_host = (float*)malloc(nH * nW * sizeof(float));
	for (int i = 0; i < nH; ++i)
	{
		cudaMemcpy((unsigned char*)pDst_host + nW * i * sizeof(float), (unsigned char*)pDst_dev + i * nLineStep_npp,  nW * sizeof(float), cudaMemcpyDeviceToHost);
	}
	//cudaMemcpy(pDst_host, pDst_dev, nH * nW * sizeof(float), cudaMemcpyDeviceToHost);
	for (int i = 0; i < 8; ++i)
	{
		printf("%f ", pDst_host[i]);
	}

	// call api to free memory
	std::cout << "\nhello world \n" << std::endl;
	return 0;

参考文献：

CUDA库之NPP入门(一)：NVIDIA 2D Image and Signal Processing Performance Primitives：

gpu显存（全局内存）在使用时数据对齐的问题