數據採集部分不料輸出了「0」

Question

編輯
在最初發布的代碼片段（見下文）我沒有正確發送struct到device，這已得到修復，但結果仍然是一樣的。在我的完整代碼中，這個錯誤並不存在。 （在我最初發布的命令中有兩個錯誤 - 一個是從HostToDevice複製的結構，但實際上是相反的，並且副本的大小也是錯誤的，抱歉;兩個錯誤都是固定的，但重新編譯的代碼仍顯示下面描述的零現象一樣，我的全部代碼。）

EDIT 2
在我的代碼去proprietarization重寫的急速我做這dalekchef好心向我指出了幾個錯誤（ struct到設備的副本在設備上進行分配之前執行，在我重寫的代碼中，設備cudaMalloc調用沒有乘以sizeof(...)陣列e的類型lements。我添加了這些修補程序，重新編譯並重新測試，但它沒有解決問題。還仔細檢查了我的原始代碼 - 它沒有出現這些錯誤。再次抱歉，爲了混淆。

我想從一個大的模擬程序轉儲統計數據。下面顯示一個類似的簡化代碼。這兩個代碼都表現出同樣的問題 - 它們輸出零時，它們應該輸出平均值。

#include "stdio.h" 

struct __align__(8) DynamicVals 
{ 
    double a; 
    double b; 
    int n1; 
    int n2; 
    int perDump; 
}; 

__device__ int *dev_arrN1, *dev_arrN2; 
__device__ double *dev_arrA, *dev_arrB; 
__device__ DynamicVals *dev_myVals; 
__device__ int stepsA, stepsB; 
__device__ double sumA, sumB; 
__device__ int stepsN1, stepsN2; 
__device__ int sumN1, sumN2; 

__global__ void TEST 
(int step, double dev_arrA[], double dev_arrB[], 
int dev_arrN1[], int dev_arrN2[],DynamicVals *dev_myVals) 
{ 
    if (step % dev_myVals->perDump) 
    { 
     dev_arrN1[step/dev_myVals->perDump] = 0; 
     dev_arrN2[step/dev_myVals->perDump] = 0; 
     dev_arrA[step/dev_myVals->perDump] = 0.0; 
     dev_arrB[step/dev_myVals->perDump] = 0.0; 
     stepsA = 0; 
     stepsB = 0; 
     stepsN1 = 0; 
     stepsN2 = 0; 
     sumA = 0.0; 
     sumB = 0.0; 
     sumN1 = 0; 
     sumN2 = 0; 
    } 

    sumA += dev_myVals->a; 
    sumB += dev_myVals->b; 
    sumN1 += dev_myVals->n1; 
    sumN2 += dev_myVals->n2; 
    stepsA++; 
    stepsB++; 
    stepsN1++; 
    stepsN2++; 

    if (sumA > 100000000) 
    { 
     dev_arrA[step/dev_myVals->perDump] += 
    sumA/stepsA; 
     sumA = 0.0; 
     stepsA = 0; 
    } 
    if (sumB > 100000000) 
    { 
     dev_arrB[step/dev_myVals->perDump] += 
    sumB/stepsB; 
     sumB = 0.0; 
     stepsB = 0; 
    } 
    if (sumN1 > 1000000) 
    { 
     dev_arrN1[step/dev_myVals->perDump] += 
    sumN1/stepsN1; 
     sumN1 = 0; 
     stepsN1 = 0; 
    } 
    if (sumN2 > 1000000) 
    { 
     dev_arrN2[step/dev_myVals->perDump] += 
    sumN2/stepsN2; 
     sumN2 = 0; 
     stepsN2 = 0; 
    } 

    if ((step+1) % dev_myVals->perDump) 
    { 
     dev_arrA[step/dev_myVals->perDump] += 
    sumA/stepsA; 
     dev_arrB[step/dev_myVals->perDump] += 
    sumB/stepsB; 
     dev_arrN1[step/dev_myVals->perDump] += 
    sumN1/stepsN1; 
     dev_arrN2[step/dev_myVals->perDump] += 
    sumN2/stepsN2; 
    } 
} 

