0

いくつかの CUDA カーネルを使用するプログラムを持っていますが、これは実行に 50 ～ 100 ミリ秒かかりますが、他のプログラムは 0 ～ 5 ミリ秒かかります。これはすべての分岐と関係があると思いますが、それを減らす方法はよくわかりません。計算機能 2.1 デバイス用にコンパイルしています。誰かが私を正しい方向に向けることができれば、それは素晴らしいことです.

// chosen using occupancy spreadsheet
#define SCORE_THREADS_PER_BLOCK 448

__device__ double ScoringMatrixVal(double *scoring_matrix, size_t pitch, unsigned int row, unsigned int column) {
  return *((double*)((char*) scoring_matrix + row * pitch) + column);
}

__global__ void ScoreBindingSites(char *input_sequence, unsigned long is_length, unsigned int *rvd_sequence, unsigned int rs_len, double cutoff, unsigned int rvd_num, double *scoring_matrix, size_t sm_pitch, unsigned char *results) {

  int block_seq_index = SCORE_THREADS_PER_BLOCK * (blockIdx.y * gridDim.x + blockIdx.x);
  int thread_id = (blockDim.x * threadIdx.y) + threadIdx.x;
  int seq_index = block_seq_index + thread_id;

  if (seq_index < 1 || seq_index >= is_length || seq_index + rs_len >= is_length - 1) return;

  if (input_sequence[seq_index - 1] == 'T' || input_sequence[seq_index - 1] == 't') {

    double thread_result = 0;

    for (int i = 0; i < rs_len; i++) {

      int rvd_index = i;

      int sm_col = 4;

      char base = input_sequence[seq_index + i];

      if (base == 'A' || base == 'a')    
        sm_col = 0;
      if (base == 'C' || base == 'c')
        sm_col = 1;
      if (base == 'G' || base == 'g')
        sm_col = 2;
      if (base == 'T' || base == 't')
        sm_col = 3;

      thread_result += ScoringMatrixVal(scoring_matrix, sm_pitch, rvd_sequence[rvd_index], sm_col);

    }

    results[seq_index] |= (thread_result < cutoff ? 1UL : 0UL) << (2 * rvd_num);

  } 

  if (input_sequence[seq_index + rs_len] == 'A' || input_sequence[seq_index + rs_len] == 'a') {

    double thread_result = 0;

    for (int i = 0; i < rs_len; i++) {

      int rvd_index = rs_len - i - 1;

      int sm_col = 4;

      char base = input_sequence[seq_index + i];

      if (base == 'A' || base == 'a')    
        sm_col = 3;
      if (base == 'C' || base == 'c')
        sm_col = 2;
      if (base == 'G' || base == 'g')
        sm_col = 1;
      if (base == 'T' || base == 't')
        sm_col = 0;

      thread_result += ScoringMatrixVal(scoring_matrix, sm_pitch, rvd_sequence[rvd_index], sm_col);

    }

    results[seq_index] |= (thread_result < cutoff ? 1UL : 0UL) << (2 * rvd_num + 1);

  }

}

ScoreBindingSites は、ブロックごとに (32, 14) スレッドと、入力シーケンスをカバーするのに十分なブロックで起動されます。役立つ場合は、完全なソースをここで見つけることができます。

4

2 に答える 2

1

このコードを改善するためにできることがいくつかあります。

  • 上で提案したように、との2つのループをマージし'T'ます'A'。ifループ内のステートメントの小さなカスケードが述語命令としてコンパイルされる可能性が非常に高いため、これはおそらくブランチ分岐の最大の原因です（ NVidia CUDA Cプログラミングガイドのセクション5.4.2を参照）。

  • バイトサイズのグローバルメモリアクセスはひどい考えです。代わりに、メインループの各反復で、、、、およびの各値に対して、input_sequenceおよびresultsをbase宣言することをお勧めします。char4base.xbase.ybase.zbase.w

  • また、何が行われているのかを詳しく調べたい場合もありますScoringMatrixVal。メモリから値を読み取るだけですか？もしそうなら、あなたはそれを一定のメモリと交換できますか？またはテクスチャ？

アップデート

リクエストに応じて、2番目のポイントで私が意味したことは次のとおりです。ただし、コードはテストしていません。バグやタイプミスを見つけた場合は、遠慮なく保管してください。rs_len簡単にするために、それは4の倍数であると仮定したことに注意してください。

// chosen using occupancy spreadsheet
#define SCORE_THREADS_PER_BLOCK 448

__device__ double ScoringMatrixVal(double *scoring_matrix, size_t pitch, unsigned int row, unsigned int column) {
  return scoring_matrix[ row*pitch/sizeof(double) + column ];
}

