| /external/mesa3d/src/amd/compiler/ |
| D | aco_lower_phis.cpp | 56 return Operand(Temp(it->second, program->lane_mask)); in get_ssa()
66 return Operand(program->lane_mask); in get_ssa()
77 return Operand(program->lane_mask); in get_ssa()
79 Temp res = Temp(program->allocateTmp(program->lane_mask)); in get_ssa()
203 std::fill(state->latest.begin(), state->latest.end(), Operand(program->lane_mask)); in lower_divergent_bool_phi()
210 state->writes[block->logical_preds[i]] = program->allocateId(program->lane_mask); in lower_divergent_bool_phi()
230 Temp new_cur = {state->writes.at(pred->index), program->lane_mask}; in lower_divergent_bool_phi()
288 if (phi->definitions[0].regClass() == program->lane_mask) in lower_phis()
|
| D | aco_live_var_analysis.cpp | 115 assert(new_demand.sgpr >= (int16_t) program->lane_mask.size()); in process_live_temps_per_block()
116 …] = RegisterDemand(new_demand.vgpr, new_demand.sgpr - (exec_live ? program->lane_mask.size() : 0)); in process_live_temps_per_block()
183 assert(new_demand.sgpr >= (int16_t) program->lane_mask.size()); in process_live_temps_per_block()
184 new_demand.sgpr -= exec_live ? program->lane_mask.size() : 0; in process_live_temps_per_block()
|
| D | aco_ir.cpp | 94 program->lane_mask = program->wave_size == 32 ? s1 : s2; in init_program()
|
| D | aco_instruction_selection_setup.cpp | 318 return RegClass(RegType::sgpr, ctx->program->lane_mask.size() * components); in get_reg_class()
577 unsigned lane_mask_size = ctx->program->lane_mask.size(); in init_context()
|
| D | aco_optimizer.cpp | 1621 if (instr->definitions[0].regClass() != ctx.program->lane_mask) in combine_ordering_test()
1721 if (instr->definitions[0].regClass() != ctx.program->lane_mask) in combine_comparison_ordering()
1818 if (instr->definitions[0].regClass() != ctx.program->lane_mask) in combine_constant_comparison_ordering()
2285 Definition(ctx.program->allocateTmp(ctx.program->lane_mask)); in combine_add_sub_b2i()
|
| D | aco_ir.h | 1627 RegClass lane_mask; variable
|
| D | aco_instruction_selection.cpp | 938 assert(dst.regClass() == ctx->program->lane_mask); in emit_comparison()
9307 assert(instr->dest.ssa.bit_size != 1 || dst.regClass() == ctx->program->lane_mask); in visit_phi()
9796 assert(cond.regClass() == ctx->program->lane_mask); in begin_divergent_if_then()
10066 assert(cond.regClass() == ctx->program->lane_mask); in visit_if()
10879 …nt] = Definition{ctx->program->allocateId(ctx->program->lane_mask), exec, ctx->program->lane_mask}; in add_startpgm()
|
| /external/tensorflow/tensorflow/core/util/ |
| D | gpu_device_functions.h | 276 __device__ inline unsigned GpuShuffleXorGetSrcLane(int lane_mask, int width) { in GpuShuffleXorGetSrcLane() argument
278 int src_lane = lane_id ^ lane_mask; in GpuShuffleXorGetSrcLane()
464 __device__ T GpuShuffleXorSync(unsigned mask, T value, int lane_mask,
468 mask, detail::GpuShuffleXorGetSrcLane(lane_mask, width)));
471 return __shfl_xor_sync(mask, value, lane_mask, width);
473 return __shfl_xor(value, lane_mask, width);
477 return __shfl_xor(static_cast<int>(value), lane_mask, width);
484 int lane_mask,
488 mask, detail::GpuShuffleXorGetSrcLane(lane_mask, width)));
492 __shfl_xor(static_cast<float>(value), lane_mask, width));
[all …]
|
| D | gpu_kernel_helper.h | 192 unsigned mask, Eigen::half value, int lane_mask, int width = warpSize) {
194 GpuShuffleXorSync(mask, static_cast<uint16>(value), lane_mask, width));
|
| /external/vixl/src/aarch64/ |
| D | simulator-aarch64.cc | 823 uint16_t lane_mask, in PrintRegisterValueFPAnnotations() argument
831 bool access = (lane_mask & (1 << (i * lane_size))) != 0; in PrintRegisterValueFPAnnotations()
1333 uint16_t lane_mask = GetPrintRegLaneMask(format); in PrintVStructAccess() local
1334 PrintVRegistersForStructuredAccess(rt_code, reg_count, lane_mask, format); in PrintVStructAccess()
1340 VIXL_ASSERT((lane_mask & access_mask) != 0); in PrintVStructAccess()
1341 lane_mask = PrintPartialAccess(access_mask, in PrintVStructAccess()
1342 lane_mask, in PrintVStructAccess()
1364 uint16_t lane_mask = 1 << (lane * lane_size_in_bytes); in PrintVSingleStructAccess() local
1365 PrintVRegistersForStructuredAccess(rt_code, reg_count, lane_mask, format); in PrintVSingleStructAccess()
1366 PrintPartialAccess(lane_mask, 0, reg_count, lane_size_in_bytes, op, address); in PrintVSingleStructAccess()
[all …]
|
| D | simulator-aarch64.h | 2378 uint16_t lane_mask,
2382 uint16_t lane_mask,
2384 PrintRegisterValueFPAnnotations(sim_register.GetBytes(), lane_mask, format);
|
| /external/mesa3d/src/amd/compiler/tests/ |
| D | helpers.cpp | 116 exec_input = bld.tmp(program->lane_mask); in setup_cs()
|