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

        h.x = (h.x>>3) + dH.x;
        h.y += dH.y;
        h.z += dH.z;
        uint hash = ((h.x ^ h.y ^ h.z) & gridParams.cellMask)<<1;
        if (hc != (x&7)) ++hash;
        particleHash8[index0_8 + l] = hash;
    }

// 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];

