天天看點
傳回

CUDA By Example(八)——流頁鎖定主機記憶體CUDA流使用單個CUDA流使用多個CUDA流GPU的工作排程機制高效地使用多個CUDA流遇到的問題(未解決)

文章目錄

  • 頁鎖定主機記憶體
    • 可分頁記憶體函數
    • 頁鎖定記憶體函數
  • CUDA流
  • 使用單個CUDA流
  • 使用多個CUDA流
  • GPU的工作排程機制
  • 高效地使用多個CUDA流
  • 遇到的問題(未解決)

頁鎖定主機記憶體

在之前的各個示例中，都是通過

cudaMalloc()

在GPU上配置設定記憶體，以及通過标準的C庫函數

malloc()

在主機上配置設定記憶體。除此之外，CUDA運作時還提供了自己獨有的機制來配置設定主機記憶體：

cudaHostAlloc()

。如果

malloc()

已經能很好地滿足C程式員的需求，那麼為什麼還要使用這個函數？

事實上，

malloc()

配置設定的記憶體與

cudaHostAlloc()

配置設定的記憶體之間存在着一個重要差異。C庫函數

malloc()

将配置設定标準的、可分頁的(Pagable) 主機記憶體，而

cudaHostAlloc()

将配置設定頁鎖定的主機記憶體。頁鎖定記憶體也稱為固定記憶體(Pinned Memory)或者不可分頁記憶體，它有一個重要的屬性：作業系統将不會對這塊記憶體分頁并交換到磁盤上，進而確定了該記憶體始終駐留在實體記憶體中。是以，作業系統能夠安全地使某個應用程式通路該記憶體的實體位址，因為這塊記憶體将不會被破壞或者重新定位。

由于GPU知道記憶體的實體位址，是以可以通過 “直接記憶體通路(Direct Memory Access, DMA)” 技術來在GPU和主機之間複制資料。由于DMA在執行複制時無需CPU的介入，這也就同樣意味着，CPU很可能在DMA的執行過程中将目标記憶體交換到磁盤上，或者通過更新作業系統的可分頁表來重新定位目标記憶體的實體位址。CPU可能會移動可分頁的資料，這就可能對DMA操作造成延遲。是以，在DMA複制過程中使用固定記憶體是非常重要的。事實上，當使用可分頁記憶體進行複制時，CUDA驅動程式仍然會通過DMA把資料傳輸給GPU。是以，複制操作将執行兩遍：

  • 第一遍從可分頁記憶體複制到一塊 ”臨時的“ 頁鎖定記憶體
  • 然後再從這個頁鎖定記憶體複制到GPU上

是以，每當從可分頁記憶體中執行複制操作時，複制速度将受限于PCIE傳輸速度和系統前端總線速度相對較低的一方。在某些系統中，這些總線在帶寬上有着巨大的差異。是以當在GPU和主機間複制資料時，這種差異會使頁鎖定主機記憶體的性能比标準可分頁記憶體的性能要高大約2倍。即使PCIE的速度與前端總線的速度相等，由于可分頁記憶體需要更多一次由CPU參與的複制操作，是以會帶來額外的開銷。

然而，你也不能進入另一個極端：查找每個

malloc

調用并将其替換為

cudaHostAlloc()

調用。固定記憶體是一把雙刃劍。當使用固定記憶體時，你将失去虛拟記憶體的所有功能。特别是，在應用程式中使用每個頁鎖定記憶體時都需要配置設定實體記憶體，因為這些記憶體不能交換到磁盤上。這意味着，與使用标準的

malloc()

調用相比，系統将更快地耗盡記憶體。是以，應用程式在實體記憶體較少的機器上會運作失敗，而且意味着應用程式将影響在系統上運作的其他應用程式的性能。

這些情況并不是說不使用

cudaHostAlloc()

，而是提醒你應該清楚頁鎖定記憶體得到隐含作用。我們建議，僅對

cudaMemcpy()

調用中的源記憶體或者目标記憶體，才使用頁鎖定記憶體，并且在不再需要使用它們時立即施放，而不是等到應用程式關閉時才施放。

cudaHostAlloc()

與到目前為止學習的其他内容一樣簡單，下面通過一個示例，說明如何配置設定固定記憶體，以及它對于标準可分頁記憶體的性能優勢。

這裡要做的就是配置設定一個GPU緩沖區，以及一個大小相等的主機緩沖區，然後在這兩個緩沖區之間執行一些複制操作。我們允許使用者指定複制的方向，例如為 “上”（從主機到裝置）或者為 “下”（從裝置到主機）。為了獲得精确的時間統計，我們為複制操作的起始時刻和結束時刻分别設定了CUDA事件。

可分頁記憶體函數

首先為

size

個整數分别配置設定主機緩沖區和GPU緩沖區

float cuda_malloc_test(int size, bool up) {
	cudaEvent_t start, stop;
	int* a, * dev_a;
	float elapsedTime;

	HANDLE_ERROR(cudaEventCreate(&start));
	HANDLE_ERROR(cudaEventCreate(&stop));

	a = (int*)malloc(size * sizeof(*a));
	HANDLE_NULL(a);
	HANDLE_ERROR(cudaMalloc((void**)&dev_a, size * sizeof(*dev_a)));

然後執行100次複制操作，并由參數

up

來指定複制方向，在完成複制操作後停止計時器。

HANDLE_ERROR(cudaEventRecord(start, 0));
	for (int i = 0; i < 100; i++) {
		if (up)
			HANDLE_ERROR(cudaMemcpy(dev_a, a, size * sizeof(*dev_a), cudaMemcpyHostToDevice));
		else
			HANDLE_ERROR(cudaMemcpy(a, dev_a, size * sizeof(*dev_a), cudaMemcpyDeviceToHost));
	}
	HANDLE_ERROR(cudaEventRecord(stop, 0));
	HANDLE_ERROR(cudaEventSynchronize(stop));
	HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop));

