Проверка, содержит ли матрица nans или бесконечные значения в CUDA

Question

Какой эффективный способ проверить большую матрицу для inf/nan элементы в CUDA (C++)? Матрица хранится как float* в памяти графического процессора. Мне не нужно расположение этих элементов, просто логический ответ да / нет, если присутствует хотя бы одна неверная запись.

Варианты:

• чтобы одно ядро ​​проверило весь массив (легко реализовать, но, вероятно, медленно)
• несколько ядер проверяют, например, строки и объединяют вывод с OR (есть ли встроенные функции CUDA для эффективного выполнения этого?)
• ..другие идеи?

Спасибо!

Answer 1

Для этого есть свои особенности, но функции, доступные для C99, должны быть в порядке:

isnan()

Чтобы проверить на inf, вы можете использовать:

isinf()

Редко быстрее, когда несколько ядер выполняют одну и ту же работу с одним хорошо написанным ядром, поэтому я не уверен, почему вы думаете, что иметь одно ядро ​​будет медленно. Этот алгоритм, скорее всего, связан с памятью, поэтому вы хотите сосредоточиться на эффективности доступа к чтению данных, т.е. на объединении. В CUDA простой способ пройти через матрицу состоит в том, чтобы каждый поток обрабатывал столбец. Это может быть эффективно реализовано с помощью цикла for и приводит к идеально слитному чтению.

Поскольку вы заботитесь только об одном результате без индексов, у нас может быть несколько потоков, записывающих в (булево) результат без атомики, для дальнейшей эффективности, поскольку все потоки, которые могут записывать в результат, будут записывать одно и то же значение.

Другой стратегией оптимизации, которую можно рассмотреть, может быть стратегия раннего выхода, но она не оптимизирует временные рамки наихудшего случая, но фактически увеличивает ее, поэтому я бы обошелся без нее, если только средняя пропускная способность не является большой проблемой.

Вот полный проработанный пример (используя тест для nan в качестве примера):

$ cat t383.cu
#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#define DSIZEW 10000
#define DSIZEH 2000
#define nTPB 256
#define BLKS 16

__global__ void isnan_test(float *data, int width, int height, bool *result){

  int idx = threadIdx.x+blockDim.x*blockIdx.x;

  while (idx < width){
    for (int i = 0; i < height; i++)
      if (isnan(data[(i*width) + idx])) *result = false;
    idx += gridDim.x+blockDim.x;
    }
}

int main(){

  float *d_data, *h_data;
  bool  *d_result, h_result=true;
  const char type = '0';

  cudaMalloc((void **)&d_data, sizeof(float)*DSIZEW*DSIZEH);
  cudaMalloc((void **)&d_result, sizeof (bool));
  h_data=(float *)malloc(sizeof(float)*DSIZEW*DSIZEH);
  for (int i=0; i<DSIZEH*DSIZEW; i++)
    h_data[i] = rand()/RAND_MAX;
  cudaMemcpy(d_data, h_data, sizeof(float)*DSIZEW*DSIZEH, cudaMemcpyHostToDevice);
  cudaMemcpy(d_result, &h_result, sizeof(bool), cudaMemcpyHostToDevice);
  isnan_test<<<BLKS,nTPB>>>(d_data, DSIZEW, DSIZEH, d_result);
  cudaMemcpy(&h_result, d_result, sizeof(bool), cudaMemcpyDeviceToHost);
  if (!h_result) {printf("error in no-NAN check\n"); return 1;}
  float my_nan = nanf(&type); // create a NAN value
  cudaMemcpy(d_data, &my_nan, sizeof(float), cudaMemcpyHostToDevice);
  isnan_test<<<BLKS,nTPB>>>(d_data, DSIZEW, DSIZEH, d_result);
  cudaMemcpy(&h_result, d_result, sizeof(bool), cudaMemcpyDeviceToHost);
  if (h_result) {printf("error in NAN check\n"); return 1;}
  printf("Success\n");
  return 0;
}


$ nvcc -arch=sm_20 -o t383 t383.cu
$ ./t383
Success
$

Обратите внимание, что я обошелся без надлежащей проверки ошибок cuda для ясности / краткости, но это всегда рекомендуется.

Для дальнейшей оптимизации вы можете поиграть с блоками по параметру сетки (BLKS) и количество потоков в параметре блока (nTPB), однако, в некоторой степени оптимальные значения будут зависеть от того, на каком графическом процессоре вы работаете.

Answer 2

Ваша проблема может быть преобразована как операция сокращения. Это может быть эффективно реализовано с помощью CUDA Thrust. Вы можете преобразовать исходный массив в логический массив с помощью CUDA isnan или же isinf а затем уменьшая преобразованный массив. Все что можно сделать, выставив thrust::transform_reduce,

Ниже приведен пример, построенный вокруг того, который Роберт Кровелла уже представил вам. Код ниже реализует в CUDA эквивалент Matlab's sum(isnan(array)),

#include <thrust\device_vector.h>
#include <thrust\reduce.h>

#define DSIZEW 10000
#define DSIZEH 2000

// --- Operator for testing nan values
struct isnan_test { 
    __host__ __device__ bool operator()(const float a) const {
        return isnan(a);
    }
};

void main(){

    thrust::host_vector<float> h_data(DSIZEW*DSIZEH);
    for (int i=0; i<DSIZEH*DSIZEW; i++)
        h_data[i] = rand()/RAND_MAX;

    const char type = '0';
    float my_nan = nanf(&type); // create a NAN value
    h_data[0] = my_nan;

    thrust::device_vector<float> d_data(h_data);

    bool h_result = thrust::transform_reduce(d_data.begin(), d_data.end(), isnan_test(), 0, thrust::plus<bool>());
    printf("Result = %d\n",h_result);

    getchar();

}