cuda內核不能同時執行

Question

我試圖探索我的Nvidia Quadro 4000的併發內核執行屬性，它具有2.0的能力。

我用2個不同的數據流，其運行相同如下：

1. 複製H2D 2個鎖定的存儲
2. 運行內核
3. Copyt D2H兩個不同的區塊，以固定內存的不同塊。

兩個流的內核完全相同，每個執行時間爲190毫秒。

在Visual分析器（版本5.0）我預期兩邊的內核以同時開始執行，然而，它們由20毫秒僅重疊。 這裏是代碼示例：

enter code here 

//initiate the streams 
     cudaStream_t stream0,stream1; 
     CHK_ERR(cudaStreamCreate(&stream0)); 
     CHK_ERR(cudaStreamCreate(&stream1)); 
     //allocate the memory on the GPU for stream0 
     CHK_ERR(cudaMalloc((void **)&def_img0, width*height*sizeof(char))); 
     CHK_ERR(cudaMalloc((void **)&ref_img0, width*height*sizeof(char))); 
     CHK_ERR(cudaMalloc((void **)&outY_img0,width_size_for_out*height_size_for_out*sizeof(char))); 
     CHK_ERR(cudaMalloc((void **)&outX_img0,width_size_for_out*height_size_for_out*sizeof(char))); 
     //allocate the memory on the GPU for stream1 
     CHK_ERR(cudaMalloc((void **)&def_img1, width*height*sizeof(char))); 
     CHK_ERR(cudaMalloc((void **)&ref_img1, width*height*sizeof(char))); 
     CHK_ERR(cudaMalloc((void **)&outY_img1,width_size_for_out*height_size_for_out*sizeof(char))); 
     CHK_ERR(cudaMalloc((void **)&outX_img1,width_size_for_out*height_size_for_out*sizeof(char))); 

     //allocate page-locked memory for stream0 
     CHK_ERR(cudaHostAlloc((void**)&host01, width*height*sizeof(char), cudaHostAllocDefault)); 
     CHK_ERR(cudaHostAlloc((void**)&host02, width*height*sizeof(char), cudaHostAllocDefault)); 
     CHK_ERR(cudaHostAlloc((void**)&host03, width_size_for_out*height_size_for_out*sizeof(char), cudaHostAllocDefault)); 
     CHK_ERR(cudaHostAlloc((void**)&host04, width_size_for_out*height_size_for_out*sizeof(char), cudaHostAllocDefault)); 

     //allocate page-locked memory for stream1 
     CHK_ERR(cudaHostAlloc((void**)&host11, width*height*sizeof(char), cudaHostAllocDefault)); 
     CHK_ERR(cudaHostAlloc((void**)&host12, width*height*sizeof(char), cudaHostAllocDefault)); 
     CHK_ERR(cudaHostAlloc((void**)&host13, width_size_for_out*height_size_for_out*sizeof(char), cudaHostAllocDefault)); 
     CHK_ERR(cudaHostAlloc((void**)&host14, width_size_for_out*height_size_for_out*sizeof(char), cudaHostAllocDefault)); 


     memcpy(host01,in1,width*height*sizeof(char)); 
     memcpy(host02,in2,width*height*sizeof(char)); 

     memcpy(host11,in1,width*height*sizeof(char)); 
     memcpy(host12,in2,width*height*sizeof(char)); 



     cudaEvent_t start, stop; 
     float time; 
     cudaEventCreate(&start); 
     cudaEventCreate(&stop); 

     dim3 dimBlock(CUDA_BLOCK_DIM, CUDA_BLOCK_DIM); 
     dim3 Grid((width-SEARCH_RADIUS*2-1)/(dimBlock.x*4)+1, (height-SEARCH_RADIUS*2-1)/(dimBlock.y*4)+1); 

     cudaEventRecord(start,0); 
     // -------------------- 
     // Copy images to device 
     // -------------------- 
     //enqueue copies of def stream0 and stream1 
     CHK_ERR(cudaMemcpyAsync(def_img0, host01,width*height*sizeof(char), cudaMemcpyHostToDevice,stream0)); 
     CHK_ERR(cudaMemcpyAsync(def_img1, host11,width*height*sizeof(char), cudaMemcpyHostToDevice,stream1)); 
     //enqueue copies of ref stream0 and stream1 
     CHK_ERR(cudaMemcpyAsync(ref_img0, host02,width*height*sizeof(char), cudaMemcpyHostToDevice,stream0)); 
     CHK_ERR(cudaMemcpyAsync(ref_img1, host12,width*height*sizeof(char), cudaMemcpyHostToDevice,stream1)); 

     CHK_ERR(cudaStreamSynchronize(stream0)); 
     CHK_ERR(cudaStreamSynchronize(stream1)); 

     //CALLING KERNEL 
     //enqueue kernel in stream0 and stream1 
     TIME_KERNEL((exhaustiveSearchKernel<CUDA_BLOCK_DIM*4,CUDA_BLOCK_DIM*4,SEARCH_RADIUS><<< Grid, dimBlock,0,stream0>>>(def_img0+SEARCH_RADIUS*width+SEARCH_RADIUS,ref_img0,outX_img0,outY_img0,width,width_size_for_out)),"exhaustiveSearchKernel stream0"); 
     TIME_KERNEL((exhaustiveSearchKernel<CUDA_BLOCK_DIM*4,CUDA_BLOCK_DIM*4,SEARCH_RADIUS><<< Grid, dimBlock,0,stream1>>>(def_img1+SEARCH_RADIUS*width+SEARCH_RADIUS,ref_img1,outX_img1,outY_img1,width,width_size_for_out)),"exhaustiveSearchKernel stream1"); 


     //Copy result back 
     CHK_ERR(cudaMemcpyAsync(host03, outX_img0, width_size_for_out*height_size_for_out*sizeof(char), cudaMemcpyDeviceToHost,stream0)); 
     CHK_ERR(cudaMemcpyAsync(host13, outX_img1, width_size_for_out*height_size_for_out*sizeof(char), cudaMemcpyDeviceToHost,stream1)); 

     CHK_ERR(cudaMemcpyAsync(host04, outY_img0, width_size_for_out*height_size_for_out*sizeof(char), cudaMemcpyDeviceToHost,stream0)); 
     CHK_ERR(cudaMemcpyAsync(host14, outY_img1, width_size_for_out*height_size_for_out*sizeof(char), cudaMemcpyDeviceToHost,stream1)); 


     CHK_ERR(cudaStreamSynchronize(stream0)); 
     CHK_ERR(cudaStreamSynchronize(stream1)); 

     cudaEventRecord(stop, 0); 
     cudaEventSynchronize(stop); 
     cudaEventElapsedTime(&time, start, stop); 
     printf("Elapsed time=%f ms\n",time); 

     memcpy(outX,host03,width_size_for_out*height_size_for_out*sizeof(char)); 
     memcpy(outY,host04,width_size_for_out*height_size_for_out*sizeof(char)); 


     cudaEventDestroy(start); 
     cudaEventDestroy(stop); 
     CHK_ERR(cudaStreamDestroy(stream0)); 
     CHK_ERR(cudaStreamDestroy(stream1)); 

     CHK_ERR(cudaDeviceReset()); 


    } 

