Commit 4ac1e6a0 authored by Victor Oliveira's avatar Victor Oliveira

cl: support for R'G'B' u8 and R'G'B'A u8

parent a5c7770e
......@@ -337,4 +337,69 @@ static const char* kernel_color_source =
" float4 out_v = (float4) (in_v.x, in_v.x, in_v.x, in_v.y); \n"
" \n"
" out[gid] = convert_uchar4_sat_rte(255.0f * out_v); \n"
"} \n"
/* R'G'B'A u8 */
/* rgba float -> r'g'b'a u8 */
"__kernel void rgbaf_to_rgba_gamma_u8 (__global const float4 * in, \n"
" __global uchar4 * out) \n"
"{ \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
" float4 out_v; \n"
" out_v = (float4)(linear_to_gamma_2_2(in_v.x), \n"
" linear_to_gamma_2_2(in_v.y), \n"
" linear_to_gamma_2_2(in_v.z), \n"
" in_v.w); \n"
" out[gid] = convert_uchar4_sat_rte(255.0f * out_v); \n"
"} \n"
/* r'g'b'a u8 -> rgba float */
"__kernel void rgba_gamma_u8_to_rgbaf (__global const uchar4 * in, \n"
" __global float4 * out) \n"
"{ \n"
" int gid = get_global_id(0); \n"
" float4 in_v = convert_float4(in[gid]) / 255.0f; \n"
" float4 out_v; \n"
" out_v = (float4)(gamma_2_2_to_linear(in_v.x), \n"
" gamma_2_2_to_linear(in_v.y), \n"
" gamma_2_2_to_linear(in_v.z), \n"
" in_v.w); \n"
" out[gid] = out_v; \n"
"} \n"
/* R'G'B' u8 */
/* rgba float -> r'g'b' u8 */
"__kernel void rgbaf_to_rgb_gamma_u8 (__global const float4 * in, \n"
" __global uchar * out) \n"
"{ \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
" float4 tmp_v; \n"
" uchar3 out_v; \n"
" tmp_v = (float4)(linear_to_gamma_2_2(in_v.x), \n"
" linear_to_gamma_2_2(in_v.y), \n"
" linear_to_gamma_2_2(in_v.z), \n"
" in_v.w); \n"
" out_v = convert_uchar3_sat_rte(255.0f * tmp_v.w * tmp_v.xyz); \n"
" vstore3 (out_v, gid, out); \n"
"} \n"
/* r'g'b' u8 -> rgba float */
"__kernel void rgb_gamma_u8_to_rgbaf (__global const uchar * in, \n"
" __global float4 * out) \n"
"{ \n"
" int gid = get_global_id(0); \n"
" uchar3 in_v; \n"
" float3 tmp_v; \n"
" float4 out_v; \n"
" in_v = vload3 (gid, in); \n"
" tmp_v = convert_float3(in_v) / 255.0f; \n"
" out_v = (float4)(gamma_2_2_to_linear(tmp_v.x), \n"
" gamma_2_2_to_linear(tmp_v.y), \n"
" gamma_2_2_to_linear(tmp_v.z), \n"
" 1.0f); \n"
" out[gid] = out_v; \n"
"} \n";
......@@ -8,7 +8,7 @@
static gegl_cl_run_data *kernels_color = NULL;
#define CL_FORMAT_N 9
#define CL_FORMAT_N 11
static const Babl *format[CL_FORMAT_N];
......@@ -41,6 +41,12 @@ CL_RGBAF_TO_YAF = 17,
CL_YAF_TO_RGBAF = 18,
CL_RGBAU8_TO_YAF = 19,
CL_YAF_TO_RGBAU8 = 20,
CL_RGBAF_TO_RGBA_GAMMA_U8 = 21,
CL_RGBA_GAMMA_U8_TO_RGBAF = 22,
CL_RGBAF_TO_RGB_GAMMA_U8 = 23,
CL_RGB_GAMMA_U8_TO_RGBAF = 24,
};
void
......@@ -74,6 +80,12 @@ gegl_cl_color_compile_kernels(void)
"rgbau8_to_yaf", /* 19 */
"yaf_to_rgbau8", /* 20 */
"rgbaf_to_rgba_gamma_u8", /* 21 */
"rgba_gamma_u8_to_rgbaf", /* 22 */
"rgbaf_to_rgb_gamma_u8", /* 23 */
"rgb_gamma_u8_to_rgbaf", /* 24 */
NULL};
format[0] = babl_format ("RGBA u8");
......@@ -85,6 +97,8 @@ gegl_cl_color_compile_kernels(void)
format[6] = babl_format ("Y u8");
format[7] = babl_format ("Y float");
format[8] = babl_format ("YA float");
format[9] = babl_format ("R'G'B'A u8");
format[10] = babl_format ("R'G'B' u8");
kernels_color = gegl_cl_compile_and_build (kernel_color_source, kernel_name);
}
......@@ -99,6 +113,8 @@ choose_kernel (const Babl *in_format, const Babl *out_format)
{
if (out_format == babl_format ("RGBA u8")) kernel = CL_RGBAF_TO_RGBAU8;
else if (out_format == babl_format ("RaGaBaA float")) kernel = CL_RGBAF_TO_RAGABAF;
else if (out_format == babl_format ("R'G'B' u8")) kernel = CL_RGBAF_TO_RGB_GAMMA_U8;
else if (out_format == babl_format ("R'G'B'A u8")) kernel = CL_RGBAF_TO_RGBA_GAMMA_U8;
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;
......@@ -140,6 +156,14 @@ choose_kernel (const Babl *in_format, const Babl *out_format)
if (out_format == babl_format ("RGBA float")) kernel = CL_YAF_TO_RGBAF;
else if (out_format == babl_format ("RGBA u8")) kernel = CL_YAF_TO_RGBAU8;
}
else if (in_format == babl_format ("R'G'B'A u8"))
{
if (out_format == babl_format ("RGBA float")) kernel = CL_RGBA_GAMMA_U8_TO_RGBAF;
}
else if (in_format == babl_format ("R'G'B' u8"))
{
if (out_format == babl_format ("RGBA float")) kernel = CL_RGB_GAMMA_U8_TO_RGBAF;
}
return kernel;
}
......@@ -171,6 +195,10 @@ gegl_cl_color_babl (const Babl *buffer_format, size_t *bytes)
*bytes = sizeof (cl_float);
else if (buffer_format == babl_format ("YA float"))
*bytes = sizeof (cl_float2);
else if (buffer_format == babl_format("R'G'B'A u8"))
*bytes = sizeof (cl_uchar4);
else if (buffer_format == babl_format("R'G'B' u8"))
*bytes = sizeof (cl_uchar3);
else
*bytes = sizeof (cl_float4);
}
......
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