1/* 2 * Copyright (c) 2017 Arm Limited. 3 * 4 * SPDX-License-Identifier: MIT 5 * 6 * Permission is hereby granted, free of charge, to any person obtaining a copy 7 * of this software and associated documentation files (the "Software"), to 8 * deal in the Software without restriction, including without limitation the 9 * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or 10 * sell copies of the Software, and to permit persons to whom the Software is 11 * furnished to do so, subject to the following conditions: 12 * 13 * The above copyright notice and this permission notice shall be included in all 14 * copies or substantial portions of the Software. 15 * 16 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR 17 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, 18 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE 19 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER 20 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, 21 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE 22 * SOFTWARE. 23 */ 24#include "helpers.h" 25#include "warp_helpers.h" 26 27/** Performs a remapping of an input image to an output given two remapping image using nearest neighbor as interpolation. 28 * 29 * This kernel performs remapping with this method of pixel coordinate translation: 30 * out(x,y) = in(mapx(x,y), mapy(x,y)); 31 * 32 * @param[in] in_ptr Pointer to the source image. Supported data types: U8. 33 * @param[in] in_stride_x Stride of the source image in X dimension (in bytes) 34 * @param[in] in_step_x in_stride_x * number of elements along X processed per work item (in bytes) 35 * @param[in] in_stride_y Stride of the source image in Y dimension (in bytes) 36 * @param[in] in_step_y in_stride_y * number of elements along Y processed per work item (in bytes) 37 * @param[in] in_offset_first_element_in_bytes Offset of the first element in the source image 38 * @param[out] out_ptr Pointer to the destination image. Supported data types: U8. 39 * @param[in] out_stride_x Stride of the destination image in X dimension (in bytes) 40 * @param[in] out_step_x out_stride_x * number of elements along X processed per work item (in bytes) 41 * @param[in] out_stride_y Stride of the destination image in Y dimension (in bytes) 42 * @param[in] out_step_y out_stride_y * number of elements along Y processed per work item (in bytes) 43 * @param[in] out_offset_first_element_in_bytes Offset of the first element in the destination image 44 * @param[in] mapx_ptr Pointer to the x remapping image. Supported data types: F32. 45 * @param[in] mapx_stride_x Stride of the remapping image in X dimension (in bytes) 46 * @param[in] mapx_step_x mapx_stride_x * number of elements along X processed per work item (in bytes) 47 * @param[in] mapx_stride_y Stride of the remapping image in Y dimension (in bytes) 48 * @param[in] mapx_step_y mapy_stride_y * number of elements along Y processed per work item (in bytes) 49 * @param[in] mapx_offset_first_element_in_bytes Offset of the first element in the remapping image 50 * @param[in] mapy_ptr Pointer to the x remapping image. Supported data types: F32. 51 * @param[in] mapy_stride_x Stride of the remapping image in X dimension (in bytes) 52 * @param[in] mapy_step_x mapy_stride_x * number of elements along X processed per work item (in bytes) 53 * @param[in] mapy_stride_y Stride of the remapping image in Y dimension (in bytes) 54 * @param[in] mapy_step_y mapy_stride_y * number of elements along Y processed per work item (in bytes) 55 * @param[in] mapy_offset_first_element_in_bytes Offset of the first element in the remapping image 56 * @param[in] width Width of the input image 57 * @param[in] height Height of the input image 58 */ 59__kernel void remap_nearest_neighbour( 60 IMAGE_DECLARATION(in), 61 IMAGE_DECLARATION(out), 62 IMAGE_DECLARATION(mapx), 63 IMAGE_DECLARATION(mapy), 64 const float width, 65 const float height) 66{ 67 Image in = CONVERT_TO_IMAGE_STRUCT_NO_STEP(in); 68 Image out = CONVERT_TO_IMAGE_STRUCT(out); 69 Image mapx = CONVERT_TO_IMAGE_STRUCT(mapx); 70 Image mapy = CONVERT_TO_IMAGE_STRUCT(mapy); 71 72 float4 mapx_coords = vload4(0, (__global float *)mapx.ptr); 73 float4 mapy_coords = vload4(0, (__global float *)mapy.ptr); 74 float8 map_coords = (float8)(mapx_coords.s0, mapy_coords.s0, mapx_coords.s1, mapy_coords.s1, 75 mapx_coords.s2, mapy_coords.s2, mapx_coords.s3, mapy_coords.s3); 76 map_coords += (float8)(0.5f); 77 78 vstore4(read_texels4(&in, convert_int8(clamp_to_border(map_coords, width, height))), 0, out.ptr); 79} 80 81/** Performs a remapping of an input image to an output given two remapping image using bilinear as interpolation. 82 * 83 * This kernel performs remapping with this method of pixel coordinate translation: 84 * out(x,y) = in(mapx(x,y), mapy(x,y)); 85 * 86 * @param[in] in_ptr Pointer to the source image. Supported data types: U8. 87 * @param[in] in_stride_x Stride of the source image in X dimension (in bytes) 88 * @param[in] in_step_x in_stride_x * number of elements along X processed per work item (in bytes) 89 * @param[in] in_stride_y Stride of the source image in Y dimension (in bytes) 90 * @param[in] in_step_y in_stride_y * number of elements along Y processed per work item (in bytes) 91 * @param[in] in_offset_first_element_in_bytes Offset of the first element in the source image 92 * @param[out] out_ptr Pointer to the destination image. Supported data types: U8. 93 * @param[in] out_stride_x Stride of the destination image in X dimension (in bytes) 94 * @param[in] out_step_x out_stride_x * number of elements along X processed per work item (in bytes) 95 * @param[in] out_stride_y Stride of the destination image in Y dimension (in bytes) 96 * @param[in] out_step_y out_stride_y * number of elements along Y processed per work item (in bytes) 97 * @param[in] out_offset_first_element_in_bytes Offset of the first element in the destination image 98 * @param[in] mapx_ptr Pointer to the x remapping image. Supported data types: F32. 99 * @param[in] mapx_stride_x Stride of the remapping image in X dimension (in bytes) 100 * @param[in] mapx_step_x mapx_stride_x * number of elements along X processed per work item (in bytes) 101 * @param[in] mapx_stride_y Stride of the remapping image in Y dimension (in bytes) 102 * @param[in] mapx_step_y mapy_stride_y * number of elements along Y processed per work item (in bytes) 103 * @param[in] mapx_offset_first_element_in_bytes Offset of the first element in the remapping image 104 * @param[in] mapy_ptr Pointer to the x remapping image. Supported data types: F32. 105 * @param[in] mapy_stride_x Stride of the remapping image in X dimension (in bytes) 106 * @param[in] mapy_step_x mapy_stride_x * number of elements along X processed per work item (in bytes) 107 * @param[in] mapy_stride_y Stride of the remapping image in Y dimension (in bytes) 108 * @param[in] mapy_step_y mapy_stride_y * number of elements along Y processed per work item (in bytes) 109 * @param[in] mapy_offset_first_element_in_bytes Offset of the first element in the remapping image 110 * @param[in] width Width of the input image 111 * @param[in] height Height of the input image 112 */ 113__kernel void remap_bilinear( 114 IMAGE_DECLARATION(in), 115 IMAGE_DECLARATION(out), 116 IMAGE_DECLARATION(mapx), 117 IMAGE_DECLARATION(mapy), 118 const float width, 119 const float height) 120{ 121 Image in = CONVERT_TO_IMAGE_STRUCT_NO_STEP(in); 122 Image out = CONVERT_TO_IMAGE_STRUCT(out); 123 Image mapx = CONVERT_TO_IMAGE_STRUCT(mapx); 124 Image mapy = CONVERT_TO_IMAGE_STRUCT(mapy); 125 126 float4 mapx_coords = vload4(0, (__global float *)mapx.ptr); 127 float4 mapy_coords = vload4(0, (__global float *)mapy.ptr); 128 float8 map_coords = (float8)(mapx_coords.s0, mapy_coords.s0, mapx_coords.s1, mapy_coords.s1, 129 mapx_coords.s2, mapy_coords.s2, mapx_coords.s3, mapy_coords.s3); 130 131 vstore4(bilinear_interpolate(&in, clamp_to_border(map_coords, width, height), width, height), 0, out.ptr); 132} 133