Estoy tratando de escribir un núcleo de histograma en OpenCL para calcular 256 histogramas R, G y B de bin de una imagen de entrada RGBA32F. Mi kernel es el siguiente:Histograma de imagen OpenCL
const sampler_t mSampler = CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP|
CLK_FILTER_NEAREST;
__kernel void computeHistogram(read_only image2d_t input, __global int* rOutput,
__global int* gOutput, __global int* bOutput)
{
int2 coords = {get_global_id(0), get_global_id(1)};
float4 sample = read_imagef(input, mSampler, coords);
uchar rbin = floor(sample.x * 255.0f);
uchar gbin = floor(sample.y * 255.0f);
uchar bbin = floor(sample.z * 255.0f);
rOutput[rbin]++;
gOutput[gbin]++;
bOutput[bbin]++;
}
Cuando lo ejecuto en una imagen 2100 x 894 (1,877,400 píxeles) i tienden a ver sólo en o alrededor de 1.870.000 valores totales se registran cuando resumo los valores del histograma para cada canal . También es un número diferente cada vez. Lo esperaba, ya que de vez en cuando, dos núcleos probablemente toman el mismo valor de la matriz de salida y lo incrementan, cancelando efectivamente una operación de incremento (¿lo asumo?).
La salida de 1,870,000 es para un {1,1} tamaño de grupo de trabajo (que es lo que parece establecerse por defecto si no especifico lo contrario). Si forzo un tamaño de grupo de trabajo más grande como {10,6}, obtengo una suma drásticamente menor en mi histograma (proporcional al cambio en el tamaño del grupo de trabajo). Esto me pareció extraño, pero supongo que lo que sucede es que todos los elementos de trabajo en el grupo incrementan el valor de la matriz de salida al mismo tiempo, por lo que solo cuenta como un incremento individual.
De todos modos, he leído en la especificación que OpenCL no tiene sincronización de memoria global, solo sincroniza dentro de los grupos de trabajo locales usando su __ memoria local. El ejemplo de histograma de nVidia divide la carga de trabajo del histograma en un conjunto de subproblemas de un tamaño específico, calcula sus histogramas parciales y luego fusiona los resultados en un solo histograma después. Esto no parece funcionar bien para imágenes de tamaño arbitrario. Supongo que podría rellenar los datos de imagen con valores ficticios ...
Siendo nuevo en OpenCL, supongo que me pregunto si hay una manera más directa de hacerlo (ya que parece que debería ser relativamente sencillo). Problema GPGPU).
Gracias!