是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优化此代码不好还是我错了？

回答

相关问题