cudaでのビオニックソートはいくつかの値を狂わせます

Question

私はより大きなプロジェクトのためにCUDAでソートアルゴリズムを作っています。私はBitonicソートを実装することに決めました。ソートする要素の数は常に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 
   

   私はそこに閉じ括弧を追加しました。それは、これらのcudaMemcpy文で配列のアドレスを取得する必要はありません

2. ：

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

   配列名は、すでに配列のアドレスですので、私はアンパサンドを削除しました。動的に割り当てられた配列を使用すると、そのような使用法は壊れてしまいます。ここ__syncthreads()の

3. ご利用が壊れている：__syncthreads()の

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

   使用を条件文の内側に条件文がthreadblock全体に均一に評価されていない限り、一般的に間違っています。これはdocumentationに記載されています。私たちは、わずかな変化で所望の効果を達成することができます：上記の変更により

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

を、あなたのコードは、ほとんどの場合のために、私のために正しく実行するに表示されます。有効性チェックでFLT_MINを使用しても、0（または任意の負の値）を正しく並べ替えることを意図している場合は疑問があります。一般に、FLT_MINはvery small, close to zeroの数字です。これが大きな負数であると考えていた場合、それはそうではありません。その結果、ゼロは乱数ジェネレータの出力であり、正しくソートされません。私はこの問題を解決するために残しておきますが、それは簡単ですが、最終的に達成したいことに依存します。（正のゼロ以外の浮動小数点値だけをソートする場合は、テストは正常ですが、この場合、乱数ジェネレータは0を返すことができます）。