在執行了 100 次複制操作後，釋放主機緩沖區和GPU緩沖區，并且銷毀計時事件。

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

	return elapsedTime;
}

頁鎖定記憶體函數

與可分頁記憶體函數的差別就在于，使用

cudaHostAlloc()

配置設定記憶體，使用

cudaFreeHost()

施放記憶體

float cuda_host_alloc_test(int size, bool up) {
	cudaEvent_t start, stop;
	int* a, * dev_a;
	float elapsedTime;

	HANDLE_ERROR(cudaEventCreate(&start));
	HANDLE_ERROR(cudaEventCreate(&stop));

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

	HANDLE_ERROR(cudaEventRecord(start, 0));
	for (int i = 0; i < 100; i++) {
		if (up)
			HANDLE_ERROR(cudaMemcpy(dev_a, a, size * sizeof(*dev_a), cudaMemcpyHostToDevice));
		else
			HANDLE_ERROR(cudaMemcpy(a, dev_a, size * sizeof(*dev_a), cudaMemcpyDeviceToHost));
	}
	HANDLE_ERROR(cudaEventRecord(stop, 0));
	HANDLE_ERROR(cudaEventSynchronize(stop));
	HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop));

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

	return elapsedTime;
}

完整代碼

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "../../common/book.h"

#include <stdio.h>
#include <iostream>

#define SIZE (10*1024*1024)

float cuda_malloc_test(int size, bool up) {
	cudaEvent_t start, stop;
	int* a, * dev_a;
	float elapsedTime;

	HANDLE_ERROR(cudaEventCreate(&start));
	HANDLE_ERROR(cudaEventCreate(&stop));

	a = (int*)malloc(size * sizeof(*a));
	HANDLE_NULL(a);
	HANDLE_ERROR(cudaMalloc((void**)&dev_a, size * sizeof(*dev_a)));

	HANDLE_ERROR(cudaEventRecord(start, 0));
	for (int i = 0; i < 100; i++) {
		if (up)
			HANDLE_ERROR(cudaMemcpy(dev_a, a, size * sizeof(*dev_a), cudaMemcpyHostToDevice));
		else
			HANDLE_ERROR(cudaMemcpy(a, dev_a, size * sizeof(*dev_a), cudaMemcpyDeviceToHost));
	}
	HANDLE_ERROR(cudaEventRecord(stop, 0));
	HANDLE_ERROR(cudaEventSynchronize(stop));
	HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop));

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

	return elapsedTime;
}

float cuda_host_alloc_test(int size, bool up) {
	cudaEvent_t start, stop;
	int* a, * dev_a;
	float elapsedTime;

	HANDLE_ERROR(cudaEventCreate(&start));
	HANDLE_ERROR(cudaEventCreate(&stop));

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

	HANDLE_ERROR(cudaEventRecord(start, 0));
	for (int i = 0; i < 100; i++) {
		if (up)
			HANDLE_ERROR(cudaMemcpy(dev_a, a, size * sizeof(*dev_a), cudaMemcpyHostToDevice));
		else
			HANDLE_ERROR(cudaMemcpy(a, dev_a, size * sizeof(*dev_a), cudaMemcpyDeviceToHost));
	}
	HANDLE_ERROR(cudaEventRecord(stop, 0));
	HANDLE_ERROR(cudaEventSynchronize(stop));
	HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop));

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

	return elapsedTime;
}

int main(void) {
	float elapsedTime;
	float MB = (float)100 * SIZE * sizeof(int) / 1024 / 1024;
	elapsedTime = cuda_malloc_test(SIZE, true);
	std::cout << "Time using cudaMalloc: " << elapsedTime << " ms\n";
	std::cout << "\tMB/s during copy up: " << MB / (elapsedTime / 1000) << std::endl;

	elapsedTime = cuda_malloc_test(SIZE, false);
	std::cout << "Time using cudaMalloc: " << elapsedTime << " ms\n";
	std::cout << "\tMB/s during copy down: " << MB / (elapsedTime / 1000) << std::endl;

	elapsedTime = cuda_host_alloc_test(SIZE, true);
	std::cout << "Time using cudaHostAlloc: " << elapsedTime << " ms\n";
	std::cout << "\tMB/s during copy up: " << MB / (elapsedTime / 1000) << std::endl;

	elapsedTime = cuda_host_alloc_test(SIZE, false);
	std::cout << "Time using cudaHostAlloc: " << elapsedTime << " ms\n";
	std::cout << "\tMB/s during copy down: " << MB / (elapsedTime / 1000) << std::endl;
}

運作結果

CUDA By Example(八)——流頁鎖定主機記憶體CUDA流使用單個CUDA流使用多個CUDA流GPU的工作排程機制高效地使用多個CUDA流遇到的問題(未解決)

可以發現使用頁鎖定記憶體比使用可分頁記憶體的讀寫速度快了2倍多。

CUDA流

在之前的文章中，我們引入了CUDA事件的概念。當時并沒有介紹

cudaEventRecord()

的第二個參數，而隻是簡要地指出這個參數用于指定插入事件的流(Stream)。

cudaEvent_t start;
cudaEventCreate(&start);
cudaEventRecord(start, 0);

CUDA流在加速應用程式方面起着重要的作用。CUDA流表示一個GPU操作隊列，并且該隊列中的操作将以指定的順序執行。我們可以在流中添加一些操作，例如核函數啟動、記憶體複制、以及事件的啟動和結束等。将這些操作添加到流的順序也就是它們的執行順序。你可以将每個流視為GPU上的一個任務，并且這些任務可以并行執行。

下面将首先介紹如何使用流，然後介紹如何使用流來加速應用程式。

使用單個CUDA流

下面首先通過在應用程式中使用單個流來說明流的用法。假設有一個CUDA C核函數，該函數帶有兩個輸入資料緩沖區，a 和 b。核函數将對這些緩沖區中相應位置上的值執行某種計算，并将生成的結果儲存到輸出緩沖區 c。下面這個示例中，将計算 a 中三個值和 b 中三個值的平均值：

