欢迎您访问程序员文章站本站旨在为大家提供分享程序员计算机编程知识!
您现在的位置是: 首页

CUDA 数据分为大小两组

程序员文章站 2022-06-03 17:58:10
...



//
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include <stdio.h>


typedef struct __align__(128) qsortAtomicData_t
{
	volatile unsigned int lt_offset;    // Current output offset for < pivot
	volatile unsigned int gt_offset;    // Current output offset for > pivot
	volatile unsigned int sorted_count; // Total count sorted, for deciding when to launch next wave
	volatile unsigned int index;        // Ringbuf tracking index. Can be ignored if not using ringbuf.
} qsortAtomicData;


////////////////////////////////////////////////////////////////////////////////////////

__global__ void vote_ballot(int *a, int *b, int n, qsortAtomicData *atomicData)
{
	unsigned int thread_id = threadIdx.x;
	unsigned int lane_id = threadIdx.x & (warpSize-1);
    if (thread_id > n)
    {
       return;
    }

    unsigned pivot = 1000;
    unsigned data = a[thread_id];

    unsigned int greater = ( data > pivot );
    __syncthreads();
    unsigned int gt_mask = __ballot(greater);
    unsigned int lt_mask = __ballot(!greater);
    __syncthreads();

    if (gt_mask == 0)
    {
        greater = (data >= pivot);
        gt_mask = __ballot(greater);    // Must re-ballot for adjusted comparator
    }

    //找到比pivot小的数字的位置

    unsigned int gt_count = __popc(gt_mask);//返回个数,比pivot大的数的个数
    unsigned int lt_count = __popc(lt_mask);

    //__ballot(int predicate):指的是当前线程所在的Wrap中第N个线程对应的predicate值不为0,则将整数0的第N位进行置位1。
    // 第32位是符号位值

    // Atomically adjust the lt_ and gt_offsets by this amount. Only one thread need do this. Share the result using shfl
    unsigned int lt_offset, gt_offset;

    //只有当threadIdx.x是warpSize的整数倍时候,lane_id=0
    if (lane_id == 0)
    {


        if (lt_count > 0)
        {
        	//获取data < pivot 的元素的位置
        	//先赋值后加
            lt_offset = atomicAdd((unsigned int *) &atomicData->lt_offset, lt_count);
            printf("lt_offset=%d   thread_id=%d   \n", lt_offset, thread_id);
        }

        if (gt_count > 0)
        {
            gt_offset = n - (atomicAdd((unsigned int *) &atomicData->gt_offset, gt_count) + gt_count);
            printf("gt_offset=%d   thread_id=%d  \n",  gt_offset, thread_id);

        }
        __syncthreads();
        printf("\n\n\n");

    }


    //printf("lt_offset=%d  gt_offset=%d  thread_id=%d  \n", lt_offset, gt_offset, thread_id );
    //lane_id1到lane_id31会获取lane_id0的数据lt_offset
    lt_offset = __shfl((int)lt_offset, 0);   // Everyone pulls the offsets from lane 0
    gt_offset = __shfl((int)gt_offset, 0);

    //printf("lt_offset=%d   thread_id=%d   \n", lt_offset, thread_id);
    //printf("gt_offset=%d   thread_id=%d  \n",  gt_offset, thread_id);


    __syncthreads();


    //获得线程在warp内的位置的掩码
    //此位置前的二进制都置1 若lane_id=3,则lane_mask_lt=(111)2=7
    unsigned lane_mask_lt;
    asm(  "mov.u32 %0, %%lanemask_lt;" : "=r"(lane_mask_lt)  );
    //printf("lane_mask_lt=%d   thread_id=%d   \n", lane_mask_lt, thread_id);

    //根据greater是否为真选择gt_mask 或者 lt_mask 赋值给my_mask
    unsigned int my_mask = greater ? gt_mask : lt_mask;
    //先与运算(都是1结果得1,其它得0),  后调用__popc()计算1的个数
    //my_offsetORI代表当前wrap中,当前thread之前有my_offset个data满足data > pivot
    unsigned int my_offset = __popc(my_mask & lane_mask_lt);

    unsigned int my_offsetORI = my_offset;
    //printf("greater=%d  my_mask=%u   my_offset=%u   thread_id=%d  \n",greater, my_mask,  my_offset, thread_id);


    // Move data.
    //所有数据分成大于和小于pivot的两部分,存储在outdata
    my_offset += greater ? gt_offset : lt_offset;



    b[my_offset] = data;




	printf("thread_id=%d,   lane_id=%d,   data=%d,   greater=%d,  lane_mask_lt=%u,   gt_mask=%u,   gt_count=%d,   gt_offset=%d,   lt_mask=%u,   lt_count=%d,   lt_offset=%d,   my_mask=%u   my_offsetORI=%d    my_offset=%d\n",
	thread_id, lane_id, data, greater, lane_mask_lt, gt_mask, gt_count, gt_offset, lt_mask, lt_count, lt_offset, my_mask, my_offsetORI, my_offset );



}



