0

基本的に、デバイス内からストリームを同期する方法を探しています。ストリームを使用して同時に実行したいカーネルの実行をシリアル化するため、cudaDeviceSynchronize() の使用は避けたいです。

より詳細な説明: 安定化された双共役勾配ソルバーであるカーネルを作成しました。ストリームを使用して、このカーネルを異なるデータで同時に起動したいと考えています。

このカーネルは cublas 関数を使用します。それらはカーネル内から呼び出されます。

ソルバーが必要とする操作の 1 つは、2 つのベクトルの内積の計算です。これは 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ため、奇妙に名前が付けられています)。したがって、関数内で への呼び出しの後に発行された cuda カーネルまたは API 呼び出しは、関数に関連付けられた cuda アクティビティが完了するまで待機する必要があります。への呼び出しの後に無害な cuda API 呼び出し、またはダミーのカーネル呼び出しを挿入すると、それが完了するまで待機する必要があります。これにより、目的のスレッドレベルの同期が得られるはずです。呼び出しの後に呼び出しを使用することもできます。floatSgemvcublasSgemv_v2cublasSgemv_v2cublasSgemv_v2cudaEventRecordcudaStreamWaitEvent

暗黙的なストリーム同期アプローチを示す例を次に示します。

#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 に答える