检查如果一个矩阵包含在CUDA

Question

什么是检查在CUDA（C++）为inf/nan元件大的矩阵的有效方式的NaN或无穷大的值？该矩阵在GPU存储器中存储为float*。如果至少存在一个错误条目，我不需要这些元素的位置，只需布尔是/否回答。

的选项有：

• 有一个内核检查整个阵列（容易实现，但可能慢）
• 有多个内核检查如行和结合输出与OR（有没有任何CUDA builtins有效地做到这一点？）
• ..其他想法？

谢谢！

Answer 1

有此instrinsics，但可用于C99的功能应该是罚款：

isnan() 

为了测试INF，您可以使用：

isinf() 

这是很少更快地拥有多个内核做一个单一编写好的内核的相同工作，所以我不确定为什么你认为单个内核会很慢。这种算法很可能是内存限制的，所以您需要关注读取数据访问效率，即合并。在CUDA中，穿过矩阵的简单方法是让每个线程处理一列。这可以通过for-loop高效地实现，并导致完美的合并读取。因为你只关心没有索引的单个结果，所以我们可以有多个线程写入（布尔）结果而不是原子，以提高效率，因为任何可能写入结果的线程都将是写同样的价值。

另一种可能考虑的优化策略是早期退出策略，但这并不能优化最坏情况下的时间，但实际上时间更长，所以我会放弃，除非平均吞吐量是一个大问题。

下面是一个完整的工作示例（使用测试楠为例）：

$ 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 
$ 

注意，我与proper cuda error checking为了清晰/简洁分配，但始终建议。

进一步优化，你可以用每格参数块（BLKS）和每块参数（nTPB）螺纹玩，但是，在一定程度上，这些最优值将取决于GPU在运行上。

Answer 2

您的问题可以改写为减少操作。这可以通过使用CUDA Thrust来有效实施。您可以通过使用CUDA的isnan或isinf，然后还原转化阵列的原始数组转换为布尔数组。所有这些都可以通过expoiting来执行thrust::transform_reduce。

下面是一个例子，围绕Robert Crovella已经介绍给你的那个例子来构建。下面的代码在CUDA中实现相当于Matlab的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(); 

}