saber - Separación de números pares e impares en CUDA

numeros pares e impares del 1 al 100

Question

Tengo una matriz de números como {1,2,3,4,5,6,7,8,9,10} y quiero separar los números pares e impares como:

even = {2,4,6,8}

y:

odd = {1,3,5,7}

Estoy al tanto de las operaciones atómicas en CUDA, y también soy consciente de que no se espera que la producción sufra condiciones de carrera. No quiero usar operaciones atómicas. ¿Cómo puedo lograr esto sin usar palabras clave atómicas?

CÓDIGO:

#include <stdio.h>
#include <cuda.h>
// Kernel that executes on the CUDA device
__global__ void square_array(float *total,float *even,float *odd, int N)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int a=total[idx];
if ((a%2)==0)
{
for (int i=0;i<=idx;i++)
{
int b = even[i];
if(b==0)
{
even[i] = total[idx];
break;
}
}
}
else
{
for (int i=0;i<idx;i++)
{
int c = odd[i];
odd[i] = total[idx];
break;
}
}
}
// main routine that executes on the host
int main(void)
{
float *total_h,*even_h, *odd_h,*total_d, *even_d,*odd_d; // Pointer to host & device arrays
const int N = 10; // Number of elements in arrays
size_t size = N * sizeof(float);
total_h = (float *)malloc(size); // Allocate array on host
even_h = (float *)malloc(size); // Allocate array on host
odd_h = (float *)malloc(size); // Allocate array on host
cudaMalloc((void **) &total_d, size);
cudaMalloc((void **) &even_d, size);
cudaMemset(even_d,0,size);
cudaMalloc((void **) &odd_d, size); // Allocate array on device
cudaMemset(odd_d,0,size);
// Initialize host array and copy it to CUDA device
for (int i=0; i<N; i++) total_h[i] = (float)i+1;
cudaMemcpy(total_d, total_h, size, cudaMemcpyHostToDevice);
// Do calculation on device:
square_array <<< 1,10 >>> (total_d,even_d,odd_d, N);
// Retrieve result from device and store it in host array
cudaMemcpy(even_h, even_d, sizeof(float)*N, cudaMemcpyDeviceToHost);
cudaMemcpy(odd_h, odd_d, sizeof(float)*N, cudaMemcpyDeviceToHost);
// Print results
printf("total Numbers/n");
for (int i=0; i<N; i++) printf("%f/n",total_h[i]);
printf("EVEN Numbers/n");
for (int i=0; i<N; i++) printf("%f/n",even_h[i]);
printf("ODD Numbers/n");
for (int i=0; i<N; i++) printf("%f/n",odd_h[i]);
// Cleanup
free(total_h);
free(even_h);
free(odd_h);
cudaFree(total_d);
cudaFree(even_d);
cudaFree(odd_d);
}

SALIDA:

Answer 1

Según lo sugerido por Jared Hoberock, sería mucho más fácil usar el algoritmo de partición eficiente disponible en CUDA Thrust en lugar de comenzar el desarrollo de una rutina de partición propia. A continuación, encontrará un ejemplo completo trabajado.

#include <thrust/device_vector.h>
#include <thrust/partition.h>
#include <thrust/execution_policy.h>
struct is_even { __host__ __device__ bool operator()(const int &x) { return (x % 2) == 0; } };
void main() {
const int N = 10;
thrust::host_vector<int> h_data(N);
for (int i=0; i<N; i++) h_data[i] = i;
thrust::device_vector<int> d_data(h_data);
thrust::device_vector<int> d_evens(N/2);
thrust::device_vector<int> d_odds(N/2);
thrust::partition_copy(d_data.begin(), d_data.end(), d_evens.begin(), d_odds.begin(), is_even());
printf("Even numbers/n");
for (int i=0; i<N/2; i++) {
int val = d_evens[i];
printf("evens[%i] = %i/n",i,val);
}
printf("Odd numbers/n");
for (int i=0; i<N/2; i++) {
int val = d_odds[i];
printf("odds[%i] = %i/n",i,val);
}
}