Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Enable image and buffer data interoperability #57

Closed
StRigaud opened this issue Oct 26, 2021 · 5 comments · Fixed by #65
Closed

Enable image and buffer data interoperability #57

StRigaud opened this issue Oct 26, 2021 · 5 comments · Fixed by #65
Assignees
@StRigaud @haesleinhuepf
Labels
request New feature or request WIP Work In Progress

Comments

@StRigaud
Copy link
Member

StRigaud commented Oct 26, 2021 •
edited
Loading

Current preamble does not support cl_image1d even though the function read_image1d and write_image1d are provided by OpenCL.
I mainly suspect an issue with the coordinate system of CLIJ that is always defined as a 2d coord (x,y) with y=0 in case of 1d.

it does, however, seems to support cl_image2d with 1d shape. Current fix would be to only implement cl_image3d and cl_image2d. And manage 1d and 2d with cl_image2d.
Though I assume a possible efficiency or results impact when applying very specific image operation? Need to be tested.

EDIT: Only implementing cl_image3d for managing all dimension compile and execute but return wrong output for 1d data.

Sources and log for information purposes (@haesleinhuepf if you wanna decrypt this).

#defines

#define GET_IMAGE_WIDTH(image_key) IMAGE_SIZE_ ## image_key ## _WIDTH
#define GET_IMAGE_HEIGHT(image_key) IMAGE_SIZE_ ## image_key ## _HEIGHT
#define GET_IMAGE_DEPTH(image_key) IMAGE_SIZE_ ## image_key ## _DEPTH

#define CONVERT_dst_PIXEL_TYPE clij_convert_float_sat
#define IMAGE_dst_PIXEL_TYPE float
#define POS_dst_TYPE int2
#define POS_dst_INSTANCE(pos0,pos1,pos2,pos3) (int2)(pos0, 0)
#define IMAGE_dst_TYPE  __write_only image1d_t
#define READ_dst_IMAGE(a,b,c) read_imagef(a,b,c)
#define WRITE_dst_IMAGE(a,b,c) write_imagef(a,b,c)
#define IMAGE_SIZE_dst_WIDTH 10
#define IMAGE_SIZE_dst_HEIGHT 1
#define IMAGE_SIZE_dst_DEPTH 1

#define CONVERT_src_PIXEL_TYPE clij_convert_float_sat
#define IMAGE_src_PIXEL_TYPE float
#define POS_src_TYPE int2
#define POS_src_INSTANCE(pos0,pos1,pos2,pos3) (int2)(pos0, 0)
#define IMAGE_src_TYPE  __read_only image1d_t
#define READ_src_IMAGE(a,b,c) read_imagef(a,b,c)
#define WRITE_src_IMAGE(a,b,c) write_imagef(a,b,c)
#define IMAGE_SIZE_src_WIDTH 10
#define IMAGE_SIZE_src_HEIGHT 1
#define IMAGE_SIZE_src_DEPTH 1

error log

<kernel>:553:61: error: no matching function for call to 'read_imagef'
  const IMAGE_dst_PIXEL_TYPE value = CONVERT_dst_PIXEL_TYPE(READ_src_IMAGE(src, sampler, pos).x + scalar);
                                                            ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
<kernel>:27:31: note: expanded from macro 'READ_src_IMAGE'
#define READ_src_IMAGE(a,b,c) read_imagef(a,b,c)
                              ^~~~~~~~~~~
cl_kernel.h:18583:25: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_t' for 1st argument
float4 __OVERLOADABLE__ read_imagef(image2d_t image, sampler_t sampler, int2 coord);
                        ^
cl_kernel.h:18601:25: note: candidate function not viable: no known conversion from 'const __attribute__((address_space(16776963))) int2' to '__attribute__((address_space(16776963))) int' for 3rd argument
float4 __OVERLOADABLE__ read_imagef(image1d_t image, sampler_t sampler, int coord);
                        ^
cl_kernel.h:18602:25: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image1d_array_t' for 1st argument
float4 __OVERLOADABLE__ read_imagef(image1d_array_t image, sampler_t sampler, int2 coord);
                        ^
