CUDA并行排序(1)——双调排序(unsigned int型数据)

最新推荐文章于 2024-03-04 23:16:55 发布
最新推荐文章于 2024-03-04 23:16:55 发布
阅读量1.2k 收藏 1
点赞数 1
分类专栏： C++ 基础知识 CUDA零基础
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/zlf19910726/article/details/78754354
版权 



#include <stdio.h>
#include<iostream>
using namespace std;
#define CHECK(res) if(res!=cudaSuccess){exit(-1);}
__global__ void helloCUDA(unsigned *indata, unsigned int len)
{

	unsigned int offset = 0;
    __shared__ unsigned sortbuf[8];     // Max of 1024 elements - TODO: make this dynamic


    // First copy data into shared memory.
    unsigned int inside = (threadIdx.x < len);
    sortbuf[threadIdx.x] = inside ? indata[threadIdx.x + offset] : 0xffffffffu;
    __syncthreads();


    // Now the sort loops
    // Here, "k" is the sort level (remember bitonic does a multi-level butterfly style sort)
    // and "j" is the partner element in the butterfly.
    // Two threads each work on one butterfly, because the read/write needs to happen
    // simultaneously
    for (unsigned int k=2; k<=16; k*=2) // Butterfly stride increments in powers of 2
    {
        printf("AAAAAAAAAAAAAAAAAAAAAA\n");

        for (unsigned int j=k>>1; j>0; j>>=1) // Strides also in powers of to, up to <k
        {
            unsigned int swap_idx = threadIdx.x ^ j; // Index of element we're compare-and-swapping with
            unsigned my_elem = sortbuf[threadIdx.x];
            unsigned swap_elem = sortbuf[swap_idx];

            __syncthreads();

            // The k'th bit of my threadid (and hence my sort item ID)
            // determines if we sort ascending or descending.
            // However, since threads are reading from the top AND the bottom of
            // the butterfly, if my ID is > swap_idx, then ascending means mine<swap.
            // Finally, if either my_elem or swap_elem is out of range, then it
            // ALWAYS acts like it's the largest number.
            // Confusing? It saves us two writes though.
            unsigned int ascend = k * (swap_idx < threadIdx.x);
            unsigned int descend = k * (swap_idx > threadIdx.x);
            bool swap = false;



            if ((threadIdx.x & k) == ascend)
            {
                if (my_elem > swap_elem)
                    swap = true;
            }

            if ((threadIdx.x & k) == descend)
            {
                if (my_elem < swap_elem)
                    swap = true;
            }

            // If we had to swap, then write my data to the other element's position.
            // Don't forget to track out-of-range status too!
            if (swap)
            {
                sortbuf[swap_idx] = my_elem;
            }

            printf("BBBBBBBBBBBBBBBBBBBBBBBBB\n");
        	printf("sortbuf addr= %p  sortbuf= %d  my_elem addr= %p  threadIdx.x= %d \n",
        			&sortbuf, sortbuf[threadIdx.x], &my_elem, threadIdx.x);

            if(k>=0&&k<=16 && j>=0&&j<=8 && threadIdx.x>=0&&threadIdx.x<=7)
            {
            	printf("k= %d  j= %d  threadIdx.x= %d  swap_idx= %d  my_elem= %d  swap_elem= %d  ascend= %d  descend= %d  swap= %d\n",
            			k, j, threadIdx.x, swap_idx, my_elem, swap_elem, ascend, descend, swap );

            }


            __syncthreads();
        }//for()
    }//for()

}

///
int main()
{
	printf("Hello main()\n");

	cudaError_t res;
	///
	unsigned int len = 16;
    unsigned *data_test = new unsigned[len];
    unsigned data_test1[] =  {6,7,16,3, 14,1,2,15, 10,9,12,13, 4,5,8,11};
    //data_test1[0]=10;
    data_test=data_test1;
    cout<<"data_test[i]: ";

    for(int i=0; i<len; i++)
    {
    	if(1)
    	{
        	cout<<data_test[i]<<"  ";
    	}
    }
    cout<<endl;
	///

    unsigned *indata =NULL;

    res = cudaMalloc((void**)(&indata), len*sizeof(unsigned));CHECK(res)

    res = cudaMemcpy((void*)(indata), (void*)(data_test), len*sizeof(unsigned), cudaMemcpyHostToDevice);CHECK(res)

    helloCUDA<<<1, len>>>(indata, len);
    cudaDeviceSynchronize();



    printf("Goodbye main()\n");
    return 0;
}





