最小縮減cuda不起作用

我寫了一個代碼來找到最小化縮減。但結果始終爲零。我不知道是什麼問題。請幫幫我。最小縮減cuda不起作用

這裏是內核代碼：我修改了Nvidia的總和縮減代碼。

#include <limits.h> 

#define NumThread 128 
#define NumBlock 32 

__global__ void min_reduce(int* In, int* Out, int n){ 
    __shared__ int sdata[NumThread]; 
    unsigned int i = blockIdx.x * NumThread + threadIdx.x; 
    unsigned int tid = threadIdx.x; 
    unsigned int gridSize = NumBlock * NumThread; 
    int myMin = INT_MAX; 

    while (i < n){ 
    if(In[i] < myMin) 
    myMin = In[i]; 
    i += gridSize; 
    } 
    sdata[tid] = myMin; 
    __syncthreads(); 

    if (NumThread >= 1024){ 
    if (tid < 512) 
    if(sdata[tid] > sdata[tid + 512]) sdata[tid] = sdata[tid + 512]; 
    __syncthreads(); 
    } 
    if (NumThread >= 512){ 
    if(sdata[tid] > sdata[tid + 256]) sdata[tid] = sdata[tid + 256]; 
    __syncthreads(); 
    } 
    if (NumThread >= 256){ 
    if(sdata[tid] > sdata[tid + 128] && sdata[tid + 128] !=0) sdata[tid] = sdata[tid + 128]; 
    __syncthreads(); 
    } 
    if (NumThread >= 128){ 
    if(sdata[tid] > sdata[tid + 64]) sdata[tid] = sdata[tid + 64]; 
    __syncthreads(); 
    } 
    //the following practice is deprecated 
    if (tid < 32){ 
    volatile int *smem = sdata; 
    if (NumThread >= 64) if(smem[tid] > smem[tid + 32]) smem[tid] = smem[tid+32]; 
    if (NumThread >= 32) if(smem[tid] > smem[tid + 16]) smem[tid] = smem[tid+16]; 
    if (NumThread >= 16) if(smem[tid] > smem[tid + 8]) smem[tid] = smem[tid+8]; 
    if (NumThread >= 8) if(smem[tid] > smem[tid + 4]) smem[tid] = smem[tid+4]; 
    if (NumThread >= 4) if(smem[tid] > smem[tid + 2]) smem[tid] = smem[tid+2]; 
    if (NumThread >= 2) if(smem[tid] > smem[tid + 1])  smem[tid] = smem[tid+1]; 
    } 
    if (tid == 0) 
    if(sdata[0] < sdata[1]) Out[blockIdx.x] = sdata[0]; 
    else Out[blockIdx.x] = sdata[1];  
}

在這裏，這是我的主要代碼：

#include <stdio.h> 
#include <stdlib.h> 

#include "min_reduction.cu" 

int main(int argc, char* argv[]){ 
    unsigned int length = 1048576; 
    int i, Size, min; 
    int *a, *out, *gpuA, *gpuOut; 

    cudaSetDevice(0); 
    Size = length * sizeof(int); 
    a = (int*)malloc(Size); 
    out = (int*)malloc(NumBlock*sizeof(int)); 
    for(i=0;i<length;i++) a[i] = (i + 10); 

    cudaMalloc((void**)&gpuA,Size); 
    cudaMalloc((void**)&gpuOut,NumBlock*sizeof(int)); 
    cudaMemcpy(gpuA,a,Size,cudaMemcpyHostToDevice); 
    min_reduce<<<NumBlock,NumThread>>>(gpuA,gpuOut,length); 
    cudaDeviceSynchronize(); 
    cudaMemcpy(out,gpuOut,NumBlock*sizeof(int),cudaMemcpyDeviceToHost); 

    min = out[0]; 
    for(i=1;i<NumBlock;i++) if(min < out[i]) min = out[i]; 
    return 0; 
}

來源

2016-02-12 Hamid_UMB