cl_kernel.h:18603:25: note: candidate function not viable: no known conversion from 'const __attribute__((address_space(16776963))) int2' to '__attribute__((address_space(16776963))) float' for 3rd argument
float4 __OVERLOADABLE__ read_imagef(image1d_t image, sampler_t sampler, float coord);
                        ^
cl_kernel.h:18626:24: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_depth_t' for 1st argument
float __OVERLOADABLE__ read_imagef(image2d_depth_t image, sampler_t sampler, int2 coord);
                       ^
cl_kernel.h:18584:25: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_t' for 1st argument
float4 __OVERLOADABLE__ read_imagef(image2d_t image, sampler_t sampler, float2 coord);
                        ^
cl_kernel.h:18585:25: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image3d_t' for 1st argument
float4 __OVERLOADABLE__ read_imagef(image3d_t image, sampler_t sampler, int4 coord);
                        ^
cl_kernel.h:18586:25: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image3d_t' for 1st argument
float4 __OVERLOADABLE__ read_imagef(image3d_t image, sampler_t sampler, float4 coord);
                        ^
cl_kernel.h:18604:25: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image1d_array_t' for 1st argument
float4 __OVERLOADABLE__ read_imagef(image1d_array_t image, sampler_t sampler, float2 coord);
                        ^
cl_kernel.h:18605:25: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_array_t' for 1st argument
float4 __OVERLOADABLE__ read_imagef(image2d_array_t image, sampler_t sampler, int4 coord);
                        ^
cl_kernel.h:18606:25: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_array_t' for 1st argument
float4 __OVERLOADABLE__ read_imagef(image2d_array_t image, sampler_t sampler, float4 coord);
                        ^
cl_kernel.h:18627:24: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_depth_t' for 1st argument
float __OVERLOADABLE__ read_imagef(image2d_depth_t image, sampler_t sampler, float2 coord);
                       ^
cl_kernel.h:18628:24: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_array_depth_t' for 1st argument
float __OVERLOADABLE__ read_imagef(image2d_array_depth_t image, sampler_t sampler, int4 coord);
                       ^
cl_kernel.h:18629:24: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_array_depth_t' for 1st argument
float __OVERLOADABLE__ read_imagef(image2d_array_depth_t image, sampler_t sampler, float4 coord);
                       ^
cl_kernel.h:18646:25: note: candidate function not viable: requires 2 arguments, but 3 were provided
float4 __OVERLOADABLE__ read_imagef(image2d_t image, int2 coord);
                        ^
cl_kernel.h:18636:25: note: candidate function not viable: requires 2 arguments, but 3 were provided
float4 __OVERLOADABLE__ read_imagef(image1d_t image, int coord);
                        ^
cl_kernel.h:18637:25: note: candidate function not viable: requires 2 arguments, but 3 were provided
float4 __OVERLOADABLE__ read_imagef(image1d_array_t image, int2 coord);
                        ^
cl_kernel.h:18638:25: note: candidate function not viable: requires 2 arguments, but 3 were provided
float4 __OVERLOADABLE__ read_imagef(image1d_buffer_t image, int coord);
                        ^
