高效复制跨步数据（到和从CUDA设备）的方法？

Question

有没有可能高效地复制具有常量(甚至非常量)步幅的数据到和从CUDA设备？

我想对一个大的对称矩阵进行对角化。

使用Jacobi算法，每次迭代中都会使用两行和两列进行一系列操作。

由于矩阵本身太大而无法完全复制到设备中，因此我正在寻找一种将两行和列复制到设备的方法。

使用三角形矩阵形式存储数据是不错的选择，但附加缺点如下：

- 行长度不是常数（不是个大问题）。 - 列值的步幅不是常数（每行步幅增加1）。

即使使用三角形形式，仍然无法将整个矩阵存储在GPU上。

我查看了一些时间，并意识到逐个复制跨越值非常慢（同步和异步都如此）。

//编辑：删除解决方案——添加一个答案

- Pixelchemist

---

也许你指的是C++，而不是C？ - wildplasser

两者皆是。我使用 C++ 替换了 2d 标签。 - Pixelchemist

在C语言中，字符串不能作为“<<”运算符的操作数，因此您应该同时删除C标签。 - wildplasser

您可能对使用cudaMemcpy2D(或[async版本](http://docs.nvidia.com/cuda/cuda-runtime-api/index.html#group__CUDART__ MEMORY_1g7f182f1a8b8750c7fa9e18aeb280d31c))的方法感兴趣，如@njuffa在这篇帖子中所述。 - Robert Crovella

您可以通过将主机内存映射到设备，然后使用自定义内核进行复制来复制任何形状的数据。 - tera

Answer 1

感谢Robert Crovella提供正确的提示使用cudamemcpy2d。 我会附上我的测试代码，让每个人都有可能理解...

如果有人提出解决行优先级三角形矩阵复制问题的建议，请随时编写另一个答案。

__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;

}

- Pixelchemist