在 cuda 中使用共享内存会导致内存写入错误答案

【问题标题】：using shared memory in cuda gives memory write error在 cuda 中使用共享内存会导致内存写入错误
【发布时间】：2014-08-26 08:08:33
【问题描述】：

我有一个可以正常工作的内核

__global__ static void  CalcSTLDistance_Kernel(Integer ComputeParticleNumber)
{
    const Integer TID = CudaGetTargetID();
    const Integer ID  = TID;
    if(ID >= ComputeParticleNumber)
    {
        return ;
    }
    CDistance NearestDistance;
    Integer NearestID = -1;
    NearestDistance.Magnitude = 1e8;
    NearestDistance.Direction = make_Scalar3(0,0,0);
    if(c_daOutputParticleID[ID] < -1)
    {
        c_daSTLDistance[ID] = NearestDistance;
        c_daSTLID[ID] = NearestID;
        return;
    }
    Scalar3 TargetPosition = c_daParticlePosition[ID];

    Integer TriangleID;     
    Integer CIDX, CIDY, CIDZ;
    Integer CID = GetCellID(&CONSTANT_BOUNDINGBOX,&TargetPosition,CIDX, CIDY, CIDZ);
    Integer Range = 1;
    if(CID >=0 && CID < c_CellNum)
    {
        for(Integer k = -Range; k <= Range; ++k)
        {
            for(Integer j = -Range; j <= Range; ++j)
            {
                for(Integer i = -Range; i <= Range; ++i)
                {
                    Integer MCID = GetCellID(&CONSTANT_BOUNDINGBOX,CIDX +i, CIDY + j,CIDZ + k);
                    if(MCID < 0 || MCID >= c_CellNum)
                    {
                        continue;
                    }
                    unsigned int TriangleNum = c_daCell[MCID].m_TriangleNum;
                    for(unsigned int l = 0; l < TriangleNum; ++l)
                    {
                        TriangleID = c_daCell[MCID].m_TriangleID[l];
                        if( TriangleID >= 0 && TriangleID < c_TriangleNum && TriangleID != NearestID)// No need to calculate again for the same triangle
                        {
                            CDistance Distance ;
                            Distance.Magnitude = CalcDistance(&c_daTriangles[TriangleID], &TargetPosition, &Distance.Direction);
                            if(Distance.Magnitude < NearestDistance.Magnitude)
                            {
                                NearestDistance = Distance;
                                NearestID = TriangleID;
                            }
                        }
                    }   
                }
            }
        }
    }
    c_daSTLDistance[ID] = NearestDistance;
    c_daSTLID[ID] = NearestID;
}

这里 c_daParticlePosition 是常量内存 float3 数据类型。所以在这里我想使用共享内存所以我尝试创建float3类型的共享内存并尝试将常量日期复制到共享内存但是它显示未知错误并且使用cuda-memcheck它说

这里的线程数是 255，2 块大小

shared_memory 代码

__global__ static void CalcSTLDistance_Kernel(Integer ComputeParticleNumber)
{
    //const Integer TID = CudaGetTargetID();
    const Integer ID  =CudaGetTargetID(); 
    extern __shared__ float3 s[];
    /*if(ID >= ComputeParticleNumber)
    {
        return ;
    }*/
    s[ID] = c_daParticlePosition[ID];
    __syncthreads();

    CDistance NearestDistance;
    Integer NearestID = -1;
    NearestDistance.Magnitude = 1e8;
    NearestDistance.Direction.x = 0;
    NearestDistance.Direction.y = 0;
    NearestDistance.Direction.z = 0;//make_Scalar3(0,0,0);
    //if(c_daOutputParticleID[ID] < -1)
    //{
    //  c_daSTLDistance[ID] = NearestDistance;
    //  c_daSTLID[ID] = NearestID;
    //  return;
    //}

    //Scalar3 TargetPosition = c_daParticlePosition[ID];

    Integer TriangleID;     
    Integer CIDX, CIDY, CIDZ;
    Integer CID = GetCellID(&CONSTANT_BOUNDINGBOX,&s[ID],CIDX, CIDY, CIDZ);
    if(CID >=0 && CID < c_CellNum)
    {
        //Integer Range = 1;
        for(Integer k = -1; k <= 1; ++k)
        {
            for(Integer j = -1; j <= 1; ++j)
            {
                for(Integer i = -1; i <= 1; ++i)
                {
                    Integer MCID = GetCellID(&CONSTANT_BOUNDINGBOX,CIDX +i, CIDY + j,CIDZ + k);
                    if(MCID < 0 || MCID >= c_CellNum)
                    {
                        continue;
                    }
                    unsigned int TriangleNum = c_daCell[MCID].m_TriangleNum;
                    for(unsigned int l = 0; l < TriangleNum; ++l)
                    {
                        TriangleID = c_daCell[MCID].m_TriangleID[l];
                        /*if(c_daTrianglesParameters[c_daTriangles[TriangleID].ModelIDNumber].isDrag)
                        {
                            continue;
                        }*/

                        if( TriangleID >= 0 && TriangleID < c_TriangleNum && TriangleID != NearestID)// No need to calculate again for the same triangle
                        {
                        CDistance Distance ;
                            Distance.Magnitude = CalcDistance(&c_daTriangles[TriangleID], &s[ID], &Distance.Direction);
                            if(Distance.Magnitude < NearestDistance.Magnitude)
                            {
                                NearestDistance = Distance;
                                NearestID = TriangleID;
                            }
                        }
                    }   
                }
            }
        }
    }
    c_daSTLDistance[ID] = NearestDistance;
    c_daSTLID[ID] = NearestID;
}

错误

  Invalid __shared__ write of size 4
    =========     at 0x00000128 in CalcSTLDistance_Kernel(int)
    =========     by thread (159,0,0) in block (0,0,0)
    =========     Address 0x0000077c is out of bounds

【问题讨论】：

最后一个代码是给出错误的当前代码，共享内存大小是默认值。我正在使用 GTX 750 Ti，它的计算能力为 5.0

标签： cuda shared-memory

【解决方案1】：

您可以在this article 中找到有关如何使用共享内存的有用信息。特别关注静态共享内存和动态共享内存部分。

根据上面的文章，您应该会发现您只是在写入超出数组s 的范围，正如错误消息所说的那样。要解决此问题，您可以：

在编译时指定共享内存数组s的大小，如果你提前知道的话，比如__shared__ float3 s[123456];
或使用动态大小的s 数组，这基本上是您目前正在做的事情，但是ALSO 将第三个内核启动参数指定为CalcSTLDistance_Kernel<<<gridSize, blockSize, sharedMemorySizeInBytes>>>。如果您将使用 123456 float3s 的数组，请使用 int sharedMemorySizeInBytes = 123456 * sizeof(float3)

【讨论】：

感谢您提到第三个内核启动参数需要是字节大小。在我的研究中，没有人提到这个小细节，我错误地认为它是元素的数量。