| /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 | 45 // Returns the bytes read by fusion parameter 'param', by returning the byte
83 // Returns the bytes read by all fusion parameters of instruction 'fusion'.
84 double CalculateBytesReadByFusionInstruction(HloInstruction* fusion) {
86 for (auto* fused_instruction : fusion->fused_instructions()) {
95 // Returns the flops to bytes transferred ratio of instruction 'fusion'.
96 double CalculateFlopsToBytesRatio(HloInstruction* fusion) {
97 CHECK_EQ(HloOpcode::kFusion, fusion->opcode());
99 double bytes = CalculateBytesReadByFusionInstruction(fusion);
101 if (fusion->IsMultiOutputFusion())
[all...] |
| gpu_compiler.cc | 228 // assignment, fusion would already have run, and the gte(customcall, 0)
229 // would probably already be into a fusion node. We can't simplify across
245 HloPassFix<HloPassPipeline> fusion("fusion");
246 fusion.AddInvariantChecker<HloVerifier>();
247 fusion.AddPass<GpuInstructionFusion>(/*may_duplicate=*/false);
248 fusion.AddPass<GpuInstructionFusion>(/*may_duplicate=*/true);
249 fusion.AddPass<FusionMerger>();
250 TF_RETURN_IF_ERROR(fusion.Run(hlo_module).status());
261 // Do another fusion pass, with the expectation that we may be able t
[all...] |
| ir_emitter_unnested.cc | 492 Status IrEmitterUnnested::HandleFusion(HloInstruction* fusion) {
493 HloInstruction* root = fusion->fused_expression_root();
497 if (HloInstruction::FusionKind::kInput == fusion->fusion_kind()) {
500 VLOG(3) << "Emitting fused reduction to vector: " << fusion->ToString();
502 thunks.emplace_back(BuildKernelThunk(fusion));
504 fusion, static_cast<KernelThunk*>(thunks.back().get())));
506 thunks.emplace_back(BuildKernelThunk(fusion));
508 MakeUnique<SequentialThunk>(std::move(thunks), fusion));
510 for (HloInstruction* operand : fusion->operands()) {
511 parameter_arrays.push_back(GetIrArray(*operand, *fusion));
[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/ |
| ops.h | 38 // Checks if the given fusion node is amenable to being implemented by
41 HloInstruction* fusion, const BufferAssignment& assignment) {
42 CHECK_EQ(fusion->opcode(), HloOpcode::kFusion);
43 HloInstruction* fused_root = fusion->fused_expression_root();
45 fusion->fusion_kind() != HloInstruction::FusionKind::kLoop) {
57 auto* operand = fusion->operand(fusion_operand->parameter_number());
59 assignment.HasAllocationAt(fusion, {}) &&
60 assignment.SharesSliceAtIndex(fusion, {}, operand, index);
71 // Given a loop-fusion node whose root is a dynamic-update-slice op whose
73 // (sequential) code for a fusion node that does the dynamic-update-slice i
[all...] |
| ops.cc | 116 HloInstruction* fusion,
121 CHECK_EQ(fusion->opcode(), HloOpcode::kFusion);
123 << fusion->ToShortString();
125 auto* dynamic_update_slice = fusion->fused_expression_root();
135 // update_shape is inside a fusion node -- it's never materialized in memory
137 // fusion node for iteration, since that corresponds to the order in memory of
143 // fusion elsewhere.)
145 LayoutUtil::CopyLayoutBetweenShapes(fusion->shape(), &update_shape));
156 fusion_output_array, launch_dimensions, IrName(fusion), ir_builder);
160 HloInstruction* fusion,
[all...] |
| /external/tensorflow/tensorflow/compiler/xla/service/ |
| hlo_instruction_test.cc | 625 // Create a fusion instruction containing a single unary operation.
632 auto* fusion = computation->CreateFusionInstruction( local
635 EXPECT_THAT(fusion->operands(), ElementsAre(constant));
636 EXPECT_THAT(constant->users(), ElementsAre(fusion));
641 // Create a fusion instruction containing a single binary operation.
650 auto* fusion = computation->CreateFusionInstruction( local
653 EXPECT_THAT(fusion->operands(), ElementsAre(constant1, constant2));
654 EXPECT_THAT(constant1->users(), ElementsAre(fusion));
655 EXPECT_THAT(constant2->users(), ElementsAre(fusion));
672 auto* fusion = computation->CreateFusionInstruction local
695 auto* fusion = computation->CreateFusionInstruction( local
772 auto* fusion = computation->CreateFusionInstruction( local
819 auto* fusion = computation->CreateFusionInstruction( local
1018 HloInstruction* fusion = computation->CreateFusionInstruction( local
1061 HloInstruction* fusion = computation->CreateFusionInstruction( local
1104 HloInstruction* fusion = computation->CreateFusionInstruction( local
1132 auto* fusion = computation->CreateFusionInstruction( local
1174 auto fusion = computation->CreateFusionInstruction( local
1249 HloInstruction* fusion = computation->CreateFusionInstruction( local
[all...] |
| liveness_util_test.cc | 105 auto fusion = computation_->CreateFusionInstruction( local
110 // The fusion instruction never uses tuple element 0, but does use element 1.
112 DoesNotUseOperandBuffer(tuple, {0}, fusion, *points_to_analysis_));
114 DoesNotUseOperandBuffer(tuple, {1}, fusion, *points_to_analysis_));
116 EXPECT_TRUE(DoesNotUseOperandBuffer(tuple, {0}, fusion, *dataflow_analysis_));
118 DoesNotUseOperandBuffer(tuple, {1}, fusion, *dataflow_analysis_));
219 auto fusion = computation_->CreateFusionInstruction( local
224 // The fusion instruction can share with tuple element 1.
225 EXPECT_FALSE(CanShareOperandBufferWithUser(tuple, {0}, fusion, {},
227 EXPECT_TRUE(CanShareOperandBufferWithUser(tuple, {1}, fusion, {},
294 auto fusion = computation_->CreateFusionInstruction( local
336 auto fusion = computation_->CreateFusionInstruction( local
367 auto fusion = computation_->CreateFusionInstruction( local
[all...] |
| hlo_verifier.cc | 617 Status HloVerifier::CheckFusionInstruction(HloInstruction* fusion) const {
618 // The parent fusion instruction of the fusion computation must be 'fusion'.
619 HloComputation* fused_computation = fusion->fused_instructions_computation();
620 if (fusion != fused_computation->FusionInstruction()) {
624 fusion->ToString().c_str());
627 // Fused root instruction and fused parameters must all be owned by the fusion
631 fusion->fused_parameters();
632 const HloInstruction* fused_root = fusion->fused_expression_root()
[all...] |
| graphviz_example.cc | 135 // Create a fusion instruction containing the chain of map instructions.
136 auto fusion = builder.AddInstruction(HloInstruction::CreateFusion( local
138 fusion->FuseInstruction(map2);
139 fusion->FuseInstruction(map1);
148 HloInstruction::CreateCall(fusion->shape(), {fusion}, call_computation));
|
| liveness_util.cc | 43 // Find fusion parameter associated with 'operand'.
51 // points-to set of fusion parameter at 'index'.
80 // Find fusion parameter associated with 'operand'.
83 // Iterate through all users of all uses of the fusion parameter value.
133 // in 'fusion.fused_instructions', where the singleton use is the fused
136 // REQUIRES: 'fusion' opcode is a kFusion instruction.
139 HloInstruction* fusion, const int64 use_operand_index,
141 CHECK_EQ(HloOpcode::kFusion, fusion->opcode());
142 // Check that 'operand' is unique in the operand list of 'fusion'.
143 if (fusion->OperandIndices(operand).size() > 1)
[all...] |
| transpose_folding_test.cc | 79 // Instructions after folding: x, y, and the fusion.
87 HloInstruction* fusion = *instruction_set.begin(); local
88 EXPECT_EQ(HloOpcode::kFusion, fusion->opcode());
90 // The fusion instruction should contain two parameters, one transpose and
92 EXPECT_EQ(4, fusion->fused_instruction_count());
130 // The created fusion instruction should contain two parameters, two
194 // Instructions after folding: x, y, and the fusion.
204 HloInstruction* fusion = local
206 EXPECT_EQ(HloOpcode::kFusion, fusion->opcode());
208 // The fusion instruction should contain two parameters, one transpose an
[all...] |
| hlo_cost_analysis_test.cc | 337 // bottleneck time on fusion nodes.
373 auto* fusion = computation->CreateFusionInstruction( local
384 ASSERT_IS_OK(fusion->Accept(&fusion_analysis));
415 auto* fusion = computation->CreateFusionInstruction( local
419 ASSERT_IS_OK(fusion->Accept(&fusion_analysis));
|
| /external/boringssl/src/util/fipstools/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
|
| out.s | 107 # WAS addis 31,2,.LC0@toc@ha # gpr load fusion, type long
121 # WAS addis 19,2,.LC3@toc@ha # gpr load fusion, type long
148 # WAS addis 20,2,.LC6@toc@ha # gpr load fusion, type long
277 # WAS addis 31,2,.LC11@toc@ha # gpr load fusion, type long
327 # WAS addis 6,2,.LC12@toc@ha # gpr load fusion, type long
392 # WAS addis 6,2,.LC13@toc@ha # gpr load fusion, type long
|
| /device/google/contexthub/firmware/variant/argonkey/ |
| argonkey.mk | 35 # Fusion algorithms
36 SRCS_os += os/algos/fusion.c \
|
| /device/google/contexthub/firmware/variant/lunchbox/ |
| lunchbox.mk | 37 # Fusion algorithms
38 SRCS_os += os/algos/fusion.c \
|
| /device/google/contexthub/firmware/variant/neonkey/ |
| neonkey.mk | 68 # Fusion algorithm
69 SRCS_os += os/algos/fusion.c \
|
| /device/google/contexthub/firmware/variant/nucleo/ |
| nucleo.mk | 42 # Fusion algorithm
43 SRCS_os += os/algos/fusion.c \
|
| /device/google/contexthub/firmware/ |
| argonkey_aux_variant_config.mk | 67 os/algos/fusion.c \
|
| lunchbox_aux_variant_config.mk | 76 os/algos/fusion.c \
|
| neonkey_aux_variant_config.mk | 68 os/algos/fusion.c \
|
| nucleo_aux_variant_config.mk | 68 os/algos/fusion.c \
|
| /device/google/contexthub/firmware/os/drivers/orientation/ |
| orientation.c | 30 #include <algos/fusion.h>
113 struct Fusion fusion; member in struct:FusionTask
114 struct Fusion game;
306 osLog(LOG_INFO, "Accel sample has been processed by fusion sensor %d\n",
341 // returns false if addSample fails for any fusion sensor
400 if (fusionHasEstimate(&mTask.fusion)) {
401 fusionGetAttitude(&mTask.fusion, &attitude);
404 fusionGetRotationMatrix(&mTask.fusion, &R);
464 // Keep draining raw samples and producing fusion samples only i
[all...] |
