Как изменить подматрицу разреженной матрицы на устройстве CUDA

Question

У меня есть разреженная матричная структура, которую я использую совместно с 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]], не возвращая весь массив валов из системной памяти в массив памяти устройства CUDA d_val?

Answer 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.

Примечания:

1. cudaMemcpy2D медленнее, чем можно было бы ожидать по сравнению с cudaMemcpy операции на же числа элементов.
2. У вызовов API CUDA есть определенные накладные расходы. Если большая часть матрицы должна обновляться разбросанным образом, все же можно быстрее просто передать весь массив d_val, воспользовавшись тем, что это можно сделать, используя одну операцию cudaMemcpy.
3. Описанный здесь метод не может быть использован, если ненулевые значения изменяют местоположение в разреженной матрице. В этом случае я не могу дать общий ответ о том, как хирургическое обновление разреженной матрицы CSR на устройстве. И некоторые относительно простые изменения могут потребовать обновления большинства данных массива (3 вектора) в любом случае.