幫你一個忙，並用[適當的錯誤檢查]包圍所有你的CUDA調用（http://stackoverflow.com/questions/14038589/what-is-the-canonical-way-to-check-for-錯誤 - 使用最CUDA的運行時API）。你的代碼有不正確的同步，'__syncthreads'應該被執行[由塊中的所有線程]（http://stackoverflow.com/questions/6666382/can-i-use-syncthreads-after-having-dropped-threads）。此外，在這裏使用'volatile'是不正確的，歪曲同步編程被破壞**，請不要使用它。 –

我不知道我的一切，@HubertApplebaum表示同意，但我可以用proper cuda error checking的建議達成一致。正如你在代碼中提到的，warp同步編程可以被認爲是已棄用，但我不能支持它是中斷（還）的說法。但我不想爲此爭論;這不是你的問題的核心。

另一個有用的調試建議將遵循步驟here編譯您的代碼與-lineinfo並運行您的代碼與cuda-memcheck。如果你這樣做，你會看到很多這樣的報道：

========= Invalid __shared__ read of size 4 
=========  at 0x000001e0 in /home/bob/misc/t1074.cu:39:min_reduce(int*, int*, int) 
=========  by thread (64,0,0) in block (24,0,0) 
=========  Address 0x00000200 is out of bounds 
=========  Saved host backtrace up to driver entry point at kernel launch time 
=========  Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x2cd) [0x15859d] 
=========  Host Frame:./t1074 [0x16dc1] 
=========  Host Frame:./t1074 [0x315d3] 
=========  Host Frame:./t1074 [0x28f5] 
=========  Host Frame:./t1074 [0x2623] 
=========  Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x21d65] 
=========  Host Frame:./t1074 [0x271d]

這將表明雙方是在你的代碼的主要問題是，你是不正確索引到你的__shared__存儲器陣列以及具體線路代碼在哪裏發生。整齊！（對我而言，這是第39行，但在您的情況下可能會有不同的行）。如果再鑽入那行，你將要學習的這部分代碼：

#define NumThread 128 
    ... 
    __shared__ int sdata[NumThread]; 
    ... 
    if (NumThread >= 128){ 
    if(sdata[tid] > sdata[tid + 64]) sdata[tid] = sdata[tid + 64]; //line 39 in my case 
    __syncthreads(); 
    }

您已經定義NumThread在128，並且有靜態分配的，許多int數量的共享存儲陣列。一切都很好。 if語句中的代碼怎麼樣？ if條件將得到滿足，這意味着塊中的所有128個線程將執行該if語句的主體。但是，您正在從共享內存中讀取sdata[tid + 64]，而對於tid大於63（即每個塊中的線程數量的一半）的線程，這會生成一個大於127的共享內存索引（這是超出範圍即非法）。

（你已經顯示了這具體代碼）的修補程序相當簡單，只需添加另一種，如果測試：

if (NumThread >= 128){ 
    if (tid < 64) 
     if(sdata[tid] > sdata[tid + 64]) sdata[tid] = sdata[tid + 64]; 
    __syncthreads(); 
    }

如果你作出這樣的修改你的代碼，並重新運行cuda-memcheck測試，你會看到所有運行時報告的錯誤都消失了。好極了！

但是代碼仍然沒有產生正確的答案。你在這裏做另一個錯誤：

for(i=1;i<NumBlock;i++) if(min < out[i]) min = out[i];

如果你想找到最低值，並仔細想想這個邏輯，你會意識到你應該這樣做：

for(i=1;i<NumBlock;i++) if(min > out[i]) min = out[i]; 
           ^
           | 
           greater than

有了這兩個變化，你的代碼產生正確的結果對我來說：

$ cat t1074.cu 
#include <stdio.h> 
#include <stdlib.h> 


#include <limits.h> 

#define NumThread 128 
#define NumBlock 32 

