cuda 中的双音排序错误地处理了一些值

Question

我正在为一个更大的项目在 CUDA 上制作排序算法，我决定实施双音排序。我将要排序的元素数量总是 2 的幂，实际上是 512。我需要一个具有最终位置的数组，因为此方法将用于对代表另一个质量矩阵的数组进行排序解决方案。

fitness 是我要排序的数组，numElements 是元素的数量，orden 最初是一个空数组，其中包含 numElements 个位置，将在最开始以这种方式填充：orden[i]=i。实际上 orden 与这个问题无关，但我保留了它。

我的问题是某些值没有正确排序，直到现在我一直无法弄清楚我有什么问题。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <ctime>
#include <cuda.h>
#include <curand.h>
#include <curand_kernel.h>
#include <device_functions.h>
#include "float.h"


__global__ void sorting(int * orden, float * fitness, int numElements);

// Populating array with random values for testing purposes
__global__ void populate( curandState * state, float * fitness{

    curandState localState = state[threadIdx.x];
    int a = curand(&localState) % 500;
    fitness[threadIdx.x] = a;
}

//Curand setup for the populate method 
__global__ void setup_cuRand(curandState * state, unsigned long seed)
{
    int id = threadIdx.x;
    curand_init(seed, id, 0, &state[id]);
}

int main()
{
    float * arrayx;
    int numelements = 512;
    int * orden;
    float arrayCPU[512] = { 0 };
    curandState * state;

    cudaDeviceReset();
    cudaSetDevice(0);
    cudaMalloc(&state, numelements * sizeof(curandState));
    cudaMalloc((void **)&arrayx, numelements*sizeof(float));
    cudaMalloc((void **)&orden, numelements*sizeof(int));






    setup_cuRand << <1, numelements >> >(state, unsigned(time(NULL)));

    populate << <1, numelements >> > (state, arrayx);
    cudaMemcpy(&arrayCPU, arrayx, numelements * sizeof(float), cudaMemcpyDeviceToHost);
    for (int i = 0; i < numelements; i++)
        printf("fitness[%i] = %f\n", i, arrayCPU[i]);

    sorting << <1, numelements >> >(orden, arrayx, numelements);
    printf("\n\n");

    cudaMemcpy(&arrayCPU, arrayx, numelements * sizeof(float), cudaMemcpyDeviceToHost);
    for (int i = 0; i < numelements; i++)
        printf("fitness[%i] = %f\n", i, arrayCPU[i]);



    cudaDeviceReset();


    return 0;
}
__device__ bool isValid(float n){
    return !(isnan(n) || isinf(n) || n != n || n <= FLT_MIN || n >= FLT_MAX);

}

__global__ void sorting(int * orden, float * fitness, int numElements){
    int i = 0;
    int j = 0;
    float f = 0.0;
    int aux = 0;

    //initial orden registered (1, 2, 3...)
    orden[threadIdx.x] = threadIdx.x;
    //Logarithm on base 2 of numElements
    for (i = 2; i <= numElements; i = i * 2){
        // descending from i reducing to half each iteration
        for (j = i; j >= 2; j = j / 2){

            if (threadIdx.x % j  < j / 2){
                __syncthreads();
                // ascending or descending consideration using (threadIdx.x % (i*2) < i) 
                if ((threadIdx.x % (i * 2) < i) && (fitness[threadIdx.x] >  fitness[threadIdx.x + j / 2] || !isValid(fitness[threadIdx.x])) ||
                    ((threadIdx.x % (i * 2) >= i) && (fitness[threadIdx.x] <= fitness[threadIdx.x + j / 2] || !isValid(fitness[threadIdx.x + j / 2])))){

                    aux = orden[threadIdx.x];
                    orden[threadIdx.x] = orden[threadIdx.x + j / 2];
                    orden[threadIdx.x + j / 2] = aux;
                    //Se reubican los fitness
                    f = fitness[threadIdx.x];
                    fitness[threadIdx.x] = fitness[threadIdx.x + j / 2];
                    fitness[threadIdx.x + j / 2] = f;
                }
            }
        }
    }
}

例如，我在随机执行时得到的输出：

A random execution

这是我的双调排序的表示：

Bitonic sorting Schema，箭头指向最差值比较的位置

Answer 1

以下是我发现的问题：

1. 在您发布的代码中，这不会编译：

   __global__ void populate( curandState * state, float * fitness{
                                                                 ^
                                                      missing close parenthesis
   

   我在那里添加了一个右括号。

2. 在这些cudaMemcpy 语句中不必取数组的地址：

   cudaMemcpy(&arrayCPU, arrayx, numelements * sizeof(float), cudaMemcpyDeviceToHost);
   ....
   cudaMemcpy(&arrayCPU, arrayx, numelements * sizeof(float), cudaMemcpyDeviceToHost);
   

   数组名已经是数组的地址，所以我去掉了和号。如果你使用动态分配的数组，这种用法会被破坏。

3. 您在此处对__syncthreads() 的使用已损坏：

   for (j = i; j >= 2; j = j / 2){
   
       if (threadIdx.x % j  < j / 2){
           __syncthreads();
   

   在条件语句中使用__syncthreads() 通常是不正确的，除非条件语句在线程块中统一计算。这在documentation 中有介绍。稍微改动一下就可以达到想要的效果：

   for (j = i; j >= 2; j = j / 2){
       __syncthreads();
       if (threadIdx.x % j  < j / 2){
   

通过上述更改，在大多数情况下，您的代码对我来说似乎可以正常运行。如果您打算正确排序 0（或任何负值），您在有效性检查中对 FLT_MIN 的使用也是有问题的。一般来说，FLT_MIN 是一个数字，即very small, close to zero。如果您认为这是一个很大的负数，那不是。因此，零是您的随机数生成器的可能输出，它不会被正确排序。我会把这个留给你来解决，它应该很简单，但这取决于你最终想要实现的目标。 （如果您只想对正的非零浮点值进行排序，测试可能没问题，但在这种情况下，您的随机数生成器可以返回 0。）