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。所以编译器可以足够聪明地优化一个完整的空子内核调用。