Low theoretical occupancy for kernel

I have a kernel which, according to the Nsight analysis tool, could only ever achieve 50% occupancy. Anyone have any ideas why? Thanks.

__global__ void restriction(double *d_u, double *d_rhs, double *d_res ,int d_dimUcoarse, int d_dimUFine){

	//Thread Indices

	int x = blockIdx.x*blockDim.x+threadIdx.x+1;

	int y = blockIdx.y*blockDim.y+threadIdx.y+1;

	if(x<d_dimUcoarse&&y<d_dimUcoarse){

		//Use half the defect

		d_rhs[y*d_dimUcoarse+x]=d_res[2*y*d_dimUFine+2*x]/2.0;

		d_u[y*d_dimUcoarse+x]=0;

	}

	

}

Never mind, it was the /2. replaced with *0.5 and all is well. I would’ve thought the compiler would handle that!