7

フォローアップ Q から: CUDA: Calling a __device__ function from a kernel

ソート操作を高速化しようとしています。簡略化された疑似バージョンは次のとおりです。

// some costly swap operation
__device__ swap(float* ptrA, float* ptrB){
  float saveData;         // swap some 
  saveData= *Adata;       //   big complex
  *Adata= *Bdata          //     data chunk
  *Bdata= saveData;
}

// a rather simple sort operation
__global__ sort(float data[]){
  for (i=0; i<limit: i++){
  find left swap point
  find right swap point
  swap<<<1,1>>>(left, right);
  }
}

(注: この単純なバージョンでは、ブロック内のリダクション テクニックは示されていません。) アイデアは、スワップ ポイントを簡単 (高速) に特定できるようにすることです。スワップ操作はコストがかかります (遅い)。したがって、1 つのブロックを使用して、スワップ ポイントを見つけて識別します。他のブロックを使用してスワップ操作を行います。つまり、実際のスワッピングを並行して行います。これはまともな計画のように聞こえます。ただし、コンパイラがデバイス呼び出しをインライン化する場合、並列スワッピングは行われません。デバイス呼び出しをインライン化しないようにコンパイラに指示する方法はありますか?

4

2 に答える 2

8

この質問は久しぶりです。同じ問題をグーグルで検索したところ、このページにたどり着きました。私は解決策を得たようです。

解決：

何とか[ここ][1]にたどり着き、別のカーネル内からカーネルを起動するクールなアプローチを見ました。

__global__ void kernel_child(float *var1, int N){
    //do data operations here
}


__global__ void kernel_parent(float *var1, int N)
{
    kernel_child<<<1,2>>>(var1,N);
}

cuda 5.0 以降の動的並列処理により、これが可能になりました。また、実行中は必ずcompute_35 以上のアーキテクチャを使用してください。

ターミナルの方法 上記の親カーネル (最終的には子カーネルを実行します) をターミナルから実行できます。Linux マシンで検証済み。

$ nvcc -arch=sm_35 -rdc=true yourFile.cu
$ ./a.out

それが役に立てば幸い。ありがとうございました！[1]: http://developer.download.nvidia.com/assets/cuda/docs/TechBrief_Dynamic_Parallelism_in_CUDA_v2.pdf

于 2016-03-11T06:02:51.517 に答える
5

編集 (2016):

動的並列処理は、第 2 世代の Kepler アーキテクチャ GPU で導入されました。デバイスでのカーネルの起動は、コンピューティング機能 3.5 以降のデバイスでサポートされています。

元の回答:

次世代のハードウェアが利用可能になる年末まで待つ必要があります。現在の CUDA デバイスは他のカーネルからカーネルを起動できません - 現在サポートされていません。

于 2012-07-31T19:15:38.423 に答える