thrust count occurence

Question

Possible Duplicate:
Counting occurences of numbers in cuda array

Is there a way to use Thrust or CUDA to count occurrences for the duplicates in an array?

For example:
If I have a device vector

{11, 11, 9, 1, 3, 11, 1, 2, 9, 1, 11}

I should get

1:3, 2:1, 3:1, 9:2, 11:4

If Thrust cannot do that, how can I use a kernel to do that?

Thanks! I am doing a concentration calculation. That's why I am asking this question.
Assume that there are 100000 particles in the domain which has nx * ny * nz cells, I need to calculate the concentration of each cell (how many particles in each cell).

My kernel is this

__global__ void concentration_kernel(float3* posPtr, uint* device_cons) 
{
    __shared__ uint cache[256];
    uint x = threadIdx.x + blockIdx.x * blockDim.x;
    uint y = threadIdx.y + blockIdx.y * blockDim.y;
    uint offset = x + y * blockDim.x * gridDim.x; 

    float3 posf3 = posPtr[offset];//make_float3(43.5,55,0.66);//
    uint cellIndex = (uint)(posf3.z+1)*153*110 + (uint)(posf3.y)*153 + (uint)posf3.x;

    cache[threadIdx.x] = device_cons[cellIndex];
    __syncthreads();
    uint a = cache[threadIdx.x];
    a++;
    cache[threadIdx.x] = a;
    __syncthreads();

    device_cons[cellIndex] = cache[threadIdx.x]; 
}

Answer 1

You can first sort the vector using thrust::sort and then use thrust::reduce_by_key. However, you also need to create a new vector (called values) of 1's (and of the same length as your sorted vector) after sort. These values will be added up to get the counts:

reduce_by_key is a generalization of reduce to key-value pairs. For each group of consecutive keys in the range [keys_first, keys_last) that are equal, reduce_by_key copies the first element of the group to the keys_output. The corresponding values in the range are reduced using the plus and the result copied to values_output.