int main() 
{ 
    const int TOTAL_STEPS = 10000000; 
    DynamicVals vals; 
    int *arrN1, *arrN2; 
    double *arrA, *arrB; 
    int statCnt; 

    vals.perDump = TOTAL_STEPS/10; 
    statCnt = TOTAL_STEPS/vals.perDump+1; 
    vals.a = 30000.0; 
    vals.b = 60000.0; 
    vals.n1 = 10000; 
    vals.n2 = 20000; 

    cudaMalloc((void**)&dev_arrA, statCnt*sizeof(double)); 
    cudaMalloc((void**)&dev_arrB, statCnt*sizeof(double)); 
    cudaMalloc((void**)&dev_arrN1, statCnt*sizeof(int)); 
    cudaMalloc((void**)&dev_arrN2, statCnt*sizeof(int)); 
    cudaMalloc((void**)&dev_myVals, sizeof(DynamicVals)); 
    cudaMemcpy(dev_myVals, &vals, sizeof(DynamicVals), 
      cudaMemcpyHostToDevice); 

    arrA = (double *)malloc(statCnt * sizeof(double)); 
    arrB = (double *)malloc(statCnt * sizeof(double)); 
    arrN1 = (int *)malloc(statCnt * sizeof(int)); 
    arrN2 = (int *)malloc(statCnt * sizeof(int)); 

    for (int i=0; i< TOTAL_STEPS; i++) 
     TEST<<<1,1>>>(i, dev_arrA,dev_arrB,dev_arrN1,dev_arrN2,dev_myVals); 

    cudaMemcpy(arrA,dev_arrA,statCnt * sizeof(double),cudaMemcpyDeviceToHost); 
    cudaMemcpy(arrB,dev_arrB,statCnt * sizeof(double),cudaMemcpyDeviceToHost); 
    cudaMemcpy(arrN1,dev_arrN1,statCnt * sizeof(int),cudaMemcpyDeviceToHost); 
    cudaMemcpy(arrN2,dev_arrN2,statCnt * sizeof(int),cudaMemcpyDeviceToHost); 

    for (int i=0; i< statCnt; i++) 
    { 
     printf("Step: %d ; A=%g B=%g N1=%d N2=%d\n", 
     i*vals.perDump, 
     arrA[i], arrB[i], arrN1[i], arrN2[i]); 
    } 
} 

輸出：

Step: 0 ; A=0 B=0 N1=0 N2=0 
Step: 1000000 ; A=0 B=0 N1=0 N2=0 
Step: 2000000 ; A=0 B=0 N1=0 N2=0 
Step: 3000000 ; A=0 B=0 N1=0 N2=0 
Step: 4000000 ; A=0 B=0 N1=0 N2=0 
Step: 5000000 ; A=0 B=0 N1=0 N2=0 
Step: 6000000 ; A=0 B=0 N1=0 N2=0 
Step: 7000000 ; A=0 B=0 N1=0 N2=0 
Step: 8000000 ; A=0 B=0 N1=0 N2=0 
Step: 9000000 ; A=0 B=0 N1=0 N2=0 
Step: 10000000 ; A=0 B=0 N1=0 N2=0 

現在，如果我用一個小週期爲我的垃圾場或者如果我的＃分別是較小的，我可以逃脫只需直接

1. add
2. 按期劃分和期末

...算法，但我使用臨時總和，否則我的int會溢出（double不會溢出，但我擔心它會丟失精度）。

如果我使用上面的直接算法獲得更小的值，我會得到正確的非零值，但第二個我使用中間值（例如stepsA,sumA等），值將變爲零。 我知道我在這裏做些傻事......我錯過了什麼？