__global__ void kernel(int* a, int* b, int* c) {
	int idx = threadIdx.x + blockIdx.x * blockDim.x;
	if (idx < N) {
		int idx1 = (idx + 1) % 256;
		int idx2 = (idx + 2) % 256;
		float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
		float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
		c[idx] = (as + bs) / 2;
	}
}

這個核函數很簡單，下面重要的是函數

main()

中與流相關的代碼

int main(void) {
	cudaDeviceProp prop;
	int whichDevice;
	HANDLE_ERROR(cudaGetDevice(&whichDevice));
	HANDLE_ERROR(cudaGetDeviceProperties(&prop, whichDevice));
	if (!prop.deviceOverlap) {
		std::cout << "Device will not handle overlaps, so no speed up from streams" << std::endl;

		return 0;
	}

首先選擇一個支援裝置重疊(Device Overlap)功能的裝置。支援裝置重疊功能的GPU能夠在執行一個CUDA C核函數的同時，還能在裝置與主機之間執行複制操作。

正如前面提到的，我們将使用多個流來實作這種計算與資料傳輸的重疊，但首先來看看如何建立和使用一個流。與其他需要測量性能提升(或者降低)的示例一樣，首先建立和啟動一個事件計時器：

cudaEvent_t start, stop;
	float elapsedTime;

	// 啟動計時器
	HANDLE_ERROR(cudaEventCreate(&start));
	HANDLE_ERROR(cudaEventCreate(&stop));
	HANDLE_ERROR(cudaEventRecord(start, 0));

啟動計時器之後，建立在應用程式中使用的流：

// 初始化流
	cudaStream_t stream;
	HANDLE_ERROR(cudaStreamCreate(&stream));

這就是建立流需要的全部工作，接下來是資料配置設定操作

int* host_a, * host_b, * host_c;
	int* dev_a, * dev_b, * dev_c;

	// 在GPU上配置設定記憶體
	HANDLE_ERROR(cudaMalloc((void**)&dev_a, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_b, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_c, N * sizeof(int)));
	
	// 配置設定由流使用的頁鎖定記憶體
	HANDLE_ERROR(cudaHostAlloc((void**)&host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));
	HANDLE_ERROR(cudaHostAlloc((void**)&host_b, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));
	HANDLE_ERROR(cudaHostAlloc((void**)&host_c, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));

	for (int i = 0; i < FULL_DATA_SIZE; i++) {
		host_a[i] = rand();
		host_b[i] = rand();
	}

我們在GPU和主機上分别配置設定好了輸入記憶體和輸出記憶體。注意，由于程式将使用主機上的固定記憶體，是以調用

cudaHostAlloc()

來執行記憶體配置設定操作。

使用固定記憶體的原因并不隻在于使複制操作執行得更快，還存在另外一個好處。會在後面進行詳細地分析，我們将使用一種新的

cudaMemcpy()

函數，并且在這個新函數中需要頁鎖定主機記憶體。在配置設定完輸入記憶體後，調用C的庫函數

rand()

并用随機整數填充主機記憶體。

在建立了流和計時事件，并且配置設定了裝置記憶體和主機記憶體後，就準備好了執行一些計算。通常，我們會将這個階段一帶而過，隻是将兩個輸入緩沖區複制到GPU，啟動核函數，然後将輸出緩沖區複制回主機。我們将再次沿用這種模式，隻是進行了一些小修改。

首先，我們不将輸入緩沖區整體都複制到GPU，而是将輸入緩沖區劃分為更小的塊，并在每個塊上執行一個包含三個步驟的過程。我們将一部分輸入緩沖區複制到GPU，在這部分緩沖區上運作核函數，然後将輸出緩沖區中的這部分結果複制回主機。

想象一下需要使用這種方法的一種情形：GPU的記憶體遠少于主機記憶體，由于整個緩沖區無法一次性填充到GPU，是以需要分塊進行計算。

執行"分塊"計算的代碼如下所示：

//在整體資料上循環，每個資料塊的大小為N
	for (int i = 0; i < FULL_DATA_SIZE; i += N) {
		// 将鎖定記憶體以異步方式複制到裝置上
		HANDLE_ERROR(cudaMemcpyAsync(dev_a, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream));
		HANDLE_ERROR(cudaMemcpyAsync(dev_b, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream));

		kernel << <N / 256, 256, 0, stream >> > (dev_a, dev_b, dev_c);

		// 将資料從裝置複制到鎖定記憶體
		HANDLE_ERROR(cudaMemcpyAsync(host_c + i, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost, stream));
	}

注意到，代碼沒有使用熟悉的

cudaMemcpy()

，而是通過一個新函數

cudaMemcpyAsync()

在GPU與主機之間複制資料。這些函數之間的差異雖然很小，但卻很重要。

cudaMemcpy()

的行為類似于C庫函數

memcpy()

。尤其是，這個函數将以同步方式執行，這意味着，當函數傳回時，複制操作就已經完成，并且在輸出緩沖區中包含了複制進去的内容。

異步函數的行為與同步函數相反，通過名字

cudaMemcpyAsync()

就可以知道。在調用

cudaMemcpyAsync()

時，隻是放置一個請求，表示在流中執行一次記憶體複制操作，這個流是通過參數

stream

來指定的。當函數傳回時，我們無法確定複制操作是否已經啟動，更無法保證它是否已經結束。我們能夠得到的保證是，複制操作肯定會當下一個被放入流中的操作之前執行。任何傳遞給

cudaMemcpyAsync()

的主機記憶體指針都必須已經通過

cudaHostAlloc()

配置設定好記憶體。也就是，你隻能以異步方式對頁鎖定記憶體進行複制操作。

注意，在核函數調用的尖括号中還可以帶有一個流參數。此時核函數調用将是異步的，就像之前與GPU之間的記憶體複制操作一樣。從技術上來說，當循環疊代完一次時，有可能不會啟動任何記憶體複制或核函數執行。

這裡隻能確定的是：第一次放入流中的複制操作将在第二次複制操作之前執行。第二個複制操作将在核函數啟動之前完成，而核函數将在第三次複制操作開始之前完成。流就像一個有序的工作隊列，GPU從該隊列中依次取出工作并執行。

當

for()

循環結束時，在隊列中應包含了許多等待GPU執行的工作。如果想要確定GPU執行完了計算和記憶體複制等操作，那麼就需要将GPU與主機同步。也就是說，主機在繼續執行之前，要首先等待GPU執行完成。可以調用

cudaStreamSynchronize()

并指定想要等待的流：

//将計算結果從頁鎖定記憶體複制到主機記憶體
	HANDLE_ERROR(cudaStreamSynchronize(stream));

當程式執行到

stream

與主機同步之後的代碼時，所有的計算和複制操作都已經完成，是以可以停止計時器，收集性能資料，并釋放輸入緩沖區和輸出緩沖區。

HANDLE_ERROR(cudaEventRecord(stop, 0));

	HANDLE_ERROR(cudaEventSynchronize(stop));
	HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop));
	std::cout << "Time taken: " << elapsedTime << " ms" << std::endl;

	//釋放流和記憶體
	HANDLE_ERROR(cudaFreeHost(host_a));
	HANDLE_ERROR(cudaFreeHost(host_b));
	HANDLE_ERROR(cudaFreeHost(host_c));
	HANDLE_ERROR(cudaFree(dev_a));
	HANDLE_ERROR(cudaFree(dev_b));
	HANDLE_ERROR(cudaFree(dev_c));

