| 36 | } // namespace |
| 37 | |
| 38 | std::string Add::GetElementWiseCode( |
| 39 | const TensorDescriptor& src_descriptor, |
| 40 | const TensorDescriptor& dst_descriptor, CalculationsPrecision precision, |
| 41 | const std::vector<ElementwiseOperation*>& linked_operations) { |
| 42 | TensorCodeGenerator src_tensor("src_data", "dst_size", src_descriptor); |
| 43 | TensorCodeGenerator dst_tensor("dst_data", "dst_size", dst_descriptor); |
| 44 | |
| 45 | std::string c = GetCommonDefines(precision); |
| 46 | |
| 47 | c += "__kernel void main_function(\n"; |
| 48 | c += src_tensor.GetDeclaration(AccessType::READ); |
| 49 | c += GetArgsDeclaration(); |
| 50 | c += ::tflite::gpu::cl::GetArgsDeclaration(linked_operations); |
| 51 | c += dst_tensor.GetDeclaration(AccessType::WRITE) + ",\n"; |
| 52 | c += " int4 dst_size\n"; |
| 53 | c += ") {\n"; |
| 54 | c += " int X = get_global_id(0);\n"; |
| 55 | c += " int Y = get_global_id(1);\n"; |
| 56 | c += " int Z = get_global_id(2);\n"; |
| 57 | c += " if (X >= dst_size.x || Y >= dst_size.y) { \n"; |
| 58 | c += " return; \n"; |
| 59 | c += " } \n"; |
| 60 | c += " FLT4 src = (FLT4)(0.0);\n"; |
| 61 | c += " " + dst_tensor.GetAddress("address", "X", "Y", "Z") + "\n"; |
| 62 | if (src_depthes_[0] != dst_depth_) { |
| 63 | c += " if (Z < " + std::to_string(src_depthes_[0]) + ") {\n"; |
| 64 | if (src_descriptor.storage_type == TensorStorageType::TEXTURE_2D) { |
| 65 | c += " float t_y = address.y - Z; \n"; |
| 66 | c += " int ti_y = (t_y + 0.5) * " + inv_divisor_name_ + "; \n"; |
| 67 | c += " int2 tmp_add = (int2)(address.x, ti_y * " + |
| 68 | std::to_string(src_depthes_[0]) + " + Z);\n"; |
| 69 | c += " src += " + |
| 70 | src_tensor.Read3D("tmp_add", TextureAddressMode::DONT_CARE) + ";\n"; |
| 71 | } else { |
| 72 | c += " src += " + |
| 73 | src_tensor.Read3D("address", TextureAddressMode::DONT_CARE) + ";\n"; |
| 74 | } |
| 75 | c += " }\n"; |
| 76 | } else { |
| 77 | c += " src += " + |
| 78 | src_tensor.Read3D("address", TextureAddressMode::DONT_CARE) + ";\n"; |
| 79 | } |
| 80 | c += " " + GetCoreCode("src", "Z", "address"); |
| 81 | c += PostProcess(linked_operations, "src", "Z", "address"); |
| 82 | c += " " + dst_tensor.Write3D("src", "address") + "\n"; |
| 83 | c += "} \n"; |
| 84 | return c; |
| 85 | } |
| 86 | |
| 87 | Add::Add(const OperationDef& definition, const std::vector<int>& channels, |
| 88 | int dst_channels) |
nothing calls this directly
no test coverage detected