胡思乱想一个gpu问题

加速计算 CUDA
1

今天偶然间遇到一个问题，不知道如何解决才好，所以来这里求助下。。。
问题是这样的：
假设有一个32*32的矩阵A
然后A矩阵的元素值只可能为固定种，就假设为3种好了，取值为0，1，2
然后A矩阵内的元素0，1，2随机分布，大概意思如下
A:
0 2 1 2 0 1…1 1 1 2 2 0
2 2 1… 1 2 0
…
2 2 1… 1 2 0
0 2 1 2 0 1…1 1 1 2 2 0

遍历过程由左到右（行0-15），由上到下（列0-31），然后继续（行16-31），列（0-31）。。做成排列的效果，如下：
B：
0 0 0 …0 0 0
0 1 1…1 1 1
…
1 1 2 …2 2 2
2 2 2…2 2 2

这种情况用gpu能实现么？
如果能实现。。大概怎么操作好呢？
拜谢。。

2

[

当然能实现了…
(1)最简单的第一种想法, 我觉得这就是一个排序…
(2)第二种, 因为你只有3种值, 先来个kernel统计出0的个数,1的个数,2的个数. 然后再来个kernel去赋值得到结果矩阵. 第二种可能会比较快…特别是矩阵规模巨大的情况下.

3

[

首先统计各种情况的数量
count[ 3 ]={ 0, 0, 0 };
prefix[ 3 ];


for( int i=0; i<n; ++i ){
   ++count[ a[ i ] ];
}
prefix[ 0 ]=0;
prefix[ 1 ]=count[ 0 ];
prefix[ 2 ]=count[ 0 ]+count[ 1 ];

对于中小规模问题直接CPU更好：


  for( int i=0; i<n; ++i ){
   dst[ prefix[ a[ i ] ]++ ]=a[ i ];
  }

CUDA实现：
先在每个block内得到排好的分布，并统计各种情况的数量，
然后再对各个block的结果进行合并


step1:
__global__
void kLocalOrdering( int* d, int* nLocal,const int* s, int n )
{
   extern __shared__ int* SMEM;
   int* sp[ 3 ];
   sp[0]=SMEM;
   sp[1]=&sp0[ blockDim.x ];
   sp[2]=&sp1[ blockDim.x ];

   unsigned int tidx=blockIdx.x*blockDim.x+threadIdx.x;
   unsigned int active=tidx<n;

   sp[ 0 ][ threadIdx.x ]=~0;
   sp[ 1 ][ threadIdx.x ]=~0;
   sp[ 2 ][ threadIdx.x ]=~0;

   if( active )
   {
   e=s[ tidx ];
   sp[ e ][ threadIdx.x ]=e;
   } __syncthreads();

   //用scan或者warp vote函数统计,统计结果放到count[]中（int count[ 3 ] )

   ( sp[ 1 ][ threadIdx.x ]<~0 )?( sp[ 0 ][ threadIdx.x ]=sp[ 1 ][ threadIdx.x ] );
   ( sp[ 2 ][ threadIdx.x ]<~0 )?( sp[ 0 ][ threadIdx.x ]=sp[ 2 ][ threadIdx.x ] );
   
   if( active )
   {
   d[ tidx ]=sp[ 0 ][ threadIdx.x ];
   if( threadIdx.x<3 ){
   nLocal[ blockIdx.x+threadIdx.x*n ]=count[ threadIdx.x ];
   }
   }
}
   
step2:
对上一步统计的count用scan操作得到各种情况相对于数组的全局偏移，然后写入即可（类似radixSort过程）

也许还有更好的方法，LZ自己想想吧

4

thrust::sort，就可以满足要求。

5

如果只有数量很少的几种分布，只要做一个直方图统计出各种数量
然后直接对结果矩阵赋值从多少到多少是几就行了