| 26 | namespace { |
| 27 | |
| 28 | std::string GetPaddingCode( |
| 29 | const TensorDescriptor& src_descriptor, |
| 30 | const TensorDescriptor& dst_descriptor, CalculationsPrecision precision, |
| 31 | const std::vector<ElementwiseOperation*>& linked_operations) { |
| 32 | TensorCodeGenerator src_tensor("src_data", "src_size", src_descriptor); |
| 33 | TensorCodeGenerator dst_tensor("dst_data", "dst_size", dst_descriptor); |
| 34 | |
| 35 | std::string code = GetCommonDefines(precision); |
| 36 | const std::string channels[] = {".x", ".y", ".z", ".w"}; |
| 37 | |
| 38 | code += "__kernel void main_function(\n"; |
| 39 | code += src_tensor.GetDeclaration(AccessType::READ); |
| 40 | code += GetArgsDeclaration(linked_operations); |
| 41 | code += dst_tensor.GetDeclaration(AccessType::WRITE) + ",\n"; |
| 42 | code += " int4 src_size, \n"; |
| 43 | code += " int4 dst_size, \n"; |
| 44 | code += " int4 prepended \n"; |
| 45 | code += ") {\n"; |
| 46 | code += " int X = get_global_id(0);\n"; |
| 47 | code += " int Y = get_global_id(1);\n"; |
| 48 | code += " int Z = get_global_id(2);\n"; |
| 49 | code += " if (X >= dst_size.x || Y >= dst_size.y) return; \n"; |
| 50 | code += " FLT4 result = (FLT4)(0.0);\n"; |
| 51 | code += " int s_x = X - prepended.x;\n"; |
| 52 | code += " int s_y = Y - prepended.y;\n"; |
| 53 | code += " bool inside_x = s_x >= 0 && s_x < src_size.x;\n"; |
| 54 | code += " bool inside_y = s_y >= 0 && s_y < src_size.y;\n"; |
| 55 | code += " if (inside_x && inside_y) {\n"; |
| 56 | code += " int start_channel = Z * 4;\n"; |
| 57 | for (int i = 0; i < 4; ++i) { |
| 58 | const auto& s = channels[i]; |
| 59 | code += " {\n"; |
| 60 | code += " int channel = start_channel + " + std::to_string(i) + ";\n"; |
| 61 | code += " int s_z = channel - prepended.z;\n"; |
| 62 | code += " if (s_z >= 0 && s_z < src_size.z) {\n"; |
| 63 | code += " FLT4 t = " + |
| 64 | src_tensor.Read3D("s_x", "s_y", "s_z / 4", |
| 65 | TextureAddressMode::DONT_CARE) + |
| 66 | ";\n"; |
| 67 | code += " FLT t_ar[4] = {t.x, t.y, t.z, t.w};\n"; |
| 68 | code += " result" + s + " = t_ar[s_z % 4];\n"; |
| 69 | code += " }\n"; |
| 70 | code += " }\n"; |
| 71 | } |
| 72 | code += " }\n"; |
| 73 | code += " " + dst_tensor.GetAddress("address", "X", "Y", "Z") + "\n"; |
| 74 | code += PostProcess(linked_operations, "result", "Z", "address"); |
| 75 | code += " " + dst_tensor.Write3D("result", "address"); |
| 76 | code += "}\n"; |
| 77 | |
| 78 | return code; |
| 79 | } |
| 80 | } // namespace |
| 81 | |
| 82 | Padding::Padding(const OperationDef& definition, const PadAttributes& attr) |
no test coverage detected