最後，在退出應用程式之前，記得銷毀對GPU操作進行排隊的流。

HANDLE_ERROR(cudaStreamDestroy(stream));

	return 0;
}

這個示例并沒有充分說明流的強大功能。當然，如果當主機正在執行一些工作時，GPU也正忙于處理填充到流的工作，那麼即使使用單個流也有助于應用程式速度的提升。但即使不需要在主機上做太多的工作，我們仍然可以通過使用流來加速應用程式。

完整代碼

#include "../../common/book.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <iostream>

#define N (1024*1024)
#define FULL_DATA_SIZE (N*20)

__global__ void kernel(int* a, int* b, int* c) {
	int idx = threadIdx.x + blockIdx.x * blockDim.x;
	if (idx < N) {
		int idx1 = (idx + 1) % 256;
		int idx2 = (idx + 2) % 256;
		float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
		float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
		c[idx] = (as + bs) / 2;
	}
}

int main(void) {
	cudaDeviceProp prop;
	int whichDevice;
	HANDLE_ERROR(cudaGetDevice(&whichDevice));
	HANDLE_ERROR(cudaGetDeviceProperties(&prop, whichDevice));
	if (!prop.deviceOverlap) {
		std::cout << "Device will not handle overlaps, so no speed up from streams" << std::endl;

		return 0;
	}
	
	cudaEvent_t start, stop;
	float elapsedTime;

	// 啟動計時器
	HANDLE_ERROR(cudaEventCreate(&start));
	HANDLE_ERROR(cudaEventCreate(&stop));
	HANDLE_ERROR(cudaEventRecord(start, 0));

	// 初始化流
	cudaStream_t stream;
	HANDLE_ERROR(cudaStreamCreate(&stream));

	int* host_a, * host_b, * host_c;
	int* dev_a, * dev_b, * dev_c;

	// 在GPU上配置設定記憶體
	HANDLE_ERROR(cudaMalloc((void**)&dev_a, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_b, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_c, N * sizeof(int)));
	
	// 配置設定由流使用的頁鎖定記憶體
	HANDLE_ERROR(cudaHostAlloc((void**)&host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));
	HANDLE_ERROR(cudaHostAlloc((void**)&host_b, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));
	HANDLE_ERROR(cudaHostAlloc((void**)&host_c, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));

	for (int i = 0; i < FULL_DATA_SIZE; i++) {
		host_a[i] = rand();
		host_b[i] = rand();
	}

	//在整體資料上循環，每個資料塊的大小為N
	for (int i = 0; i < FULL_DATA_SIZE; i += N) {
		// 将鎖定記憶體以異步方式複制到裝置上
		HANDLE_ERROR(cudaMemcpyAsync(dev_a, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream));
		HANDLE_ERROR(cudaMemcpyAsync(dev_b, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream));

		kernel << <N / 256, 256, 0, stream >> > (dev_a, dev_b, dev_c);

		// 将資料從裝置複制到鎖定記憶體
		HANDLE_ERROR(cudaMemcpyAsync(host_c + i, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost, stream));
	}

	//将計算結果從頁鎖定記憶體複制到主機記憶體
	HANDLE_ERROR(cudaStreamSynchronize(stream));

	HANDLE_ERROR(cudaEventRecord(stop, 0));

	HANDLE_ERROR(cudaEventSynchronize(stop));
	HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop));
	std::cout << "Time taken: " << elapsedTime << " ms" << std::endl;

	//釋放流和記憶體
	HANDLE_ERROR(cudaFreeHost(host_a));
	HANDLE_ERROR(cudaFreeHost(host_b));
	HANDLE_ERROR(cudaFreeHost(host_c));
	HANDLE_ERROR(cudaFree(dev_a));
	HANDLE_ERROR(cudaFree(dev_b));
	HANDLE_ERROR(cudaFree(dev_c));

	HANDLE_ERROR(cudaStreamDestroy(stream));

	return 0;
}

運作結果

CUDA By Example(八)——流頁鎖定主機記憶體CUDA流使用單個CUDA流使用多個CUDA流GPU的工作排程機制高效地使用多個CUDA流遇到的問題(未解決)

使用多個CUDA流

