您當前位置：首頁 > 服務器 > 《GPU高性能編程CUDA實戰》學習筆記（十）

《GPU高性能編程CUDA實戰》學習筆記（十）

來源：程序員人生發布時間：2016-10-05 09:04:59 閱讀次數：3679次

第10章流

GPU上履行大范圍數據并行計算性能高于cpu上履行，另外，NVIDIA 圖形處理器還支持1種并行性（Parallelism）。這類并行性類似于cpu多線程利用程序中的任務并行性（Task Parallelism）。任務并行性是指并行履行兩個或多個不同的任務，而不是在大量數據上履行同1個任務。

在并行環境中，任務可以是任意操作。如1個線程繪制GUI，另外一個線程通過網絡下載更新包。本章介紹 CUDA 流，和如何通過流在GPU上同時履行多個任務。

10.1 本章目標

了解如何分配頁鎖定（Page-Locked）類型的主機內存。
了解CUDA流的概念。
了解如何使用CUDA 流來加速利用程序。

10.2 頁鎖定主機內存

C 庫分配主機內存使用 malloc() ；CUDA也能夠分配主機內存，使用 cudaHostAlloc() 。

差異：malloc()分配的內存是標準的、可分頁的（Pagable）主機內存；cudaHostAlloc()分配的內存是頁鎖定的主機內存。

頁鎖定內存：也稱固定內存（Pinned Memory）或不可分頁內存，它有1個重要屬性------操作系統將不會對這塊內存分頁并交換到磁盤上，從而確保了該內存始終駐留在物理內存中。因此，操作系統能夠安全地使某個利用程序訪問該內存的物理地址，由于這塊內存將不會被破壞或重新定位。

由于gpu知道內存的物理地址，因此可以通過“直接內存訪問（Direct Memory Access，DMA）” 技術來在gpu和主機之間復制數據。由于DMA 在履行復制時無需cpu參與，這就說明 cpu極可能在dma的履行進程中將目標內存交換到磁盤上，或通過更新操作系統的可分頁表來重新定位目標內存的物理地址。cpu可能會移動可分頁的數據，這便可能對dma操作造成延遲。因此，固定內存很重要。

固定內存是雙刃劍，所以建議：僅對 cudaMemcpy() 調用中的源內存或目標內存，才使用頁鎖定內存，并且不再需要時立即釋放。

#include "../common/book.h" #define SIZE (64*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( *a ), cudaMemcpyHostToDevice ) ); else HANDLE_ERROR( cudaMemcpy( a, dev_a, size * sizeof( *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; // try it with cudaMalloc elapsedTime = cuda_malloc_test( SIZE, true ); printf( "Time using cudaMalloc: %3.1f ms\n", elapsedTime ); printf( "\tMB/s during copy up: %3.1f\n", MB/(elapsedTime/1000) ); elapsedTime = cuda_malloc_test( SIZE, false ); printf( "Time using cudaMalloc: %3.1f ms\n", elapsedTime ); printf( "\tMB/s during copy down: %3.1f\n", MB/(elapsedTime/1000) ); // now try it with cudaHostAlloc elapsedTime = cuda_host_alloc_test( SIZE, true ); printf( "Time using cudaHostAlloc: %3.1f ms\n", elapsedTime ); printf( "\tMB/s during copy up: %3.1f\n", MB/(elapsedTime/1000) ); elapsedTime = cuda_host_alloc_test( SIZE, false ); printf( "Time using cudaHostAlloc: %3.1f ms\n", elapsedTime ); printf( "\tMB/s during copy down: %3.1f\n", MB/(elapsedTime/1000) ); }

cudaHostAlloc() --- cudaFreeHost()

通過履行上面代碼，可以看到性能提升。

10.3 CUDA 流

cudaEventRecord( , ) CUDA事件的第2個參數就是用于指定插入事件的流（Stream）。

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

CUDA 流在加速利用程序方面起側重要的作用，CUDA 流表示1個GPU 操作隊列，并且該隊列中的操作將以指定的順序履行。我們可以在流中添加1些操作，例如核函數的啟動、內存復制，和事件的啟動和結束等。將這些操作添加到流的順序也是它們的履行順序。可以將每一個流視為GPU上的1個任務，并且這些任務可以并行履行。下面介紹使用和加速利用程序。

10.4 使用單個CUDA流

#include "../common/book.h" #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) { printf( "Device will not handle overlaps, so no speed up from streams\n" ); return 0; } cudaEvent_t start, stop; float elapsedTime; cudaStream_t stream; int *host_a, *host_b, *host_c; int *dev_a, *dev_b, *dev_c; // start the timers HANDLE_ERROR( cudaEventCreate( &start ) ); HANDLE_ERROR( cudaEventCreate( &stop ) ); // initialize the stream HANDLE_ERROR( cudaStreamCreate( &stream ) ); // allocate the memory on the 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) ) ); // allocate host locked memory, used to stream 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(); } HANDLE_ERROR( cudaEventRecord( start, 0 ) ); // now loop over full data, in bite-sized chunks for (int i=0; i<FULL_DATA_SIZE; i+= N) { // copy the locked memory to the device, async 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 ); // copy the data from device to locked memory HANDLE_ERROR( cudaMemcpyAsync( host_c+i, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost, stream ) ); } // copy result chunk from locked to full buffer HANDLE_ERROR( cudaStreamSynchronize( stream ) ); HANDLE_ERROR( cudaEventRecord( stop, 0 ) ); HANDLE_ERROR( cudaEventSynchronize( stop ) ); HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime, start, stop ) ); printf( "Time taken: %3.1f ms\n", elapsedTime ); // cleanup the streams and memory 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; }

