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

How to change sub-matrix of a sparse matrix on CUDA device

本文关键字：CUDA      更新时间：2023-10-16

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

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

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

矩阵数据结构将以阵列形式驻留在CUDA设备上：d_col、d_row和d_val

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

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

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

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

本质上：我如何更改CUDA设备端阵列（d_val）中标记updateInds[1]、updateInds[2]…处的条目，。。。，updateInds[n]到一组新的值val[updatInds[1]]，val[updateInds[2]。。。，val[updateInds[3]，不将整个val阵列从系统内存重新复制到CUDA设备内存阵列d_val？

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

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

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;

则也可以使用对CCD_ 8的单个调用来执行该传送。方法概述如下。

注：

  1. 与对相同数量的元素执行cudaMemcpy操作相比，cudaMemcpy2D比您预期的要慢
  2. CUDA API调用有一些固有的开销。如果矩阵的很大一部分要以分散的方式更新，那么利用可以使用单个cudaMemcpy操作来完成的事实，仅仅传输整个d_val阵列实际上可能更快
  3. 如果非零值改变其在稀疏矩阵中的位置，则不能使用此处描述的方法。在这种情况下，我无法提供如何在设备上通过手术更新CSR稀疏矩阵的一般答案。而且某些相对简单的更改可能需要更新大部分数组数据（3个矢量）
相关文章：