注：
答）是的，我知道在其上面形式不平行此代碼，本身並不能保證並行。它是更長代碼的一個小統計數據收集部分的一部分。在該代碼中，它被封裝在一個線程索引特定的條件邏輯中，以防止衝突（使其並行）並用作模擬程序（保證並行化）的數據收集。希望你能理解上面的代碼來自哪裏，並避免對它缺乏線程安全性的評論。 （這種免責聲明是從過去的經驗中得到的，這些經歷從不明白我發佈摘錄而不是完整的代碼的人那裏收到了非生產性的評論，儘管我這樣寫得不那麼直截了當。）

B.）是的，我知道變量的名稱是不明確的。這就是我想說的。我正在處理的代碼是專有的，儘管它最終會被公開。我只寫這篇文章，因爲我過去發佈了類似的匿名代碼，並收到了關於我的命名約定的粗魯的評論。

C.）是的，我已經多次閱讀CUDA manual，雖然我確實犯了錯誤，但我承認有些功能我不明白。我在這裏沒有使用共享內存，但我在完整的代碼中使用共享內存（OF COURSE）。

D.）是的，上面的代碼確實代表了與我的非工作代碼的數據轉儲部分完全相同的功能，其中與此特定問題無關的邏輯被刪除，並且具有線程安全條件。變量名稱已被更改，但算法上它應該保持不變，並且由完全相同的非工作輸出（零）驗證。

E.）我確實意識到上述片段中的「動態」struct具有非動態值。我將結構命名爲，因爲在完整的代碼中，這個struct包含模擬數據，並且是動態的。簡化代碼中的靜態特性不應使統計信息收集代碼失敗，這僅表示每次轉儲的平均值應該是常數（並且不爲零）。

Answer 1

我在這裏看到的最大問題是範圍之一。這段代碼的寫法讓我得出結論：你可能不瞭解C++中的變量範圍是如何工作的，以及設備和主機代碼範圍在CUDA中的工作原理。一些觀察：

1. 當你在代碼中這種類型的事情：

   __device__ double *dev_arrA, *dev_arrB; 
   __global__ void TEST(int step, double dev_arrA[], double dev_arrB[], ....)

   你有一個變量範圍的問題。在編譯單元作用域和函數作用域都聲明瞭dev_arrA。這兩個聲明不引用相同的變量 - 函數單元範圍聲明（在內核中）優先於內核中的編譯單元範圍聲明。您修改該變量，您正在修改內核範圍聲明，而不是__device__變量。這可能會導致各種微妙和未預料的行爲。避免有具有在多個作用域中聲明的相同變量會好得多。

2. 在聲明使用__device__說明符的變量，它意在是排他性設備上下文符號，只應在設備代碼直接使用。所以像這樣：

   __device__ double *dev_arrA; 
   int main() 
   { 
   .... 
   cudaMalloc((void**)&dev_arrA, statCnt*sizeof(double)); 
   .... 
   }

   是非法的。您無法直接在__device__變量上調用像cudaMalloc這樣的API函數。即使它會編譯（因爲主機和設備代碼中涉及CUDA編譯軌跡的hackery），但這樣做是不正確的。在上面的例子中，dev_arrA是一個設備符號。您可以通過API符號操作調用與它進行交互，但這在技術上是合法的。在你的代碼中，用於保存設備指針並作爲內核參數傳遞的變量（如dev_arrA）應該在main()範圍內聲明，並按值傳遞給內核。

這是上述兩件事情的組合，可能會導致您的問題。

但是，難點在於您選擇將粗糙的150行代碼（其中很多是冗餘的）作爲repro格式發佈。我懷疑任何人都會關心你的問題，用精細的梳子去仔細閱讀那些代碼，並確定問題的具體位置。此外，您在問題中進行這些討厭的「頂級編輯」的習慣會迅速將可能已經合理編寫起點的內容變成難以理解的難以理解的虛擬變更日誌，並且不可能對任何人有所幫助。此外，輕度被動攻擊性票據部分沒有真正的目的 - 它沒有增加任何有價值的問題。

