CUDA Dynamic Parallelizm;從設備流同步

Question

我基本上尋找一種方法來同步設備內的流。我想避免使用cudaDeviceSynchronize（），因爲它會序列化我想要使用流同時執行的內核的執行;

更詳細的描述：我寫了一個內核，這是一個穩定的雙共軛梯度解算器。我想在不同的數據流上同時吃這個內核。

該內核使用cublas函數。它們在內核中被調用。

解算器需要的操作之一是計算兩個向量的點積。這可以用cublasdot（）完成。但是由於這個調用是同步的，因此不同流中的內核的執行會被序列化。我不用調用點積函數，而是使用cublasspmv（）來計算點積，這是異步調用的。問題是這個函數在結果計算之前返回。因此，我想要同步來自設備的流 - 我正在尋找相當於cudaStreamSynchronize（）但可從設備調用的流。

__device__ float _cDdot(cublasHandle_t & cublasHandle, const int n, real_t * x, real_t * y) { 
     float *norm; norm = new float; 
     float alpha = 1.0f; float beta = 0.0f; 

     cublasSgemv_v2(cublasHandle, CUBLAS_OP_N ,1 , n, &alpha, x, 1, y, 1, &beta, norm, 1); 

     return *norm; 
} 

我能做些什麼來確保結果是在函數返回之前計算的？當然，cudaDeviceSynchronize（）的插入工作，但正如我所提到的，它串行化我的內核跨流的執行。

感謝， 馬辛

Answer 1

也許如果你讀the programming guide dynamic parallelism section仔細（特別是流，事件和同步），你可能會得到一些想法。以下是我想出了：

沒有與調用您_cDdot函數的執行順序（相關的隱式空流（設備）名字古怪的，恕我直言，因爲你與float大量的工作在這種情況下，即使用Sgemv）。因此，在您的函數中調用cublasSgemv_v2後發出的任何cuda內核或API調用都應該等到與該函數關聯的任何cuda活動完成。如果在調用cublasSgemv_v2之後插入無害的cuda API調用或虛擬內核調用，它應該等待完成。這應該會給你以後的線程級同步。您也可以使用呼叫cudaEventRecord，然後撥打cudaStreamWaitEvent。

下面是一個例子來說明隱含流同步的方法：

nvcc -arch=sm_35 -rdc=true -o t302 t302.cu -lcudadevrt -lcublas -lcublas_device 

結果：

#include <stdio.h> 
#include <cublas_v2.h> 
#define SZ 16 

__global__ void dummy_kernel(float *in, float *out){ 
    *out = *in; 
} 

__device__ float _cDdot(cublasHandle_t & cublasHandle, const int n, float * x, float * y, const int wait) { 
     float *norm; norm = new float; 
     float alpha = 1.0f; float beta = 0.0f; 
     *norm = 0.0f; 
     cublasSgemv_v2(cublasHandle, CUBLAS_OP_N ,1 , n, &alpha, x, 1, y, 1, &beta, norm, 1); 
     if (wait){ 
     dummy_kernel<<<1,1>>>(norm, norm); 
     } 
     return *norm; 
} 


__global__ void compute(){ 
    cublasHandle_t my_h; 
    cublasStatus_t status; 
    status = cublasCreate(&my_h); 
    if (status != CUBLAS_STATUS_SUCCESS) printf("cublasCreate fail\n"); 
    float *x, *y; 
    x = new float[SZ]; 
    y = new float[SZ]; 
    for (int i = 0; i < SZ; i++){ 
    x[i] = 1.0f; 
    y[i] = 1.0f;} 
    float result = _cDdot(my_h, SZ, x, y, 0); 
    printf("result with no wait = %f\n", result); 
    result = _cDdot(my_h, SZ, x, y, 1); 
    printf("result with wait = %f\n", result); 
} 

int main(){ 

    compute<<<1,1>>>(); 
    cudaDeviceSynchronize(); 
    return 0; 
} 

與編譯

$ ./t302 
result with no wait = 0.000000 
result with wait = 16.000000 
$ 

不幸的是我嘗試了完全空dummy_kernel;這不起作用，除非我編譯-G。所以編譯器可以足夠聰明地優化一個完整的空子內核調用。