cl_kernel.h:18696:1: note: candidate function not viable: requires 4 arguments, but 3 were provided
read_imagef(image1d_t image, sampler_t sampler, float coord, float lod);
^
cl_kernel.h:18779:1: note: candidate function not viable: requires 5 arguments, but 3 were provided
read_imagef(image2d_array_t image, sampler_t sampler, float4 coord,
^
cl_kernel.h:18775:1: note: candidate function not viable: requires 4 arguments, but 3 were provided
read_imagef(image2d_array_depth_t image, sampler_t sampler, float4 coord,
^
cl_kernel.h:18763:1: note: candidate function not viable: requires 4 arguments, but 3 were provided
read_imagef(image2d_array_t image, sampler_t sampler, float4 coord,
^
cl_kernel.h:18751:1: note: candidate function not viable: requires 5 arguments, but 3 were provided
read_imagef(image1d_array_t image, sampler_t sampler, float2 coord,
^
cl_kernel.h:18739:1: note: candidate function not viable: requires 4 arguments, but 3 were provided
read_imagef(image1d_array_t image, sampler_t sampler, float2 coord,
^
cl_kernel.h:18727:1: note: candidate function not viable: requires 5 arguments, but 3 were provided
read_imagef(image3d_t image, sampler_t sampler, float4 coord,
^
cl_kernel.h:18718:1: note: candidate function not viable: requires 4 arguments, but 3 were provided
read_imagef(image3d_t image, sampler_t sampler, float4 coord, float lod);
^
cl_kernel.h:18706:1: note: candidate function not viable: requires 5 arguments, but 3 were provided
read_imagef(image1d_t image, sampler_t sampler, float coord,
^
cl_kernel.h:18791:1: note: candidate function not viable: requires 5 arguments, but 3 were provided
read_imagef(image2d_array_depth_t image, sampler_t sampler, float4 coord,
^
cl_kernel.h:18692:1: note: candidate function not viable: requires 5 arguments, but 3 were provided
read_imagef(image2d_depth_t image, sampler_t sampler, float2 coord,
^
cl_kernel.h:18680:1: note: candidate function not viable: requires 5 arguments, but 3 were provided
read_imagef(image2d_t image, sampler_t sampler, float2 coord,
^
cl_kernel.h:18677:1: note: candidate function not viable: requires 4 arguments, but 3 were provided
read_imagef(image2d_depth_t image, sampler_t sampler, float2 coord, float lod);
^
cl_kernel.h:18668:1: note: candidate function not viable: requires 4 arguments, but 3 were provided
read_imagef(image2d_t image, sampler_t sampler, float2 coord, float lod);
^
cl_kernel.h:18650:25: note: candidate function not viable: requires 2 arguments, but 3 were provided
float4 __OVERLOADABLE__ read_imagef(image3d_t image, int4 coord);
                        ^
cl_kernel.h:18649:24: note: candidate function not viable: requires 2 arguments, but 3 were provided
float __OVERLOADABLE__ read_imagef(image2d_depth_t image, int2 coord);
                       ^
cl_kernel.h:18648:24: note: candidate function not viable: requires 2 arguments, but 3 were provided
float __OVERLOADABLE__ read_imagef(image2d_array_depth_t image, int4 coord);
                       ^
cl_kernel.h:18647:25: note: candidate function not viable: requires 2 arguments, but 3 were provided
float4 __OVERLOADABLE__ read_imagef(image2d_array_t image, int4 coord);
                        ^
<kernel>:555:3: error: no matching function for call to 'write_imagef'
  WRITE_dst_IMAGE (dst, pos, value);
  ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
<kernel>:14:32: note: expanded from macro 'WRITE_dst_IMAGE'
#define WRITE_dst_IMAGE(a,b,c) write_imagef(a,b,c)
                               ^~~~~~~~~~~~
cl_kernel.h:18818:23: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_t' for 1st argument
void __OVERLOADABLE__ write_imagef(image2d_t image, int2 coord, float4 val);
                      ^
cl_kernel.h:18825:23: note: candidate function not viable: no known conversion from 'const __attribute__((address_space(16776963))) int2' to '__attribute__((address_space(16776963))) int' for 2nd argument
void __OVERLOADABLE__ write_imagef(image1d_t, int coord, float4 color);
                      ^
cl_kernel.h:18828:23: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image1d_array_t' for 1st argument
void __OVERLOADABLE__ write_imagef(image1d_array_t, int2 coord, float4 color);
                      ^
cl_kernel.h:18846:23: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_array_depth_t' for 1st argument
void __OVERLOADABLE__ write_imagef(image2d_array_depth_t image, int4 coord, float depth);
                      ^
cl_kernel.h:18831:23: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image1d_buffer_t' for 1st argument
void __OVERLOADABLE__ write_imagef(image1d_buffer_t, int coord, float4 color);
                      ^
cl_kernel.h:18835:23: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_array_t' for 1st argument
void __OVERLOADABLE__ write_imagef(image2d_array_t image, int4 coord, float4 color);
                      ^
cl_kernel.h:18839:23: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image3d_t' for 1st argument
void __OVERLOADABLE__ write_imagef(image3d_t image, int4 coord, float4 color);
@StRigaud
Copy link
Member Author

StRigaud commented Nov 5, 2021 •
edited
Loading

After a bit of investigation I think the issues is related to the coordinate coding which in the current defines is
int2 with position as (pos0, 0) in 1D
int2 with position as (pos0, pos1) in 2D
int4 with position as (pos0, pos1, pos2, 0) in 3D

cl_image3d works with int4 with position as (pos0, pos1, pos2, 0)
cl_image2d works with int2 with position as (pos0, pos1)
but
cl_image1d requires int with position as (pos0)

The last case is not treated because for buffer we only use the int2 or int4 options and never int. A special case for cl_image1d to use int with position as (pos0) could solve it.

To be tested (and hoping it doesn't need more ...).

EDIT: apparently it's not that simple :/

#define CONVERT_dst_PIXEL_TYPE clij_convert_float_sat
#define IMAGE_dst_PIXEL_TYPE float
#define POS_dst_TYPE int
#define POS_dst_INSTANCE(pos0,pos1,pos2,pos3) (int)(pos0)
#define IMAGE_dst_TYPE __write_only image1d_t
#define READ_dst_IMAGE(a,b,c) read_imagef(a,b,c)
#define WRITE_dst_IMAGE(a,b,c) write_imagef(a,b,c)
#define IMAGE_SIZE_dst_WIDTH 10
#define IMAGE_SIZE_dst_HEIGHT 1
#define IMAGE_SIZE_dst_DEPTH 1

#define CONVERT_src_PIXEL_TYPE clij_convert_float_sat
#define IMAGE_src_PIXEL_TYPE float
#define POS_src_TYPE int
#define POS_src_INSTANCE(pos0,pos1,pos2,pos3) (int)(pos0)
#define IMAGE_src_TYPE __read_only image1d_t
#define READ_src_IMAGE(a,b,c) read_imagef(a,b,c)
#define WRITE_src_IMAGE(a,b,c) write_imagef(a,b,c)
#define IMAGE_SIZE_src_WIDTH 10
#define IMAGE_SIZE_src_HEIGHT 1
#define IMAGE_SIZE_src_DEPTH 1

errors:

Kernel : Fail to build program "add_image_and_scalar" from source.
	Exception caught! clBuildProgram error code -11
build log:
<kernel>:36:26: warning: unknown OpenCL extension 'cl_amd_printf' - ignoring
#pragma OPENCL EXTENSION cl_amd_printf : enable
                         ^
<kernel>:553:61: error: no matching function for call to 'read_imagef'
  const IMAGE_dst_PIXEL_TYPE value = CONVERT_dst_PIXEL_TYPE(READ_src_IMAGE(src, sampler, pos).x + scalar);
                                                            ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
<kernel>:27:31: note: expanded from macro 'READ_src_IMAGE'
#define READ_src_IMAGE(a,b,c) read_imagef(a,b,c)
                              ^~~~~~~~~~~
cl_kernel.h:18583:25: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_t' for 1st argument
float4 __OVERLOADABLE__ read_imagef(image2d_t image, sampler_t sampler, int2 coord);
                        ^
cl_kernel.h:18601:25: note: candidate function not viable: no known conversion from 'const __attribute__((address_space(16776963))) int2' to '__attribute__((address_space(16776963))) int' for 3rd argument
float4 __OVERLOADABLE__ read_imagef(image1d_t image, sampler_t sampler, int coord);
                        ^
cl_kernel.h:18602:25: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image1d_array_t' for 1st argument
float4 __OVERLOADABLE__ read_imagef(image1d_array_t image, sampler_t sampler, int2 coord);
                        ^
cl_kernel.h:18603:25: note: candidate function not viable: no known conversion from 'const __attribute__((address_space(16776963))) int2' to '__attribute__((address_space(16776963))) float' for 3rd argument
float4 __OVERLOADABLE__ read_imagef(image1d_t image, sampler_t sampler, float coord);
                        ^
cl_kernel.h:18626:24: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_depth_t' for 1st argument
float __OVERLOADABLE__ read_imagef(image2d_depth_t image, sampler_t sampler, int2 coord);
                       ^
cl_kernel.h:18584:25: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_t' for 1st argument
float4 __OVERLOADABLE__ read_imagef(image2d_t image, sampler_t sampler, float2 coord);
                        ^
cl_kernel.h:18585:25: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image3d_t' for 1st argument
float4 __OVERLOADABLE__ read_imagef(image3d_t image, sampler_t sampler, int4 coord);
                        ^
cl_kernel.h:18586:25: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image3d_t' for 1st argument
float4 __OVERLOADABLE__ read_imagef(image3d_t image, sampler_t sampler, float4 coord);
                        ^
cl_kernel.h:18604:25: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image1d_array_t' for 1st argument
float4 __OVERLOADABLE__ read_imagef(image1d_array_t image, sampler_t sampler, float2 coord);
                        ^
cl_kernel.h:18605:25: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_array_t' for 1st argument
float4 __OVERLOADABLE__ read_imagef(image2d_array_t image, sampler_t sampler, int4 coord);
                        ^
cl_kernel.h:18606:25: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_array_t' for 1st argument
float4 __OVERLOADABLE__ read_imagef(image2d_array_t image, sampler_t sampler, float4 coord);
                        ^
cl_kernel.h:18627:24: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_depth_t' for 1st argument
float __OVERLOADABLE__ read_imagef(image2d_depth_t image, sampler_t sampler, float2 coord);
                       ^
cl_kernel.h:18628:24: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_array_depth_t' for 1st argument
float __OVERLOADABLE__ read_imagef(image2d_array_depth_t image, sampler_t sampler, int4 coord);
                       ^
cl_kernel.h:18629:24: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_array_depth_t' for 1st argument
float __OVERLOADABLE__ read_imagef(image2d_array_depth_t image, sampler_t sampler, float4 coord);
                       ^
cl_kernel.h:18646:25: note: candidate function not viable: requires 2 arguments, but 3 were provided
float4 __OVERLOADABLE__ read_imagef(image2d_t image, int2 coord);
                        ^
cl_kernel.h:18636:25: note: candidate function not viable: requires 2 arguments, but 3 were provided
float4 __OVERLOADABLE__ read_imagef(image1d_t image, int coord);
                        ^
cl_kernel.h:18637:25: note: candidate function not viable: requires 2 arguments, but 3 were provided
float4 __OVERLOADABLE__ read_imagef(image1d_array_t image, int2 coord);
                        ^
cl_kernel.h:18638:25: note: candidate function not viable: requires 2 arguments, but 3 were provided
float4 __OVERLOADABLE__ read_imagef(image1d_buffer_t image, int coord);
                        ^
cl_kernel.h:18696:1: note: candidate function not viable: requires 4 arguments, but 3 were provided
read_imagef(image1d_t image, sampler_t sampler, float coord, float lod);
^
cl_kernel.h:18779:1: note: candidate function not viable: requires 5 arguments, but 3 were provided
read_imagef(image2d_array_t image, sampler_t sampler, float4 coord,
^
cl_kernel.h:18775:1: note: candidate function not viable: requires 4 arguments, but 3 were provided
read_imagef(image2d_array_depth_t image, sampler_t sampler, float4 coord,
^
cl_kernel.h:18763:1: note: candidate function not viable: requires 4 arguments, but 3 were provided
read_imagef(image2d_array_t image, sampler_t sampler, float4 coord,
^
cl_kernel.h:18751:1: note: candidate function not viable: requires 5 arguments, but 3 were provided
read_imagef(image1d_array_t image, sampler_t sampler, float2 coord,
^
cl_kernel.h:18739:1: note: candidate function not viable: requires 4 arguments, but 3 were provided
read_imagef(image1d_array_t image, sampler_t sampler, float2 coord,
^
cl_kernel.h:18727:1: note: candidate function not viable: requires 5 arguments, but 3 were provided
read_imagef(image3d_t image, sampler_t sampler, float4 coord,
^
cl_kernel.h:18718:1: note: candidate function not viable: requires 4 arguments, but 3 were provided
read_imagef(image3d_t image, sampler_t sampler, float4 coord, float lod);
^
cl_kernel.h:18706:1: note: candidate function not viable: requires 5 arguments, but 3 were provided
read_imagef(image1d_t image, sampler_t sampler, float coord,
^
cl_kernel.h:18791:1: note: candidate function not viable: requires 5 arguments, but 3 were provided
read_imagef(image2d_array_depth_t image, sampler_t sampler, float4 coord,
^
cl_kernel.h:18692:1: note: candidate function not viable: requires 5 arguments, but 3 were provided
read_imagef(image2d_depth_t image, sampler_t sampler, float2 coord,
^
cl_kernel.h:18680:1: note: candidate function not viable: requires 5 arguments, but 3 were provided
read_imagef(image2d_t image, sampler_t sampler, float2 coord,
^
cl_kernel.h:18677:1: note: candidate function not viable: requires 4 arguments, but 3 were provided
read_imagef(image2d_depth_t image, sampler_t sampler, float2 coord, float lod);
^
cl_kernel.h:18668:1: note: candidate function not viable: requires 4 arguments, but 3 were provided
read_imagef(image2d_t image, sampler_t sampler, float2 coord, float lod);
^
cl_kernel.h:18650:25: note: candidate function not viable: requires 2 arguments, but 3 were provided
float4 __OVERLOADABLE__ read_imagef(image3d_t image, int4 coord);
                        ^
cl_kernel.h:18649:24: note: candidate function not viable: requires 2 arguments, but 3 were provided
float __OVERLOADABLE__ read_imagef(image2d_depth_t image, int2 coord);
                       ^
cl_kernel.h:18648:24: note: candidate function not viable: requires 2 arguments, but 3 were provided
float __OVERLOADABLE__ read_imagef(image2d_array_depth_t image, int4 coord);
                       ^
cl_kernel.h:18647:25: note: candidate function not viable: requires 2 arguments, but 3 were provided
float4 __OVERLOADABLE__ read_imagef(image2d_array_t image, int4 coord);
                        ^
<kernel>:555:3: error: no matching function for call to 'write_imagef'
  WRITE_dst_IMAGE (dst, pos, value);
  ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
<kernel>:14:32: note: expanded from macro 'WRITE_dst_IMAGE'
#define WRITE_dst_IMAGE(a,b,c) write_imagef(a,b,c)
                               ^~~~~~~~~~~~
cl_kernel.h:18818:23: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_t' for 1st argument
void __OVERLOADABLE__ write_imagef(image2d_t image, int2 coord, float4 val);
                      ^
cl_kernel.h:18825:23: note: candidate function not viable: no known conversion from 'const __attribute__((address_space(16776963))) int2' to '__attribute__((address_space(16776963))) int' for 2nd argument
void __OVERLOADABLE__ write_imagef(image1d_t, int coord, float4 color);
                      ^
cl_kernel.h:18828:23: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image1d_array_t' for 1st argument
void __OVERLOADABLE__ write_imagef(image1d_array_t, int2 coord, float4 color);
                      ^
cl_kernel.h:18846:23: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_array_depth_t' for 1st argument
void __OVERLOADABLE__ write_imagef(image2d_array_depth_t image, int4 coord, float depth);
                      ^
cl_kernel.h:18831:23: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image1d_buffer_t' for 1st argument
void __OVERLOADABLE__ write_imagef(image1d_buffer_t, int coord, float4 color);
                      ^
cl_kernel.h:18835:23: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image2d_array_t' for 1st argument
void __OVERLOADABLE__ write_imagef(image2d_array_t image, int4 coord, float4 color);
                      ^
cl_kernel.h:18839:23: note: candidate function not viable: no known conversion from '__attribute__((address_space(16776963))) image1d_t' to '__attribute__((address_space(16776963))) image3d_t' for 1st argument
void __OVERLOADABLE__ write_imagef(image3d_t image, int4 coord, float4 color);

@StRigaud
Copy link
Member Author

StRigaud commented Nov 8, 2021 •
edited
Loading

So, to follow up on this.

It find a read function for image1d_t but fail to use it because of conversion between int and int2 even though we are using 'int':

cl_kernel.h:18601:25: note: candidate function not viable: no known conversion from 'const __attribute__((address_space(16776963))) int2' to '__attribute__((address_space(16776963))) int' for 3rd argument
float4 __OVERLOADABLE__ read_imagef(image1d_t image, sampler_t sampler, int coord);

This is due to the kernel _2d version containing this:

const int2 pos = (int2){x,y};

Same but with int4 would be found in _3d version.

Hence, we face an incompatibility issue from CLIJ kernel to image1d_t. To be discussed if there is a work around or not 🙃

EDIT:
Creating a _1d version of the kernel by replacing

const int2 pos = (int2){x,y};

with

const int pos = (int){x};

Solve the issue, kernel is executing properly with image1d_t

Still need to test if this solves the issue of running a kernel with dimension reduction (e.g. MaximumOfAllPIxelsKernel)

EDIT2: Spoiler alert

....

it works with MaximumOfAllPixelsKernel!

@StRigaud
Copy link
Member Author

StRigaud commented Nov 10, 2021 •
edited
Loading

So, to finish this thread, to allows image1d_t full compatibility and enable Image and Buffer interoperability in CLIc, we need to do the following modifications.

In generating #defines:

  • 3d - int4 (pos0, pos1, pos2, 0)
  • 2d - int2 (pos0, 0)
  • 1d - int (pos0)

In dimensional specific kernel (aka named with _2d or _3d at the end), create a _1d version:

  • 3d - add_image_and_scalar_3d_x.cl
const int x = get_global_id(0);
const int y = get_global_id(1);
const int z = get_global_id(2);
const int4 pos = (int4){x,y,z,0};
  • 2d - add_image_and_scalar_2d_x.cl
const int x = get_global_id(0);
const int y = get_global_id(1);
const int2 pos = (int4){x,y};
  • 1d - add_image_and_scalar_1d_x.cl
const int x = get_global_id(0);
const int pos = (int){0};

In preamble.cl, introduce 1d buffer read and write functions:

inline char2 read_buffer1dc(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global char * buffer_var, sampler_t sampler, int position )
{
    int pos = (int){position};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos = max((MINMAX_TYPE)pos, (MINMAX_TYPE)0);
        pos = min((MINMAX_TYPE)pos, (MINMAX_TYPE)read_buffer_width - 1);
    }
    int pos_in_buffer = pos;
    if (pos < 0 || pos >= read_buffer_width ) {
        return (char2){0,0};
    }
    return (char2){buffer_var[pos_in_buffer],0};
}

inline uchar2 read_buffer1duc(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global uchar * buffer_var, sampler_t sampler, int position )
{
    int pos = (int){position};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos = max((MINMAX_TYPE)pos, (MINMAX_TYPE)0);
        pos = min((MINMAX_TYPE)pos, (MINMAX_TYPE)read_buffer_width - 1);
    }
    int pos_in_buffer = pos;
    if (pos < 0 || pos >= read_buffer_width) {
        return (uchar2){0,0};
    }
    return (uchar2){buffer_var[pos_in_buffer],0};
}

inline short2 read_buffer1di(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global short * buffer_var, sampler_t sampler, int position )
{
    int pos = (int){position};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos = max((MINMAX_TYPE)pos, (MINMAX_TYPE)0);
        pos = min((MINMAX_TYPE)pos, (MINMAX_TYPE)read_buffer_width - 1);
    }
    int pos_in_buffer = pos;
    if (pos < 0 || pos >= read_buffer_width) {
        return (short2){0,0};
    }
    return (short2){buffer_var[pos_in_buffer],0};
}

inline ushort2 read_buffer1dui(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global ushort * buffer_var, sampler_t sampler, int position )
{
    int pos = (int){position};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos = max((MINMAX_TYPE)pos, (MINMAX_TYPE)0);
        pos = min((MINMAX_TYPE)pos, (MINMAX_TYPE)read_buffer_width - 1);
    }
    int pos_in_buffer = pos;
    if (pos < 0 || pos >= read_buffer_width) {
        return (ushort2){0,0};
    }
    return (ushort2){buffer_var[pos_in_buffer],0};
}

inline float2 read_buffer1df(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global float* buffer_var, sampler_t sampler, int position )
{
    int pos = (int){position};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos = max((MINMAX_TYPE)pos, (MINMAX_TYPE)0);
        pos = min((MINMAX_TYPE)pos, (MINMAX_TYPE)read_buffer_width - 1);
    }
    int pos_in_buffer = pos;
    if (pos < 0 || pos >= read_buffer_width) {
        return (float2){0,0};
    }
    return (float2){buffer_var[pos_in_buffer],0};
}

inline void write_buffer1dc(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global char * buffer_var, int pos, char value )
{
    int pos_in_buffer = pos;
    if (pos < 0 || pos >= write_buffer_width) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer1duc(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global uchar * buffer_var, int pos, uchar value )
{
    int pos_in_buffer = pos;
    if (pos < 0 || pos >= write_buffer_width) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer1di(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global short * buffer_var, int pos, short value )
{
    int pos_in_buffer = pos;
    if (pos < 0 || pos >= write_buffer_width) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer1dui(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global ushort * buffer_var, int pos, ushort value )
{
    int pos_in_buffer = pos;
    if (pos < 0 || pos >= write_buffer_width) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer1df(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global float* buffer_var, int pos, float value )
{
    int pos_in_buffer = pos;
    if (pos < 0 || pos >= write_buffer_width) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

The only question left is the cost and impact of adding _1d kernel version in CLIJ kernels. I do not think it will have any impact though I prefer to ask.

EDIT:
After a quick overview of clij kernel, not all dimensional specific are fit to 1d, this would mainly concern basic operation not related to image them selves, such as set_2d_x,cl , or equal_2d_x.cl. Those specific to image, like draw_sphere_2d_x.cl should not be extended. This raise the need of input/output testing in some kernel to avoid miss use.
Also some, once in 1d, become repetitive, such as set_column_2d_x.cl and set_row_2d_x.cl which become the same operation once in 1d.

StRigaud added a commit that referenced this issue Nov 15, 2021
@StRigaud
WIP: require solving issue [#57](#57)
dfa068f
StRigaud added a commit that referenced this issue Nov 15, 2021
@StRigaud
WIP: require solving issue [#57]
541a7df
StRigaud added a commit that referenced this issue Nov 15, 2021
@StRigaud
WIP: Preamble does not support cl_image1d #57
6197393
@StRigaud
Copy link
Member Author

The merging of the branch #ocl_buffer_image_interoperability which is solving this issue is directly linked to the resolution of issues #11 and #15 from CLIJ kernel repository.

@StRigaud StRigaud self-assigned this Dec 10, 2021
@StRigaud StRigaud added request New feature or request WIP Work In Progress labels Dec 10, 2021
@StRigaud StRigaud changed the title Preamble does not support cl_image1d Enable image and buffer data interoperability Dec 10, 2021
@StRigaud StRigaud linked a pull request Dec 10, 2021 that will close this issue
Ocl buffer image interoperability #65
Merged
@StRigaud
Copy link
Member Author

StRigaud commented Feb 9, 2022

#65 is solving the issue by introducing an Object class holding a cl::Memory which is the super-class for both cl::Buffer, cl::Image1D, cl::Image2D, cl::Image3D.

Kernels will take as input and output an Object, no matter the memory type that it is holding and process it in device.
Both type Buffer and Image should interoperate seamlessly, exception to kernels requiring a specific memory type.

The memory type is to be precised when calling create or push, default being Buffer.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees

@StRigaud StRigaud

@haesleinhuepf haesleinhuepf

Labels
request New feature or request WIP Work In Progress
Projects
None yet
Milestone
No milestone
Development

Successfully merging a pull request may close this issue.

Ocl buffer image interoperability clEsperanto/CLIc
2 participants
@StRigaud @haesleinhuepf