Hello main()
data_test[i]: 6  7  16  3  14  1  2  15  10  9  12  13  4  5  8  11  
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
sortbuf addr= 0x7fdc7f000000  sortbuf= 6  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 0 
sortbuf addr= 0x7fdc7f000000  sortbuf= 7  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 1 
sortbuf addr= 0x7fdc7f000000  sortbuf= 16  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 2 
sortbuf addr= 0x7fdc7f000000  sortbuf= 3  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 3 
sortbuf addr= 0x7fdc7f000000  sortbuf= 1  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 4 
sortbuf addr= 0x7fdc7f000000  sortbuf= 14  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 5 
sortbuf addr= 0x7fdc7f000000  sortbuf= 15  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 6 
sortbuf addr= 0x7fdc7f000000  sortbuf= 2  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 7 
sortbuf addr= 0x7fdc7f000000  sortbuf= 9  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 8 
sortbuf addr= 0x7fdc7f000000  sortbuf= 10  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 9 
sortbuf addr= 0x7fdc7f000000  sortbuf= 13  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 10 
sortbuf addr= 0x7fdc7f000000  sortbuf= 12  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 11 
sortbuf addr= 0x7fdc7f000000  sortbuf= 4  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 12 
sortbuf addr= 0x7fdc7f000000  sortbuf= 5  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 13 
sortbuf addr= 0x7fdc7f000000  sortbuf= 11  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 14 
sortbuf addr= 0x7fdc7f000000  sortbuf= 8  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 15 
k= 2  j= 1  threadIdx.x= 0  swap_idx= 1  my_elem= 6  swap_elem= 7  ascend= 0  descend= 2  swap= 0
k= 2  j= 1  threadIdx.x= 1  swap_idx= 0  my_elem= 7  swap_elem= 6  ascend= 2  descend= 0  swap= 0
k= 2  j= 1  threadIdx.x= 2  swap_idx= 3  my_elem= 16  swap_elem= 3  ascend= 0  descend= 2  swap= 0
k= 2  j= 1  threadIdx.x= 3  swap_idx= 2  my_elem= 3  swap_elem= 16  ascend= 2  descend= 0  swap= 0
k= 2  j= 1  threadIdx.x= 4  swap_idx= 5  my_elem= 14  swap_elem= 1  ascend= 0  descend= 2  swap= 1
k= 2  j= 1  threadIdx.x= 5  swap_idx= 4  my_elem= 1  swap_elem= 14  ascend= 2  descend= 0  swap= 1
k= 2  j= 1  threadIdx.x= 6  swap_idx= 7  my_elem= 2  swap_elem= 15  ascend= 0  descend= 2  swap= 1
k= 2  j= 1  threadIdx.x= 7  swap_idx= 6  my_elem= 15  swap_elem= 2  ascend= 2  descend= 0  swap= 1
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
sortbuf addr= 0x7fdc7f000000  sortbuf= 6  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 0 
sortbuf addr= 0x7fdc7f000000  sortbuf= 3  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 1 
sortbuf addr= 0x7fdc7f000000  sortbuf= 16  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 2 
sortbuf addr= 0x7fdc7f000000  sortbuf= 7  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 3 
sortbuf addr= 0x7fdc7f000000  sortbuf= 15  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 4 
sortbuf addr= 0x7fdc7f000000  sortbuf= 14  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 5 
sortbuf addr= 0x7fdc7f000000  sortbuf= 1  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 6 
sortbuf addr= 0x7fdc7f000000  sortbuf= 2  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 7 
sortbuf addr= 0x7fdc7f000000  sortbuf= 9  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 8 
sortbuf addr= 0x7fdc7f000000  sortbuf= 10  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 9 
sortbuf addr= 0x7fdc7f000000  sortbuf= 13  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 10 
sortbuf addr= 0x7fdc7f000000  sortbuf= 12  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 11 
sortbuf addr= 0x7fdc7f000000  sortbuf= 11  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 12 
sortbuf addr= 0x7fdc7f000000  sortbuf= 8  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 13 
sortbuf addr= 0x7fdc7f000000  sortbuf= 4  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 14 
sortbuf addr= 0x7fdc7f000000  sortbuf= 5  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 15 
k= 4  j= 2  threadIdx.x= 0  swap_idx= 2  my_elem= 6  swap_elem= 16  ascend= 0  descend= 4  swap= 0
k= 4  j= 2  threadIdx.x= 1  swap_idx= 3  my_elem= 7  swap_elem= 3  ascend= 0  descend= 4  swap= 1
k= 4  j= 2  threadIdx.x= 2  swap_idx= 0  my_elem= 16  swap_elem= 6  ascend= 4  descend= 0  swap= 0
k= 4  j= 2  threadIdx.x= 3  swap_idx= 1  my_elem= 3  swap_elem= 7  ascend= 4  descend= 0  swap= 1
k= 4  j= 2  threadIdx.x= 4  swap_idx= 6  my_elem= 1  swap_elem= 15  ascend= 0  descend= 4  swap= 1
k= 4  j= 2  threadIdx.x= 5  swap_idx= 7  my_elem= 14  swap_elem= 2  ascend= 0  descend= 4  swap= 0
k= 4  j= 2  threadIdx.x= 6  swap_idx= 4  my_elem= 15  swap_elem= 1  ascend= 4  descend= 0  swap= 1
k= 4  j= 2  threadIdx.x= 7  swap_idx= 5  my_elem= 2  swap_elem= 14  ascend= 4  descend= 0  swap= 0
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
sortbuf addr= 0x7fdc7f000000  sortbuf= 3  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 0 
sortbuf addr= 0x7fdc7f000000  sortbuf= 6  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 1 
sortbuf addr= 0x7fdc7f000000  sortbuf= 7  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 2 
sortbuf addr= 0x7fdc7f000000  sortbuf= 16  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 3 
sortbuf addr= 0x7fdc7f000000  sortbuf= 15  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 4 
sortbuf addr= 0x7fdc7f000000  sortbuf= 14  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 5 
sortbuf addr= 0x7fdc7f000000  sortbuf= 2  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 6 
sortbuf addr= 0x7fdc7f000000  sortbuf= 1  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 7 
sortbuf addr= 0x7fdc7f000000  sortbuf= 9  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 8 
sortbuf addr= 0x7fdc7f000000  sortbuf= 10  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 9 
sortbuf addr= 0x7fdc7f000000  sortbuf= 12  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 10 
sortbuf addr= 0x7fdc7f000000  sortbuf= 13  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 11 
sortbuf addr= 0x7fdc7f000000  sortbuf= 11  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 12 
sortbuf addr= 0x7fdc7f000000  sortbuf= 8  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 13 
sortbuf addr= 0x7fdc7f000000  sortbuf= 5  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 14 
sortbuf addr= 0x7fdc7f000000  sortbuf= 4  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 15 
k= 4  j= 1  threadIdx.x= 0  swap_idx= 1  my_elem= 6  swap_elem= 3  ascend= 0  descend= 4  swap= 1
k= 4  j= 1  threadIdx.x= 1  swap_idx= 0  my_elem= 3  swap_elem= 6  ascend= 4  descend= 0  swap= 1
k= 4  j= 1  threadIdx.x= 2  swap_idx= 3  my_elem= 16  swap_elem= 7  ascend= 0  descend= 4  swap= 1
k= 4  j= 1  threadIdx.x= 3  swap_idx= 2  my_elem= 7  swap_elem= 16  ascend= 4  descend= 0  swap= 1
k= 4  j= 1  threadIdx.x= 4  swap_idx= 5  my_elem= 15  swap_elem= 14  ascend= 0  descend= 4  swap= 0
k= 4  j= 1  threadIdx.x= 5  swap_idx= 4  my_elem= 14  swap_elem= 15  ascend= 4  descend= 0  swap= 0
k= 4  j= 1  threadIdx.x= 6  swap_idx= 7  my_elem= 1  swap_elem= 2  ascend= 0  descend= 4  swap= 1
k= 4  j= 1  threadIdx.x= 7  swap_idx= 6  my_elem= 2  swap_elem= 1  ascend= 4  descend= 0  swap= 1
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
sortbuf addr= 0x7fdc7f000000  sortbuf= 3  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 0 
sortbuf addr= 0x7fdc7f000000  sortbuf= 6  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 1 
sortbuf addr= 0x7fdc7f000000  sortbuf= 2  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 2 
sortbuf addr= 0x7fdc7f000000  sortbuf= 1  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 3 
sortbuf addr= 0x7fdc7f000000  sortbuf= 15  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 4 
sortbuf addr= 0x7fdc7f000000  sortbuf= 14  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 5 
sortbuf addr= 0x7fdc7f000000  sortbuf= 7  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 6 
sortbuf addr= 0x7fdc7f000000  sortbuf= 16  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 7 
sortbuf addr= 0x7fdc7f000000  sortbuf= 11  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 8 
sortbuf addr= 0x7fdc7f000000  sortbuf= 10  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 9 
sortbuf addr= 0x7fdc7f000000  sortbuf= 12  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 10 
sortbuf addr= 0x7fdc7f000000  sortbuf= 13  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 11 
sortbuf addr= 0x7fdc7f000000  sortbuf= 9  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 12 
sortbuf addr= 0x7fdc7f000000  sortbuf= 8  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 13 
sortbuf addr= 0x7fdc7f000000  sortbuf= 5  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 14 
sortbuf addr= 0x7fdc7f000000  sortbuf= 4  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 15 
k= 8  j= 4  threadIdx.x= 0  swap_idx= 4  my_elem= 3  swap_elem= 15  ascend= 0  descend= 8  swap= 0
k= 8  j= 4  threadIdx.x= 1  swap_idx= 5  my_elem= 6  swap_elem= 14  ascend= 0  descend= 8  swap= 0
k= 8  j= 4  threadIdx.x= 2  swap_idx= 6  my_elem= 7  swap_elem= 2  ascend= 0  descend= 8  swap= 1
k= 8  j= 4  threadIdx.x= 3  swap_idx= 7  my_elem= 16  swap_elem= 1  ascend= 0  descend= 8  swap= 1
k= 8  j= 4  threadIdx.x= 4  swap_idx= 0  my_elem= 15  swap_elem= 3  ascend= 8  descend= 0  swap= 0
k= 8  j= 4  threadIdx.x= 5  swap_idx= 1  my_elem= 14  swap_elem= 6  ascend= 8  descend= 0  swap= 0
k= 8  j= 4  threadIdx.x= 6  swap_idx= 2  my_elem= 2  swap_elem= 7  ascend= 8  descend= 0  swap= 1
k= 8  j= 4  threadIdx.x= 7  swap_idx= 3  my_elem= 1  swap_elem= 16  ascend= 8  descend= 0  swap= 1
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
sortbuf addr= 0x7fdc7f000000  sortbuf= 2  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 0 
sortbuf addr= 0x7fdc7f000000  sortbuf= 1  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 1 
sortbuf addr= 0x7fdc7f000000  sortbuf= 3  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 2 
sortbuf addr= 0x7fdc7f000000  sortbuf= 6  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 3 
sortbuf addr= 0x7fdc7f000000  sortbuf= 7  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 4 
sortbuf addr= 0x7fdc7f000000  sortbuf= 14  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 5 
sortbuf addr= 0x7fdc7f000000  sortbuf= 15  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 6 
sortbuf addr= 0x7fdc7f000000  sortbuf= 16  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 7 
sortbuf addr= 0x7fdc7f000000  sortbuf= 12  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 8 
sortbuf addr= 0x7fdc7f000000  sortbuf= 13  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 9 
sortbuf addr= 0x7fdc7f000000  sortbuf= 11  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 10 
sortbuf addr= 0x7fdc7f000000  sortbuf= 10  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 11 
sortbuf addr= 0x7fdc7f000000  sortbuf= 9  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 12 
sortbuf addr= 0x7fdc7f000000  sortbuf= 8  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 13 
sortbuf addr= 0x7fdc7f000000  sortbuf= 5  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 14 
sortbuf addr= 0x7fdc7f000000  sortbuf= 4  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 15 
k= 8  j= 2  threadIdx.x= 0  swap_idx= 2  my_elem= 3  swap_elem= 2  ascend= 0  descend= 8  swap= 1
k= 8  j= 2  threadIdx.x= 1  swap_idx= 3  my_elem= 6  swap_elem= 1  ascend= 0  descend= 8  swap= 1
k= 8  j= 2  threadIdx.x= 2  swap_idx= 0  my_elem= 2  swap_elem= 3  ascend= 8  descend= 0  swap= 1
k= 8  j= 2  threadIdx.x= 3  swap_idx= 1  my_elem= 1  swap_elem= 6  ascend= 8  descend= 0  swap= 1
k= 8  j= 2  threadIdx.x= 4  swap_idx= 6  my_elem= 15  swap_elem= 7  ascend= 0  descend= 8  swap= 1
k= 8  j= 2  threadIdx.x= 5  swap_idx= 7  my_elem= 14  swap_elem= 16  ascend= 0  descend= 8  swap= 0
k= 8  j= 2  threadIdx.x= 6  swap_idx= 4  my_elem= 7  swap_elem= 15  ascend= 8  descend= 0  swap= 1
k= 8  j= 2  threadIdx.x= 7  swap_idx= 5  my_elem= 16  swap_elem= 14  ascend= 8  descend= 0  swap= 0
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
sortbuf addr= 0x7fdc7f000000  sortbuf= 1  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 0 
sortbuf addr= 0x7fdc7f000000  sortbuf= 2  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 1 
sortbuf addr= 0x7fdc7f000000  sortbuf= 3  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 2 
sortbuf addr= 0x7fdc7f000000  sortbuf= 6  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 3 
sortbuf addr= 0x7fdc7f000000  sortbuf= 7  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 4 
sortbuf addr= 0x7fdc7f000000  sortbuf= 14  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 5 
sortbuf addr= 0x7fdc7f000000  sortbuf= 15  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 6 
sortbuf addr= 0x7fdc7f000000  sortbuf= 16  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 7 
sortbuf addr= 0x7fdc7f000000  sortbuf= 13  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 8 
sortbuf addr= 0x7fdc7f000000  sortbuf= 12  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 9 
sortbuf addr= 0x7fdc7f000000  sortbuf= 11  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 10 
sortbuf addr= 0x7fdc7f000000  sortbuf= 10  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 11 
sortbuf addr= 0x7fdc7f000000  sortbuf= 9  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 12 
sortbuf addr= 0x7fdc7f000000  sortbuf= 8  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 13 
sortbuf addr= 0x7fdc7f000000  sortbuf= 5  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 14 
sortbuf addr= 0x7fdc7f000000  sortbuf= 4  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 15 
k= 8  j= 1  threadIdx.x= 0  swap_idx= 1  my_elem= 2  swap_elem= 1  ascend= 0  descend= 8  swap= 1
k= 8  j= 1  threadIdx.x= 1  swap_idx= 0  my_elem= 1  swap_elem= 2  ascend= 8  descend= 0  swap= 1
k= 8  j= 1  threadIdx.x= 2  swap_idx= 3  my_elem= 3  swap_elem= 6  ascend= 0  descend= 8  swap= 0
k= 8  j= 1  threadIdx.x= 3  swap_idx= 2  my_elem= 6  swap_elem= 3  ascend= 8  descend= 0  swap= 0
k= 8  j= 1  threadIdx.x= 4  swap_idx= 5  my_elem= 7  swap_elem= 14  ascend= 0  descend= 8  swap= 0
k= 8  j= 1  threadIdx.x= 5  swap_idx= 4  my_elem= 14  swap_elem= 7  ascend= 8  descend= 0  swap= 0
k= 8  j= 1  threadIdx.x= 6  swap_idx= 7  my_elem= 15  swap_elem= 16  ascend= 0  descend= 8  swap= 0
k= 8  j= 1  threadIdx.x= 7  swap_idx= 6  my_elem= 16  swap_elem= 15  ascend= 8  descend= 0  swap= 0
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
AAAAAAAAAAAAAAAAAAAAAA
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
sortbuf addr= 0x7fdc7f000000  sortbuf= 1  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 0 
sortbuf addr= 0x7fdc7f000000  sortbuf= 2  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 1 
sortbuf addr= 0x7fdc7f000000  sortbuf= 3  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 2 
sortbuf addr= 0x7fdc7f000000  sortbuf= 6  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 3 
sortbuf addr= 0x7fdc7f000000  sortbuf= 7  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 4 
sortbuf addr= 0x7fdc7f000000  sortbuf= 8  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 5 
sortbuf addr= 0x7fdc7f000000  sortbuf= 5  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 6 
sortbuf addr= 0x7fdc7f000000  sortbuf= 4  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 7 
sortbuf addr= 0x7fdc7f000000  sortbuf= 13  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 8 
sortbuf addr= 0x7fdc7f000000  sortbuf= 12  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 9 
sortbuf addr= 0x7fdc7f000000  sortbuf= 11  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 10 
sortbuf addr= 0x7fdc7f000000  sortbuf= 10  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 11 
sortbuf addr= 0x7fdc7f000000  sortbuf= 9  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 12 
sortbuf addr= 0x7fdc7f000000  sortbuf= 14  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 13 
sortbuf addr= 0x7fdc7f000000  sortbuf= 15  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 14 
sortbuf addr= 0x7fdc7f000000  sortbuf= 16  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 15 
k= 16  j= 8  threadIdx.x= 0  swap_idx= 8  my_elem= 1  swap_elem= 13  ascend= 0  descend= 16  swap= 0
k= 16  j= 8  threadIdx.x= 1  swap_idx= 9  my_elem= 2  swap_elem= 12  ascend= 0  descend= 16  swap= 0
k= 16  j= 8  threadIdx.x= 2  swap_idx= 10  my_elem= 3  swap_elem= 11  ascend= 0  descend= 16  swap= 0
k= 16  j= 8  threadIdx.x= 3  swap_idx= 11  my_elem= 6  swap_elem= 10  ascend= 0  descend= 16  swap= 0
k= 16  j= 8  threadIdx.x= 4  swap_idx= 12  my_elem= 7  swap_elem= 9  ascend= 0  descend= 16  swap= 0
k= 16  j= 8  threadIdx.x= 5  swap_idx= 13  my_elem= 14  swap_elem= 8  ascend= 0  descend= 16  swap= 1
k= 16  j= 8  threadIdx.x= 6  swap_idx= 14  my_elem= 15  swap_elem= 5  ascend= 0  descend= 16  swap= 1
k= 16  j= 8  threadIdx.x= 7  swap_idx= 15  my_elem= 16  swap_elem= 4  ascend= 0  descend= 16  swap= 1
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
sortbuf addr= 0x7fdc7f000000  sortbuf= 1  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 0 
sortbuf addr= 0x7fdc7f000000  sortbuf= 2  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 1 
sortbuf addr= 0x7fdc7f000000  sortbuf= 3  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 2 
sortbuf addr= 0x7fdc7f000000  sortbuf= 4  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 3 
sortbuf addr= 0x7fdc7f000000  sortbuf= 7  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 4 
sortbuf addr= 0x7fdc7f000000  sortbuf= 8  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 5 
sortbuf addr= 0x7fdc7f000000  sortbuf= 5  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 6 
sortbuf addr= 0x7fdc7f000000  sortbuf= 6  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 7 
sortbuf addr= 0x7fdc7f000000  sortbuf= 9  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 8 
sortbuf addr= 0x7fdc7f000000  sortbuf= 12  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 9 
sortbuf addr= 0x7fdc7f000000  sortbuf= 11  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 10 
sortbuf addr= 0x7fdc7f000000  sortbuf= 10  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 11 
sortbuf addr= 0x7fdc7f000000  sortbuf= 13  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 12 
sortbuf addr= 0x7fdc7f000000  sortbuf= 14  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 13 
sortbuf addr= 0x7fdc7f000000  sortbuf= 15  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 14 
sortbuf addr= 0x7fdc7f000000  sortbuf= 16  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 15 
k= 16  j= 4  threadIdx.x= 0  swap_idx= 4  my_elem= 1  swap_elem= 7  ascend= 0  descend= 16  swap= 0
k= 16  j= 4  threadIdx.x= 1  swap_idx= 5  my_elem= 2  swap_elem= 8  ascend= 0  descend= 16  swap= 0
k= 16  j= 4  threadIdx.x= 2  swap_idx= 6  my_elem= 3  swap_elem= 5  ascend= 0  descend= 16  swap= 0
k= 16  j= 4  threadIdx.x= 3  swap_idx= 7  my_elem= 6  swap_elem= 4  ascend= 0  descend= 16  swap= 1
k= 16  j= 4  threadIdx.x= 4  swap_idx= 0  my_elem= 7  swap_elem= 1  ascend= 16  descend= 0  swap= 0
k= 16  j= 4  threadIdx.x= 5  swap_idx= 1  my_elem= 8  swap_elem= 2  ascend= 16  descend= 0  swap= 0
k= 16  j= 4  threadIdx.x= 6  swap_idx= 2  my_elem= 5  swap_elem= 3  ascend= 16  descend= 0  swap= 0
k= 16  j= 4  threadIdx.x= 7  swap_idx= 3  my_elem= 4  swap_elem= 6  ascend= 16  descend= 0  swap= 1
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
sortbuf addr= 0x7fdc7f000000  sortbuf= 1  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 0 
sortbuf addr= 0x7fdc7f000000  sortbuf= 2  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 1 
sortbuf addr= 0x7fdc7f000000  sortbuf= 3  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 2 
sortbuf addr= 0x7fdc7f000000  sortbuf= 4  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 3 
sortbuf addr= 0x7fdc7f000000  sortbuf= 5  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 4 
sortbuf addr= 0x7fdc7f000000  sortbuf= 6  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 5 
sortbuf addr= 0x7fdc7f000000  sortbuf= 7  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 6 
sortbuf addr= 0x7fdc7f000000  sortbuf= 8  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 7 
sortbuf addr= 0x7fdc7f000000  sortbuf= 9  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 8 
sortbuf addr= 0x7fdc7f000000  sortbuf= 10  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 9 
sortbuf addr= 0x7fdc7f000000  sortbuf= 11  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 10 
sortbuf addr= 0x7fdc7f000000  sortbuf= 12  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 11 
sortbuf addr= 0x7fdc7f000000  sortbuf= 13  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 12 
sortbuf addr= 0x7fdc7f000000  sortbuf= 14  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 13 
sortbuf addr= 0x7fdc7f000000  sortbuf= 15  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 14 
sortbuf addr= 0x7fdc7f000000  sortbuf= 16  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 15 
k= 16  j= 2  threadIdx.x= 0  swap_idx= 2  my_elem= 1  swap_elem= 3  ascend= 0  descend= 16  swap= 0
k= 16  j= 2  threadIdx.x= 1  swap_idx= 3  my_elem= 2  swap_elem= 4  ascend= 0  descend= 16  swap= 0
k= 16  j= 2  threadIdx.x= 2  swap_idx= 0  my_elem= 3  swap_elem= 1  ascend= 16  descend= 0  swap= 0
k= 16  j= 2  threadIdx.x= 3  swap_idx= 1  my_elem= 4  swap_elem= 2  ascend= 16  descend= 0  swap= 0
k= 16  j= 2  threadIdx.x= 4  swap_idx= 6  my_elem= 7  swap_elem= 5  ascend= 0  descend= 16  swap= 1
k= 16  j= 2  threadIdx.x= 5  swap_idx= 7  my_elem= 8  swap_elem= 6  ascend= 0  descend= 16  swap= 1
k= 16  j= 2  threadIdx.x= 6  swap_idx= 4  my_elem= 5  swap_elem= 7  ascend= 16  descend= 0  swap= 1
k= 16  j= 2  threadIdx.x= 7  swap_idx= 5  my_elem= 6  swap_elem= 8  ascend= 16  descend= 0  swap= 1
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
BBBBBBBBBBBBBBBBBBBBBBBBB
sortbuf addr= 0x7fdc7f000000  sortbuf= 1  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 0 
sortbuf addr= 0x7fdc7f000000  sortbuf= 2  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 1 
sortbuf addr= 0x7fdc7f000000  sortbuf= 3  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 2 
sortbuf addr= 0x7fdc7f000000  sortbuf= 4  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 3 
sortbuf addr= 0x7fdc7f000000  sortbuf= 5  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 4 
sortbuf addr= 0x7fdc7f000000  sortbuf= 6  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 5 
sortbuf addr= 0x7fdc7f000000  sortbuf= 7  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 6 
sortbuf addr= 0x7fdc7f000000  sortbuf= 8  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 7 
sortbuf addr= 0x7fdc7f000000  sortbuf= 9  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 8 
sortbuf addr= 0x7fdc7f000000  sortbuf= 10  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 9 
sortbuf addr= 0x7fdc7f000000  sortbuf= 11  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 10 
sortbuf addr= 0x7fdc7f000000  sortbuf= 12  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 11 
sortbuf addr= 0x7fdc7f000000  sortbuf= 13  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 12 
sortbuf addr= 0x7fdc7f000000  sortbuf= 14  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 13 
sortbuf addr= 0x7fdc7f000000  sortbuf= 15  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 14 
sortbuf addr= 0x7fdc7f000000  sortbuf= 16  my_elem addr= 0x7fdc81fffc50  threadIdx.x= 15 
k= 16  j= 1  threadIdx.x= 0  swap_idx= 1  my_elem= 1  swap_elem= 2  ascend= 0  descend= 16  swap= 0
k= 16  j= 1  threadIdx.x= 1  swap_idx= 0  my_elem= 2  swap_elem= 1  ascend= 16  descend= 0  swap= 0
k= 16  j= 1  threadIdx.x= 2  swap_idx= 3  my_elem= 3  swap_elem= 4  ascend= 0  descend= 16  swap= 0
k= 16  j= 1  threadIdx.x= 3  swap_idx= 2  my_elem= 4  swap_elem= 3  ascend= 16  descend= 0  swap= 0
k= 16  j= 1  threadIdx.x= 4  swap_idx= 5  my_elem= 5  swap_elem= 6  ascend= 0  descend= 16  swap= 0
k= 16  j= 1  threadIdx.x= 5  swap_idx= 4  my_elem= 6  swap_elem= 5  ascend= 16  descend= 0  swap= 0
k= 16  j= 1  threadIdx.x= 6  swap_idx= 7  my_elem= 7  swap_elem= 8  ascend= 0  descend= 16  swap= 0
k= 16  j= 1  threadIdx.x= 7  swap_idx= 6  my_elem= 8  swap_elem= 7  ascend= 16  descend= 0  swap= 0
Goodbye main()