這里是單個流來講明流的用法，主要看main函數。

第1件事：選擇1個支持裝備堆疊（Device Overlap）功能的裝備。支持裝備堆疊功能的GPU能夠在履行1個CUDA C核函數的同時，還能在裝備與主機之間履行復制操作。正如前面說的，我們將使用多個流來實現這類計算與數據傳輸的堆疊。

（1）創建流，

cudaStream_t stream; // initialize the stream HANDLE_ERROR( cudaStreamCreate( &stream ) );

（2）數據分配，

int *host_a, *host_b, *host_c; int *dev_a, *dev_b, *dev_c; // allocate the memory on the 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) ) ); // allocate host locked memory, used to stream 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(); }

cudaHostAlloc() 使用主機固定內存，除復制快，其他好處1會分析。

1般情況，我們是將輸入緩沖區復制到GPU，啟動核函數，然后將輸出緩沖區復制回主機；不過這次我們有些改動。

1.將輸入緩沖區劃分為更小的塊，并在每一個塊上履行1個包括3個步驟的進程；

2.將1部份輸入緩沖區復制到GPU，這部份緩沖區運行核函數；

3.將輸出緩沖區中的這部份結果復制回主機。

使用情形： GPU的內存遠小于主機內存，全部緩沖區沒法1次性填充到GPU，因此需要分塊進行計算。

// now loop over full data, in bite-sized chunks for (int i=0; i<FULL_DATA_SIZE; i+= N) { // copy the locked memory to the device, async 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 ); // copy the data from device to locked memory HANDLE_ERROR( cudaMemcpyAsync( host_c+i, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost, stream ) ); }

cudaMemcpy() 這個函數將以同步方式履行，說明，當函數返回時，復制操作已完成，并且在輸出緩沖區中包括了復制進去的內容。

cudaMemcpyAsync() 異步，只是放置1個要求，表示在流中履行1次內存復制操作，這個流是通過參數stream來指定的。任何傳遞給cudaMemcpyAsync() 的主機內存指針必須已通過cudaHostAlloc() 分配好內存。只能以異步方式對頁鎖定內存進行復制操作。

10.5 使用多個CUDA 流

在任何支持內存復制和核函數的履行相互堆疊的裝備上，當使用多個流時，利用程序的整體性能都會提示。

#include "../common/book.h" #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) { printf( "Device will not handle overlaps, so no speed up from streams\n" ); return 0; } cudaEvent_t start, stop; float elapsedTime; cudaStream_t stream0, stream1; int *host_a, *host_b, *host_c; int *dev_a0, *dev_b0, *dev_c0; int *dev_a1, *dev_b1, *dev_c1; // start the timers HANDLE_ERROR( cudaEventCreate( &start ) ); HANDLE_ERROR( cudaEventCreate( &stop ) ); // initialize the streams HANDLE_ERROR( cudaStreamCreate( &stream0 ) ); HANDLE_ERROR( cudaStreamCreate( &stream1 ) ); // allocate the memory on the 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) ) ); // allocate host locked memory, used to stream 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(); } HANDLE_ERROR( cudaEventRecord( start, 0 ) ); // now loop over full data, in bite-sized chunks for (int i=0; i<FULL_DATA_SIZE; i+= N*2) { // copy the locked memory to the device, async 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 ); // copy the data from device to locked memory HANDLE_ERROR( cudaMemcpyAsync( host_c+i, dev_c0, N * sizeof(int), cudaMemcpyDeviceToHost, stream0 ) ); // copy the locked memory to the device, async 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 ); // copy the data from device to locked memory 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 ) ); printf( "Time taken: %3.1f ms\n", elapsedTime ); // cleanup the streams and memory 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; }

10.6 GPU的工作調度機制

流可以看作是：有序的操作序列，其中包括內存復制操作、核函數調用。硬件中沒有流的概念，而是包括1個或多個引擎來履行內存復制操作，和1個引擎來履行核函數。這些引擎彼此獨立的對操作進行排隊。

10.7 高效地使用多個CUDA流

如果同時調度某個流的所有操作，那末容易在無意中阻塞另外一個流的復制操作或核函數履行。解決這個問題，在將操作放入流的隊列時應當采取寬度優先方式，而非深度優先方式。也就是說，將這兩個流之間的操作交叉添加。

#include "../common/book.h" #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) { printf( "Device will not handle overlaps, so no speed up from streams\n" ); return 0; } cudaEvent_t start, stop; float elapsedTime; cudaStream_t stream0, stream1; int *host_a, *host_b, *host_c; int *dev_a0, *dev_b0, *dev_c0; int *dev_a1, *dev_b1, *dev_c1; // start the timers HANDLE_ERROR( cudaEventCreate( &start ) ); HANDLE_ERROR( cudaEventCreate( &stop ) ); // initialize the streams HANDLE_ERROR( cudaStreamCreate( &stream0 ) ); HANDLE_ERROR( cudaStreamCreate( &stream1 ) ); // allocate the memory on the 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) ) ); // allocate host locked memory, used to stream 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(); } HANDLE_ERROR( cudaEventRecord( start, 0 ) ); // now loop over full data, in bite-sized chunks for (int i=0; i<FULL_DATA_SIZE; i+= N*2) { // enqueue copies of a in stream0 and 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 ) ); // enqueue copies of b in stream0 and 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 ) ); // enqueue kernels in stream0 and 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 ); // enqueue copies of c from device to locked memory 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 ) ); printf( "Time taken: %3.1f ms\n", elapsedTime ); // cleanup the streams and memory 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; }

這個速度更快。

生活不易，碼農辛苦
如果您覺得本網站對您的學習有所幫助,可以手機掃描二維碼進行捐贈
程序員人生