Commit f51dcd2c authored by Victor Oliveira's avatar Victor Oliveira Committed by Øyvind "pippin" Kolås

Support for RGBu8

And minor changes in color conversion
parent b923b265
......@@ -88,7 +88,7 @@ gegl_buffer_cl_cache_merge (GeglBuffer *buffer,
if (!roi)
roi = &buffer->extent;
gegl_cl_color_babl (buffer->format, NULL, &size);
gegl_cl_color_babl (buffer->format, &size);
if (G_UNLIKELY (!cache_entries))
{
......@@ -213,8 +213,8 @@ gegl_buffer_cl_cache_from (GeglBuffer *buffer,
gint i;
gegl_cl_color_op conv = gegl_cl_color_supported (buffer->format, format);
gegl_cl_color_babl (buffer->format, NULL, &buf_size);
gegl_cl_color_babl (format, NULL, &dest_size);
gegl_cl_color_babl (buffer->format, &buf_size);
gegl_cl_color_babl (format, &dest_size);
if (G_UNLIKELY (!cache_entries))
{
......
......@@ -94,8 +94,8 @@ gegl_buffer_cl_iterator_add (GeglBufferClIterator *iterator,
else
i->conv[self] = gegl_cl_color_supported (buffer->format, format);
gegl_cl_color_babl (buffer->format, NULL, &i->buf_cl_format_size[self]);
gegl_cl_color_babl (format, NULL, &i->op_cl_format_size [self]);
gegl_cl_color_babl (buffer->format, &i->buf_cl_format_size[self]);
gegl_cl_color_babl (format, &i->op_cl_format_size [self]);
if (self!=0)
{
......
......@@ -234,4 +234,28 @@ static const char* kernel_color_source =
" linear_to_gamma_2_2(rgb.z), \n"
" in_v.w); \n"
" out[gid] = convert_uchar4_sat_rte(255.0f * out_v); \n"
"} \n"
/* -- RGB u8 -- */
/* RGB u8 -> RGBA float */
"__kernel void rgbu8_to_rgbaf (__global const uchar3 * in, \n"
" __global float4 * out) \n"
"{ \n"
" int gid = get_global_id(0); \n"
" float3 in_v = convert_float3(in[gid]) / 255.0f; \n"
" float4 out_v; \n"
" out_v.xyz = in_v; \n"
" out_v.w = 1.0f; \n"
" out[gid] = out_v; \n"
"} \n"
/* RGBA float -> RGB u8 */
"__kernel void rgbaf_to_rgbu8 (__global const float4 * in, \n"
" __global uchar3 * out) \n"
"{ \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
" float4 out_v = in_v; \n"
" out[gid] = convert_uchar3_sat_rte(255.0f * out_v.w * out_v.xyz); \n"
"} \n";
......@@ -6,7 +6,7 @@
static gegl_cl_run_data *kernels_color = NULL;
#define CL_FORMAT_N 5
#define CL_FORMAT_N 6
static const Babl *format[CL_FORMAT_N];
......@@ -29,6 +29,9 @@ CL_RGBAF_TO_YCBCRAF = 10,
CL_YCBCRAF_TO_RGBAF = 11,
CL_RGBAU8_TO_YCBCRAF = 12,
CL_YCBCRAF_TO_RGBAU8 = 13,
CL_RGBU8_TO_RGBAF = 14,
CL_RGBAF_TO_RGBU8 = 15,
};
void
......@@ -52,6 +55,9 @@ gegl_cl_color_compile_kernels(void)
"rgbau8_to_ycbcraf", /* 12 */
"ycbcraf_to_rgbau8", /* 13 */
"rgbu8_to_rgbaf", /* 14 */
"rgbaf_to_rgbu8", /* 15 */
NULL};
format[0] = babl_format ("RGBA u8"),
......@@ -59,6 +65,7 @@ gegl_cl_color_compile_kernels(void)
format[2] = babl_format ("RaGaBaA float"),
format[3] = babl_format ("R'G'B'A float"),
format[4] = babl_format ("Y'CbCrA float"),
format[5] = babl_format ("RGB u8"),
kernels_color = gegl_cl_compile_and_build (kernel_color_source, kernel_name);
}
......@@ -75,6 +82,7 @@ choose_kernel (const Babl *in_format, const Babl *out_format)
else if (out_format == babl_format ("RaGaBaA float")) kernel = CL_RGBAF_TO_RAGABAF;
else if (out_format == babl_format ("R'G'B'A float")) kernel = CL_RGBAF_TO_RGBA_GAMMA_F;
else if (out_format == babl_format ("Y'CbCrA float")) kernel = CL_RGBAF_TO_YCBCRAF;
else if (out_format == babl_format ("RGB u8")) kernel = CL_RGBAF_TO_RGBU8;
}
else if (in_format == babl_format ("RGBA u8"))
{
......@@ -98,12 +106,16 @@ choose_kernel (const Babl *in_format, const Babl *out_format)
if (out_format == babl_format ("RGBA float")) kernel = CL_YCBCRAF_TO_RGBAF;
else if (out_format == babl_format ("RGBA u8")) kernel = CL_YCBCRAF_TO_RGBAU8;
}
else if (in_format == babl_format ("RGB u8"))
{
if (out_format == babl_format ("RGBA float")) kernel = CL_RGBU8_TO_RGBAF;
}
return kernel;
}
gboolean
gegl_cl_color_babl (const Babl *buffer_format, cl_image_format *cl_format, size_t *bytes)
gegl_cl_color_babl (const Babl *buffer_format, size_t *bytes)
{
int i;
gboolean supported_format = FALSE;
......@@ -114,23 +126,16 @@ gegl_cl_color_babl (const Babl *buffer_format, cl_image_format *cl_format, size_
if (!supported_format)
return FALSE;
if (cl_format)
if (bytes)
{
if (buffer_format == babl_format ("RGBA u8"))
{
cl_format->image_channel_order = CL_RGBA;
cl_format->image_channel_data_type = CL_UNORM_INT8;
}
*bytes = sizeof (cl_uchar4);
else if (buffer_format == babl_format ("RGB u8"))
*bytes = sizeof (cl_uchar3);
else
{
cl_format->image_channel_order = CL_RGBA;
cl_format->image_channel_data_type = CL_FLOAT;
}
*bytes = sizeof (cl_float4);
}
if (bytes)
*bytes = (buffer_format == babl_format ("RGBA u8"))? sizeof (cl_uchar4) : sizeof (cl_float4);
return TRUE;
}
......@@ -160,7 +165,7 @@ gegl_cl_color_conv (cl_mem in_tex, cl_mem out_tex, const size_t size,
if (in_format == out_format)
{
size_t s;
gegl_cl_color_babl (in_format, NULL, &s);
gegl_cl_color_babl (in_format, &s);
/* just copy in_tex to out_tex */
errcode = gegl_clEnqueueCopyBuffer (gegl_cl_get_command_queue(),
......
......@@ -13,7 +13,7 @@ typedef enum
void gegl_cl_color_compile_kernels(void);
gboolean gegl_cl_color_babl (const Babl *buffer_format, cl_image_format *cl_format, size_t *bytes);
gboolean gegl_cl_color_babl (const Babl *buffer_format, size_t *bytes);
gegl_cl_color_op gegl_cl_color_supported (const Babl *in_format, const Babl *out_format);
......
......@@ -88,8 +88,8 @@ gegl_operation_point_filter_cl_process (GeglOperation *operation,
gboolean err;
/* non-texturizable format! */
if (!gegl_cl_color_babl (in_format, NULL, NULL) ||
!gegl_cl_color_babl (out_format, NULL, NULL))
if (!gegl_cl_color_babl (in_format, NULL) ||
!gegl_cl_color_babl (out_format, NULL))
{
g_warning ("[OpenCL] Non-texturizable input of output format!");
return FALSE;
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment