CUDA 优化此代码很糟糕还是我错了?

Is CUDA optimising this Code Badly or am I wrong?

本文关键字：我错了 优化 代码 CUDA      更新时间：2023-10-16

这是我一直在处理的代码块，并收到了我意想不到的结果。我已经削减了我的完整代码块，只突出了问题。我希望在这个块的末尾，spID应该是一个tid值的块，除了那些lbBooltruespID应该_CCL_SHARED_MEM_MAX_VALUE (255)的像素。但是，如果我使用 NSight 调试__syncthreads()处的数据，我发现所有spID值都等同于lbBooltrue为 0。

我的块由 16 x 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 以lbBooltrue？ 如果为零，此 Cuda 是否优化了对共享内存的写入？ 有没有办法让布尔检查将值设置为 255？

您的共享内存分配已中断。__shared__ float lbSpecific;分配一个浮点值。 然后，将spID设置为该地址，并使用远远超出单个浮点分配的位置。

只需使用正确的大小和类型分配所需的共享内存，然后跳过类型转换。

__shared__ _CCL_SHARED_MEM_TYPE spID[TOTAL_BLOCK_SIZE];
相关文章：