如何在 CUDA 设备上更改稀疏矩阵的子矩阵答案

【问题标题】：How to change sub-matrix of a sparse matrix on CUDA device如何在 CUDA 设备上更改稀疏矩阵的子矩阵
【发布时间】：2015-09-08 09:48:37
【问题描述】：

我有一个稀疏矩阵结构，我将它与 CUBLAS 结合使用来实现线性求解器类。我预计我将要解决的稀疏矩阵的维度将相当大（大约为 10 ^ 7 x 10 ^ 7）。我还预计求解器将需要多次使用，并且该矩阵的一部分也需要多次更新（在计算解决方案之间）。

将整个矩阵结构从系统内存复制到设备内存可能会成为性能瓶颈，因为在给定时间只需要更改一小部分矩阵条目。

我希望能够做的是有一种方法只更新特定的子集/子矩阵，而不是每次我需要更改矩阵时将整个矩阵结构从系统内存重新复制到设备内存.

矩阵数据结构将驻留在数组中的 CUDA 设备上： d_col、d_row 和 d_val

在系统端，我会有相应的数组 I、J 和 val。

因此，理想情况下，我只想更改与系统数组 val 中已更改的值相对应的 d_val 子集。

请注意，我预计不会在矩阵中添加或删除任何条目，只有现有条目的值会发生变化。

我天真地认为要实现这一点，我会在主机端有一个整数数组或向量，例如updateInds ，它将跟踪 val 中已更改的条目的索引，但我不确定如何有效地告诉 CUDA 设备更新 d_val 的相应值。

本质上：如何将 CUDA 设备端数组 (d_val) 中 indicies updateInds[1],updateInds[2],...,updateInds[n] 中的条目更改为一组新值 val[updatInds [1]], val[updateInds[2]], ..., val[updateInds[3]], 不将整个 val 数组从系统内存重新复制到 CUDA 设备内存数组 d_val?

【问题讨论】：

您有具体的 CUDA 编程问题要问吗？
它归结为：如何将 CUDA 设备端数组 (d_val) 中 indicies updateInds[1],updateInds[2],...,updateInds[n] 中的条目更改为新的一组值 val[updatInds[1]], val[updateInds[2]], ..., val[updateInds[3]，没有将整个 val 数组从系统内存重新复制到 CUDA 设备内存数组 d_val？跨度>
您是指主机代码还是设备代码？你的矩阵真的是 COO 格式吗？
我对 CUDA 还是很陌生，所以我基本上遵循 CUDA 6.5 SDK 示例库中的共轭梯度求解器示例。矩阵由调用定义：cusaprseCreateMatDescr(&descr)cusparseSetMatType(descr, CUSPARSE_MATRIX_TYPE_GENERAL); 矩阵通过调用加载到设备内存中：cudaMemcpy(d_col,J,nz*sizeof(int), cudaMemcpyHostToDevice)cudaMemcpy(d_row,I,(N+1)*sizeof(int), cudaMemcpyHostToDevice)cudaMemcpy(d_val, val, nz*sizeof(float), cudaMemcpyHostToDevice)
这对于密集矩阵来说非常有意义。我认为稀疏矩阵没有意义。目前尚不清楚您是否了解 CSR 稀疏矩阵存储的工作原理（CSR 是您引用的 cuda 示例中使用的格式）。我可以将稀疏矩阵中的单个值从零更改为非零，这可能需要更改 3 个数组中的 2 个（csrRowPtr、csrValA）中的几乎每个值。所以我认为一般情况很难处理，并且可能需要更新大部分数据。（与密集病例不同，您可以进行手术改变）。

标签： c++ arrays cuda sparse-matrix

【解决方案1】：

只要你只想改变与CSR（或CSC，或COO）稀疏矩阵表示相关的值数组的数值，过程并不复杂。

假设我有这样的代码（摘自CUDA共轭梯度sample）：

checkCudaErrors(cudaMalloc((void **)&d_val, nz*sizeof(float)));
...
cudaMemcpy(d_val, val, nz*sizeof(float), cudaMemcpyHostToDevice);

现在，在代码中的这一点之后，假设我需要更改d_val 数组中的一些值，对应于我在val 中所做的更改：

for (int i = 10; i < 25; i++)
  val[i] = 4.0f;

移动这些特定更改的过程在概念上与使用memcpy 更新数组相同，但我们将使用cudaMemcpy 更新设备上的d_val 数组：

cudaMemcpy(d_val+10, val+10, 15*sizeof(float), cudaMempcyHostToDevice);

由于这些值都是连续的，我可以使用单个 cudaMemcpy 调用来实现传输。

如果我有几个与上述类似的不相交区域，则需要多次调用cudaMemcpy，每个区域调用一次。如果碰巧这些区域等距且等长：

for (int i = 10; i < 5; i++)
  val[i] = 1.0f;
for (int i = 20; i < 5; i++)
  val[i] = 2.0f;
for (int i = 30; i < 5; i++)
  val[i] = 4.0f;

那么也可以使用对cudaMemcpy2D 的一次调用来执行此传输。方法概述here。

注意事项：

与 相同数量的元素上的 cudaMemcpy 操作相比，cudaMemcpy2D 比您预期的要慢。
CUDA API 调用有一些固有的开销。如果矩阵的很大一部分要以分散的方式更新，实际上传输整个d_val 数组可能仍然更快，因为这可以使用单个cudaMemcpy 操作来完成。
如果非零值改变了它们在稀疏矩阵中的位置，则不能使用此处描述的方法。在这种情况下，我无法就如何通过手术方式更新设备上的 CSR 稀疏矩阵提供一般性答案。并且某些相对简单的更改可能需要更新大部分数组数据（3 个向量）。

【讨论】：