kann nicht funktionieren OpenCL Sum Reduktion von AMD Probe Reduktion, wenn Workgroups Dimension

Question

Der folgende Code stammt aus the amd website

__kernel 
void reduce(__global float* buffer, 
      __local float* scratch, 
      __const int length, 
      __global float* result) { 

    int global_index = get_global_id(0); 
    float accumulator = INFINITY; 
    // Loop sequentially over chunks of input vector 
    while (global_index < length) { 
    float element = buffer[global_index]; 
    accumulator = (accumulator < element) ? accumulator : element; 
    global_index += get_global_size(0); 
    } 

    // Perform parallel reduction 
    int local_index = get_local_id(0); 
    scratch[local_index] = accumulator; 
    barrier(CLK_LOCAL_MEM_FENCE); 
    for(int offset = get_local_size(0)/2; 
     offset > 0; 
     offset = offset/2) { 
    if (local_index < offset) { 
     float other = scratch[local_index + offset]; 
     float mine = scratch[local_index]; 
     scratch[local_index] = (mine < other) ? mine : other; 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 
    } 
    if (local_index == 0) { 
    result[get_group_id(0)] = scratch[0]; 
    } 
} 

Ändern ich es angepasst, um es als eine Summe Reduktion funktioniert:

__kernel 
void reduce(__global float* buffer, 
      __local float* scratch, 
      __const int length, 
      __global float* result) { 

    int global_index = get_global_id(0); 
    float accumulator = 0.0; 
    // Loop sequentially over chunks of input vector 
    while (global_index < length) { 
    float element = buffer[global_index]; 
    accumulator = accumulator + element; 
    global_index += get_global_size(0); 
    } 

    // Perform parallel reduction 
    int local_index = get_local_id(0); 
    scratch[local_index] = accumulator; 
    barrier(CLK_LOCAL_MEM_FENCE); 
    for(int offset = get_local_size(0)/2; 
     offset > 0; 
     offset = offset/2) { 
    if (local_index < offset) { 
     float other = scratch[local_index + offset]; 
     float mine = scratch[local_index]; 
     scratch[local_index] = mine + other; 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 
    } 
    if (local_index == 0) { 
    result[get_group_id(0)] = scratch[0]; 
    } 
} 

Und es funktioniert wie ein Charme, wenn ich nur eine Arbeitsgruppe verwende (dh ich gebe NULL als local_work_size bis clEnqueueNDRangeKernel()), aber Dinge geraten außer Kontrolle, wenn ich versuche, die Arbeitsgruppe Dimension zu ändern. (Ich würde sagen, ich bin ein absoluter Neuling in OpenCL)

Was ich tue, ist wie folgt

#define GLOBAL_DIM 600 
#define WORK_DIM 60 

size_t global_1D[3] = {GLOBAL_DIM,1,1}; 
size_t work_dim[3] = {WORK_DIM,1,1}; 
err = clEnqueueNDRangeKernel(commands, av_velocity_kernel, 1, NULL, global_1D, work_dim, 0, NULL, NULL); //TODO CHECK THIS LINE 
if (err) { 
    printf("Error: Failed to execute av_velocity_kernel!\n");   printf("\n%s",err_code(err)); fflush(stdout);  return EXIT_FAILURE; } 

Muss ich es tun der falsche Weg ??

Außerdem habe ich festgestellt, dass wenn ich #define GLOBAL_DIM 60000 (was ich brauche) ich keinen lokalen Speicher haben. Bekomme ich "mehr" lokalen Speicher, wenn ich mehrere Arbeitsgruppen verwende oder der lokale Speicher gleichmäßig zwischen Arbeitsgruppen verteilt ist?

Answer 1

Zunächst funktionieren diese Reduktionskerne nur korrekt, wenn die Arbeitsgruppengröße eine Potenz von zwei ist. Das bedeutet, dass Sie anstelle von 60 etwas 64 verwenden sollten. Es gibt auch keine Möglichkeit, dass Sie durch Ändern von GLOBAL_DIM keinen lokalen Speicher mehr haben: Sie tun wahrscheinlich etwas falsch, wenn Sie den Kernel aufrufen.