下面将單個流的版本改為使用兩個不同的流。改進這個程式的思想很簡單：分塊計算以及記憶體複制和核函數執行的重疊。

即在第 0 個流執行核函數的同時，第一個流将輸入緩沖區複制到GPU。然後，在第 0 個流将計算結果複制回主機的同時，第 1 個流将執行核函數…

如下圖所示，這裡假設記憶體複制操作和核函數執行的時間大緻相，且GPU可以同時執行一個記憶體複制操作和一個核函數，是以空的方框表示一個流正在等待執行哦某個操作的時刻，這個操作不能與其他流的操作互相重疊。

CUDA By Example(八)——流頁鎖定主機記憶體CUDA流使用單個CUDA流使用多個CUDA流GPU的工作排程機制高效地使用多個CUDA流遇到的問題(未解決)

事實上，實際的執行時間線可能比上圖給出的更好看，在一些新的 NVIDIA GPU 中同時支援核函數和兩次記憶體複制操作，一次是從主機到裝置，另一次是從裝置到主機。在任何支援記憶體複制和核函數的執行互相重疊的裝置上，當使用多個流時，應用程式的整體性能都會提升。

核函數代碼保持不變：

__global__ void kernel(int* a, int* b, int* c) {
	int idx = threadIdx.x + blockIdx.x * blockDim.x;
	if (idx < N) {
		int idx1 = (idx + 1) % 256;
		int idx2 = (idx + 2) % 256;
		float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
		float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
		c[idx] = (as + bs) / 2;
	}
}

與使用單個流的版本一樣，我們将判斷裝置是否支援計算與記憶體複制操作的重疊。如果裝置支援重疊，那麼就像前面一樣建立CUDA事件并對應用程式計時。

int main( void ) {
	cudaDeviceProp prop;
	int whichDevice;
	HANDLE_ERROR(cudaGetDevice(&whichDevice));
	HANDLE_ERROR(cudaGetDeviceProperties(&prop, whichDevice));
	if (!prop.deviceOverlap) {
		std::cout << "Device will not handle overlaps, so no speed up from streams" << std::endl;

		return 0;
	}

	cudaEvent_t start, stop;
	float elapsedTime;

	// 啟動計時器
	HANDLE_ERROR(cudaEventCreate(&start));
	HANDLE_ERROR(cudaEventCreate(&stop));
	HANDLE_ERROR(cudaEventRecord(start, 0));

接下來建立兩個流，建立方式與前面單個流的版本完全一樣。

// 初始化流
	cudaStream_t stream0, stream1;
	HANDLE_ERROR(cudaStreamCreate(&stream0));
	HANDLE_ERROR(cudaStreamCreate(&stream1));

假設在主機上仍然是兩個輸入緩沖區和一個輸出緩沖區。輸入緩沖區中填充的是随機資料，與使用單個流的應用程式采樣的方式一樣。然而，現在我們将使用兩個流來處理資料，配置設定兩組相同的GPU緩沖區，這樣每個流都可以獨立地在輸入資料塊上執行工作。

int* host_a, * host_b, * host_c;
	int* dev_a0, * dev_b0, * dev_c0; // 為第0個流配置設定的GPU記憶體
	int* dev_a1, * dev_b1, * dev_c1; // 為第1個流配置設定的GPU記憶體

	// 在GPU上配置設定記憶體
	HANDLE_ERROR(cudaMalloc((void**)&dev_a0, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_b0, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_c0, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_a1, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_b1, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_c1, N * sizeof(int)));

	// 配置設定由流使用的頁鎖定記憶體
	HANDLE_ERROR(cudaHostAlloc((void**)&host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));
	HANDLE_ERROR(cudaHostAlloc((void**)&host_b, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));
	HANDLE_ERROR(cudaHostAlloc((void**)&host_c, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));

	for (int i = 0; i < FULL_DATA_SIZE; i++) {
		host_a[i] = rand();
		host_b[i] = rand();
	}

然後，程式在輸入資料塊上循環。然而，由于現在使用了兩個流，是以在

for()

循環的疊代中需要處理的資料量也是原來的兩倍。在

stream()

中，我們首先将

a

和

b

的異步複制操作放入GPU的隊列，然後将一個核函數執行放入隊列，接下來再将一個複制回

c

的操作放入隊列：

//在整體資料上循環，每個資料塊的大小為N
	for (int i = 0; i < FULL_DATA_SIZE; i += N*2) {
		// 将鎖定記憶體以異步方式複制到裝置上
		HANDLE_ERROR(cudaMemcpyAsync(dev_a0, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));
		HANDLE_ERROR(cudaMemcpyAsync(dev_b0, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));

		kernel << <N / 256, 256, 0, stream0 >> > (dev_a0, dev_b0, dev_c0);

		// 将資料從裝置複制到鎖定記憶體
		HANDLE_ERROR(cudaMemcpyAsync(host_c + i, dev_c0, N * sizeof(int), cudaMemcpyDeviceToHost, stream0));

在将這些操作放入

stream0

的隊列後，再把下一個資料塊上的相同操作放入

stream1

的隊列中。

// 将鎖定記憶體以異步方式複制到裝置上
		HANDLE_ERROR(cudaMemcpyAsync(dev_a1, host_a + i + N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));
		HANDLE_ERROR(cudaMemcpyAsync(dev_b1, host_b + i + N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));

		kernel << <N / 256, 256, 0, stream1 >> > (dev_a1, dev_b1, dev_c1);

		// 将資料從裝置複制到鎖定記憶體
		HANDLE_ERROR(cudaMemcpyAsync(host_c + i + N, dev_c1, N * sizeof(int), cudaMemcpyDeviceToHost, stream1));
	}

這樣，在

for()

循環的疊代過程中，将交替地把每個資料塊放入這兩個流的隊列，直到所有待處理的輸入資料都被放入隊列。在結束了

for()

循環後，在停止應用程式的計時器之前，首先将 GPU 與 GPU進行同步。由于使用了兩個流，是以需要對二者都進行同步。

之後，停止計時器，顯示經曆的時間，并且執行清理工作。當然，我們要記住，現在需要銷毀兩個流，并且需要釋放兩倍的GPU記憶體。

HANDLE_ERROR(cudaEventRecord(stop, 0));

	HANDLE_ERROR(cudaEventSynchronize(stop));
	HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop));
	std::cout << "Time taken: " << elapsedTime << " ms" << std::endl;

	//釋放流和記憶體
	HANDLE_ERROR(cudaFreeHost(host_a));
	HANDLE_ERROR(cudaFreeHost(host_b));
	HANDLE_ERROR(cudaFreeHost(host_c));
	HANDLE_ERROR(cudaFree(dev_a0));
	HANDLE_ERROR(cudaFree(dev_b0));
	HANDLE_ERROR(cudaFree(dev_c0));
	HANDLE_ERROR(cudaFree(dev_a1));
	HANDLE_ERROR(cudaFree(dev_b1));
	HANDLE_ERROR(cudaFree(dev_c1));

	HANDLE_ERROR(cudaStreamDestroy(stream0));
	HANDLE_ERROR(cudaStreamDestroy(stream1));
	return 0;
}

