奇偶排序使用CUDA加速

奇偶排序

可以证明奇偶排序是一种使序列逆序数降为0(全局最优)的迭代算法，当连续一轮奇排列和偶排列完成后，交换的元素个数都为0个，则排序已收敛到全局最优。
考虑到每次奇或偶排列需要LENGTH/2组比较，且在每次排序中，组之间没有依赖关系，于是在序列长度较大时，使用CUDA对其加速成为可能，将每一组比较对作为一个CUDA线程执行，于是奇排序有如下代码:

int curLocation = *tid % *length;
if(!(curLocation % 2) && curLocation != *length-1)
	if(array[curLocation] > array[curLocation+1])
	{
        int temp = array[curLocation];
        atomicExch(&array[curLocation], array[curLocation+1]);
        atomicExch(&array[curLocation+1], temp);
        atomicAdd(vOdd, 1);
	 }

偶排序有如下代码:

int curLocation = *tid % *length;
if(curLocation % 2 == 1 && curLocation != *length-1)
    if(array[curLocation] > array[curLocation+1])
    {
        int temp = array[curLocation];
        atomicExch(&array[curLocation], array[curLocation+1]);
        atomicExch(&array[curLocation+1], temp);
        atomicAdd(vEven, 1);
    }

每个线程执行一组奇、偶排序，考虑到偶排序需要在奇排序完成之后才能执行，因此需要使用线程同步函数，由于GRID中只设置了一个BLOCK，于是线程同步函数就可以保证所有分组比较全部完成，代码如下:

//printf("4\n");
int tId = threadIdx.x + blockIdx.x * blockDim.x;
//printf("tID: %d\n", tId);
oddSort(array, length, &tId, vOdd);
// 等待Block内所有线程完成奇排序
__syncthreads();  
evenSort(array, length, &tId, vEven);
// 等待Block内所有线程完成偶排序
__syncthreads();

于是整个CUDA并行加速奇偶排序算法如下所示:

#include <cuda.h>
#include <cuda_runtime_api.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <iostream>

using namespace std;

#define LENGTH 1024

__device__ void oddSort(int *array, int *length, int* tid, int *vOdd)
{
    int curLocation = *tid % *length;
    if(!(curLocation % 2) && curLocation != *length-1)
        if(array[curLocation] > array[curLocation+1])
        {
            int temp = array[curLocation];
            atomicExch(&array[curLocation], array[curLocation+1]);
            atomicExch(&array[curLocation+1], temp);
            atomicAdd(vOdd, 1);
        }
}

__device__ void evenSort(int *array, int *length, int* tid, int *vEven)
{
    int curLocation = *tid % *length;
    if(curLocation % 2 == 1 && curLocation != *length-1)
        if(array[curLocation] > array[curLocation+1])
        {
            int temp = array[curLocation];
            atomicExch(&array[curLocation], array[curLocation+1]);
            atomicExch(&array[curLocation+1], temp);
            atomicAdd(vEven, 1);
        }
}


__global__ void oddEvenSort(int *array, int *length, int *vOdd, int* vEven)
{
    //printf("4\n");
    int tId = threadIdx.x + blockIdx.x * blockDim.x;
    //printf("tID: %d\n", tId);
    oddSort(array, length, &tId, vOdd);
    // 等待Block内所有线程完成奇排序
    __syncthreads();  
    evenSort(array, length, &tId, vEven);
    // 等待Block内所有线程完成偶排序
    __syncthreads();
}


extern "C" void hostOddEvenSort(int *array, int length, int odd, int even)
{
    int *arrayDev, *lengthDev, *vOdd, *vEven;
    cudaMalloc(&arrayDev, length*sizeof(int));
    cudaMemcpy(arrayDev, array, length*sizeof(int), cudaMemcpyHostToDevice);
    cudaMalloc(&lengthDev, sizeof(int));
    cudaMemcpy(lengthDev, &length, sizeof(int), cudaMemcpyHostToDevice);
    cudaMalloc(&vOdd, sizeof(int));
    cudaMalloc(&vEven, sizeof(int));

    int *arrayHost, *vOddHost, *vEvenHost;
    arrayHost = (int*)malloc(length*sizeof(int));
    vOddHost = (int*)malloc(sizeof(int));
    vEvenHost = (int*)malloc(sizeof(int));
    do
    {
        cudaMemcpy(vOdd, &odd, sizeof(int), cudaMemcpyHostToDevice);
        cudaMemcpy(vEven, &even, sizeof(int), cudaMemcpyHostToDevice);
        //cout<< 2<< endl;
        oddEvenSort<<<1, LENGTH>>>(arrayDev, lengthDev, vOdd, vEven);
        //cout<< 3<< endl;
        cudaMemcpy(vOddHost, vOdd, sizeof(int), cudaMemcpyDeviceToHost);
        cudaMemcpy(vEvenHost, vEven, sizeof(int), cudaMemcpyDeviceToHost);
        cout<< "odd: "<< *vOddHost<< "even: "<< *vEvenHost<< endl;
    }while(*vOddHost!=0 || *vEvenHost!=0);
    cudaMemcpy(arrayHost, arrayDev, length*sizeof(int), cudaMemcpyDeviceToHost);
    for(int i=0; i<length; i++)
        cout<< arrayHost[i]<< ",";
    cout<< "."<< endl;
}


int main()
{
    cout<< "begin: "<< endl;
    int *array = new int[LENGTH];
    // cout<< -1<< endl;
    // 逆序数输入
    for(int i=0; i<LENGTH; i++)
        array[i] = LENGTH - i;
    // cout<< 0<< endl;
    // 奇偶排序后顺序输出
    hostOddEvenSort(array, LENGTH, 0, 0);
    return 0;
}