6月28日更新--“如何并行化我的应用”方案有奖征集活动--获奖名单公布

图腾部落

原帖由 cyrosly 于 2010-6-8 06:33 发表
基于点集的Balance kd-tree的并行构建：

1首先将点集的坐标数据转换成无符号整型
2将数据按照X轴方向进行排序
3将排序号的数据进行分割，如果slice中的元素个数是偶数则分成的2个slice正好相等，
如果是奇数个数则中间的元素同时即属于左子slice又属于右子slice
4对每一迭代层中的所有slice中的元素进行kdnode编码
5然后在下一层迭代中对上面分成的2个slice分别独立的按照Y轴向排序，饼尽享如上不相同的操作步骤
6这是一个递归（当然程序实现时时没有递归的，以为每个迭代层中的所有slice可以独立的并行计算
7用户可以设置一个阀值，用来当每个slice中（kdnode)的元素数量小于该值时停止树的划分。
8及时对于相当大的数据规模，实际的迭代次数没有那么多，因为当每个slice的尺寸小于或等于最大可用
的block尺寸时（为什么是最大尺寸的一半，因为要考虑slice尺寸为奇数划分时的长度扩展）对应于剩下
的迭代中的划分，编码计算可以在内核中一次完成，这样最后的几步迭代实际上只调用一次内核就可以了

    另外还需要注意的是，将浮点数转化成无符号整型后，还需要翻转符号位，因为负数的符号位是1
而正数的符号位是0，翻转后才能确保负数总是排在正数之前（假设按照从小到大的顺序）。
   还要记得，排序时XYZ坐标是关键字（每层的关键字coord[ level%3 ] )，而点ID是被排序的目标数组。

   还有一点需要注意，在准备数据前，要先根据阀值确定树最终将会构建都少层，然后再根据层数计算
出扩展空间（如果当前slice中的元素是奇数个，那么当将其划分时程2个单独的slice后，这2个slice的长度
将大于原来的slice尺寸，因此实现确定最大存储空间而不是计算中重新分配将更有效率）

伪代码可以表示如下：
   
    source_slice[];
      target_slice[];
      AXIS=0;
      PROC Create_Point_Kdtree( target slice, source_slice, AXIS )
      {
        while( slice_size>5;
    }
    inline __device__ uint inclusive( uint* smem ) const
    {
        uint pref=warpOp( smem ); __syncthreads();

        if( lane==31 ){
            smem[ warpid ]=pref;
        } __syncthreads();

        for( uint n=1; n>5 ); n0 ){
        pref+=smem[ warpid-1 ];
         } __syncthreads();

         return pref;
    }

    inline __device__ void exclusive( uint* smem )
    {
        uint pref=warpOp( smem ); __syncthreads();

        if( lane==31 ){
            smem[ warpid ]=pref;
        } __syncthreads();

        for( uint n=1; n>5 ); n0 ){
            pref+=smem[ warpid-1 ];
        } __syncthreads();

        ( threadIdx.x= 1 ){ smem[ threadIdx.x ]+=smem[ threadIdx.x- 1 ]; }
        if( lane>= 2 ){ smem[ threadIdx.x ]+=smem[ threadIdx.x- 2 ]; }
        if( lane>= 4 ){ smem[ threadIdx.x ]+=smem[ threadIdx.x- 4 ]; }
        if( lane>= 8 ){ smem[ threadIdx.x ]+=smem[ threadIdx.x- 8 ]; }
        if( lane>=16 ){ smem[ threadIdx.x ]+=smem[ threadIdx.x-16 ]; }
        return smem[ threadIdx.x ];
    }

    uint lane;
    uint warpid;
};

extern "C"
{

//从这个内核到kernel_scatter是排序用的内核，采用基数排序
__global__
void kernel_scan_blocks( uint* prefix, uint* block_suffix, const uint* list, uint slice_size )
{
    __shared__ uint smem[ MAX_CTA_DIM ];

                const uint tidx=IM( MAX_CTA_DIM, blockIdx.x )+threadIdx.x;
                const uint active=tidx

稍后代码功能会调整

yyfn风辰

图腾部落

原帖由 yyfn风辰 于 2010-6-10 11:35 发表

我说的是论坛插代码的功能。

cyrosly

看来洋洋洒洒的一堆代码和文字算是白贴了，竟然被无视

图腾部落

原帖由 cyrosly 于 2010-6-17 11:45 发表
看来洋洋洒洒的一堆代码和文字算是白贴了，竟然被无视

呵，修宇，已经把你的名字已经添加进名单了，我的失误，很抱歉。

[ 本帖最后由 图腾部落 于 2010-6-17 13:46 编辑 ]

cyrosly

话说最近几个月钱赚的很尴尬，一张显卡卖了2K,不到半个月清零，接了个小项目2.5k奢侈了几天，没了。又接了个小项目,报酬是N个果冻，结果牌子还错了，都分给同事了，等于白做。再次接到项目，1K的定金，结果朋友这时来广州出差，要待大约半个月，不到三天又没了，所以期待这500奖金再帮我撑几天...

图腾部落

原帖由 cyrosly 于 2010-7-5 22:54 发表
话说最近几个月钱赚的很尴尬，一张显卡卖了2K,不到半个月清零，接了个小项目2.5k奢侈了几天，没了。又接了个小项目,报酬是N个果冻，结果牌子还错了，都分给同事了，等于白做。再次接到项目，1K的定金，结果朋友这时来广州出差，要待大约半个月，不到三天又没了，所以期待这500奖金再帮我撑几天...

够奢侈的你，1个月不到，用了6K