Quelle est la meilleure pratique à faire réduire dans OpenCL?

Question

Imaginez une opération binaire (nommez-la "+") avec une propriété associative. Lorsque vous pouvez calculer a1 + a2 + a3 + a4 + ... en parallèle, premier calcul

b1 = a1 + a2 
b2 = a3 + a4 

puis

c1 = b1 + b2 
c2 = b3 + b4 

faire alors la même chose pour les résultats de l'étape précédente, et ainsi de suite, jusqu'à ce qu'il y est un élément gauche. J'apprençais OpenCL et j'essayais d'implémenter cette approche pour récapituler tous les éléments du tableau. Je suis un débutant total dans cette technologie, donc le programme pourrait sembler quelque chose d'étrange.

Ceci est le noyau:

__kernel void reduce (__global float *input, __global float *output) 
{ 
    size_t gl = get_global_id (0); 
    size_t s = get_local_size (0); 
    int i; 
    float accum = 0; 

    for (i=0; i<s; i++) { 
     accum += input[s*gl+i]; 
    } 

    output[gl] = accum; 
} 

C'est le programme principal:

#include <stdio.h> 
#include <stdlib.h> 
#include <fcntl.h> 
#include <unistd.h> 
#include <sys/mman.h> 
#include <sys/stat.h> 
#include <CL/cl.h> 

#define N (64*64*64*64) 

#include <sys/time.h> 
#include <stdlib.h> 

double gettime() 
{ 
    struct timeval tv; 
    gettimeofday (&tv, NULL); 
    return (double)tv.tv_sec + (0.000001 * (double)tv.tv_usec); 
} 

int main() 
{ 
    int i, fd, res = 0; 
    void* kernel_source = MAP_FAILED; 

    cl_context context; 
    cl_context_properties properties[3]; 
    cl_kernel kernel; 
    cl_command_queue command_queue; 
    cl_program program; 
    cl_int err; 
    cl_uint num_of_platforms=0; 
    cl_platform_id platform_id; 
    cl_device_id device_id; 
    cl_uint num_of_devices=0; 
    cl_mem input, output; 
    size_t global, local; 

    cl_float *array = malloc (sizeof (cl_float)*N); 
    cl_float *array2 = malloc (sizeof (cl_float)*N); 
    for (i=0; i<N; i++) array[i] = i; 

    fd = open ("kernel.cl", O_RDONLY); 
    if (fd == -1) { 
     perror ("Cannot open kernel"); 
     res = 1; 
     goto cleanup; 
    } 
    struct stat s; 

    res = fstat (fd, &s); 
    if (res == -1) { 
     perror ("Cannot stat() kernel"); 
     res = 1; 
     goto cleanup; 
    } 

    kernel_source = mmap (NULL, s.st_size, PROT_READ, MAP_PRIVATE, fd, 0); 
    if (kernel_source == MAP_FAILED) { 
     perror ("Cannot map() kernel"); 
     res = 1; 
     goto cleanup; 
    } 

    if (clGetPlatformIDs (1, &platform_id, &num_of_platforms) != CL_SUCCESS) { 
     printf("Unable to get platform_id\n"); 
     res = 1; 
     goto cleanup; 
    } 

    if (clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, 
         &num_of_devices) != CL_SUCCESS) 
    { 
     printf("Unable to get device_id\n"); 
     res = 1; 
     goto cleanup; 
    } 
    properties[0]= CL_CONTEXT_PLATFORM; 
    properties[1]= (cl_context_properties) platform_id; 
    properties[2]= 0; 
    context = clCreateContext(properties,1,&device_id,NULL,NULL,&err); 
    command_queue = clCreateCommandQueue(context, device_id, 0, &err); 
    program = clCreateProgramWithSource(context, 1, (const char**)&kernel_source, NULL, &err); 


    if (clBuildProgram(program, 0, NULL, NULL, NULL, NULL) != CL_SUCCESS) { 
     char buffer[4096]; 
     size_t len; 

     printf("Error building program\n"); 
     clGetProgramBuildInfo (program, device_id, CL_PROGRAM_BUILD_LOG, sizeof (buffer), buffer, &len); 
     printf ("%s\n", buffer); 
     res = 1; 
     goto cleanup; 
    } 

    kernel = clCreateKernel(program, "reduce", &err); 
    if (err != CL_SUCCESS) { 
     printf("Unable to create kernel\n"); 
     res = 1; 
     goto cleanup; 
    } 

    // create buffers for the input and ouput 
    input = clCreateBuffer(context, CL_MEM_READ_ONLY, 
          sizeof(cl_float) * N, NULL, NULL); 
    output = clCreateBuffer(context, CL_MEM_WRITE_ONLY, 
          sizeof(cl_float) * N, NULL, NULL); 

    // load data into the input buffer 
    clEnqueueWriteBuffer(command_queue, input, CL_TRUE, 0, 
          sizeof(cl_float) * N, array, 0, NULL, NULL); 

    size_t size = N; 
    cl_mem tmp; 
    double time = gettime(); 
    while (size > 1) 
    { 
     // set the argument list for the kernel command 
     clSetKernelArg(kernel, 0, sizeof(cl_mem), &input); 
     clSetKernelArg(kernel, 1, sizeof(cl_mem), &output); 
     global = size; 
     local = 64; 

     // enqueue the kernel command for execution 
     clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global, 
          &local, 0, NULL, NULL); 
     clFinish(command_queue); 
     size = size/64; 
     tmp = output; 
     output = input; 
     input = tmp; 
    } 
    cl_float answer[1]; 
    clEnqueueReadBuffer(command_queue, tmp, CL_TRUE, 0, 
         sizeof(cl_float), array, 0, NULL, NULL); 
    time = gettime() - time; 
    printf ("%f %f\n", array[0], time); 

