检查矩阵是否包含 CUDA 中的 nans 或无限值

Question

在 CUDA (C++) 中检查inf/nan元素的大矩阵的有效方法是什么？矩阵存储float*在 GPU 内存中。我不需要这些元素的位置，如果至少存在一个错误条目，则只需要一个布尔值是/否答案。

选项是：

• 让一个内核检查整个数组（易于实现但可能很慢）
• 有多个内核检查例如行并将输出与 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)，但是，在某种程度上，这些的最佳值将取决于您正在运行的 GPU。