完整代碼

#include "../../common/book.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <iostream>

#define N (1024*1024)
#define FULL_DATA_SIZE (N*20)

__global__ void kernel(int* a, int* b, int* c) {
	int idx = threadIdx.x + blockIdx.x * blockDim.x;
	if (idx < N) {
		int idx1 = (idx + 1) % 256;
		int idx2 = (idx + 2) % 256;
		float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
		float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
		c[idx] = (as + bs) / 2;
	}
}

int main(void) {
	cudaDeviceProp prop;
	int whichDevice;
	HANDLE_ERROR(cudaGetDevice(&whichDevice));
	HANDLE_ERROR(cudaGetDeviceProperties(&prop, whichDevice));
	if (!prop.deviceOverlap) {
		std::cout << "Device will not handle overlaps, so no speed up from streams" << std::endl;

		return 0;
	}

	cudaEvent_t start, stop;
	float elapsedTime;

	// 啟動計時器
	HANDLE_ERROR(cudaEventCreate(&start));
	HANDLE_ERROR(cudaEventCreate(&stop));
	HANDLE_ERROR(cudaEventRecord(start, 0));

	// 初始化流
	cudaStream_t stream0, stream1;
	HANDLE_ERROR(cudaStreamCreate(&stream0));
	HANDLE_ERROR(cudaStreamCreate(&stream1));

	int* host_a, * host_b, * host_c;
	int* dev_a0, * dev_b0, * dev_c0; // 為第0個流配置設定的GPU記憶體
	int* dev_a1, * dev_b1, * dev_c1; // 為第1個流配置設定的GPU記憶體

	// 在GPU上配置設定記憶體
	HANDLE_ERROR(cudaMalloc((void**)&dev_a0, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_b0, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_c0, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_a1, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_b1, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_c1, N * sizeof(int)));

	// 配置設定由流使用的頁鎖定記憶體
	HANDLE_ERROR(cudaHostAlloc((void**)&host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));
	HANDLE_ERROR(cudaHostAlloc((void**)&host_b, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));
	HANDLE_ERROR(cudaHostAlloc((void**)&host_c, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));

	for (int i = 0; i < FULL_DATA_SIZE; i++) {
		host_a[i] = rand();
		host_b[i] = rand();
	}

	//在整體資料上循環，每個資料塊的大小為N
	for (int i = 0; i < FULL_DATA_SIZE; i += N*2) {
		// 将鎖定記憶體以異步方式複制到裝置上
		HANDLE_ERROR(cudaMemcpyAsync(dev_a0, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));
		HANDLE_ERROR(cudaMemcpyAsync(dev_b0, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));

		kernel << <N / 256, 256, 0, stream0 >> > (dev_a0, dev_b0, dev_c0);

		// 将資料從裝置複制到鎖定記憶體
		HANDLE_ERROR(cudaMemcpyAsync(host_c + i, dev_c0, N * sizeof(int), cudaMemcpyDeviceToHost, stream0));

		// 将鎖定記憶體以異步方式複制到裝置上
		HANDLE_ERROR(cudaMemcpyAsync(dev_a1, host_a + i + N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));
		HANDLE_ERROR(cudaMemcpyAsync(dev_b1, host_b + i + N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));

		kernel << <N / 256, 256, 0, stream1 >> > (dev_a1, dev_b1, dev_c1);

		// 将資料從裝置複制到鎖定記憶體
		HANDLE_ERROR(cudaMemcpyAsync(host_c + i + N, dev_c1, N * sizeof(int), cudaMemcpyDeviceToHost, stream1));
	}

	HANDLE_ERROR(cudaStreamSynchronize(stream0));
	HANDLE_ERROR(cudaStreamSynchronize(stream1));

	HANDLE_ERROR(cudaEventRecord(stop, 0));

	HANDLE_ERROR(cudaEventSynchronize(stop));
	HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop));
	std::cout << "Time taken: " << elapsedTime << " ms" << std::endl;

	//釋放流和記憶體
	HANDLE_ERROR(cudaFreeHost(host_a));
	HANDLE_ERROR(cudaFreeHost(host_b));
	HANDLE_ERROR(cudaFreeHost(host_c));
	HANDLE_ERROR(cudaFree(dev_a0));
	HANDLE_ERROR(cudaFree(dev_b0));
	HANDLE_ERROR(cudaFree(dev_c0));
	HANDLE_ERROR(cudaFree(dev_a1));
	HANDLE_ERROR(cudaFree(dev_b1));
	HANDLE_ERROR(cudaFree(dev_c1));

	HANDLE_ERROR(cudaStreamDestroy(stream0));
	HANDLE_ERROR(cudaStreamDestroy(stream1));
	return 0;
}

運作結果

