使用cuda选择基数

我一直在使用CUDA开发基数选择，它利用k个最小元素对给定数量的元素进行排序。这种基数选择背后的主要思想是扫描从MSB到LSB的32位整数。它将左侧的所有0位和右侧的所有1位进行分区。包含k个最小元素的边被递归求解。我的分区过程运行得很好，但我在处理递归函数调用时遇到了问题。我无法停止递归。请帮帮我！我的内核函数如下：这是kernel.h

  #include "header.h"
  #define WARP_SIZE 32
  #define BLOCK_SIZE 32
__device__ int Partition(int *d_DataIn, int firstidx, int lastidx, int k, int N, int bit)
{
int threadID = threadIdx.x + BLOCK_SIZE * blockIdx.x;
int WarpID = threadID >> 5;
int LocWarpID = threadID - 32 * WarpID;
int NumWarps = N / WARP_SIZE;
int pivot;
__shared__ int DataPartition[BLOCK_SIZE];
__shared__ int DataBinary[WARP_SIZE];
for(int i = 0; i < NumWarps; i++)
{
    if(LocWarpID >= firstidx && LocWarpID <=lastidx)
    {
        int r = d_DataIn[i * WARP_SIZE + LocWarpID];
        int p = (r>>(31-bit))&1;
        unsigned int B = __ballot(p);
        unsigned int B_flip = ~B;
        if(p==1)
        {
            int b = B << (32-LocWarpID);
            int RightLoc = __popc(b);
            DataPartition[lastidx - RightLoc] = r;
        }
        else
        {
            int b_flip = B_flip << (32 - LocWarpID);
            int LeftLoc = __popc(b_flip);
            DataPartition[LeftLoc] = r;
        }
    if(LocWarpID <= lastidx - __popc(B))
        {
            d_DataIn[LocWarpID] = DataPartition[LocWarpID];
        }
    else
        {
            d_DataIn[LocWarpID] = DataPartition[LocWarpID];
        }
        pivot = lastidx - __popc(B); 
        return pivot+1;
    }
}   
}

__device__ int RadixSelect(int *d_DataIn, int firstidx, int lastidx, int k, int N, int bit)
{
if(firstidx == lastidx)
    return *d_DataIn;
int q = Partition(d_DataIn, firstidx, lastidx, k, N, bit);
int length = q - firstidx;
if(k == length)
    return *d_DataIn; 
else if(k < length)
    return RadixSelect(d_DataIn, firstidx, q-1, k, N, bit+1);
else
    return RadixSelect(d_DataIn, q, lastidx, k-length, N, bit+1);
}
 __global__ void radix(int *d_DataIn, int firstidx, int lastidx, int k, int N, int bit)
{
RadixSelect(d_DataIn, firstidx, lastidx, k, N, bit);
}

主机代码是main.cu，看起来像：

#include "header.h"
#include <iostream>
#include <fstream>
#include "kernel.h"
#define BLOCK_SIZE 32
using namespace std;
int main()
{
int N = 32;
thrust::host_vector<float>h_HostFloat(N);
thrust::counting_iterator <unsigned int> Numbers(0);
thrust::transform(Numbers, Numbers + N, h_HostFloat.begin(),         RandomFloatNumbers(1.f, 100.f));
thrust::host_vector<int>h_HostInt(N);
thrust::transform(h_HostFloat.begin(), h_HostFloat.end(), h_HostInt.begin(), FloatToInt());
thrust::device_vector<float>d_DeviceFloat = h_HostFloat;
thrust::device_vector<int>d_DeviceInt(N);   
thrust::transform(d_DeviceFloat.begin(), d_DeviceFloat.end(), d_DeviceInt.begin(), FloatToInt());
int *d_DataIn = thrust::raw_pointer_cast(d_DeviceInt.data());
int *h_DataOut;
float *h_DataOut1;
int fsize = N * sizeof(float);
int size = N * sizeof(int);
h_DataOut = new int[size];
h_DataOut1 = new float[fsize];
int firstidx = 0;
int lastidx = BLOCK_SIZE-1;
int k = 20;
int bit = 1;
int NUM_BLOCKS = N / BLOCK_SIZE;
radix <<< NUM_BLOCKS, BLOCK_SIZE >>> (d_DataIn, firstidx, lastidx, k, N, bit);
cudaMemcpy(h_DataOut, d_DataIn, size, cudaMemcpyDeviceToHost);  
WriteData(h_DataOut1, h_DataOut, 10, N);

return 0;
   }

我使用的标题列表：

   #include "cuda.h"
   #include "cuda_runtime_api.h"
   #include "device_launch_parameters.h"
   #include <thrust/host_vector.h>
   #include <thrust/device_vector.h>
   #include <thrust/transform.h>
   #include <thrust/generate.h>
   #include "functor.h"
   #include <thrust/iterator/counting_iterator.h>
   #include <thrust/copy.h>
   #include <thrust/device_ptr.h>

另一个头文件"functor.h"，用于将浮点数转换为int类型并生成随机浮点数。

  #include <thrust/random.h>
  #include <sstream>
  #include <fstream>
  #include <iomanip>
  struct RandomFloatNumbers
   {
float a, b;
__host__ __device__
    RandomFloatNumbers(float _a, float _b) : a(_a), b(_b) {};  
__host__ __device__
    float operator() (const unsigned int n) const{
        thrust::default_random_engine rng;
        thrust::uniform_real_distribution<float> dist(a,b);
        rng.discard(n);
        return dist(rng);
}
    };
   struct FloatToInt
    {
__host__ __device__
    int operator() (const float &x)
const {
    union {
        float f_value;
        int i_value;
    } value;
    value.f_value = x;
    return value.i_value;
}
    };
   float IntToFloat(int &x)
     {
    union{
          float f_value;
          int i_value;
        }value;
        value.i_value = x;
        return value.f_value;
     }
   bool WriteData(float *h_DataOut1, int *h_DataOut, int bit, int N)
    {
std::ofstream data;
std::stringstream file;
file << "out\Partition_";
file << std::setfill('0') <<std::setw(2) << bit;
file << ".txt";
data.open((file.str()).c_str());
if(data.is_open() ==  false)
{
    std::cout << "File is not open" << std::endl;
    return false;
}
    for(int i = 0; i < N; i++)
    {
        h_DataOut1[i] = IntToFloat(h_DataOut[i]);
        //cout << h_HostFloat[i] << " t" <<  h_DataOut1[i] << endl;
        //std::bitset<32>bitshift(h_DataOut[i]&1<<31-bit);
        //data << bitshift[31-bit] << "t" <<h_DataOut1[i] <<std::endl;
        data << h_DataOut1[i] << std::endl;
    }
    data << std::endl;
data.close();
std::cout << "Partition=" <<bit <<"n"; 
return true;
    }