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;


		//Use half the defect






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