4

一定の（または一定でない）値でストライドされたデータをCUDAデバイスとの間で効率的にコピーする可能性はありますか？

大きな対称行列を対角化したい。

ヤコビアルゴリズムを使用すると、各反復内で2つの行と2つの列を使用する一連の操作があります。

マトリックス自体が大きすぎてデバイスに完全にコピーできないため、2つの行と列をデバイスにコピーする方法を探しています。

三角行列形式を使用してデータを保存すると便利ですが、次のような追加の欠点があります。

  • 一定でない行の長さ[そのような問題ではありません]
  • 列値のストライドが一定でない[ストライドは行ごとに1ずつ増加します。]

発生します。[編集：三角形のフォームを使用しても、マトリックス全体をGPUに保存することはまだ不可能です。]

いくつかのタイミングを調べて、ストライド値を1つずつコピーするのは非常に遅い（同期および非同期）ことを認識しました。

//編集：解決策を削除-回答を追加

4

1 に答える 1

3

cudamemcpy2dを使用するための正しいヒントを与えてくれたRobertCrovellaに感謝します。誰もが理解できるように、テストコードを追加します...

行優先の三角行列を使用してコピーの問題を解決するための提案を誰かが思いついた場合は、別の答えを自由に書いてください。

__global__ void setValues (double *arr, double value)
{
  arr[blockIdx.x] = value;
}

int main( void ) 
{
  // define consts
  static size_t const R = 10, C = 10, RC = R*C;

  // create matrices and initialize
  double * matrix = (double*) malloc(RC*sizeof(double)), 
    *final_matrix = (double*) malloc(RC*sizeof(double));
  for (size_t i=0; i<RC; ++i) matrix[i] = rand()%R+10;
  memcpy(final_matrix, matrix, RC*sizeof(double));

  // create vectors on the device
  double *dev_col, *dev_row, 
    *h_row = (double*) malloc(C*sizeof(double)), 
    *h_col = (double*) malloc(R*sizeof(double));
  cudaMalloc((void**)&dev_row, C * sizeof(double));
  cudaMalloc((void**)&dev_col, R * sizeof(double));

  // choose row / col to copy
  size_t selected_row = 7, selected_col = 3;

  // since we are in row-major order we can copy the row at once 
  cudaMemcpy(dev_row, &matrix[selected_row*C], 
    C * sizeof(double), cudaMemcpyHostToDevice);
  // the colum needs to be copied using cudaMemcpy2D 
  // with Columnsize*sizeof(type) as source pitch
  cudaMemcpy2D(dev_col, sizeof(double), &matrix[selected_col], 
    C*sizeof(double), sizeof(double), R, cudaMemcpyHostToDevice);

  // copy back to host to check whether we got the right column and row
  cudaMemcpy(h_row, dev_row, C * sizeof(double), cudaMemcpyDeviceToHost);
  cudaMemcpy(h_col, dev_col, R * sizeof(double), cudaMemcpyDeviceToHost);
  // change values to evaluate backcopy
  setValues<<<R, 1>>>(dev_col, 88.0); // column should be 88
  setValues<<<C, 1>>>(dev_row, 99.0); // row should be 99
  // backcopy
  cudaMemcpy(&final_matrix[selected_row*C], dev_row, 
    C * sizeof(double), cudaMemcpyDeviceToHost);
  cudaMemcpy2D(&final_matrix[selected_col], C*sizeof(double), dev_col, 
    sizeof(double), sizeof(double), R, cudaMemcpyDeviceToHost);

  cudaDeviceSynchronize();
  // output for checking functionality

  printf("Initial Matrix:\n");
  for (size_t i=0; i<R; ++i)
  {
    for (size_t j=0; j<C; ++j) printf(" %lf", matrix[i*C+j]);
    printf("\n");
  }
  printf("\nRow %u values: ", selected_row);
  for (size_t i=0; i<C; ++i) printf(" %lf", h_row[i]);
  printf("\nCol %u values: ", selected_col);
  for (size_t i=0; i<R; ++i) printf(" %lf", h_col[i]);
  printf("\n\n");

  printf("Final Matrix:\n");
  for (size_t i=0; i<R; ++i)
  {
    for (size_t j=0; j<C; ++j) printf(" %lf", final_matrix[i*C+j]);
    printf("\n");
  }

  cudaFree(dev_col);
  cudaFree(dev_row);
  free(matrix);
  free(final_matrix);
  free(h_row);
  free(h_col);
  cudaDeviceReset();
  return 0;

}
于 2013-06-14T15:40:06.483 に答える