affine transform debug

[iionichi]

Context and Problem

This is code which I have extended for affine transformation to include interpolation at the software level. However the desired output is not getting generated.

#ifndef SAMPLER_FILTER
#define SAMPLER_FILTER CLK_FILTER_NEAREST
#endif

#ifndef SAMPLER_ADDRESS
#define SAMPLER_ADDRESS CLK_ADDRESS_CLAMP
#endif

__kernel void affine_transform_2d(
    IMAGE_input_TYPE input,
    IMAGE_output_TYPE output,
    IMAGE_mat_TYPE mat)
{

  const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE|
      SAMPLER_ADDRESS |	SAMPLER_FILTER;

  uint i = get_global_id(0);
  uint j = get_global_id(1);

  uint Nx = GET_IMAGE_WIDTH(input);
  uint Ny = GET_IMAGE_HEIGHT(input);

  float x = i+0.5f;
  float y = j+0.5f;

  float y2 = (mat[3]*y+mat[4]*x+mat[5]);
  float x2 = (mat[0]*y+mat[1]*x+mat[2]);

  int2 coord_norm = (int2)(x2,y2);

  float2 frac_coord = (float2)(x2-floor(x2),y2-floor(y2));

  float pix = 0.0;

  if (x2 >= 0 && y2 >= 0 &&
  x2 < GET_IMAGE_WIDTH(input) && y2 < GET_IMAGE_HEIGHT(input)) {
    
    float2 top_left_coord = (float2)(coord_norm.x,coord_norm.y);
    float2 top_right_coord = (float2)((coord_norm.x+1),(coord_norm.y));
    float2 bottom_left_coord = (float2)((coord_norm.x),(coord_norm.y+1));
    float2 bottom_right_coord = (float2)((coord_norm.x+1),(coord_norm.y+1));

    float top_left_color = (float) (READ_input_IMAGE(input, sampler, POS_input_INSTANCE(top_left_coord.x, top_left_coord.y, 0, 0)).x);
    float top_right_color = (float) (READ_input_IMAGE(input, sampler, POS_input_INSTANCE(top_right_coord.x, top_right_coord.y, 0, 0)).x);
    float bottom_left_color = (float) (READ_input_IMAGE(input, sampler, POS_input_INSTANCE(bottom_left_coord.x, bottom_left_coord.y, 0, 0)).x);
    float bottom_right_color = (float) (READ_input_IMAGE(input, sampler, POS_input_INSTANCE(bottom_right_coord.x, bottom_right_coord.y, 0, 0)).x);

    pix = (float) mix(mix(top_left_color, top_right_color, frac_coord.x), mix(bottom_left_color, bottom_right_color, frac_coord.x), frac_coord.y);
  }
  printf("This is sparta: %d,%d",i,j);
  WRITE_output_IMAGE(output, POS_output_INSTANCE(i, j, 0, 0), CONVERT_output_PIXEL_TYPE(pix));
}

Input image

Following is the output from the above code.

Following is the output from interpolation done by openCL

What I would like to do

I want to debug at the kernel level to understand what is happening and how the calculations are being performed. People on the internet said to run the kernel in a debugger to closely monitor the kernel. If someone has already done it I would be happy to learn how to do it.

Something I found

In the book OpenCL in Action it says in page 137 that:

With this I understand that OpenCL has in-built functions to deal with interpolations when it is not supported in some systems. In this case I don't think it would be necessary to add this interpolation at the kernel level.

[haesleinhuepf]

Hi @iionichi ,

awesome! Great to see progress here :-)

Would you mind uploading your code and the (notebook?) that produces the output shown above to a branch or fork on github? In that way, others could reproduce your experiment and potentially help.

Have a great weekend!

Best,
Robert

[iionichi]

@haesleinhuepf
I have pushed the code to the repository.

[haesleinhuepf]

I have pushed the code to the repository.

Where? Can you provide a link? Thanks!

[iionichi]

I have pushed the code to the repository.

Where? Can you provide a link? Thanks!

Since I cannot commit in the this repository I had forked it to push the code in the affine branch : https://github.com/iionichi/pyclesperanto_prototype/tree/affine

Here is the pull request to merge the update to the affine branch of this repository : #311

[haesleinhuepf]

Just for completeness, the notebook is located there:
https://github.com/clEsperanto/pyclesperanto_prototype/blob/1343678a103fa4cc4335c1b904fb93661caa41bb/pyclesperanto_prototype/_tier8/cle_test.ipynb

[haesleinhuepf]

printing in opencl:

int i = 5;
if (x == 5 && y == 6) {
    printf("%d", i);
}

[haesleinhuepf]

In case you want to make sure that pixels outside the image are considered == 0, feel free to adapt the strategy used here:
031d524

[iionichi]

@haesleinhuepf can you share how you showed the color beyond the edge of the image?

[iionichi]

I have updated the code.
Instead of considering pixels outside the image as zero or same as the boundary pixel I have considered pixel previous to the current pixel.
With this the interpolation is very similar to the one which openCL does with hardware acceleration. The shifting problem still stays unfortunately.
The updated kernel: https://github.com/iionichi/pyclesperanto_prototype/blob/affine/pyclesperanto_prototype/_tier8/affine_transform_2d_x.cl
Notebook link: https://github.com/iionichi/pyclesperanto_prototype/blob/affine/pyclesperanto_prototype/_tier8/cle_test.ipynb