本文通過使用malloc分配記憶體和cudaHostAlloc分配頁鎖定記憶體，說明使用頁鎖定記憶體可提高運算效率，並指出哪些場合適合使用頁鎖定記憶體。

malloc分配的是標準的可分頁的(pagable)的主機記憶體，作業系統在對記憶體進行排程的時候可能會將這種記憶體分頁或者交換到磁碟上，需要的時候再調回記憶體，這樣就會增加運算時間。而cudaHostAlloc分配的是頁鎖定的(page-locked)主機記憶體，作業系統不會對這塊記憶體分頁和交換到磁碟上，確保該記憶體始終駐留在實體記憶體中。

下面通過100M資料在主機和裝置上的交換說明二者的差異。貼上程式碼：

/********************************************************************
*  PageLockedMem.cu
*  Compare the performance of general mem and page locked mem.
*********************************************************************/

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cutil_inline.h>

#define _SIZE 100*1024*1024

/************************************************************************/
/* Init CUDA                                                            */
/************************************************************************/
bool InitCUDA(void)
{
    ......
}

float cudaMallocTest(int size, bool dir)
{
	cudaEvent_t start, stop;
	float elapsedTime;

	int *a, *dev_a;

	a = (int*)malloc(size*sizeof(int));
	if (!a)
	{
		printf("Mem error!\n");
	}
	cutilSafeCall(cudaMalloc((void**)&dev_a, size*sizeof(int)));

	cutilSafeCall(cudaEventCreate(&start));
	cutilSafeCall(cudaEventCreate(&stop));
	cutilSafeCall(cudaEventRecord(start, 0));

	for (int i=0; i<10; i++)
	{
		if (dir)
		{
			cutilSafeCall(cudaMemcpy(dev_a, a, size*sizeof(int), cudaMemcpyHostToDevice));
		}
		else
		{
			cutilSafeCall(cudaMemcpy(a, dev_a, size*sizeof(int), cudaMemcpyDeviceToHost));
		}
	}

	cutilSafeCall(cudaEventRecord(stop, 0));
	cudaEventSynchronize(stop);
	cutilSafeCall(cudaEventElapsedTime(&elapsedTime, start, stop));

	free(a);
	cutilSafeCall(cudaFree(dev_a));
	cutilSafeCall(cudaEventDestroy(start));
	cutilSafeCall(cudaEventDestroy(stop));

	return elapsedTime;
}

float cudaHostAllocTest(int size, bool dir)
{
	cudaEvent_t start, stop;
	float elapsedTime;

	int *a, *dev_a;

	cutilSafeCall(cudaHostAlloc((void**)&a, size*sizeof(int), cudaHostAllocDefault));
	cutilSafeCall(cudaMalloc((void**)&dev_a, size*sizeof(int)));

	cutilSafeCall(cudaEventCreate(&start));
	cutilSafeCall(cudaEventCreate(&stop));
	cutilSafeCall(cudaEventRecord(start, 0));

	for (int i=0; i<10; i++)
	{
		if (dir)
		{
			cutilSafeCall(cudaMemcpy(dev_a, a, size*sizeof(int), cudaMemcpyHostToDevice));
		}
		else
		{
			cutilSafeCall(cudaMemcpy(a, dev_a, size*sizeof(int), cudaMemcpyDeviceToHost));
		}
	}

	cutilSafeCall(cudaEventRecord(stop, 0));
	cudaEventSynchronize(stop);
	cutilSafeCall(cudaEventElapsedTime(&elapsedTime, start, stop));

	cutilSafeCall(cudaFreeHost(a));
	cutilSafeCall(cudaFree(dev_a));
	cutilSafeCall(cudaEventDestroy(start));
	cutilSafeCall(cudaEventDestroy(stop));

	return elapsedTime;
}



int main(int argc, char* argv[])
{

	if(!InitCUDA()) {
		return 0;
	}

	float elapsedTime;
	float MB = (float)100*_SIZE*sizeof(int)/1024/1024;

	elapsedTime = cudaMallocTest(_SIZE, true);

	printf("Time using cudaMalloc: %3.1f ms\n", elapsedTime);
	printf("\tMB/s during copy up: %3.1f \n", MB/(elapsedTime/1000));

	elapsedTime = cudaMallocTest(_SIZE, true);

	printf("Time using cudaMalloc: %3.1f ms\n", elapsedTime);
	printf("\tMB/s during copy down: %3.1f \n", MB/(elapsedTime/1000));

	elapsedTime = cudaHostAllocTest(_SIZE, true);

	printf("Time using cudaHostAlloc: %3.1f ms\n", elapsedTime);
	printf("\tMB/s during copy up: %3.1f \n", MB/(elapsedTime/1000));

	elapsedTime = cudaHostAllocTest(_SIZE, true);

	printf("Time using cudaHostAlloc: %3.1f ms\n", elapsedTime);
	printf("\tMB/s during copy down: %3.1f \n", MB/(elapsedTime/1000));
	return 0;
}
 

可以看出運算時間縮短了約2倍。

但是並不是所有的場合都適合用頁鎖定記憶體，因為使用固定記憶體時，將失去虛擬記憶體的所有功能，即需要為每個頁鎖定記憶體分配實體記憶體，系統將更快耗盡記憶體(跟使用普通記憶體相比)。所以要根據需要進行選擇。