cleanup: 
    free (array); 
    free (array2); 
    clReleaseMemObject(input); 
    clReleaseMemObject(output); 
    clReleaseProgram(program); 
    clReleaseKernel(kernel); 
    clReleaseCommandQueue(command_queue); 
    clReleaseContext(context); 

    if (kernel_source != MAP_FAILED) munmap (kernel_source, s.st_size); 
    if (fd != -1) close (fd); 

    _Exit (res); // Kludge 
    return res; 
} 

Je noyau Réexécutez jusqu'à ce qu'il n'y a qu'un seul élément dans le tampon. Cette approche est-elle correcte pour calculer la somme des éléments dans OpenCL? Le temps que je mesure avec gettime est environ 10 fois plus lent lors de l'exécution d'une simple boucle sur le CPU (Clang compilé 4.0.0 et -O2 -ffast-math drapeaux). Matériel que j'utilise: Amd Ryzen 5 1600X et Amd Radeon HD 6950.

Answer 1

Il y a deux choses que vous pouvez faire pour essayer d'améliorer les performances.

Premièrement, débarrassez-vous de l'appel clFinish à l'intérieur de votre boucle. Cela oblige les exécutions individuelles des noyaux à dépendre de l'état entier de la file d'attente de commande atteignant un point de synchronisation avec l'hôte avant de continuer, ce qui est inutile. La seule synchronisation requise est que les noyaux s'exécutent dans l'ordre, et même si vous avez une file d'attente en désordre (que votre programme ne demande pas de toute façon), vous pouvez le garantir en utilisant simplement des objets événement.

size_t size = N; 
size_t total_expected_events = 0; 
for(size_t event_count = size; event_count > 1; event_count /= 64) 
    total_expected_events++; 
cl_event * events = malloc(total_expected_events * sizeof(cl_event)); 
cl_mem tmp; 
double time = gettime(); 
size_t event_index = 0; 
while (size > 1) 
{ 
    // set the argument list for the kernel command 
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &input); 
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &output); 
    global = size; 
    local = 64; 

    if(event_index == 0) 
     // enqueue the kernel command for execution 
     clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global, 
          &local, 0, NULL, events); 
    else 
     clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global, 
          &local, 1, events + (event_index - 1), events + event_index); 
    size = size/64; 
    tmp = output; 
    output = input; 
    input = tmp; 
    event_index++; 
} 
clFinish(command_queue); 
for(; event_index > 0; event_index--) 
    clReleaseEvent(events[event_index-1]); 
free(events); 
cl_float answer[1]; 
clEnqueueReadBuffer(command_queue, tmp, CL_TRUE, 0, 
        sizeof(cl_float), array, 0, NULL, NULL); 

L'autre chose à regarder potentiellement en effectue la réduction en un seul noyau, au lieu de l'étaler sur plusieurs invocations du même noyau. This is one potential exemple, même si cela peut être plus compliqué que vous ne le souhaitez.