% ctype = dtype_to_c_type(dtype)

__kernel void conv2d(const int height, const int width, const int out_height, const int out_width, __global const <%= ctype %> *images, __global const <%= ctype %> *filter, __global <%= ctype %> *output) {

// Get the index of the current element to be processed
const int batch_index = get_global_id(0);
const int h_index = get_global_id(1);
const int w_index = get_global_id(2);
const int h_index_with_stride = h_index * <%= stride[0] %> - <%= padding[0] %>;
const int w_index_with_stride = w_index * <%= stride[1] %> - <%= padding[1] %>;

const int image_index = batch_index * height * width * <%= ch %>;
const int image_row_width = width * <%= ch %>;
const int out_image_row_size = out_height * out_width * <%= out_ch %>;

for (int out_channel_index = 0; out_channel_index < <%= out_ch %>; out_channel_index++) {
  <%= ctype %> sum = 0;
  for (int channel_index = 0; channel_index < <%= ch %>; channel_index++) {
    for(int y = 0; y < <%= fh %>; y++) {
      for (int x = 0; x < <%= fw %>; x++) {
        if ( (h_index_with_stride + y) < height && (w_index_with_stride + x) < width &&
             (h_index_with_stride + y) >= 0 && (w_index_with_stride + x) >=0) {
          <%= ctype %> f = filter[y*<%= fw * ch * out_ch %> + x*<%= ch * out_ch %> + (channel_index*<%= out_ch %>) + out_channel_index];
          sum += images[image_index + (h_index_with_stride + y)*image_row_width + (w_index_with_stride + x)*<%= ch %> + channel_index] * f;
        }
      }
    }
  }
  output[batch_index * out_image_row_size  + h_index * out_width * <%= out_ch %> +  w_index * <%= out_ch %> + out_channel_index ] = sum;
}

}