| /external/mesa3d/src/amd/vulkan/ |
| D | radv_llvm_helper.cpp | 34 wave_size(arg_wave_size), passes(NULL), passes_wave32(NULL) {} in radv_llvm_per_thread_info()
64 struct ac_compiler_passes *p = wave_size == 32 ? passes_wave32 : passes; in compile_to_memory_buffer()
73 arg_wave_size == wave_size) in is_same()
81 unsigned wave_size; member in radv_llvm_per_thread_info
116 unsigned wave_size) in radv_init_llvm_compiler() argument
120 if (I.is_same(family, tm_options, wave_size)) { in radv_init_llvm_compiler()
126 radv_llvm_per_thread_list.emplace_back(family, tm_options, wave_size); in radv_init_llvm_compiler()
|
| D | radv_shader_helper.h | 33 unsigned wave_size);
|
| D | radv_shader.c | 914 (info->wave_size == 32 ? 8 : 4)) | in radv_postprocess_config()
1153 .wave_size = binary->info.wave_size, in radv_shader_variant_create()
1437 info.wave_size = 64; in radv_create_trap_handler_shader()
1529 uint8_t wave_size = variant->info.wave_size; in radv_get_max_waves() local
1544 DIV_ROUND_UP(max_workgroup_size, wave_size); in radv_get_max_waves()
1556 unsigned vgprs = align(conf->num_vgprs, wave_size == 32 ? 8 : 4); in radv_get_max_waves()
|
| D | radv_pipeline.c | 2068 wavesize = gs_info->wave_size; in gfx10_get_ngg_info()
2071 ? infos[MESA_SHADER_TESS_EVAL].wave_size in gfx10_get_ngg_info()
2072 : infos[MESA_SHADER_VERTEX].wave_size; in gfx10_get_ngg_info()
2161 unsigned wave_size = 64; in radv_pipeline_init_gs_ring_state() local
2175 wave_size, alignment); in radv_pipeline_init_gs_ring_state()
2177 unsigned esgs_ring_size = max_gs_waves * 2 * wave_size * in radv_pipeline_init_gs_ring_state()
2179 unsigned gsvs_ring_size = max_gs_waves * 2 * wave_size * in radv_pipeline_init_gs_ring_state()
2843 infos[i].wave_size = in radv_fill_shader_info()
3313 info.wave_size = 64; /* Wave32 not supported. */ in radv_create_shaders()
4874 S_0286D8_PS_W32_EN(ps->info.wave_size == 32)); in radv_pipeline_generate_fragment_shader()
[all …]
|
| /external/mesa3d/src/amd/common/ |
| D | ac_binary.c | 39 void ac_parse_shader_binary_config(const char *data, size_t nbytes, unsigned wave_size, in ac_parse_shader_binary_config() argument
54 if (wave_size == 32) in ac_parse_shader_binary_config()
136 conf->num_vgprs = align(conf->num_vgprs, wave_size == 32 ? 16 : 8); in ac_parse_shader_binary_config()
|
| D | ac_rtld.h | 59 unsigned wave_size; member
96 unsigned wave_size; member
|
| D | ac_binary.h | 53 void ac_parse_shader_binary_config(const char *data, size_t nbytes, unsigned wave_size,
|
| D | ac_rtld.c | 260 binary->wave_size = i.wave_size; in ac_rtld_open()
524 ac_parse_shader_binary_config(config_data, config_nbytes, binary->wave_size, true, info, &c); in ac_rtld_read_config()
|
| /external/mesa3d/src/amd/compiler/ |
| D | aco_ir.cpp | 93 program->wave_size = info->wave_size; in init_program()
94 program->lane_mask = program->wave_size == 32 ? s1 : s2; in init_program()
109 program->vgpr_alloc_granule = program->wave_size == 32 ? 15 : 7; in init_program()
111 program->vgpr_alloc_granule = program->wave_size == 32 ? 7 : 3; in init_program()
|
| D | aco_spill.cpp | 64 unsigned wave_size; member
71 processed(program->blocks.size(), false), wave_size(program->wave_size) {} in spill_ctx()
1337 S_008F0C_INDEX_STRIDE(ctx.program->wave_size == 64 ? 3 : 2); in load_scratch_resource()
1370 unsigned find_available_slot(std::vector<bool>& used, unsigned wave_size, in find_available_slot() argument
1373 unsigned wave_size_minus_one = wave_size - 1; in find_available_slot()
1389 if (is_sgpr && ((slot & wave_size_minus_one) > wave_size - size)) { in find_available_slot()
1390 slot = align(slot, wave_size); in find_available_slot()
1422 unsigned slot = find_available_slot(slots_used, ctx.wave_size, in assign_spill_slots_helper()
1443 unsigned slot = find_available_slot(slots_used, ctx.wave_size, in assign_spill_slots_helper()
1495 std::vector<Temp> vgpr_spill_temps((sgpr_spill_slots + ctx.wave_size - 1) / ctx.wave_size); in assign_spill_slots()
[all …]
|
| D | aco_live_var_analysis.cpp | 260 ? program->wave_size in calc_waves_per_workgroup()
263 return align(workgroup_size, program->wave_size) / program->wave_size; in calc_waves_per_workgroup()
325 if (program->wave_size == 32) in calc_min_waves()
|
| D | aco_lower_to_hw_instr.cpp | 488 assert(cluster_size == ctx->program->wave_size || op == aco_opcode::p_reduce); in emit_reduction()
489 assert(cluster_size <= ctx->program->wave_size); in emit_reduction()
653 if (ctx->program->wave_size == 64) { in emit_reduction()
706 assert(cluster_size == ctx->program->wave_size); in emit_reduction()
764 if (ctx->program->wave_size == 64) { in emit_reduction()
799 Operand(PhysReg{tmp + k}, v1), Operand(ctx->program->wave_size - 1)); in emit_reduction()
819 assert(program->wave_size == 64); in emit_gfx10_wave64_bpermute()
905 for (unsigned n = 0; n < program->wave_size; ++n) { in emit_gfx6_bpermute()
1932 else if (ctx.program->chip_class >= GFX10 && ctx.program->wave_size == 64) in lower_to_hw_instr()
1998 program->workgroup_size > program->wave_size; in lower_to_hw_instr()
|
| D | aco_print_asm.cpp | 177 if (program->chip_class >= GFX10 && program->wave_size == 64) { in print_asm()
|
| /external/mesa3d/src/amd/compiler/tests/ |
| D | helpers.h | 72 unsigned wave_size=64, enum radeon_family family=CHIP_UNKNOWN);
74 enum radeon_family family=CHIP_UNKNOWN, unsigned wave_size=64);
|
| D | helpers.cpp | 77 void create_program(enum chip_class chip_class, Stage stage, unsigned wave_size, enum radeon_family… in create_program() argument
80 info.wave_size = wave_size; in create_program()
94 enum radeon_family family, unsigned wave_size) in setup_cs() argument
106 create_program(chip_class, compute_cs, wave_size, family); in setup_cs()
|
| /external/mesa3d/src/gallium/drivers/radeonsi/ |
| D | si_debug.c | 877 gl_shader_stage stage, unsigned wave_size) in si_add_split_disasm() argument
882 .wave_size = wave_size, in si_add_split_disasm()
951 unsigned wave_size = si_get_shader_wave_size(shader); in si_print_annotated_shader() local
958 instructions, stage, wave_size); in si_print_annotated_shader()
962 &num_inst, instructions, stage, wave_size); in si_print_annotated_shader()
966 &num_inst, instructions, stage, wave_size); in si_print_annotated_shader()
969 instructions, stage, wave_size); in si_print_annotated_shader()
972 instructions, stage, wave_size); in si_print_annotated_shader()
|
| D | si_shader_llvm.c | 96 if (ac->wave_size == 32) in si_compile_llvm()
118 .wave_size = ac->wave_size, in si_compile_llvm()
130 struct ac_llvm_compiler *compiler, unsigned wave_size) in si_llvm_context_init() argument
137 AC_FLOAT_MODE_DEFAULT_OPENGL, wave_size, 64); in si_llvm_context_init()
|
| D | si_shader.c | 832 .wave_size = si_get_shader_wave_size(shader), in si_shader_binary_open()
909 gl_shader_stage stage, unsigned wave_size, in si_shader_dump_disassembly() argument
918 .wave_size = wave_size, in si_shader_dump_disassembly()
1129 unsigned wave_size = si_get_shader_wave_size(shader); in si_shader_dump() local
1134 si_shader_dump_disassembly(sscreen, &shader->prolog->binary, stage, wave_size, debug, in si_shader_dump()
1138 wave_size, debug, "previous stage", file); in si_shader_dump()
1140 si_shader_dump_disassembly(sscreen, &shader->prolog2->binary, stage, wave_size, in si_shader_dump()
1143 si_shader_dump_disassembly(sscreen, &shader->binary, stage, wave_size, debug, "main", in si_shader_dump()
1147 si_shader_dump_disassembly(sscreen, &shader->epilog->binary, stage, wave_size, debug, in si_shader_dump()
1933 unsigned wave_size = sscreen->compute_wave_size; in si_compile_shader() local
[all …]
|
| D | si_shader_llvm_gs.c | 165 LLVMConstInt(ctx->ac.i32, ctx->ac.wave_size, false), ""), in si_llvm_emit_es_epilogue()
383 num_records = ctx->ac.wave_size; in si_preload_gs_rings()
388 stream_offset += stride * ctx->ac.wave_size; in si_preload_gs_rings()
|
| D | si_compute_blit.c | 216 unsigned wave_size = sctx->screen->compute_wave_size; in si_compute_do_clear_or_copy() local
217 unsigned dwords_per_wave = dwords_per_thread * wave_size; in si_compute_do_clear_or_copy()
223 info.block[0] = MIN2(wave_size, num_instructions); in si_compute_do_clear_or_copy()
|
| D | si_shader_internal.h | 228 struct ac_llvm_compiler *compiler, unsigned wave_size);
|
| D | si_state_draw.c | 177 unsigned wave_size = sctx->screen->ge_wave_size; in si_emit_derived_tess_state() local
179 if (temp_verts_per_tg > wave_size && temp_verts_per_tg % wave_size < wave_size * 3 / 4) in si_emit_derived_tess_state()
180 *num_patches = (temp_verts_per_tg & ~(wave_size - 1)) / max_verts_per_patch; in si_emit_derived_tess_state()
186 unsigned one_wave = wave_size / max_verts_per_patch; in si_emit_derived_tess_state()
|
| D | gfx10_shader_ngg.c | 46 LLVMConstInt(ctx->ac.i32, ctx->ac.wave_size, false), ""); in get_thread_id_in_tg()
717 LLVMConstInt(ctx->ac.i32, ctx->ac.wave_size, 0), ""); in update_thread_counts()
721 ac_build_imin(&ctx->ac, *new_num_threads, LLVMConstInt(ctx->ac.i32, ctx->ac.wave_size, 0)); in update_thread_counts()
744 unsigned max_waves = ctx->ac.wave_size == 64 ? 2 : 4; in gfx10_emit_ngg_culling_epilogue()
747 if (ctx->ac.wave_size == 64) { in gfx10_emit_ngg_culling_epilogue()
1022 LLVMConstInt(ctx->ac.i32, ctx->ac.wave_size, 0), ""), in gfx10_emit_ngg_culling_epilogue()
1674 numprims = ac_build_reduce(&ctx->ac, numprims, nir_op_iadd, ctx->ac.wave_size); in gfx10_ngg_gs_emit_epilogue()
|
| /external/mesa3d/src/amd/llvm/ |
| D | ac_llvm_build.c | 60 enum ac_float_mode float_mode, unsigned wave_size, in ac_llvm_context_init() argument
67 ctx->wave_size = wave_size; in ac_llvm_context_init()
71 ac_create_module(wave_size == 32 ? compiler->tm_wave32 : compiler->tm, ctx->context); in ac_llvm_context_init()
96 ctx->iN_wavemask = LLVMIntTypeInContext(ctx->context, ctx->wave_size); in ac_llvm_context_init()
445 if (ctx->wave_size == 64) in ac_build_ballot()
471 if (ctx->wave_size == 64) in ac_get_i1_sgpr_mask()
1754 if (ctx->wave_size == 32) { in ac_get_thread_id()
1760 set_range_metadata(ctx, tid, 0, ctx->wave_size); in ac_get_thread_id()
3385 if (ctx->wave_size == 32) { in ac_build_mbcnt()
4007 result = ac_build_scan(ctx, op, result, identity, ctx->wave_size, true); in ac_build_inclusive_scan()
[all …]
|
| D | ac_llvm_build.h | 136 unsigned wave_size; member
146 enum ac_float_mode float_mode, unsigned wave_size,
|