CUDA By Example(八)——流頁鎖定主機記憶體CUDA流使用單個CUDA流使用多個CUDA流GPU的工作排程機制高效地使用多個CUDA流遇到的問題(未解決)

GPU的工作排程機制

雖然從邏輯上來看，不同的流之間是互相獨立的，但事實上這種了解并不完全符合GPU的隊列機制。程式員可以将流視為有序的操作序列，其中既包含記憶體複制操作，又包含核函數調用。

然而，在硬體中并沒有流的概念，而是包含一個或多個引擎來執行記憶體複制操作，以及一個引擎來執行核函數。這些引擎彼此獨立地對操作進行排隊，是以将導緻如下圖所示的任務排程情形。圖中的箭頭說明了硬體引擎如何排程流中隊列的操作并實際執行。

CUDA By Example(八)——流頁鎖定主機記憶體CUDA流使用單個CUDA流使用多個CUDA流GPU的工作排程機制高效地使用多個CUDA流遇到的問題(未解決)

是以，在某種程度上，使用者與硬體關于GPU工作的排隊方式有着完全不同的了解，而CUDA驅動程式則負責對使用者和硬體進行協調。首先，在操作被添加到流的順序中包含了重要的依賴性。

如上圖中，第 0 個流對 A 的記憶體複制需要在對 B 的記憶體複制之前完成，而對 B 的複制又要在核函數 A 啟動之前完成。然而，一旦這些操作放入到硬體的記憶體複制引擎和核函數執行引擎的隊列中，這些依賴性将丢失，是以CUDA驅動程式需要確定硬體的執行單元不破壞内部的依賴性。

這意味着說明？之前在代碼中，應用程式基本上是對

a

調用一次

cudaMemcpyAsync()

，對

b

調用一次

cudaMemcpyAsync()

，然後再是執行核函數以及調用

cudaMemcpyAsync()

将

c

複制回主機。應用程式首先将對第 0 個流的所有操作放入隊列，然後是第 1 個流的所有操作。CUDA 驅動程式負責按照這些操作的順序把它們排程到硬體上執行，這就維持了流内部的依賴性。下圖說明了這些依賴性，其中從複制操作到核函數的箭頭表示，複制操作要等核函數執行完成之後才能開始。

CUDA By Example(八)——流頁鎖定主機記憶體CUDA流使用單個CUDA流使用多個CUDA流GPU的工作排程機制高效地使用多個CUDA流遇到的問題(未解決)

假定了解了 GPU 的工作排程原理後，我們可以得到關于這些操作再硬體上執行的時間線，如下圖所示

CUDA By Example(八)——流頁鎖定主機記憶體CUDA流使用單個CUDA流使用多個CUDA流GPU的工作排程機制高效地使用多個CUDA流遇到的問題(未解決)

由于第 0 個流中将

c

複制回主機的操作要等待核函數執行完成，是以第 1 個流中将

a

和

b

複制到 GPU 的操作雖然是完全獨立的，但卻被阻塞了，這是因為GPU引擎是按照指定的順序來執行工作。這種情況也說明了為什麼上面使用了兩個流卻沒有獲得很大的速度提升。

這個問題的直接原因是我們沒有意識到硬體的工作方式與CUDA流程式設計模型的方式是不同的。

硬體在處理記憶體複制和核函數執行時分别采用了不同的引擎，是以我們需要知道，将操作放入流中隊列中的順序将影響着 CUDA 驅動程式排程這些操作以及執行的方式。下面，我們将看到如何幫助硬體實作記憶體複制操作與核函數執行的重疊。

高效地使用多個CUDA流

如上節所看到的，如果同時排程某個流的所有操作，那麼容易在無意中阻塞另一個流的複制操作或者核函數執行。要解決這個問題，在将操作放入流的隊列時應采用寬度優先方式，而非深度優先方式。

也就是說，不是首先添加第 0 個流的所有四個操作(即a的複制、b的複制、核函數以及c的複制)，然後再添加第 1 個流的所有四個操作。而是将這兩個流之間的操作交叉添加。首先，将 a 的複制操作添加到第 0 個流，然後将 a 的複制操作添加到第 1 個流。接着，将 b 的複制操作添加到第 0 個流，再将 b 的複制操作添加到第 1 個流。接下來，将核函數調用添加到第 0 個流，再将相同的操作添加到第 1 個流中。最後，将 c 的複制操作添加到第 0 個流中，然後将相同的操作添加到第 1 個流中。

下面給出具體的代碼，隻需要修改

for()

循環内的代碼。

for (int i = 0; i < FULL_DATA_SIZE; i += N*2) {
		// 将複制a的操作放入stream0和stream1的隊列
		HANDLE_ERROR(cudaMemcpyAsync(dev_a0, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));
		HANDLE_ERROR(cudaMemcpyAsync(dev_a1, host_a + i + N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));

		// 将複制b的操作放入stream0和stream1的隊列
		HANDLE_ERROR(cudaMemcpyAsync(dev_b0, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));
		HANDLE_ERROR(cudaMemcpyAsync(dev_b1, hos + i + N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));
		
		// 将核函數的執行放入stream0和stream1的隊列中
		kernel << <N / 256, 256, 0, stream0 >> > (dev_a0, dev_b0, dev_c0);
		kernel << <N / 256, 256, 0, stream1 >> > (dev_a1, dev_b1, dev_c1);

		// 将複制c的操作放入stream0和stream1的隊列中
		HANDLE_ERROR(cudaMemcpyAsync(host_c + i, dev_c0, N * sizeof(int), cudaMemcpyDeviceToHost, stream0));
		HANDLE_ERROR(cudaMemcpyAsync(host_c + i + N, dev_c1, N * sizeof(int), cudaMemcpyDeviceToHost, stream1));
	}

如果記憶體複制操作的時間與核函數執行的時間大緻相當，那麼新的執行時間線将如下圖所示。引擎箭的依賴性通過箭頭表示，可以看到在新的排程順序中，這些依賴性仍然能得到滿足。

