CUDA的併發性創建的流

Question

我以這種方式創建的流默認流：

cudaStream_t stream0; 
cudaStream_t stream1; 
cudaStreamCreate(&stream0); 
cudaStreamCreate(&stream1); 

我運行內核的功能，如

singlecore<<<1,1>>>(devL2,1000); 
singlecore<<<1,1,0,stream0>>>(devL2,1000); 

兩個內核目前尚未執行。但是，如果我在stream1執行第一內核：

singlecore<<<1,1,0,stream1>>>(devL2,1000); 
singlecore<<<1,1,0,stream0>>>(devL2,1000); 

，他們將目前執行。

我不知道默認流中的內核函數當前是不是可以執行的。

Answer 1

是對發佈到默認流的cuda命令有限制。參照C編程指南節上implicit synchronization： ... •任何CUDA命令：

「來自不同流的兩個命令不能同時如果以下任一項操作在兩者之間是向他們發出主機線程中運行到默認流， 「

因此，作爲一般經驗法則，對於重疊的複製和計算操作，最簡單的方法是將所有這些操作編程爲一組非默認流。有一些漏洞（你已經發現）可能與默認流（和其他流）中發佈的命令重疊，但它需要仔細瞭解默認流和其他流之間的限制，以及請注意您發佈命令的順序。在C編程指南中解釋了good example。閱讀「重疊行爲」部分的所有內容。

在第一個例子中，發給默認流的內核阻塞了發給其他流的內核的執行。在第二個示例中，您可以具有併發性，因爲發佈到非默認流的內核不會阻止發佈到默認流的內核的執行。

Answer 2

我想根據最新發布的CUDA 7.0來更新Robert Crovella的答案，該版本截至2015年3月，發佈候選版本。

隨着CUDA 7.0，默認流是常規的流在這個意義上，在默認流命令可以在非默認流命令同時運行。這一新功能的更詳細的解釋可以在

CUDA 7 Streams Simplify Concurrency

此功能被發現，可通過附加--default stream per-thread編譯選項可以方便地實現。

在上面鏈接的頁面上，可以找到Mark Harris制定的一個例子。在這裏，我想恢復我在False dependency issue for the Fermi architecture發佈的示例。特別是，在下面的新示例中，儘管我創建了3流，但我不再使用第一個流，並在其位置採用默認流。

這是時間軸製作沒有的--default stream per-thread編譯選項：

正如你所看到的，在默認流的執行沒有利用併發。

在此另一邊，這是的--default stream per-thread編譯選項生成的時間表：

正如你可以看到現在，默認流執行與其他兩個流的執行重疊。

#include <iostream> 

#include "cuda_runtime.h" 
#include "device_launch_parameters.h" 

#include <stdio.h> 

#include "Utilities.cuh" 

using namespace std; 

#define NUM_THREADS 32 
#define NUM_BLOCKS 16 
#define NUM_STREAMS 3 

__global__ void kernel(const int *in, int *out, int N) 
{ 
    int start = blockIdx.x * blockDim.x + threadIdx.x; 
    int end = N; 
    for (int i = start; i < end; i += blockDim.x * gridDim.x) 
    { 
     out[i] = in[i] * in[i]; 
    } 
} 

