1/* Copyright 2021 Huawei Technologies Co., Ltd 2 * 3 * Licensed under the Apache License, Version 2.0 (the "License"); 4 * you may not use this file except in compliance with the License. 5 * You may obtain a copy of the License at 6 * 7 * http://www.apache.org/licenses/LICENSE-2.0 8 * 9 * Unless required by applicable law or agreed to in writing, software 10 * distributed under the License is distributed on an "AS IS" BASIS, 11 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. 12 * See the License for the specific language governing permissions and 13 * limitations under the License. 14 */ 15__constant sampler_t smp_none = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; 16 17__kernel void ElementAddInt8(__read_only image2d_t input_a, __read_only image2d_t input_b, 18 __write_only image2d_t output, const int2 output_shape, float act_min, float act_max, 19 const float4 scale, const char4 zero_point) { 20 int X = get_global_id(0); 21 int Y = get_global_id(1); 22 if (X >= output_shape.x || Y >= output_shape.y) { 23 return; 24 } 25 char4 a = convert_char4(read_imagei(input_a, smp_none, (int2)(X, Y))); 26 char4 b = convert_char4(read_imagei(input_b, smp_none, (int2)(X, Y))); 27 28 float4 real_a = convert_float4(a - zero_point.x) * scale.x; 29 float4 real_b = convert_float4(b - zero_point.y) * scale.y; 30 int4 result = convert_int4(round((real_a + real_b) / scale.z)) + zero_point.z; 31 result = clamp(result, (int)(act_min), (int)(act_max)); 32 write_imagei(output, (int2)(X, Y), result); 33} 34