因此，我會留下一個您發佈的代碼的大大簡化版本，我認爲這些代碼包含您嘗試工作的所有基本功能。我把它作爲一個「讀者的練習」，將它重新轉化爲你正在嘗試做的事情。

#include "stdio.h" 

typedef float Real; 
struct __align__(8) DynamicVals 
{ 
    Real a; 
    int n1; 
    int perDump; 
}; 

__device__ int stepsA; 
__device__ Real sumA; 
__device__ int stepsN1; 
__device__ int sumN1; 

__global__ void TEST 
(int step, Real dev_arrA[], int dev_arrN1[], DynamicVals *dev_myVals) 
{ 
    if (step % dev_myVals->perDump) 
    { 
     dev_arrN1[step/dev_myVals->perDump] = 0; 
     dev_arrA[step/dev_myVals->perDump] = 0.0; 
     stepsA = 0; 
     stepsN1 = 0; 
     sumA = 0.0; 
     sumN1 = 0; 
    } 

    sumA += dev_myVals->a; 
    sumN1 += dev_myVals->n1; 
    stepsA++; 
    stepsN1++; 

    dev_arrA[step/dev_myVals->perDump] += sumA/stepsA; 
    dev_arrN1[step/dev_myVals->perDump] += sumN1/stepsN1; 
} 

inline void gpuAssert(cudaError_t code, char *file, int line, 
       bool abort=true) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), 
      file, line); 
     if (abort) exit(code); 
    } 
} 

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } 

int main() 
{ 
    const int TOTAL_STEPS = 1000; 
    DynamicVals vals; 
    int *arrN1; 
    Real *arrA; 
    int statCnt; 

    vals.perDump = TOTAL_STEPS/10; 
    statCnt = TOTAL_STEPS/vals.perDump; 
    vals.a = 30000.0; 
    vals.n1 = 10000; 

    Real *dev_arrA; 
    int *dev_arrN1; 
    DynamicVals *dev_myVals; 

    gpuErrchk(cudaMalloc((void**)&dev_arrA, statCnt*sizeof(Real))); 
    gpuErrchk(cudaMalloc((void**)&dev_arrN1, statCnt*sizeof(int))); 
    gpuErrchk(cudaMalloc((void**)&dev_myVals, sizeof(DynamicVals))); 
    gpuErrchk(cudaMemcpy(dev_myVals, &vals, sizeof(DynamicVals), 
       cudaMemcpyHostToDevice)); 

    arrA = (Real *)malloc(statCnt * sizeof(Real)); 
    arrN1 = (int *)malloc(statCnt * sizeof(int)); 

    for (int i=0; i< TOTAL_STEPS; i++) { 
     TEST<<<1,1>>>(i, dev_arrA,dev_arrN1,dev_myVals); 
     gpuErrchk(cudaPeekAtLastError()); 
    } 

    gpuErrchk(cudaMemcpy(arrA,dev_arrA,statCnt * sizeof(Real), 
       cudaMemcpyDeviceToHost)); 
    gpuErrchk(cudaMemcpy(arrN1,dev_arrN1,statCnt * sizeof(int), 
       cudaMemcpyDeviceToHost)); 

    for (int i=0; i< statCnt; i++) 
    { 
     printf("Step: %d ; A=%g N1=%d\n", 
       i*vals.perDump, arrA[i], arrN1[i]); 
    } 
} 

Answer 2

幾件事情：

好像你調用cudaMemcpy爲dev_MyVals你正在爲它調用cudaMalloc之前。這不是應該如此。

另外：當您執行cudaMalloc調用時，您不會乘以sizeof int。

您應該真的檢查所有CUDA調用cudaMalloc/cudaMemcpy的錯誤代碼。他們都應該返回錯誤或CUDA_SUCCESS。我相信CUDA的例子都展示瞭如何做到這一點。

另外，爲了將來的參考不要在CUDA中使用模運算符，它非常慢。只是谷歌的「模數CUDA」的一些替代品。

讓我知道它是怎麼回事，這可能需要幾次迭代才能解決。