__global__ void min_reduce(int* In, int* Out, int n){ 
    __shared__ int sdata[NumThread]; 
    unsigned int i = blockIdx.x * NumThread + threadIdx.x; 
    unsigned int tid = threadIdx.x; 
    unsigned int gridSize = NumBlock * NumThread; 
    int myMin = INT_MAX; 

    while (i < n){ 
    if(In[i] < myMin) 
    myMin = In[i]; 
    i += gridSize; 
    } 
    sdata[tid] = myMin; 
    __syncthreads(); 

    if (NumThread >= 1024){ 
    if (tid < 512) 
    if(sdata[tid] > sdata[tid + 512]) sdata[tid] = sdata[tid + 512]; 
    __syncthreads(); 
    } 
    if (NumThread >= 512){ 
    if(sdata[tid] > sdata[tid + 256]) sdata[tid] = sdata[tid + 256]; 
    __syncthreads(); 
    } 
    if (NumThread >= 256){ 
    if(sdata[tid] > sdata[tid + 128] && sdata[tid + 128] !=0) sdata[tid] = sdata[tid + 128]; 
    __syncthreads(); 
    } 
    if (NumThread >= 128){ 
    if (tid < 64) 
    if(sdata[tid] > sdata[tid + 64]) sdata[tid] = sdata[tid + 64]; 
    __syncthreads(); 
    } 
    //the following practice is deprecated 
    if (tid < 32){ 
    volatile int *smem = sdata; 
    if (NumThread >= 64) if(smem[tid] > smem[tid + 32]) smem[tid] = smem[tid+32]; 
    if (NumThread >= 32) if(smem[tid] > smem[tid + 16]) smem[tid] = smem[tid+16]; 
    if (NumThread >= 16) if(smem[tid] > smem[tid + 8]) smem[tid] = smem[tid+8]; 
    if (NumThread >= 8) if(smem[tid] > smem[tid + 4]) smem[tid] = smem[tid+4]; 
    if (NumThread >= 4) if(smem[tid] > smem[tid + 2]) smem[tid] = smem[tid+2]; 
    if (NumThread >= 2) if(smem[tid] > smem[tid + 1])  smem[tid] = smem[tid+1]; 
    } 
    if (tid == 0) 
    if(sdata[0] < sdata[1]) Out[blockIdx.x] = sdata[0]; 
    else Out[blockIdx.x] = sdata[1]; 
} 

int main(int argc, char* argv[]){ 
    unsigned int length = 1048576; 
    int i, Size, min; 
    int *a, *out, *gpuA, *gpuOut; 

    cudaSetDevice(0); 
    Size = length * sizeof(int); 
    a = (int*)malloc(Size); 
    out = (int*)malloc(NumBlock*sizeof(int)); 
    for(i=0;i<length;i++) a[i] = (i + 10); 
    a[10]=5; 
    cudaMalloc((void**)&gpuA,Size); 
    cudaMalloc((void**)&gpuOut,NumBlock*sizeof(int)); 
    cudaMemcpy(gpuA,a,Size,cudaMemcpyHostToDevice); 
    min_reduce<<<NumBlock,NumThread>>>(gpuA,gpuOut,length); 
    cudaDeviceSynchronize(); 
    cudaMemcpy(out,gpuOut,NumBlock*sizeof(int),cudaMemcpyDeviceToHost); 

    min = out[0]; 
    for(i=1;i<NumBlock;i++) if(min > out[i]) min = out[i]; 
    printf("min = %d\n", min); 
    return 0; 
} 
$ nvcc -o t1074 t1074.cu 
$ cuda-memcheck ./t1074 
========= CUDA-MEMCHECK 
min = 5 
========= ERROR SUMMARY: 0 errors 
$

注意，你已經在1024個線程情況下，如果檢查，你可能想要爲512和256線程案例添加一個適當的if-check，就像我爲上面的128個線程案例添加它一樣。

來源

2016-02-12 05:16:30

謝謝你的回答。代碼工作正常，但是當我改變輸入數組a，並添加這一行a [10] = 5;初始化後。代碼，找不到最小值？你能告訴我我的代碼有什麼問題嗎？ –

我加了一個[10] = 5;初始化到我在我的答案張貼的代碼後，它似乎工作正常。（我已經在我的答案中更新了完整的測試案例以證明這一點）。也許您應該針對您的新代碼提出一個新問題，而這個新問題不起作用。您是否添加了適當的cuda錯誤檢查並使用'cuda-memcheck'運行您的代碼？ –

謝謝......問題已解決 –

最小縮減cuda不起作用

回答

相關問題