int main() 
{ 
    const int N = 6000000; 

    // --- Host side input data allocation and initialization. Registering host memory as page-locked (required for asynch cudaMemcpyAsync). 
    int *h_in = new int[N]; for(int i = 0; i < N; i++) h_in[i] = 5; 
    gpuErrchk(cudaHostRegister(h_in, N * sizeof(int), cudaHostRegisterPortable)); 

    // --- Host side input data allocation and initialization. Registering host memory as page-locked (required for asynch cudaMemcpyAsync). 
    int *h_out = new int[N]; for(int i = 0; i < N; i++) h_out[i] = 0; 
    gpuErrchk(cudaHostRegister(h_out, N * sizeof(int), cudaHostRegisterPortable)); 

    // --- Host side check results vector allocation and initialization 
    int *h_checkResults = new int[N]; for(int i = 0; i < N; i++) h_checkResults[i] = h_in[i] * h_in[i]; 

    // --- Device side input data allocation. 
    int *d_in = 0;    gpuErrchk(cudaMalloc((void **)&d_in, N * sizeof(int))); 

    // --- Device side output data allocation. 
    int *d_out = 0;    gpuErrchk(cudaMalloc((void **)&d_out, N * sizeof(int))); 

    int streamSize = N/NUM_STREAMS; 
    size_t streamMemSize = N * sizeof(int)/NUM_STREAMS; 

    // --- Set kernel launch configuration 
    dim3 nThreads  = dim3(NUM_THREADS,1,1); 
    dim3 nBlocks  = dim3(NUM_BLOCKS, 1,1); 
    dim3 subKernelBlock = dim3((int)ceil((float)nBlocks.x/2)); 

    // --- Create CUDA streams 
    cudaStream_t streams[NUM_STREAMS]; 
    for(int i = 0; i < NUM_STREAMS; i++) 
     gpuErrchk(cudaStreamCreate(&streams[i])); 

    /**************************/ 
    /* BREADTH-FIRST APPROACH */ 
    /**************************/ 

    int offset = 0; 
    cudaMemcpyAsync(&d_in[offset], &h_in[offset], streamMemSize, cudaMemcpyHostToDevice,  0); 
    for(int i = 1; i < NUM_STREAMS; i++) { 
     int offset = i * streamSize; 
     cudaMemcpyAsync(&d_in[offset], &h_in[offset], streamMemSize, cudaMemcpyHostToDevice,  streams[i]); 
    } 

    kernel<<<subKernelBlock, nThreads>>>(&d_in[offset], &d_out[offset], streamSize/2); 
    kernel<<<subKernelBlock, nThreads>>>(&d_in[offset + streamSize/2], &d_out[offset + streamSize/2], streamSize/2); 

    for(int i = 1; i < NUM_STREAMS; i++) 
    { 
     int offset = i * streamSize; 
     kernel<<<subKernelBlock, nThreads, 0, streams[i]>>>(&d_in[offset], &d_out[offset], streamSize/2); 
     kernel<<<subKernelBlock, nThreads, 0, streams[i]>>>(&d_in[offset + streamSize/2], &d_out[offset + streamSize/2], streamSize/2); 
    } 

    for(int i = 1; i < NUM_STREAMS; i++) { 
     int offset = i * streamSize; 
     cudaMemcpyAsync(&h_out[offset], &d_out[offset], streamMemSize, cudaMemcpyDeviceToHost, streams[i]); 
    } 

    cudaMemcpyAsync(&h_out[offset], &d_out[offset], streamMemSize, cudaMemcpyDeviceToHost, 0); 
    for(int i = 1; i < NUM_STREAMS; i++) { 
     int offset = i * streamSize; 
     cudaMemcpyAsync(&h_out[offset], &d_out[offset], streamMemSize, cudaMemcpyDeviceToHost, 0); 
    } 

    for(int i = 0; i < NUM_STREAMS; i++) 
     gpuErrchk(cudaStreamSynchronize(streams[i])); 

    gpuErrchk(cudaDeviceSynchronize()); 

    // --- Release resources 
    gpuErrchk(cudaHostUnregister(h_in)); 
    gpuErrchk(cudaHostUnregister(h_out)); 
    gpuErrchk(cudaFree(d_in)); 
    gpuErrchk(cudaFree(d_out)); 

    for(int i = 0; i < NUM_STREAMS; i++) 
     gpuErrchk(cudaStreamDestroy(streams[i])); 

    cudaDeviceReset(); 

    // --- GPU output check 
    int sum = 0; 
    for(int i = 0; i < N; i++)  
     sum += h_checkResults[i] - h_out[i]; 

    cout << "Error between CPU and GPU: " << sum << endl; 

    delete[] h_in; 
    delete[] h_out; 
    delete[] h_checkResults; 

    return 0; 
}