David,
attached is a patch that tries another workaround for NVIDIA's OpenCL
bug. Maybe you can give it a try.
Ulrich
Am 02.12.2013 20:26, schrieb David Vincent-Jones:
After installing the latest Nvidia drivers my machine appears to go a
long way towards fully implementing the compiles needed for OpenCL but
then fails trying to compile 'blendop.cl'.
I have attached the terminal info and would appreciate if somebody can
tell me if this is a simple adjustment that I need to make.
David
diff --git a/data/kernels/blendop.cl b/data/kernels/blendop.cl
index da43f8f..f69502a 100644
--- a/data/kernels/blendop.cl
+++ b/data/kernels/blendop.cl
@@ -262,7 +262,7 @@ blendif_factor_rgb(const float4 input, const float4 output, const unsigned int b
}
__kernel void
-blendop_mask_Lab (__read_only image2d_t in_a, __read_only image2d_t in_b, __read_only image2d_t mask_in, __write_only image2d_t mask, const int width, const int height,
+blendop_mask_Lab (__read_only image2d_t mask_in, __write_only image2d_t mask, __read_only image2d_t in_a, __read_only image2d_t in_b, const int width, const int height,
const float gopacity, const int blendif, global const float *blendif_parameters, const unsigned int mask_mode, const unsigned int mask_combine, const int2 offs)
{
const int x = get_global_id(0);
@@ -284,7 +284,7 @@ blendop_mask_Lab (__read_only image2d_t in_a, __read_only image2d_t in_b, __read
#if 0
__kernel void
-blendop_mask_RAW (__read_only image2d_t in_a, __read_only image2d_t in_b, __read_only image2d_t mask_in, __write_only image2d_t mask, const int width, const int height,
+blendop_mask_RAW (__read_only image2d_t mask_in, __write_only image2d_t mask, __read_only image2d_t in_a, __read_only image2d_t in_b, const int width, const int height,
const float gopacity, const int blendif, global const float *blendif_parameters, const unsigned int mask_mode, const unsigned int mask_combine, const int2 offs)
{
const int x = get_global_id(0);
@@ -304,7 +304,7 @@ blendop_mask_RAW (__read_only image2d_t in_a, __read_only image2d_t in_b, __read
// thanks to Jens Fendler for finding this workaround.
// TODO: review after some time (May 2013) if this is still needed.
__kernel void
-blendop_mask_RAW (__read_only image2d_t in_a, __read_only image2d_t in_b, __read_only image2d_t mask_in, __write_only image2d_t mask, const int width, const int height,
+blendop_mask_RAW (__read_only image2d_t mask_in, __write_only image2d_t mask, __read_only image2d_t in_a, __read_only image2d_t in_b, const int width, const int height,
const float gopacity, const int blendif, global const float *blendif_parameters, const unsigned int mask_mode, const unsigned int mask_combine, const int2 offs)
{
const int x = get_global_id(0);
@@ -326,7 +326,7 @@ blendop_mask_RAW (__read_only image2d_t in_a, __read_only image2d_t in_b, __read
#endif
__kernel void
-blendop_mask_rgb (__read_only image2d_t in_a, __read_only image2d_t in_b, __read_only image2d_t mask_in, __write_only image2d_t mask, const int width, const int height,
+blendop_mask_rgb (__read_only image2d_t mask_in, __write_only image2d_t mask, __read_only image2d_t in_a, __read_only image2d_t in_b, const int width, const int height,
const float gopacity, const int blendif, global const float *blendif_parameters, const unsigned int mask_mode, const unsigned int mask_combine, const int2 offs)
{
const int x = get_global_id(0);
diff --git a/src/develop/blend.c b/src/develop/blend.c
index 424d4e0..a0680f6 100644
--- a/src/develop/blend.c
+++ b/src/develop/blend.c
@@ -2324,10 +2324,10 @@ dt_develop_blend_process_cl (struct dt_iop_module_t *self, struct dt_dev_pixelpi
dt_opencl_finish(devid);
size_t sizes[] = { ROUNDUPWD(width), ROUNDUPHT(height), 1 };
- dt_opencl_set_kernel_arg(devid, kernel_mask, 0, sizeof(cl_mem), (void *)&dev_in);
- dt_opencl_set_kernel_arg(devid, kernel_mask, 1, sizeof(cl_mem), (void *)&dev_out);
- dt_opencl_set_kernel_arg(devid, kernel_mask, 2, sizeof(cl_mem), (void *)&dev_mask);
- dt_opencl_set_kernel_arg(devid, kernel_mask, 3, sizeof(cl_mem), (void *)&dev_mask);
+ dt_opencl_set_kernel_arg(devid, kernel_mask, 0, sizeof(cl_mem), (void *)&dev_mask);
+ dt_opencl_set_kernel_arg(devid, kernel_mask, 1, sizeof(cl_mem), (void *)&dev_mask);
+ dt_opencl_set_kernel_arg(devid, kernel_mask, 2, sizeof(cl_mem), (void *)&dev_in);
+ dt_opencl_set_kernel_arg(devid, kernel_mask, 3, sizeof(cl_mem), (void *)&dev_out);
dt_opencl_set_kernel_arg(devid, kernel_mask, 4, sizeof(int), (void *)&width);
dt_opencl_set_kernel_arg(devid, kernel_mask, 5, sizeof(int), (void *)&height);
dt_opencl_set_kernel_arg(devid, kernel_mask, 6, sizeof(float), (void *)&opacity);
------------------------------------------------------------------------------
Sponsored by Intel(R) XDK
Develop, test and display web and hybrid apps with a single code base.
Download it for free now!
http://pubads.g.doubleclick.net/gampad/clk?id=111408631&iu=/4140/ostg.clktrk
_______________________________________________
darktable-devel mailing list
darktable-devel@lists.sourceforge.net
https://lists.sourceforge.net/lists/listinfo/darktable-devel
