0

我基本上是在寻找一种从设备内部同步流的方法。我想避免使用 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() 是可行的，但正如我所提到的，它会跨流序列化我的内核的执行。

4

1 回答 1

1

可能如果您仔细阅读编程指南动态并行部分（尤其是流、事件和同步），您可能会得到一些想法。这是我想出的：

有一个隐式 NULL 流（在设备上）与调用您的_cDdot函数的执行序列相关联（奇怪的是，恕我直言，因为在这种情况下您正在使用float数量，即使用Sgemv）。因此，在您的函数中调用之后发出的任何 cuda 内核或 API 调用cublasSgemv_v2都应该等到与该cublasSgemv_v2函数关联的任何 cuda 活动完成。如果你在调用之后插入一个无害的 cuda API 调用，或者一个虚拟内核调用cublasSgemv_v2，它应该等待它完成。这应该为您提供您所追求的线程级同步。您可能还可以使用一个cudaEventRecord呼叫，然后是一个cudaStreamWaitEvent呼叫。

这是一个显示隐式流同步方法的示例：

#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;
}

编译：

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

结果：

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

不幸的是，我尝试了一个完全空的dummy_kernel；那没有用，除非我用-G. 所以编译器可能足够聪明，可以优化出一个完整的空子内核调用。

于 2013-12-19T06:32:21.083 回答