//
#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
扫一扫
2万+
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