__global__ void ScoreBindingSites(char4 *input_sequence, unsigned long is_length, unsigned int *rvd_sequence, unsigned int rs_len, double cutoff, unsigned int rvd_num, double *scoring_matrix, size_t sm_pitch, unsigned char *results) {

  int block_seq_index = SCORE_THREADS_PER_BLOCK * (blockIdx.y * gridDim.x + blockIdx.x);
  int thread_id = (blockDim.x * threadIdx.y) + threadIdx.x;
  int seq_index = block_seq_index + thread_id;

  if (seq_index < 1 || seq_index >= is_length || seq_index + rs_len >= is_length - 1) return;

  if (input_sequence[seq_index - 1] == 'T' || input_sequence[seq_index - 1] == 't') {

    double4 thread_result = make_double4( 0 );

    for (int i = 0; i < rs_len/4; i++) {

      int rvd_index = 4*i;

      int4 sm_col = make_int4( 4 );

      char4 base = input_sequence[seq_index + i];

      if (base.x == 'A' || base.x == 'a')    
        sm_col.x = 0;
      else if (base.x == 'C' || base.x == 'c')
        sm_col.x = 1;
      else if (base.x == 'G' || base.x == 'g')
        sm_col.x = 2;
      else if (base.x == 'T' || base.x == 't')
        sm_col.x = 3;
      thread_result.x += ScoringMatrixVal(scoring_matrix, sm_pitch, rvd_sequence[rvd_index + 0], sm_col.x);

      if (base.y == 'A' || base.y == 'a')    
        sm_col.y = 0;
      else if (base.y == 'C' || base.y == 'c')
        sm_col.y = 1;
      else if (base.y == 'G' || base.y == 'g')
        sm_col.y = 2;
      else if (base.y == 'T' || base.y == 't')
        sm_col.y = 3;
      thread_result.y += ScoringMatrixVal(scoring_matrix, sm_pitch, rvd_sequence[rvd_index + 1], sm_col.y);

      if (base.z == 'A' || base.z == 'a')    
        sm_col.z = 0;
      else if (base.z == 'C' || base.z == 'c')
        sm_col.z = 1;
      else if (base.z == 'G' || base.z == 'g')
        sm_col.z = 2;
      else if (base.z == 'T' || base.z == 't')
        sm_col.z = 3;
      thread_result.z += ScoringMatrixVal(scoring_matrix, sm_pitch, rvd_sequence[rvd_index + 2], sm_col.z);

      if (base.w == 'A' || base.w == 'a')    
        sm_col.w = 0;
      else if (base.w == 'C' || base.w == 'c')
        sm_col.w = 1;
      else if (base.w == 'G' || base.w == 'g')
        sm_col.w = 2;
      else if (base.w == 'T' || base.w == 't')
        sm_col.w = 3;
      thread_result.w += ScoringMatrixVal(scoring_matrix, sm_pitch, rvd_sequence[rvd_index + 3], sm_col.w);

    }

    double acc_thread_result = thread_result.x + thread_result.y + thead_result.z + thread_result.w;

    results[seq_index] |= (acc_thread_result < cutoff ? 1UL : 0UL) << (2 * rvd_num);

  }

  if (input_sequence[seq_index + rs_len] == 'A' || input_sequence[seq_index + rs_len] == 'a') {

    ...

  }

}

いくつかの注意：

  • ScoringMatrixValポインター演算の混乱全体がコンパイラーをスローする可能性があるため、通常の配列アクセスを使用するように関数を書き直しました。
  • 相互に排他的であるように見えるため、ステートメントをステートメントifのカスケードに変換しました。if-elseifコンパイラーは述語命令を使用し、4つのif-elseifブロックをインターリーブすると思います。
  • これらすべてを、、、、、などの文字コードを除いてchar[256]すべてが設定されている場所に置き換えることを検討してください。4AaCc
  • if-elseif-statementsをテーブルルックアップに変換する場合、とに2つの異なるテーブルを使用できるためinput_sequence[seq_index - 1] == 'T'、input_sequence[seq_index + rs_len] == 'A'すべてを1つのループに保持できます。

コードをめちゃくちゃにしすぎていないことを願っています。これがお役に立てば幸いです。

于 2012-07-09T09:49:46.403 に答える
0

私があなたのカーネルで理解している限り、各スレッドは最大 32 文字を読み取り、各文字をチェックしていくつかのデータを出力します。

別のブロックアプローチと別のインデックスを使用して暗黙的にシミュレートすることで、ループを完全に削除できます (それが可能な場合)。

各ブロックには 32 個のスレッドがあり、各スレッドは 1 回のループ反復の結果を計算します。

速いかどうかはわかりませんが、テストする価値があります。

テーブルルックアップを使用して条件を置き換えるペドロの回答は、明確に考慮に入れる必要があります。

マイナーな変更：

input_sequence[seq_index - 1] == 'T' || input_sequence[seq_index - 1] == 't'1 回のメモリ読み取りのみに最適化されていますか?

一度使用された variable のみを削除して、レジスタを保存しますthreadId。

于 2012-07-10T13:19:36.990 に答える