吃我老猪一钉耙
关注
专栏目录
CUDA：并行快速排序的实例
希望我的博客，能帮上你解决学习中工作中所遇到的问题
12-21 293
CUDA：并行快速排序的实例
CUDA并行排序(1)——整数
12-06
// 1. A small-set insertion sort. We do this on any set with <=32 elements // 2. A partitioning kernel, which - given a pivot - separates an input // array into elements <=pivot, and >pivot. Two quicksorts will then // be launched to resolve each of these. // 3. A quicksort co-ordinator, which figures out what kernels to launch // and when.
CUDA动态并行实现快速排序
xiaohu的博客
11-02 2930
排序是任何应用的基本构造块的关键算法之一。有许多排序算法已经被广泛研究。，常见的排序算法时间和空间复杂度如下：一些排序算法属于分治算法的范畴。这些算法适用于并行性，并适合 GPU 等架构，在这些架构中，要排序的数据可以被划分进行排序。快速排序就是这样一种算法。如前所述，Quicksort 属于“分而治之”的范畴。与其他排序算法(如需要额外存储的合并排序)相比，快速排序的内存负载和要求更低。快速排序的更实际的实现使用随机化版本。随机化版本的预期时间复杂度为O(nlogn)。
CUDA(六). 从并行排序方法理解并行化思维——冒泡、归并、双调排序的GPU实现
热门推荐
Rachel Zhang的专栏
09-14 4万+
在第五讲中我们学习了GPU三个重要的基础并行算法： Reduce, Scan 和 Histogram，分析了 其作用与串并行实现方法。 在第六讲中，本文以冒泡排序 Bubble Sort、归并排序 Merge Sort 和排序网络中的双调排序 Bitonic Sort 为例， 讲解如何从数据结构课上学的串行并行排序方法转换到并行排序，并附GPU实现代码。 在并行方法中，我们将考虑到并行方法需要注意
cuda双调排序
tiger1334的专栏
03-28 1692
#include "cuda_runtime.h" #include "device_launch_parameters.h" #include "device_functions.h" #include #include #include #define SHARED_SIZE_LIMIT 1024U #define UMUL(a, b) __umul24((a), (b)) #def
CUDA并行排序(2)——一维Double型10000
12-15
double a,b,c; char d,e; int size=20; double arr[size][3];
CUDA 六 从并行排序方法理解并行化思维——冒泡 归并 双调排序的GPU实现
kahncc的博客
11-10 359
CUDA 六 从并行排序方法理解并行化思维——冒泡 归并 双调排序的GPU实现
NVIDIA CUDA 高度并行处理器编程（二）：数据并行执行模型
ikun的博仔的博客
05-10 769
NVIDIA CUDA 高度并行处理器编程（二）：数据并行执行模型1. CUDA的线程组织 内置变量：内置变量的值一般都由运行时系统预初始化，例如 CUDA 的 kernel 函数中， gradDim，blockDim，blockIdx，threadIdx都是内置变量，它们的值由 CUDA 运行时系统预初始化，可以在 kernel 函数中引用。其他地方应避免使用这些变量。 1. CUDA的线程组织 上一节向量加法中的线程被组织成二级的层次结构：一个网络包含一个或更多的线程块，每块包含一个或更多的线程。一个块
CUDA C编程4 - 动态并行理解
DU_YULIN的博客
03-01 1858
系列文章目录 文章目录系列文章目录前言一、动态并行概念二、嵌套执行三、动态并行的限制条件四、嵌套规约4.1 基本方法（线程同步及线程块同步）4.2 加速方法（去除线程同步以及线程块同步）4.3 进一步加速方法（待完善。。。）五、运行结果总结参考资料附录：完整代码 前言 之前学习了CUDA C编程中动态并行技术， 这里进行总结，以加深理解，并希望能够对大家有所帮助。 一、动态并行概念 之前接触的CUDA C设备函数的调用都是在CPU端进行的，即设备函数的执行在CPU的控制下。自然而然，我们就会
CUDA并行排序(4)——双调排序(二维double型数据)
A person does not really understand something until after teaching it to a computer. Donald.E.Knuth
12-09 904
#include #include #include #include using namespace std; #define CHECK(res) if(res!=cudaSuccess){exit(-1);} __global__ void helloCUDA(double **dp_out_params, unsigned int len) { //在核函数内部定义的变量，没有 _
双调排序详解
09-25
双调排序的详解，还是比较容易看懂的，可用于cuda并行算法（png）
cuda 作业 排序算法 sort
05-24
中科院 cuda 作业 排序算法 sort
基于CUDA的大量数字排序
11-26
对于大量的数据，用传统CPU排序比较慢，本资料根据GPU的CUDA平台，实现1000个数据的排序，对于想学并行计算的朋友来说是好的资料。
sequential-vs-parallel-sort:C ++中实现的顺序排序算法与CUDA中实现的并行排序算法之间的比较研究，这是硕士学位论文的一部分
05-23
信息 作为硕士论文的一部分，对C ++中实现的顺序排序算法与CUDA中实现的并行排序算法之间的比较进行了研究。 我们实现了七个算法：双音排序，多步双音排序，自适应双音排序，合并排序，快速排序，基数排序和样本排序。 顺序算法是在使用C ++的CPU上实现的，而并行算法是在使用CUDA平台的GPU上实现的。 我们改进了上述实现，并采用了它们，以便能够对任意长度的输入序列进行排序。 我们在六个不同的输入分布上比较了算法，这些分布包括32位数字，32位键值对，64位数字和64位键值对。 结果表明，基数排序是最快的顺序排序算法，而基数排序和合并排序是最快的并行排序算法（取决于输入分布）。 与顺序实现相比，通过并行实现，我们最多可实现157倍的加速。 作者：DarkoBožidar 导师：TomažDobravec博士 资料下载 结果： : 硕士学位论文（斯洛文尼亚版） ： : usp
CUDA并行排序(2)——一维Double型
12-07
Sorted 100000000 Double elems in 1803.735 ms (55.441 Melems/sec)
CUDA与大数组的双调排序
林地宁宁的树洞
05-26 1760
CUDA 与大数组的双调排序 双调排序 双调排序，又称作 Bitonic Sort，整体复杂度为 nlog⁡2(n)n\log^2(n)nlog2(n)。 归并排序劣势 相比于归并排序的 nlog⁡(n)n\log(n)nlog(n) 的复杂度，看起来优势不大，但是胜在能够比归并排序更加有效的运用多核多线程，因此在 CUDA 加持下，效率往往快于归并排序。 归并排序的缺点： 当归并排序到后期时，待归并的段主要有少数的“大段”构成，而这几个大段的归并比较难运用到多核多线程的优势。尽管可以通过一定手段优化几个大
【并行计算】Bitonic Sort（双调排序）基础
jiange_zh的博客
10-31 1万+
双调序列双调序列(Bitonic Sequence)是指由一个非严格增序列X和非严格减序列Y构成的序列，比如序列（23,10,8,3,5,7,11,78）。定义:一个序列a1,a2,…,an是双调序列(Bitonic Sequence)，如果： (1)存在一个ak(1≤k≤n), 使得a1≥…≥ak≤…≤an成立；或者 (2)序列能够循环移位满足条件(1)Batcher定理将任意一个长为2n的双
CUDA：实现双调排序奇偶合并排序(附完整源码)
最新发布
希望我的博客，能帮上你解决学习中工作中所遇到的问题
03-04 179
CUDA：实现双调排序奇偶合并排序(附完整源码)
cuda对global memery的双调排序
chanzhennan的博客
10-09 362
__device__ void swap(int& a, int& b) { int t = a; a = b; b = t; } __global__ void sort(int* a, int flag_j, int flag_i, int count) { unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x; unsigned int tid_comp = tid ^ flag_j.
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值