view src/parallel_execution/examples/bitonicSort/CUDAbitonicSwap.cu @ 420:764c92c3b181

Fix
author Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
date Fri, 06 Oct 2017 15:10:07 +0900
parents 49159fbdd1fb
children 6bb391fc9e12
line wrap: on
line source

extern "C" {
    struct Integer {
        int value;
    };
    __global__ void bitonicSwap(struct Integer* array, int* blockPtr, int* firstPtr, int* prefixPtr) {
        int block = *blockPtr;
        int first = *firstPtr;
        int prefix = *prefixPtr;
        int i = 0;
C_bitonicSwap:
        if (i < prefix) {
            int index = i + (blockIdx.x * blockDim.x + threadIdx.x) * prefix;
            int position = index/block;
            int index1 = index+block*position;
            int index2 = (first == 1)? ((block<<1)*(position+1))-(index1%block)-1 : index1+block;
            if (array[index2].value < array[index1].value) {
                struct Integer tmp = array[index1];
                array[index1] = array[index2];
                array[index2] = tmp;
            }
            i++;
            goto C_bitonicSwap;
        }
    }
}