Answer 1

計算能力2.x的-3.0

計算能力2.x的-3.0設備具有單一的硬件工作隊列。 CUDA驅動程序將命令推入工作隊列。 GPU主機讀取命令並將工作分發給複製引擎或CUDA工作分配器（CWD）。 CUDA驅動程序將同步命令插入到硬件工作隊列中，以確保同一個流上的工作不能同時運行。當主機點擊一個同步命令時，它將停止，直到相關工作完成。

同時內核執行提高GPU利用率當電網太小，不能填滿整個GPU或當網格具有尾效應（線程塊的子集執行遠長於其他線程塊）。

案例1：背靠背在一個流

內核如果一個應用推出兩款kernesl回到同一流回來CUDA驅動插入同步命令將不會派遣第二個內核CWD直到第一個內核已經完成。

案例2：背靠背內核啓動兩個流

如果一個應用程序啓動對不同的流兩個內核主機將讀取命令和調度命令CWD。 CWD將柵格化第一個網格（順序取決於體系結構）並將線程塊分派給SM。只有當所有來自第一個網格的線程塊都被分派後，CWD纔會從第二個網格分派線程塊。

計算能力3.5

計算能力3.5推出了一些新的功能，以提高GPU利用率。這些包括： - HyperQ支持多個獨立的硬件工作隊列。 - Dynamic Parallelism允許設備代碼啓動新工作。 - CWD容量增加到32個網格。

資源

• CUDA C/C++ Streams and Concurrency [webinar]
• NVIDIA Fermi Whitepaper
• NVIDIA Kepler GK110 Architecture Whitepaper