檢查如果一個矩陣包含在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(); 

}