help! kernel compile status as warning

[font=“Courier”]I compile the kernel as down and it return :

warning:latest line of file ends without a newline

#endif

kernel

#ifndef __GEN_PARTICLES_KERNEL__

#define __GEN_PARTICLES_KERNEL__

#define IMAD(X,Y,Z) (__mul24(X,Y)+Z)

#define IMUL(X,Y) __mul24(X,Y)

__global__ 

void GenParticles_kernel(float3* oPos,float4* oCol,float scale)

{

    unsigned int tidx=IMAD(blockIdx.x,blockDim.x,threadIdx.x);

    unsigned int tidy=IMAD(blockIdx.y,blockDim.y,threadIdx.y);

    unsigned int tidz=IMAD(blockIdx.z,blockDim.z,threadIdx.z);

    unsigned int threadsPerCol=IMUL(blockDim.x,gridDim.x);

    unsigned int threadsPerRow=IMUL(blockDim.y,gridDim.y);

    unsigned int oid=IMUL(IMAD(tidz,threadsPerRow,tidy),threadsPerCol)+tidx;

   float sPosX=(float)tidx;

    float sPosY=(float)tidy;

    sPosX*=scale;

    sPosY*=scale;

    oPos[oid].x=sPosX;

    oPos[oid].y=sPosY;

    oPos[oid].z=scale*gridDim.y;

   __syncthreads();

}

#endif

but when changed to this ,it as good:

//device code for generate cube mesh particles

#ifndef __GEN_PARTICLES_KERNEL__

#define __GEN_PARTICLES_KERNEL__

#define IMAD(X,Y,Z) (__mul24(X,Y)+Z)

#define IMUL(X,Y) __mul24(X,Y)

texture<float4,2,cudaReadModeElementType> f4tex;

__global__ 

void GenParticles_kernel(float3* oPos,float4* oCol,float scale)

{

    unsigned int tidx=IMAD(blockIdx.x,blockDim.x,threadIdx.x);

    unsigned int tidy=IMAD(blockIdx.y,blockDim.y,threadIdx.y);

    unsigned int tidz=IMAD(blockIdx.z,blockDim.z,threadIdx.z);

    unsigned int threadsPerCol=IMUL(blockDim.x,gridDim.x);

    unsigned int threadsPerRow=IMUL(blockDim.y,gridDim.y);

   //global index for per-particles

    //z*witdh*height+y*width+x

    unsigned int oid=IMUL(IMAD(tidz,threadsPerRow,tidy),threadsPerCol)+tidx;

   

    //caculate texcoords

    float u=tidx/(float)threadsPerCol;

    float v=tidy/(float)threadsPerRow;

    u+=0.5F;

    v+=0.5F;

    //compute out pos of per particle with per thread

    float sPosX=(float)tidx;

    float sPosY=(float)tidy;

    sPosX*=scale;

    sPosY*=scale;

   oPos[oid].x=sPosX;

    oPos[oid].y=sPosY;

    oPos[oid].z=scale*gridDim.y;//gridDim.y determinates the z-offset of per layer point mesh

    oCol[oid]=tex2D(f4tex,u,v);

    __syncthreads();

}

#endif

And this kernel is in same case:

#ifndef __UPDATE_PARTICLES_KERNEL__

#define __UPDATE_PARTICLES_KERNEL__

#define IMAD(X,Y,Z) (__mul24(X,Y)+Z)

#define IMUL(X,Y) __mul24(X,Y)

#define __UPDATE_COLOR__

__global__ void transferParticles_kernel(

float3* newPos,

float3* oldPos,

#ifdef __UPDATE_COLOR__

float4* newCol,

float4* oldCol,

float scaleR,float scaleG,float scaleB,float scaleA,

#endif

float dt,

float forceX,float forceY,float forceZ,

float attenCoeffXA,float attenCoeffXB,float attenCoeffXC,

float attenCoeffYA,float attenCoeffYB,float attenCoeffYC,

float attenCoeffZA,float attenCoeffZB,float attenCoeffZC)

{

    unsigned int tidx=IMAD(blockIdx.x,blockDim.x,threadIdx.x);

    unsigned int tidy=IMAD(blockIdx.y,blockDim.y,threadIdx.y);

    unsigned int tidz=IMAD(blockIdx.z,blockDim.z,threadIdx.z);

    unsigned int threadsPerCol=IMUL(blockDim.x,gridDim.x);

    unsigned int threadsPerRow=IMUL(blockDim.y,gridDim.y);

    unsigned int oid=IMUL(IMAD(tidz,threadsPerRow,tidy),threadsPerCol)+tidx;

   float attenuationX=1.0F/(dt*(dt*attenCoeffXA+attenCoeffXB)+attenCoeffXC);

    float attenuationY=1.0F/(dt*(dt*attenCoeffYA+attenCoeffYB)+attenCoeffYC);

    float attenuationZ=1.0F/(dt*(dt*attenCoeffZA+attenCoeffZB)+attenCoeffZC);

#ifdef __UPDATE__COLOR__

    newCol[oid].x=scaleR*oldCol[oid].x;

    newCol[oid].y=scaleG*oldCol[oid].y;

    newCol[oid].z=scaleB*oldCol[oid].z;

    newCol[oid].w=scaleA*oldCol[oid].w;

#endif

   newPos[oid].x=attenuationX*sinf(dt)*cosf(dt/2.0F)*forceX*oldPos[oid].x;

    newPos[oid].y=attenuationY*sinf(dt)*cosf(dt/2.0F)*forceY*oldPos[oid].y;

    newPos[oid].z=attenuationZ*cosf(dt)*forceZ*oldPos[oid].z;

   __syncthreads();

}

#endif

why?

Thanks!

[/font]