是CUDA優化此代碼不好還是我錯了？

這是一段代碼，我一直在努力，並得到了我沒有想到的結果。我已經減少了我的完整代碼塊來突出顯示問題。我預計在這個塊的末尾spID應該是一個tid值的塊，但那些像素的例外是lbBool是true，其中spID應該是_CCL_SHARED_MEM_MAX_VALUE (255)。但是如果我使用NSight在__syncthreads()調試數據，我覺得所有的spID值等同的情況lbBool是true爲0是CUDA優化此代碼不好還是我錯了？

我塊由16個線程組成16所以uint8足以存儲的所有值（ 0-255）。我意識到將會有一個ID爲255的有效像素和一個值爲255的壞點。這很好。

我使用unsigned long代替tOut。

在這種情況下，我的圖像是100x100，但它在我嘗試過的每個圖像尺寸上都失敗。我在GTX 580上運行，並定期使用256線程的內核。

調用內核：

#define _CCL_SHARED_MEM_TYPE uint8 
#define _CCL_SHARED_MEM_MAX_VALUE 255 

template<class tOut> tOut *nsGPUBaseClasses::IbxCCL4Link(bool *lbEdges,uint32 liImageWidth,uint32 liImageHeight,tOut *lpOut) 
{ 
dim3 liThreads(16,16); 
dim3 liBlocks((liImageWidth+liThreads.x-1)/liThreads.x,(liImageHeight+liThreads.y-1)/liThreads.y); 

if(lpOut == nullptr) _CHECK_CUDA_ERROR(cudaMalloc(&lpOut,sizeof(tOut)*liImageWidth*liImageHeight)); 

IbxCCL4LinkCUDA<<<liBlocks,liThreads,(sizeof(_CCL_SHARED_MEM_TYPE)*liThreads.x*liThreads.y+sizeof(bool)*2)>>>(lbEdges,liImageWidth,liImageHeight,lpOut); 

_CHECK_CUDA_ERROR_EMPTY(); 

return lpOut; 
}

而且內核本身：

template<class tOut> void __global__ IbxCCL4LinkCUDA(bool *lbBool,unsigned long liImageWidth,unsigned long liImageHeight,tOut *lpOut) 
{ 
    // Shared Memory 
    __shared__ float lbSpecific[]; 
    _CCL_SHARED_MEM_TYPE *spID=reinterpret_cast<_CCL_SHARED_MEM_TYPE*>(&lbSpecific); 

    //IDs for thread 
    unsigned long tid = threadIdx.x+threadIdx.y*blockDim.x; 
    unsigned long liXPos = threadIdx.x+blockIdx.x*blockDim.x; 
    unsigned long liYPos = (threadIdx.y+blockIdx.y*blockDim.y); 

    //Check if it is in image bounds 
    if(liXPos>=liImageWidth || liYPos>=liImageHeight) return; 
    unsigned long liPPos = liXPos+liYPos*liImageWidth; 

    //If Boolean is true 
    if(lbBool[liPPos]) 
    { 
     spID[tid] = _CCL_SHARED_MEM_MAX_VALUE;  
     lpOut[liPPos] =liImageWidth*liImageHeight; 
     return; 
    } 
    lpOut = &lpOut[liPPos]; 
    lpOut[0] = (blockIdx.x+blockIdx.y*gridDim.x)*(_CCL_SHARED_MEM_MAX_VALUE+1); 

    spID[tid] = tid; 
    __syncthreads(); 

    //More Processing Goes Here 

    lpOut[0] += static_cast<tOut>(spID[tid]); 
}

如果這是等同的位置輸出255或0至lbBool是true？如果它爲零，則此Cuda將寫入共享內存優化出來？有沒有一種方法可以使布爾檢查值設置爲255？

來源

2014-03-04 Thormidable

您的共享內存分配已損壞。 __shared__ float lbSpecific;分配一個浮點值。然後，您將spID設置爲該地址，並使用遠遠超出單個浮點分配的位置。

只需分配您想要的共享內存，使用正確的大小和類型，並跳過類型轉換。

__shared__ _CCL_SHARED_MEM_TYPE spID[TOTAL_BLOCK_SIZE];

來源

2014-03-04 13:50:35 Peter

我已經做了更正，如上所述。我之前在內核調用中聲明瞭一個共享內存塊，但並未將其作爲指針調用。將進行更正和測試，但我懷疑它不會解決問題，因爲所有其他值都正常工作。 – Thormidable

對不起，你完全正確。謝謝你的答案。 – Thormidable

是CUDA優化此代碼不好還是我錯了？

回答

相關問題