     {
         dim3 threads(BSIZE,1);
         dim3 grid((nbPoints+BSIZE-1)/BSIZE,1);
-        computeHashD<float4><<< grid, threads >>>((const float4*)x, (uint*)particleIndex8, (uint*)particleHash8, nbPoints);
+        computeHashD<float4><<< grid, threads >>>((const float4*)x, (unsigned int*)particleIndex8, (unsigned int*)particleHash8, nbPoints);
     }
 }
 

    {
        dim3 threads(BSIZE,1);
        dim3 grid((nbPoints+BSIZE-1)/BSIZE,1);
        computeHashD<float3><<< grid, threads >>>((const float3*)x, (uint*)particleIndex8, (uint*)particleHash8, nbPoints);
    }
}


// find start of each cell in sorted particle list by comparing with previous hash value
// one thread per particle
__global__ void
findCellRangeD(const unsigned int* particleHash,
        int2 * cellRange, int* cellGhost, int n)
{
    unsigned int i = __mul24(blockIdx.x, blockDim.x) + threadIdx.x;
    __shared__ unsigned int hash[BSIZE];
    if (i < n)
        hash[threadIdx.x] = particleHash[i];

