| /device/google/contexthub/firmware/os/algos/ |
| fusion.c | 17 // adapted from frameworks/native/services/sensorservice/Fusion.cpp
19 #include <algos/fusion.h>
28 // change to 0 to disable fusion debugging output
66 void initFusion(struct Fusion *fusion, uint32_t flags) {
67 fusion->flags = flags;
70 // normal fusion mode
71 fusion->param.gyro_var = DEFAULT_GYRO_VAR;
72 fusion->param.gyro_bias_var = DEFAULT_GYRO_BIAS_VAR;
73 fusion->param.acc_stdev = DEFAULT_ACC_STDEV
[all...] |
| /external/tensorflow/tensorflow/compiler/xla/service/gpu/ |
| fusion_merger.cc | 48 // Returns the bytes read by fusion parameter 'param', by returning the byte
87 // Returns the bytes read by all fusion parameters of instruction 'fusion'.
88 double CalculateBytesReadByFusionInstruction(HloInstruction* fusion) {
90 for (auto* fused_instruction : fusion->fused_instructions()) {
99 // Returns bytes transferred by instruction 'fusion', including the bytes
101 double GetCurrentBytesTransferred(HloInstruction* fusion) {
102 CHECK_EQ(HloOpcode::kFusion, fusion->opcode());
103 const double bytes_read = CalculateBytesReadByFusionInstruction(fusion);
105 if (fusion->IsMultiOutputFusion())
[all...] |
| multi_output_fusion_test.cc | 48 // Fusion with reduce instruction root and a sibling reduce instruction
62 fusion = f32[512] fusion(p0, p1), kind=kInput, calls=fused_computation
64 ROOT root = (f32[512]{0}, f32[512]{0}) tuple(fusion, reduce.2)
69 const HloInstruction* fusion = local
71 ASSERT_TRUE(fusion->IsMultiOutputFusion());
72 EXPECT_THAT(fusion->fused_expression_root(),
95 fusion.1 = f32[] fusion(p0, p1), kind=kInput, calls=fused_computation_1
96 fusion.2 = f32[] fusion(p0, p1), kind=kInput, calls=fused_computation_
157 const HloInstruction* fusion = local
190 const HloInstruction* fusion = local
252 const HloInstruction* fusion = local
296 const HloInstruction* fusion = local
352 const HloInstruction* fusion = local
401 const HloInstruction* fusion = root->operand(0)->operand(0); local
428 const HloInstruction* fusion = root->operand(0)->operand(0); local
469 const HloInstruction* fusion = root->operand(0)->operand(0); local
536 const HloInstruction* fusion = root->operand(0)->operand(0); local
[all...] |
| cudnn_conv_runner.cc | 75 absl::optional<FusionParams> fusion; member in struct:xla::gpu::__anon44396::CudnnConvParams
281 se::DeviceMemory<T> side_input(params.fusion->side_input_buf);
284 if (params.fusion->side_input_scale != 0) {
301 params.fusion->side_input_scale, bias_desc,
302 DeviceMemory<T>(params.fusion->bias_buf), params.fusion->mode,
371 params.fusion.emplace();
372 auto& fusion = *params.fusion; local
377 fusion.mode = static_cast<se::dnn::ActivationMode>
[all...] |
| ir_emitter.h | 64 // Fusion nodes are a special case. A fusion node is emitted using
97 Status HandleFusion(HloInstruction* fusion) override;
188 HloInstruction* fusion) {
191 ir_arrays.reserve(fusion->operand_count());
192 absl::c_transform(fusion->operands(), std::back_inserter(ir_arrays),
194 return GetIrArray(*operand, *fusion);
|
| nvptx_compiler.cc | 255 // pipeline, before layout assignment and fusion. This pass does some
321 // We pick the algorithm before fusion so we can generate better HLO. After
340 // However, if we were to run CudnnConvAlgorithmPicker after fusion
341 // the gte(customcall, 0) would probably already be into a fusion node. We
355 HloPassFix<HloPassPipeline> fusion("fusion");
358 fusion.AddPass<VariadicOpSplitter>();
361 fusion.AddInvariantChecker<HloVerifier>(
365 fusion.AddPass<GpuInstructionFusion>(/*may_duplicate=*/false);
366 fusion.AddPass<GpuInstructionFusion>(/*may_duplicate=*/true)
[all...] |
| /device/google/contexthub/firmware/os/inc/algos/ |
| fusion.h | 40 struct Fusion {
70 INITIALIZATION, // right after initialization of fusion
75 void initFusion(struct Fusion *fusion, uint32_t flags);
77 void fusionHandleGyro(struct Fusion *fusion, const struct Vec3 *w, float dT);
78 int fusionHandleAcc(struct Fusion *fusion, const struct Vec3 *a, float dT);
79 int fusionHandleMag(struct Fusion *fusion, const struct Vec3 *m, float dT)
[all...] |
| /external/tensorflow/tensorflow/compiler/xla/service/llvm_ir/ |
| dynamic_update_slice_util.h | 41 // Checks if the given fusion node is amenable to being implemented by
44 HloInstruction* fusion, const BufferAssignment& assignment) {
45 CHECK_EQ(fusion->opcode(), HloOpcode::kFusion);
46 HloInstruction* fused_root = fusion->fused_expression_root();
48 fusion->fusion_kind() != HloInstruction::FusionKind::kLoop) {
60 auto* operand = fusion->operand(fusion_operand->parameter_number());
62 assignment.HasAllocationAt(fusion, {}) &&
63 assignment.SharesSliceAtIndex(fusion, {}, operand, index);
74 // Given a loop-fusion node whose root is a dynamic-update-slice op whose
76 // (sequential) code for a fusion node that does the dynamic-update-slice i
[all...] |
| dynamic_update_slice_util.cc | 138 HloInstruction* fusion,
142 CHECK_EQ(fusion->opcode(), HloOpcode::kFusion);
144 << fusion->ToShortString();
146 auto* dynamic_update_slice = fusion->fused_expression_root();
156 // update_shape is inside a fusion node -- it's never materialized in memory
158 // fusion node for iteration, since that corresponds to the order in memory of
164 // fusion elsewhere.)
166 LayoutUtil::CopyLayoutBetweenShapes(fusion->shape(), &update_shape));
182 fusion_output_array, launch_dimensions, IrName(fusion), b);
186 HloInstruction* fusion,
[all...] |
| /external/deqp-deps/SPIRV-Tools/source/opt/ |
| loop_fusion_pass.cpp | 48 LoopFusion fusion(context(), &loop_0, &loop_1);
50 if (fusion.AreCompatible() && fusion.IsLegal()) {
56 fusion.Fuse();
|
| /external/swiftshader/third_party/SPIRV-Tools/source/opt/ |
| loop_fusion_pass.cpp | 48 LoopFusion fusion(context(), &loop_0, &loop_1);
50 if (fusion.AreCompatible() && fusion.IsLegal()) {
56 fusion.Fuse();
|
| /external/deqp-deps/SPIRV-Tools/test/opt/loop_optimizations/ |
| fusion_legal.cpp | 171 LoopFusion fusion(context.get(), loops[0], loops[1]);
173 EXPECT_TRUE(fusion.AreCompatible());
174 EXPECT_TRUE(fusion.IsLegal());
176 fusion.Fuse();
307 LoopFusion fusion(context.get(), loops[0], loops[1]);
309 EXPECT_TRUE(fusion.AreCompatible());
310 EXPECT_TRUE(fusion.IsLegal());
312 fusion.Fuse();
440 LoopFusion fusion(context.get(), loops[0], loops[1]);
442 EXPECT_TRUE(fusion.AreCompatible())
[all...] |
| fusion_illegal.cpp | 144 LoopFusion fusion(context.get(), loops[0], loops[1]);
146 EXPECT_TRUE(fusion.AreCompatible());
147 EXPECT_FALSE(fusion.IsLegal());
263 LoopFusion fusion(context.get(), loops[0], loops[1]);
265 EXPECT_TRUE(fusion.AreCompatible());
266 EXPECT_FALSE(fusion.IsLegal());
430 LoopFusion fusion(context.get(), loop_0, loop_1);
431 EXPECT_FALSE(fusion.AreCompatible());
435 LoopFusion fusion(context.get(), loop_0, loop_2);
436 EXPECT_TRUE(fusion.AreCompatible())
[all...] |
| fusion_compatibility.cpp | 110 LoopFusion fusion(context.get(), loops[0], loops[1]);
111 EXPECT_FALSE(fusion.AreCompatible());
195 LoopFusion fusion(context.get(), loops[0], loops[1]);
196 EXPECT_TRUE(fusion.AreCompatible());
281 LoopFusion fusion(context.get(), loops[0], loops[1]);
282 EXPECT_TRUE(fusion.AreCompatible());
368 LoopFusion fusion(context.get(), loops[0], loops[1]);
369 EXPECT_FALSE(fusion.AreCompatible());
456 LoopFusion fusion(context.get(), loops[0], loops[1]);
457 EXPECT_FALSE(fusion.AreCompatible())
[all...] |
| /external/swiftshader/third_party/SPIRV-Tools/test/opt/loop_optimizations/ |
| fusion_legal.cpp | 171 LoopFusion fusion(context.get(), loops[0], loops[1]);
173 EXPECT_TRUE(fusion.AreCompatible());
174 EXPECT_TRUE(fusion.IsLegal());
176 fusion.Fuse();
307 LoopFusion fusion(context.get(), loops[0], loops[1]);
309 EXPECT_TRUE(fusion.AreCompatible());
310 EXPECT_TRUE(fusion.IsLegal());
312 fusion.Fuse();
440 LoopFusion fusion(context.get(), loops[0], loops[1]);
442 EXPECT_TRUE(fusion.AreCompatible())
[all...] |
| fusion_illegal.cpp | 144 LoopFusion fusion(context.get(), loops[0], loops[1]);
146 EXPECT_TRUE(fusion.AreCompatible());
147 EXPECT_FALSE(fusion.IsLegal());
263 LoopFusion fusion(context.get(), loops[0], loops[1]);
265 EXPECT_TRUE(fusion.AreCompatible());
266 EXPECT_FALSE(fusion.IsLegal());
430 LoopFusion fusion(context.get(), loop_0, loop_1);
431 EXPECT_FALSE(fusion.AreCompatible());
435 LoopFusion fusion(context.get(), loop_0, loop_2);
436 EXPECT_TRUE(fusion.AreCompatible())
[all...] |
| fusion_compatibility.cpp | 110 LoopFusion fusion(context.get(), loops[0], loops[1]);
111 EXPECT_FALSE(fusion.AreCompatible());
195 LoopFusion fusion(context.get(), loops[0], loops[1]);
196 EXPECT_TRUE(fusion.AreCompatible());
281 LoopFusion fusion(context.get(), loops[0], loops[1]);
282 EXPECT_TRUE(fusion.AreCompatible());
368 LoopFusion fusion(context.get(), loops[0], loops[1]);
369 EXPECT_FALSE(fusion.AreCompatible());
456 LoopFusion fusion(context.get(), loops[0], loops[1]);
457 EXPECT_FALSE(fusion.AreCompatible())
[all...] |
| /external/tensorflow/tensorflow/compiler/xla/service/ |
| fusion_queue.h | 24 // A queue interface that allows implementations to choose fusion candidates in
31 // Dequeues the next fusion candidates: a consumer and the list of producers
40 // A callback passed to the queue implementation right after the fusion is
42 virtual void OnFusingInstruction(HloInstruction* fusion,
|
| multi_output_fusion.cc | 117 // Make sure that if only one of the instructions is a fusion, or if only one
118 // of the instructions is a multi-output fusion, it's what will be fused into.
151 HloInstruction* fusion = instr1; local
154 fusion = instr2;
159 for (auto use : fusion->users()) {
166 FusionCandidate& fusion_node = candidates_[get_candidate_id(fusion)];
170 UpdateReachability(fusion, fused, all_fusion_candidates_,
173 // Update the fusible list for fusion. Variable new_fusibles keeps
180 if (is_fused(instr) || is_connected(fusion, instr)) {
185 int64 profit = GetProfit(instr, fusion);
[all...] |
| bfloat16_propagation.h | 45 // fusion computation or its only caller is a while.
96 // Special handling in the opportunity-finding pass for fusion computations.
99 void DetermineFusionComputationPrecision(HloInstruction* fusion);
101 // Reverts changes to BF16 that will not propagate outside a fusion
102 // computation. This avoids BF16 casts overhead inside a fusion which won't
106 void RevertIfFusionInternalBF16Changes(HloInstruction* fusion);
|
| hlo_instruction_test.cc | 628 // Create a fusion instruction containing a single unary operation.
635 auto* fusion = computation->CreateFusionInstruction( local
638 EXPECT_THAT(fusion->operands(), ElementsAre(constant));
639 EXPECT_THAT(constant->users(), ElementsAre(fusion));
644 // Create a fusion instruction containing a single binary operation.
653 auto* fusion = computation->CreateFusionInstruction( local
656 EXPECT_THAT(fusion->operands(), ElementsAre(constant1, constant2));
657 EXPECT_THAT(constant1->users(), ElementsAre(fusion));
658 EXPECT_THAT(constant2->users(), ElementsAre(fusion));
675 auto* fusion = computation->CreateFusionInstruction local
698 auto* fusion = computation->CreateFusionInstruction( local
776 auto* fusion = computation->CreateFusionInstruction( local
823 auto* fusion = computation->CreateFusionInstruction( local
1073 HloInstruction* fusion = computation->CreateFusionInstruction( local
1114 HloInstruction* fusion = computation->CreateFusionInstruction( local
1157 HloInstruction* fusion = computation->CreateFusionInstruction( local
1198 HloInstruction* fusion = computation->CreateFusionInstruction( local
1219 auto* fusion = computation->CreateFusionInstruction( local
1261 auto fusion = computation->CreateFusionInstruction( local
1501 HloInstruction* fusion = computation->CreateFusionInstruction( local
[all...] |
| hlo_verifier.cc | 513 Status ShapeVerifier::HandleFusion(HloInstruction* fusion) {
514 auto& fused_parameters = fusion->fused_parameters();
515 if (fused_parameters.size() != fusion->operand_count()) {
518 " passed to the fusion instruction in: %s.",
519 fused_parameters.size(), fusion->operand_count(),
520 fusion->ToString().c_str());
524 if (!ShapesSame(fused_param->shape(), fusion->operand(param_no)->shape())) {
528 param_no, fusion->ToString().c_str());
[all...] |
| instruction_fusion_test.cc | 60 HloInstruction* fusion = local
63 ASSERT_THAT(fusion, op::Fusion()) << module->ToString();
64 EXPECT_THAT(fusion->fused_expression_root(),
79 ROOT fusion = f32[4,3] fusion(abs), kind=kLoop, calls=fused_computation
84 HloInstruction* fusion = local
87 ASSERT_THAT(fusion, op::Fusion()) << module->ToString();
88 EXPECT_THAT(fusion->fused_expression_root(), op::Add(op::Abs(), op::Abs())
105 HloInstruction* fusion = local
[all...] |
| /external/tensorflow/tensorflow/compiler/xla/service/cpu/tests/ |
| cpu_fusion_test.cc | 70 CpuInstructionFusion fusion; local
71 EXPECT_TRUE(fusion.Run(module.get()).ValueOrDie());
73 // The computation root instruction was fused. Verify the fusion instruction
117 CpuInstructionFusion fusion; local
118 EXPECT_TRUE(fusion.Run(module.get()).ValueOrDie());
120 // The computation root instruction was fused. Verify the fusion instruction
192 CpuInstructionFusion fusion; local
193 EXPECT_TRUE(fusion.Run(module.get()).ValueOrDie());
195 // The computation root instruction was fused. Verify the fusion instruction
203 // There should be 6 fused instructions in the root fusion instruction:
264 CpuInstructionFusion fusion; local
320 CpuInstructionFusion fusion; local
[all...] |
| /external/boringssl/src/util/fipstools/delocate/testdata/ppc64le-Sample2/ |
| in.s | 49 addis 31,2,.LC0@toc@ha # gpr load fusion, type long
51 addis 19,2,.LC3@toc@ha # gpr load fusion, type long
55 addis 20,2,.LC6@toc@ha # gpr load fusion, type long
118 addis 31,2,.LC11@toc@ha # gpr load fusion, type long
132 addis 6,2,.LC12@toc@ha # gpr load fusion, type long
148 addis 6,2,.LC13@toc@ha # gpr load fusion, type long
|
