2

I have a kernel code that executes properly runnable code

__global__ static void CalcSTLDistance_Kernel(Integer ComputeParticleNumber)
{
    //const Integer TID = CudaGetTargetID();
    const Integer ID  =CudaGetTargetID(); 
    /*if(ID >= ComputeParticleNumber)
    {
        return ;
    }*/
    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,&c_daParticlePosition[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], &c_daParticlePosition[ID], &Distance.Direction);
                            if(Distance.Magnitude < NearestDistance.Magnitude)
                            {
                                NearestDistance = Distance;
                                NearestID = TriangleID;
                            }
                        }
                    }   
                }
            }
        }
    }
    c_daSTLDistance[ID] = NearestDistance;
    c_daSTLID[ID] = NearestID;
}

and when I add any basic variables or perform any checking operation, it gives unknown error and while checking wih cuda-memcheck, it suggests memory read error.

here in the changed code, i tried to check the previously calculated part and tried to skip the redundant calculation. for this I tried to perform basic check operation in array but it throws memory error.

error raising code

__global__ static void CalcSTLDistance_Kernel(Integer ComputeParticleNumber)
{
    //const Integer TID = CudaGetTargetID();
    const Integer ID  =CudaGetTargetID(); 
    /*if(ID >= ComputeParticleNumber)
    {
        return ;
    }*/
    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,&c_daParticlePosition[ID],CIDX, CIDY, CIDZ);
    int len=0;
    int td[100];
    for(int m=0;m<100;m++)
    {
      td[m]=-1;

    }
    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;

                    
                     
                     bool flag = false;
                      
                       //len=len+TriangleNum ;
                     for(unsigned int l = 0; l < TriangleNum; ++l)
                     {
                        TriangleID = c_daCell[MCID].m_TriangleID[l]; 
                        //tem[l] = c_daCell[MCID].m_TriangleID[l];

                        for(int m=0;m<100;m++)
                        {
                            if(TriangleID ==td[m])
                            {
                                flag= true;
                                
                                
                                
                            }
                            if(flag == true)
                            break;
                            

                        }

                        

                        if(flag == true)
                            continue;
                        else 
                        {
                            td[len] = TriangleID;
                            len= len+1;

                            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], &c_daParticlePosition[ID], &Distance.Direction);
                              if(Distance.Magnitude < NearestDistance.Magnitude)
                                {
                                NearestDistance = Distance;
                                NearestID = TriangleID;
                                }
                             }
                         
                        }
                     }

                }
            }
        }
    }
    c_daSTLDistance[ID] = NearestDistance;
    c_daSTLID[ID] = NearestID;
}

this problem arises whenever I tried to add any piece of code,thus I suspects that this block of kernel is not allowing me to add any further code due to memory over use. is there any memory violation rule per block or thread??

how to find the total memory usuage per kernel ?? is there any way??

sam3376
  • 33
  • 6
Laxmi Kadariya
  • 1,103
  • 1
  • 14
  • 34
  • Please fix the code formatting in the second code example. That is extremely difficult to read as posted. – talonmies Jul 21 '14 at 07:17
  • I have an array and tried to insert the previously non calculated triangleID to the end of array and thus this array is checked to find if the present triangleID is in the array or not . if it is already there ,it skips the calculation. This code runs well in the case when there are less triangle but if the triangle is increased it gives unknown error. – Laxmi Kadariya Jul 21 '14 at 07:43
  • 1
    In your second code, each thread is allocating `100` `ints`. Depending on your kernel launch configuration and on your hardware, this can require too much memory. – Vitality Jul 21 '14 at 08:23
  • I am using gtx 780 gpu. does this allocation of 100 ints exceeds the thread memory limit. what is the maximum limit for each thread?? – Laxmi Kadariya Jul 21 '14 at 08:59
  • How many threads are you launching? – Vitality Jul 21 '14 at 09:51
  • @JackOLantern I am using 256 threads – Laxmi Kadariya Jul 21 '14 at 10:10
  • 1
    @JackOLantern: The number of threads should be irrelevent. Local memory reservations are (as of CUDA 6 still, I believe) made *a priori* for the maximum number of threads the target GPU supports for the maximum local memory footprint loaded from a particular module. – talonmies Jul 21 '14 at 12:14
  • You increase len by 1 in the inner most loop which can be as big as 3*3*3*c_daCell[MCID].m_TriangleNum. Could it be that len grows too large and goes above 100? – deathly809 Jul 21 '14 at 13:01

0 Answers0