| 25 | namespace { |
| 26 | |
| 27 | std::string GetUpsampleCode( |
| 28 | const TensorDescriptor& src_descriptor, |
| 29 | const TensorDescriptor& dst_descriptor, CalculationsPrecision precision, |
| 30 | const std::vector<ElementwiseOperation*>& linked_operations) { |
| 31 | TensorCodeGenerator src_tensor("src_data", "src_size", src_descriptor); |
| 32 | TensorCodeGenerator dst_tensor("dst_data", "dst_size", dst_descriptor); |
| 33 | |
| 34 | std::string c = GetCommonDefines(precision); |
| 35 | c += "__kernel void main_function(\n"; |
| 36 | c += src_tensor.GetDeclaration(AccessType::READ); |
| 37 | c += GetArgsDeclaration(linked_operations); |
| 38 | c += dst_tensor.GetDeclaration(AccessType::WRITE) + ",\n"; |
| 39 | c += " int4 src_size,\n"; |
| 40 | c += " int4 dst_size,\n"; |
| 41 | c += " float2 scale_factor\n"; |
| 42 | c += ") {\n"; |
| 43 | c += " int X = get_global_id(0);\n"; |
| 44 | c += " int Y = get_global_id(1);\n"; |
| 45 | c += " int Z = get_global_id(2);\n"; |
| 46 | c += " if (X >= dst_size.x || Y >= dst_size.y) { \n"; |
| 47 | c += " return; \n"; |
| 48 | c += " } \n"; |
| 49 | c += " float2 f_coords = (float2)(X, Y) * scale_factor;\n"; |
| 50 | c += " int2 borders = src_size.xy - (int2)(1, 1);\n"; |
| 51 | c += " int4 st;\n"; |
| 52 | c += " st.xy = (int2)(f_coords.x, f_coords.y);\n"; |
| 53 | c += " st.zw = min(st.xy + (int2)(1, 1), borders);\n"; |
| 54 | c += " float2 t = f_coords - (float2)(st.x, st.y);\n"; |
| 55 | c += " float4 src0 = " + |
| 56 | src_tensor.ReadAsFloat3D("st.x", "st.y", "Z", |
| 57 | TextureAddressMode::DONT_CARE) + |
| 58 | ";\n"; |
| 59 | c += " float4 src1 = " + |
| 60 | src_tensor.ReadAsFloat3D("st.z", "st.y", "Z", |
| 61 | TextureAddressMode::DONT_CARE) + |
| 62 | ";\n"; |
| 63 | c += " float4 src2 = " + |
| 64 | src_tensor.ReadAsFloat3D("st.x", "st.w", "Z", |
| 65 | TextureAddressMode::DONT_CARE) + |
| 66 | ";\n"; |
| 67 | c += " float4 src3 = " + |
| 68 | src_tensor.ReadAsFloat3D("st.z", "st.w", "Z", |
| 69 | TextureAddressMode::DONT_CARE) + |
| 70 | ";\n"; |
| 71 | c += " FLT4 r0 = TO_FLT4(mix(mix(src0, src1, t.x), mix(src2, src3, t.x), " |
| 72 | "t.y));\n"; |
| 73 | c += " " + dst_tensor.GetAddress("dst_addr", "X", "Y", "Z") + "\n"; |
| 74 | c += PostProcess(linked_operations, "r0", "Z", "dst_addr"); |
| 75 | c += " " + dst_tensor.Write3D("r0", "dst_addr"); |
| 76 | c += "}\n"; |
| 77 | return c; |
| 78 | } |
| 79 | |
| 80 | } // namespace |
| 81 |
no test coverage detected