int main()
{

    int *h_a, *h_b, *d_a, *d_b;
    int n = 100, m = 10;
    int nsize = n * sizeof(int);
    h_a = (int *)malloc(nsize);
    h_b = (int *)malloc(nsize);

    printf("h_a=:  \n");
    for (int i = 0; i < n; ++i)
    {
       h_a[i] = (i+600)%9*200+(i+500)%8*20+(i+400)%7*10;
       printf("%d  ",*(h_a+i));
       if (!((i+1) % m))
       {
           printf("\n");
       }
    }
    printf("\n\n\n");
    memset(h_b, 0, nsize);
    cudaMalloc(&d_a, nsize);
    cudaMalloc(&d_b, nsize);
    cudaMemcpy(d_a, h_a, nsize, cudaMemcpyHostToDevice);
    cudaMemset(d_b, 0, nsize);

///

    ////////////////////////////////////////////////////////////////////////////////
    unsigned int stacksize = 1024*1024;//1024*1024

    // This is the stack, for atomic tracking of each sort's status
    qsortAtomicData *atomicData;
    cudaMalloc((void **)&atomicData, stacksize * sizeof(qsortAtomicData));
    cudaMemset(atomicData, 0, sizeof(qsortAtomicData));     // Only need set first entry to 0


////
    vote_ballot<< <1, n >> >(d_a, d_b, n, atomicData);
    cudaMemcpy(h_b, d_b, nsize, cudaMemcpyDeviceToHost);
    printf("vote_ballot():");
    for (int i = 0; i < n; ++i)
    {
       if (!(i % m))
       {
           printf("\n");
       }
       printf("%d    ", h_b[i]);
    }
    printf("\n");
}




////////////////////////////////////////////////
//////////////////////////////////////////////////
thread_id   lane_id

data
   
greater

gt_mask  gt_count  gt_offset
lt_mask  lt_count  lt_offset

unsigned int greater = ( data > pivot );
unsigned int gt_mask = __ballot(greater);
若lane_id=N, 则有一个32位的二进制数字gt_mask,且lane_id=N对应的greater值不为0,则将gt_mask的第N位二进制位置位为1。
32位二进制数gt_mask中第N个二进制位为1,就代表当前wrap中第N个线程的data > pivot
举例: lane_id=3,且greater=1; 则gt_mask=100=4

unsigned int gt_count = __popc(gt_mask);//返回个数,比pivot大的数的个数

gt_mask =1008602887 ==  0011  1100  0001  1110  0000  1111  0000  0111
gt_count=15 gt_mask中有15个1,也就是当前wrap中有15个数比1000大

gt_offset = n - (atomicAdd((unsigned int *) &atomicData->gt_offset, gt_count) + gt_count);


my_mask  
unsigned int my_mask = greater ? gt_mask : lt_mask;

lane_mask_lt: 
若lane_id=N, 则有一个32位的二进制数字lane_mask_lt,lane_mask_lt前N位二进制位都置位为1。
举例: lane_id=3,则lane_mask_lt=111=7
////////////////////////////////////////////////

my_offsetORI=__popc(my_mask & lane_mask_lt);

greater:    my_offset = gt_offset + my_offsetORI;
lower:      my_offset = lt_offset + my_offsetORI;

////////////////////////////////////////////////
//////////////////////////////////////////////////





输出:

h_a=:  
1290  1520  1750  180  250  480  640  870  1100  1330  
1560  1790  60  220  450  680  910  1140  1370  1600  
1600  30  260  490  720  950  1180  1340  1410  1640  
70  300  530  760  920  1150  1220  1450  1680  110  
340  500  730  960  1030  1260  1490  1720  80  310  
540  770  840  1070  1300  1460  1690  120  350  580  
650  880  1040  1270  1500  1730  160  390  460  620  
850  1080  1310  1540  1770  200  200  430  660  890  
1120  1350  1580  1740  10  240  470  700  930  1160  
1320  1550  1620  50  280  510  740  900  1130  1360  



lt_offset=0   thread_id=0   
lt_offset=17   thread_id=64   
lt_offset=35   thread_id=32   
lt_offset=53   thread_id=96   
gt_offset=86   thread_id=64  
gt_offset=84   thread_id=96  
gt_offset=69   thread_id=0  
gt_offset=55   thread_id=32  












thread_id=96,   lane_id=0,   data=740,   greater=0,  lane_mask_lt=0,   gt_mask=12,   gt_count=2,   gt_offset=84,   lt_mask=3,   lt_count=2,   lt_offset=53,   my_mask=3   my_offsetORI=0    my_offset=53
thread_id=97,   lane_id=1,   data=900,   greater=0,  lane_mask_lt=1,   gt_mask=12,   gt_count=2,   gt_offset=84,   lt_mask=3,   lt_count=2,   lt_offset=53,   my_mask=3   my_offsetORI=1    my_offset=54
thread_id=98,   lane_id=2,   data=1130,   greater=1,  lane_mask_lt=3,   gt_mask=12,   gt_count=2,   gt_offset=84,   lt_mask=3,   lt_count=2,   lt_offset=53,   my_mask=12   my_offsetORI=0    my_offset=84
thread_id=99,   lane_id=3,   data=1360,   greater=1,  lane_mask_lt=7,   gt_mask=12,   gt_count=2,   gt_offset=84,   lt_mask=3,   lt_count=2,   lt_offset=53,   my_mask=12   my_offsetORI=1    my_offset=85
thread_id=32,   lane_id=0,   data=530,   greater=0,  lane_mask_lt=0,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=0    my_offset=35
thread_id=33,   lane_id=1,   data=760,   greater=0,  lane_mask_lt=1,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=1    my_offset=36
thread_id=34,   lane_id=2,   data=920,   greater=0,  lane_mask_lt=3,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=2    my_offset=37
thread_id=35,   lane_id=3,   data=1150,   greater=1,  lane_mask_lt=7,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=3252744312   my_offsetORI=0    my_offset=55
thread_id=36,   lane_id=4,   data=1220,   greater=1,  lane_mask_lt=15,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=3252744312   my_offsetORI=1    my_offset=56
thread_id=37,   lane_id=5,   data=1450,   greater=1,  lane_mask_lt=31,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=3252744312   my_offsetORI=2    my_offset=57
thread_id=38,   lane_id=6,   data=1680,   greater=1,  lane_mask_lt=63,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=3252744312   my_offsetORI=3    my_offset=58
thread_id=39,   lane_id=7,   data=110,   greater=0,  lane_mask_lt=127,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=3    my_offset=38
thread_id=40,   lane_id=8,   data=340,   greater=0,  lane_mask_lt=255,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=4    my_offset=39
thread_id=41,   lane_id=9,   data=500,   greater=0,  lane_mask_lt=511,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=5    my_offset=40
thread_id=42,   lane_id=10,   data=730,   greater=0,  lane_mask_lt=1023,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=6    my_offset=41
thread_id=43,   lane_id=11,   data=960,   greater=0,  lane_mask_lt=2047,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=7    my_offset=42
thread_id=44,   lane_id=12,   data=1030,   greater=1,  lane_mask_lt=4095,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=3252744312   my_offsetORI=4    my_offset=59
thread_id=45,   lane_id=13,   data=1260,   greater=1,  lane_mask_lt=8191,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=3252744312   my_offsetORI=5    my_offset=60
thread_id=46,   lane_id=14,   data=1490,   greater=1,  lane_mask_lt=16383,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=3252744312   my_offsetORI=6    my_offset=61
thread_id=47,   lane_id=15,   data=1720,   greater=1,  lane_mask_lt=32767,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=3252744312   my_offsetORI=7    my_offset=62
thread_id=48,   lane_id=16,   data=80,   greater=0,  lane_mask_lt=65535,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=8    my_offset=43
thread_id=49,   lane_id=17,   data=310,   greater=0,  lane_mask_lt=131071,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=9    my_offset=44
thread_id=50,   lane_id=18,   data=540,   greater=0,  lane_mask_lt=262143,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=10    my_offset=45
thread_id=51,   lane_id=19,   data=770,   greater=0,  lane_mask_lt=524287,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=11    my_offset=46
thread_id=52,   lane_id=20,   data=840,   greater=0,  lane_mask_lt=1048575,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=12    my_offset=47
thread_id=53,   lane_id=21,   data=1070,   greater=1,  lane_mask_lt=2097151,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=3252744312   my_offsetORI=8    my_offset=63
thread_id=54,   lane_id=22,   data=1300,   greater=1,  lane_mask_lt=4194303,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=3252744312   my_offsetORI=9    my_offset=64
thread_id=55,   lane_id=23,   data=1460,   greater=1,  lane_mask_lt=8388607,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=3252744312   my_offsetORI=10    my_offset=65
thread_id=56,   lane_id=24,   data=1690,   greater=1,  lane_mask_lt=16777215,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=3252744312   my_offsetORI=11    my_offset=66
thread_id=57,   lane_id=25,   data=120,   greater=0,  lane_mask_lt=33554431,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=13    my_offset=48
thread_id=58,   lane_id=26,   data=350,   greater=0,  lane_mask_lt=67108863,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=14    my_offset=49
thread_id=59,   lane_id=27,   data=580,   greater=0,  lane_mask_lt=134217727,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=15    my_offset=50
thread_id=60,   lane_id=28,   data=650,   greater=0,  lane_mask_lt=268435455,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=16    my_offset=51
thread_id=61,   lane_id=29,   data=880,   greater=0,  lane_mask_lt=536870911,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=1042222983   my_offsetORI=17    my_offset=52
thread_id=62,   lane_id=30,   data=1040,   greater=1,  lane_mask_lt=1073741823,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=3252744312   my_offsetORI=12    my_offset=67
thread_id=63,   lane_id=31,   data=1270,   greater=1,  lane_mask_lt=2147483647,   gt_mask=3252744312,   gt_count=14,   gt_offset=55,   lt_mask=1042222983,   lt_count=18,   lt_offset=35,   my_mask=3252744312   my_offsetORI=13    my_offset=68
thread_id=0,   lane_id=0,   data=1290,   greater=1,  lane_mask_lt=0,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=1008602887   my_offsetORI=0    my_offset=69
thread_id=1,   lane_id=1,   data=1520,   greater=1,  lane_mask_lt=1,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=1008602887   my_offsetORI=1    my_offset=70
thread_id=2,   lane_id=2,   data=1750,   greater=1,  lane_mask_lt=3,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=1008602887   my_offsetORI=2    my_offset=71
thread_id=3,   lane_id=3,   data=180,   greater=0,  lane_mask_lt=7,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=0    my_offset=0
thread_id=4,   lane_id=4,   data=250,   greater=0,  lane_mask_lt=15,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=1    my_offset=1
thread_id=5,   lane_id=5,   data=480,   greater=0,  lane_mask_lt=31,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=2    my_offset=2
thread_id=6,   lane_id=6,   data=640,   greater=0,  lane_mask_lt=63,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=3    my_offset=3
thread_id=7,   lane_id=7,   data=870,   greater=0,  lane_mask_lt=127,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=4    my_offset=4
thread_id=8,   lane_id=8,   data=1100,   greater=1,  lane_mask_lt=255,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=1008602887   my_offsetORI=3    my_offset=72
thread_id=9,   lane_id=9,   data=1330,   greater=1,  lane_mask_lt=511,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=1008602887   my_offsetORI=4    my_offset=73
thread_id=10,   lane_id=10,   data=1560,   greater=1,  lane_mask_lt=1023,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=1008602887   my_offsetORI=5    my_offset=74
thread_id=11,   lane_id=11,   data=1790,   greater=1,  lane_mask_lt=2047,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=1008602887   my_offsetORI=6    my_offset=75
thread_id=12,   lane_id=12,   data=60,   greater=0,  lane_mask_lt=4095,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=5    my_offset=5
thread_id=13,   lane_id=13,   data=220,   greater=0,  lane_mask_lt=8191,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=6    my_offset=6
thread_id=14,   lane_id=14,   data=450,   greater=0,  lane_mask_lt=16383,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=7    my_offset=7
thread_id=15,   lane_id=15,   data=680,   greater=0,  lane_mask_lt=32767,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=8    my_offset=8
thread_id=16,   lane_id=16,   data=910,   greater=0,  lane_mask_lt=65535,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=9    my_offset=9
thread_id=17,   lane_id=17,   data=1140,   greater=1,  lane_mask_lt=131071,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=1008602887   my_offsetORI=7    my_offset=76
thread_id=18,   lane_id=18,   data=1370,   greater=1,  lane_mask_lt=262143,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=1008602887   my_offsetORI=8    my_offset=77
thread_id=19,   lane_id=19,   data=1600,   greater=1,  lane_mask_lt=524287,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=1008602887   my_offsetORI=9    my_offset=78
thread_id=20,   lane_id=20,   data=1600,   greater=1,  lane_mask_lt=1048575,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=1008602887   my_offsetORI=10    my_offset=79
thread_id=21,   lane_id=21,   data=30,   greater=0,  lane_mask_lt=2097151,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=10    my_offset=10
thread_id=22,   lane_id=22,   data=260,   greater=0,  lane_mask_lt=4194303,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=11    my_offset=11
thread_id=23,   lane_id=23,   data=490,   greater=0,  lane_mask_lt=8388607,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=12    my_offset=12
thread_id=24,   lane_id=24,   data=720,   greater=0,  lane_mask_lt=16777215,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=13    my_offset=13
thread_id=25,   lane_id=25,   data=950,   greater=0,  lane_mask_lt=33554431,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=14    my_offset=14
thread_id=26,   lane_id=26,   data=1180,   greater=1,  lane_mask_lt=67108863,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=1008602887   my_offsetORI=11    my_offset=80
thread_id=27,   lane_id=27,   data=1340,   greater=1,  lane_mask_lt=134217727,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=1008602887   my_offsetORI=12    my_offset=81
thread_id=28,   lane_id=28,   data=1410,   greater=1,  lane_mask_lt=268435455,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=1008602887   my_offsetORI=13    my_offset=82
thread_id=29,   lane_id=29,   data=1640,   greater=1,  lane_mask_lt=536870911,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=1008602887   my_offsetORI=14    my_offset=83
thread_id=30,   lane_id=30,   data=70,   greater=0,  lane_mask_lt=1073741823,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=15    my_offset=15
thread_id=31,   lane_id=31,   data=300,   greater=0,  lane_mask_lt=2147483647,   gt_mask=1008602887,   gt_count=15,   gt_offset=69,   lt_mask=3286364408,   lt_count=17,   lt_offset=0,   my_mask=3286364408   my_offsetORI=16    my_offset=16
thread_id=64,   lane_id=0,   data=1500,   greater=1,  lane_mask_lt=0,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=504301443   my_offsetORI=0    my_offset=86
thread_id=65,   lane_id=1,   data=1730,   greater=1,  lane_mask_lt=1,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=504301443   my_offsetORI=1    my_offset=87
thread_id=66,   lane_id=2,   data=160,   greater=0,  lane_mask_lt=3,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=0    my_offset=17
thread_id=67,   lane_id=3,   data=390,   greater=0,  lane_mask_lt=7,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=1    my_offset=18
thread_id=68,   lane_id=4,   data=460,   greater=0,  lane_mask_lt=15,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=2    my_offset=19
thread_id=69,   lane_id=5,   data=620,   greater=0,  lane_mask_lt=31,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=3    my_offset=20
thread_id=70,   lane_id=6,   data=850,   greater=0,  lane_mask_lt=63,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=4    my_offset=21
thread_id=71,   lane_id=7,   data=1080,   greater=1,  lane_mask_lt=127,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=504301443   my_offsetORI=2    my_offset=88
thread_id=72,   lane_id=8,   data=1310,   greater=1,  lane_mask_lt=255,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=504301443   my_offsetORI=3    my_offset=89
thread_id=73,   lane_id=9,   data=1540,   greater=1,  lane_mask_lt=511,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=504301443   my_offsetORI=4    my_offset=90
thread_id=74,   lane_id=10,   data=1770,   greater=1,  lane_mask_lt=1023,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=504301443   my_offsetORI=5    my_offset=91
thread_id=75,   lane_id=11,   data=200,   greater=0,  lane_mask_lt=2047,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=5    my_offset=22
thread_id=76,   lane_id=12,   data=200,   greater=0,  lane_mask_lt=4095,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=6    my_offset=23
thread_id=77,   lane_id=13,   data=430,   greater=0,  lane_mask_lt=8191,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=7    my_offset=24
thread_id=78,   lane_id=14,   data=660,   greater=0,  lane_mask_lt=16383,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=8    my_offset=25
thread_id=79,   lane_id=15,   data=890,   greater=0,  lane_mask_lt=32767,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=9    my_offset=26
thread_id=80,   lane_id=16,   data=1120,   greater=1,  lane_mask_lt=65535,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=504301443   my_offsetORI=6    my_offset=92
thread_id=81,   lane_id=17,   data=1350,   greater=1,  lane_mask_lt=131071,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=504301443   my_offsetORI=7    my_offset=93
thread_id=82,   lane_id=18,   data=1580,   greater=1,  lane_mask_lt=262143,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=504301443   my_offsetORI=8    my_offset=94
thread_id=83,   lane_id=19,   data=1740,   greater=1,  lane_mask_lt=524287,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=504301443   my_offsetORI=9    my_offset=95
thread_id=84,   lane_id=20,   data=10,   greater=0,  lane_mask_lt=1048575,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=10    my_offset=27
thread_id=85,   lane_id=21,   data=240,   greater=0,  lane_mask_lt=2097151,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=11    my_offset=28
thread_id=86,   lane_id=22,   data=470,   greater=0,  lane_mask_lt=4194303,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=12    my_offset=29
thread_id=87,   lane_id=23,   data=700,   greater=0,  lane_mask_lt=8388607,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=13    my_offset=30
thread_id=88,   lane_id=24,   data=930,   greater=0,  lane_mask_lt=16777215,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=14    my_offset=31
thread_id=89,   lane_id=25,   data=1160,   greater=1,  lane_mask_lt=33554431,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=504301443   my_offsetORI=10    my_offset=96
thread_id=90,   lane_id=26,   data=1320,   greater=1,  lane_mask_lt=67108863,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=504301443   my_offsetORI=11    my_offset=97
thread_id=91,   lane_id=27,   data=1550,   greater=1,  lane_mask_lt=134217727,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=504301443   my_offsetORI=12    my_offset=98
thread_id=92,   lane_id=28,   data=1620,   greater=1,  lane_mask_lt=268435455,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=504301443   my_offsetORI=13    my_offset=99
thread_id=93,   lane_id=29,   data=50,   greater=0,  lane_mask_lt=536870911,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=15    my_offset=32
thread_id=94,   lane_id=30,   data=280,   greater=0,  lane_mask_lt=1073741823,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=16    my_offset=33
thread_id=95,   lane_id=31,   data=510,   greater=0,  lane_mask_lt=2147483647,   gt_mask=504301443,   gt_count=14,   gt_offset=86,   lt_mask=3790665852,   lt_count=18,   lt_offset=17,   my_mask=3790665852   my_offsetORI=17    my_offset=34
vote_ballot():
180    250    480    640    870    60    220    450    680    910    
30    260    490    720    950    70    300    160    390    460    
620    850    200    200    430    660    890    10    240    470    
700    930    50    280    510    530    760    920    110    340    
500    730    960    80    310    540    770    840    120    350    
580    650    880    740    900    1150    1220    1450    1680    1030    
1260    1490    1720    1070    1300    1460    1690    1040    1270    1290    
1520    1750    1100    1330    1560    1790    1140    1370    1600    1600    
1180    1340    1410    1640    1130    1360    1500    1730    1080    1310    
1540    1770    1120    1350    1580    1740    1160    1320    1550    1620