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

Checking if a matrix contains nans or infinite values in CUDA

本文关键字：nans 无限 中的 CUDA 是否 包含 检查      更新时间：2023-10-16
在

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

选项包括：

  • 让一个内核检查整个数组（易于实现，但可能很慢）
  • 让多个内核检查例如行并将输出与 OR 组合（是否有任何 CUDA 内置的方法来有效地做到这一点？
  • ..其他想法？

谢谢！

这有本能，但 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 checkn"); 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 checkn"); return 1;}
  printf("Successn");
  return 0;
}

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

请注意，为了清晰/简洁，我已经省去了适当的 cuda 错误检查，但始终建议这样做。

为了进一步优化，您可以使用每个网格的块参数（BLKS）和每个块的线程数参数（nTPB），但是，在某种程度上，这些的最佳值将取决于您运行的GPU。

您的问题可以重新转换为缩减操作。这可以通过使用 CUDA 推力有效地实现。您可以使用 CUDA 的isnan或isinf将原始数组转换为布尔数组，然后减少转换后的数组。所有这些都可以通过暴露thrust::transform_reduce来执行。

下面是一个例子，围绕罗伯特·克罗维拉已经向你展示的例子构建的。下面的代码在 CUDA 中实现了相当于 Matlab sum(isnan(array)) 的代码。

#include <thrustdevice_vector.h>
#include <thrustreduce.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 = %dn",h_result);
    getchar();
}
相关文章：