CUDA混合float和double導致內核不執行

使用CUDA CI有一個統計的內核，當我在VS2012中添加內核內的任何地方一個斷點，包括與前一個變量定義STDDEV行：CUDA混合float和double導致內核不執行

double mean, stddev, sumOfValues, sumOfValuesSquared; 
unsigned int n; 

// acquire greater than 0 values for: sumOfValues, sumOfValuesSquared, and n 

stddev = (float)(sqrt((double)(n) * sumOfValuesSquared - (sumOfValues*sumOfValues))/(double)(n));

斷點永遠不會到達，並且內核不會執行。當我刪除單行時，內核執行。我認爲這與sqrt有關，但它不是。我有另一行：

mean = sumOfValues/n;

當我使用該行時，它也不執行內核。我錯過了CUDA中的類型轉換（這是一個寄存器問題，還是單精度與雙精度）？

- 更新（2013年10月2日14:25 CST） -

我調整線程數爲1，然後1024第一次運行，它關係到我的斷點，第二用高線程數，內核不執行。請查看代碼，如下：

#include "stdafx.h" 

#include <stdio.h> 
#include <cuda.h> 
#include <cuda_runtime.h> 
#include <device_launch_parameters.h> 

typedef struct 
{ 
    unsigned int value; 
} ValueStruct; 

__global__ void FailsToExecute(ValueStruct *vs) 
{ 
    unsigned int numerator = vs->value; 
    unsigned int denominator= 3; 
    bool eject = false; 

    if(denominator > 0) 
    { 
     if(1.0f * numerator/denominator > 17.98f) 
      eject = true; 
     else 
      eject = false; 
    } 
} 

int _tmain(int argc, _TCHAR* argv[]) 
{ 
    ValueStruct *vsHost; 
    ValueStruct *vsDevice; 

    cudaMallocHost((void **)&vsHost, sizeof(ValueStruct)); 
    cudaMalloc((void **)&vsDevice, sizeof(ValueStruct)); 

    vsHost->value = 54; 

    cudaMemcpy(vsDevice, vsHost, sizeof(ValueStruct), cudaMemcpyKind::cudaMemcpyHostToDevice); 

    dim3 blocks(5); 
    dim3 threads(1024); 

    FailsToExecute<<<blocks, threads>>>(vsDevice); 

    return 0; 
}

如何計算/應付寄存器界限，我不很瞭解他們嗎？

來源

2013-10-02 user2712376

您是否正在構建支持'double'的體系結構的代碼，即計算能力爲1.3或更高？ – njuffa

看起來你沒有足夠的寄存器。嘗試減少每個塊的線程數以查看它是否可以運行。 – kangshiyin

我想這不是你的實際代碼。嘗試使用cuda-memcheck運行您的代碼。您的for循環中可能存在數據訪問衝突。或者提供一個完整的重現代碼。 SSCCE.org –

需要比可用資源更多的內核（例如寄存器，共享內存）將不會啓動。

https://devtalk.nvidia.com/default/topic/545591/how-to-debug-kernel-throwing-an-exception-/?offset=16

由於分配粒度效應對於每個GPU架構，對於特定的內核的組合所需的資源的精確計算不同：這可以通過適當的錯誤檢查如下所示被檢測到，例如與特定的啓動配置可以是不平凡的。出於這個原因，我建議使用包含粒度細節的CUDA佔用率計算器。你可以找到佔用的電子表格在這裏：

http://developer.download.nvidia.com/compute/cuda/CUDA_Occupancy_calculator.xls

來源

2013-10-02 20:02:46 njuffa

謝謝njuffa！正是我需要的。 – user2712376

除了降低每塊的線程數以符合可用佔用率外，還有另一個主要的限制因素是我調試我的代碼時。儘管我使用的CUDA卡不是我的主要產品，但我有2個顯示器插入其中。斷開顯示器後，它有更多的可用資源，因此有更多的寄存器。對於大多數人來說，這可能是一件容易的事情，但對於CUDA編碼而言是新手，在調試代碼時不會自動斷開監視器的連接。 cudaCheckLaunchError是跟蹤cudaGetLastError內核啓動爲零的關鍵。 – user2712376

此外，根據佔位率Excel文檔的指示，如果在編譯期間打開詳細輸出，則可以看到通過編譯器爲每個cuda內核分配的寄存器和其他資源的數量。動態內存分配使用多堆/堆棧可能會使跟蹤資源分配複雜化，但只要在調試過程中所有內核啓動和cuda函數後使用cudaGetLastError，調試都是一個簡單的過程。最後，對智者說一句話......不要將你的GPGPU用作顯卡，同時期待最大的資源可用於計算樂趣。 – user2712376

是否有可能的是，NVCC編譯器優化了內核爛掉？看看內核函數，我可以看到如何安全地優化零指令，因爲它實際上什麼都不做。

__global__ void FailsToExecute(ValueStruct *vs) 
{ 
    unsigned int numerator = vs->value; 
    unsigned int denominator= 3; 
    bool eject = false; 

    if(denominator > 0) 
    { 
     if(1.0f * numerator/denominator > 17.98f) 
      eject = true; 
     else 
      eject = false; 
    } 
}

設置eject因爲eject沒關係的就不再使用。所以我們可以抓住這兩項任務。 if（）條件中的表達式不會修改任何內容，因爲它對if語句的任一分支都沒有任何作用，它看起來像if（）可以被刪除。同樣繼續回到內核的頂端，似乎它可以全部優化爲無，並且仍然使內核產生相同的結果。

也許如果你添加了某種輸出，比如可能是一個布爾數組，然後將eject的結果保存到該數組中，那麼你會看到正在執行的內核。

來源

2013-10-02 21:17:50 MorbidFuzzball

我想到了同樣的事情，並且做了測試。如果您減少每塊的線程數，線程（1）而不是線程（1024），則代碼會生成，執行並且可追蹤 - PTX在反彙編器中可見。但是，由於您提出的主題，我遇到了編碼困難。它出現在使用NSIGHT進行調試時，'寄存器'是虛擬寄存器，由於優化而不遵循典型的C範圍概念。觀察到的變量更快地超出範圍。我確定這是必要的，我只是習慣了。我不相信這是什麼影響了我的案例中的代碼。 – user2712376

CUDA混合float和double導致內核不執行

回答

相關問題