CUDA By Example(八)——流頁鎖定主機記憶體CUDA流使用單個CUDA流使用多個CUDA流GPU的工作排程機制高效地使用多個CUDA流遇到的問題(未解決)

由于采用了寬度優先方式将操作放入各個流的隊列中，是以第0個流對c的複制操作将不會阻塞第1個流對a和b的記憶體複制操作。這使得GPU能夠并行地執行複制操作和核函數，進而使應用程式的運作速度顯著加快。

完整代碼

#include "../../common/book.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <iostream>

#define N (1024*1024)
#define FULL_DATA_SIZE (N*20)

__global__ void kernel(int* a, int* b, int* c) {
	int idx = threadIdx.x + blockIdx.x * blockDim.x;
	if (idx < N) {
		int idx1 = (idx + 1) % 256;
		int idx2 = (idx + 2) % 256;
		float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
		float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
		c[idx] = (as + bs) / 2;
	}
}

int main(void) {
	cudaDeviceProp prop;
	int whichDevice;
	HANDLE_ERROR(cudaGetDevice(&whichDevice));
	HANDLE_ERROR(cudaGetDeviceProperties(&prop, whichDevice));
	if (!prop.deviceOverlap) {
		std::cout << "Device will not handle overlaps, so no speed up from streams" << std::endl;

		return 0;
	}

	cudaEvent_t start, stop;
	float elapsedTime;

	// 啟動計時器
	HANDLE_ERROR(cudaEventCreate(&start));
	HANDLE_ERROR(cudaEventCreate(&stop));
	HANDLE_ERROR(cudaEventRecord(start, 0));

	// 初始化流
	cudaStream_t stream0, stream1;
	HANDLE_ERROR(cudaStreamCreate(&stream0));
	HANDLE_ERROR(cudaStreamCreate(&stream1));

	int* host_a, * host_b, * host_c;
	int* dev_a0, * dev_b0, * dev_c0; // 為第0個流配置設定的GPU記憶體
	int* dev_a1, * dev_b1, * dev_c1; // 為第1個流配置設定的GPU記憶體

	// 在GPU上配置設定記憶體
	HANDLE_ERROR(cudaMalloc((void**)&dev_a0, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_b0, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_c0, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_a1, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_b1, N * sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_c1, N * sizeof(int)));

	// 配置設定由流使用的頁鎖定記憶體
	HANDLE_ERROR(cudaHostAlloc((void**)&host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));
	HANDLE_ERROR(cudaHostAlloc((void**)&host_b, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));
	HANDLE_ERROR(cudaHostAlloc((void**)&host_c, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));

	for (int i = 0; i < FULL_DATA_SIZE; i++) {
		host_a[i] = rand();
		host_b[i] = rand();
	}

	//在整體資料上循環，每個資料塊的大小為N
	for (int i = 0; i < FULL_DATA_SIZE; i += N*2) {
		// 将鎖定記憶體以異步方式複制到裝置上
		HANDLE_ERROR(cudaMemcpyAsync(dev_a0, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));
		HANDLE_ERROR(cudaMemcpyAsync(dev_a1, host_a + i + N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));

		// 将鎖定記憶體以異步方式複制到裝置上
		HANDLE_ERROR(cudaMemcpyAsync(dev_b0, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));
		HANDLE_ERROR(cudaMemcpyAsync(dev_b1, host_b + i + N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));
		
		kernel << <N / 256, 256, 0, stream0 >> > (dev_a0, dev_b0, dev_c0);
		kernel << <N / 256, 256, 0, stream1 >> > (dev_a1, dev_b1, dev_c1);

		// 将資料從裝置複制到鎖定記憶體
		HANDLE_ERROR(cudaMemcpyAsync(host_c + i, dev_c0, N * sizeof(int), cudaMemcpyDeviceToHost, stream0));

		// 将資料從裝置複制到鎖定記憶體
		HANDLE_ERROR(cudaMemcpyAsync(host_c + i + N, dev_c1, N * sizeof(int), cudaMemcpyDeviceToHost, stream1));
	}

	HANDLE_ERROR(cudaStreamSynchronize(stream0));
	HANDLE_ERROR(cudaStreamSynchronize(stream1));

	HANDLE_ERROR(cudaEventRecord(stop, 0));

	HANDLE_ERROR(cudaEventSynchronize(stop));
	HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop));
	std::cout << "Time taken: " << elapsedTime << " ms" << std::endl;

	//釋放流和記憶體
	HANDLE_ERROR(cudaFreeHost(host_a));
	HANDLE_ERROR(cudaFreeHost(host_b));
	HANDLE_ERROR(cudaFreeHost(host_c));
	HANDLE_ERROR(cudaFree(dev_a0));
	HANDLE_ERROR(cudaFree(dev_b0));
	HANDLE_ERROR(cudaFree(dev_c0));
	HANDLE_ERROR(cudaFree(dev_a1));
	HANDLE_ERROR(cudaFree(dev_b1));
	HANDLE_ERROR(cudaFree(dev_c1));

	HANDLE_ERROR(cudaStreamDestroy(stream0));
	HANDLE_ERROR(cudaStreamDestroy(stream1));
	return 0;
}

運作結果

CUDA By Example(八)——流頁鎖定主機記憶體CUDA流使用單個CUDA流使用多個CUDA流GPU的工作排程機制高效地使用多個CUDA流遇到的問題(未解決)

遇到的問題(未解決)

我發現使用多個流實際并沒有産生運作速度的提升，我試了單個流、兩個流、四個流發現消耗的時間基本沒有差别，暫時不知道是什麼原因導緻的，歡迎大佬解答。

GPU高性能程式設計CUDA實戰 cuda 并行計算 GPU高性能程式設計
上一篇: 假如你對你的設計師不滿意@卡西
下一篇: CUDA By Example(七)——原子性原子操作簡介計算直方圖更多相關内容

繼續閱讀