| 26 | namespace { |
| 27 | |
| 28 | std::string GetReshapeCode( |
| 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 c = GetCommonDefines(precision); |
| 36 | c += "__kernel void main_function(\n"; |
| 37 | c += src_tensor.GetDeclaration(AccessType::READ); |
| 38 | c += GetArgsDeclaration(linked_operations); |
| 39 | c += dst_tensor.GetDeclaration(AccessType::WRITE) + ",\n"; |
| 40 | c += " int4 src_size, \n"; |
| 41 | c += " int4 dst_size, \n"; |
| 42 | c += " int2 plane_xz \n"; |
| 43 | c += ") {\n"; |
| 44 | c += " int X = get_global_id(0);\n"; |
| 45 | c += " int Y = get_global_id(1);\n"; |
| 46 | c += " int Z = get_global_id(2);\n"; |
| 47 | c += " if (X >= dst_size.x || Y >= dst_size.y) { \n"; |
| 48 | c += " return; \n"; |
| 49 | c += " } \n"; |
| 50 | c += " FLT temps[4];\n"; |
| 51 | c += " temps[0] = (FLT)(0.0f);\n"; |
| 52 | c += " temps[1] = (FLT)(0.0f);\n"; |
| 53 | c += " temps[2] = (FLT)(0.0f);\n"; |
| 54 | c += " temps[3] = (FLT)(0.0f);\n"; |
| 55 | c += " for (int i = 0; i < 4; ++i) {\n"; |
| 56 | c += " int dst_channel = Z * 4 + i;\n"; |
| 57 | c += " if (dst_channel < dst_size.z) {;\n"; |
| 58 | c += " int p = dst_channel + dst_size.z * X + plane_xz.y * Y;\n"; |
| 59 | c += " int src_y = p / plane_xz.x;\n"; |
| 60 | c += " int src_x = (p % plane_xz.x) / src_size.z;\n"; |
| 61 | c += " int src_ch = (p % plane_xz.x) % src_size.z;\n"; |
| 62 | c += " int src_z = src_ch / 4;\n"; |
| 63 | c += " int src_sub_ch = src_ch % 4;\n"; |
| 64 | c += " FLT4 t =" + |
| 65 | src_tensor.Read3D("src_x", "src_y", "src_z", |
| 66 | TextureAddressMode::DONT_CARE) + |
| 67 | ";\n"; |
| 68 | c += " FLT t_ar[4] = {t.x, t.y, t.z, t.w};\n"; |
| 69 | c += " temps[i] = t_ar[src_sub_ch];\n"; |
| 70 | c += " }\n"; |
| 71 | c += " }\n"; |
| 72 | c += " FLT4 result = (FLT4)(temps[0], temps[1], temps[2], temps[3]);\n"; |
| 73 | c += " " + dst_tensor.GetAddress("dst_adr", "X", "Y", "Z"); |
| 74 | c += PostProcess(linked_operations, "result", "Z", "dst_adr"); |
| 75 | c += " " + dst_tensor.Write3D("result", "dst_adr"); |
| 76 | c += "}\n"; |
| 77 | return c; |
| 78 | } |
| 79 | } // namespace |
| 80 | |
| 81 | Reshape::Reshape(Reshape&& operation) |
no test coverage detected