Home
last modified time | relevance | path

Searched refs:wave_size (Results 1 – 25 of 47) sorted by relevance

12

/third_party/mesa3d/src/amd/vulkan/
Dradv_llvm_helper.cpp31 : family(arg_family), tm_options(arg_tm_options), wave_size(arg_wave_size), passes(NULL) in radv_llvm_per_thread_info()
60 if (arg_family == family && arg_tm_options == tm_options && arg_wave_size == wave_size) in is_same()
69 unsigned wave_size; member in radv_llvm_per_thread_info
101 enum ac_target_machine_options tm_options, unsigned wave_size) in radv_init_llvm_compiler() argument
104 if (I.is_same(family, tm_options, wave_size)) { in radv_init_llvm_compiler()
110 radv_llvm_per_thread_list.emplace_back(family, tm_options, wave_size); in radv_init_llvm_compiler()
Dradv_shader.c1009 info->wave_size, in radv_lower_ngg()
1020 nir, info->wave_size, info->workgroup_size, in radv_lower_ngg()
1367 config_out->rsrc1 = S_00B848_VGPRS((num_vgprs - 1) / (info->wave_size == 32 ? 8 : 4)) | in radv_postprocess_config()
1601 .wave_size = binary->info.wave_size, in radv_shader_variant_create()
1906 info.wave_size = 64;
1918 upload_vs_prolog(struct radv_device *device, struct radv_prolog_binary *bin, unsigned wave_size) argument
1935 prolog->rsrc1 = S_00B848_VGPRS((bin->num_vgprs - 1) / (wave_size == 32 ? 8 : 4)) |
1954 info.wave_size = key->wave32 ? 32 : 64;
1977 struct radv_shader_prolog *prolog = upload_vs_prolog(device, binary, info.wave_size);
2084 uint8_t wave_size = variant->info.wave_size; local
[all …]
Dradv_llvm_helper.h33 enum ac_target_machine_options tm_options, unsigned wave_size);
Dradv_pipeline.c1942 pipeline->device->physical_device->rad_info.chip_class, infos[es_stage].wave_size, in gfx9_get_gs_info()
2116 wavesize = gs_info->wave_size; in gfx10_get_ngg_info()
2118 wavesize = nir[MESA_SHADER_TESS_CTRL] ? infos[MESA_SHADER_TESS_EVAL].wave_size in gfx10_get_ngg_info()
2119 : infos[MESA_SHADER_VERTEX].wave_size; in gfx10_get_ngg_info()
2223 unsigned wave_size = 64; in radv_pipeline_init_gs_ring_state() local
2237 align(gs->vgt_esgs_ring_itemsize * 4 * gs_vertex_reuse * wave_size, alignment); in radv_pipeline_init_gs_ring_state()
2240 max_gs_waves * 2 * wave_size * gs->vgt_esgs_ring_itemsize * 4 * gs_info->gs.vertices_in; in radv_pipeline_init_gs_ring_state()
2241 unsigned gsvs_ring_size = max_gs_waves * 2 * wave_size * gs_info->gs.max_gsvs_emit_size; in radv_pipeline_init_gs_ring_state()
2974 infos[i].wave_size = radv_get_wave_size(pipeline->device, pStages[i], i, &infos[i]); in radv_fill_shader_info()
2982 infos[MESA_SHADER_FRAGMENT].workgroup_size = infos[MESA_SHADER_FRAGMENT].wave_size; in radv_fill_shader_info()
[all …]
/third_party/mesa3d/src/amd/common/
Dac_binary.c39 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 || info->wave64_vgpr_alloc_granularity == 8) 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()
Dac_rtld.h59 unsigned wave_size; member
96 unsigned wave_size; member
Dac_nir.h99 unsigned wave_size,
110 unsigned wave_size,
Dac_nir_lower_ngg.c55 unsigned wave_size; member
84 unsigned wave_size; member
210 unsigned wave_size) in repack_invocations_in_workgroup() argument
220 nir_ssa_def *input_mask = nir_build_ballot(b, 1, wave_size, input_bool); in repack_invocations_in_workgroup()
1196 … nogs_state->max_num_waves, nogs_state->wave_size); in add_deferred_attribute_culling()
1268 unsigned wave_size, in ac_nir_lower_ngg_nogs() argument
1279 assert(max_num_es_vertices && max_workgroup_size && wave_size); in ac_nir_lower_ngg_nogs()
1298 .max_num_waves = DIV_ROUND_UP(max_workgroup_size, wave_size), in ac_nir_lower_ngg_nogs()
1300 .wave_size = wave_size, in ac_nir_lower_ngg_nogs()
1486 …nir_ssa_def *num_threads = nir_bit_count(b, nir_build_ballot(b, 1, s->wave_size, nir_imm_bool(b, t… in ngg_gs_shader_query()
[all …]
Dac_binary.h53 void ac_parse_shader_binary_config(const char *data, size_t nbytes, unsigned wave_size,
Dac_shader_util.h115 unsigned ac_compute_esgs_workgroup_size(enum chip_class chip_class, unsigned wave_size,
Dac_nir_lower_esgs_io_to_mem.c219 unsigned wave_size = 64u; /* GFX6-8 only support wave64 */ in lower_gs_per_vertex_input_load() local
221 return emit_split_buffer_load(b, ring, off, nir_imm_zero(b, 1, 32), 4u * wave_size, in lower_gs_per_vertex_input_load()
Dac_shader_util.c547 unsigned ac_compute_esgs_workgroup_size(enum chip_class chip_class, unsigned wave_size, in ac_compute_esgs_workgroup_size() argument
558 return wave_size; in ac_compute_esgs_workgroup_size()
Dac_rtld.c259 binary->wave_size = i.wave_size; in ac_rtld_open()
529 ac_parse_shader_binary_config(config_data, config_nbytes, binary->wave_size, true, info, &c); in ac_rtld_read_config()
/third_party/mesa3d/src/amd/compiler/
Daco_spill.cpp83 unsigned wave_size; member
90 processed(program->blocks.size(), false), wave_size(program->wave_size) in spill_ctx()
1414 S_008F0C_ADD_TID_ENABLE(1) | S_008F0C_INDEX_STRIDE(ctx.program->wave_size == 64 ? 3 : 2); in load_scratch_resource()
1447 find_available_slot(std::vector<bool>& used, unsigned wave_size, unsigned size, bool is_sgpr, in find_available_slot() argument
1450 unsigned wave_size_minus_one = wave_size - 1; in find_available_slot()
1466 if (is_sgpr && ((slot & wave_size_minus_one) > wave_size - size)) { in find_available_slot()
1467 slot = align(slot, wave_size); in find_available_slot()
1499 find_available_slot(slots_used, ctx.wave_size, ctx.interferences[vec[0]].first.size(), in assign_spill_slots_helper()
1520 find_available_slot(slots_used, ctx.wave_size, ctx.interferences[id].first.size(), in assign_spill_slots_helper()
1574 std::vector<Temp> vgpr_spill_temps((sgpr_spill_slots + ctx.wave_size - 1) / ctx.wave_size); in assign_spill_slots()
[all …]
Daco_live_var_analysis.cpp271 program->workgroup_size == UINT_MAX ? program->wave_size : program->workgroup_size; in calc_waves_per_workgroup()
273 return align(workgroup_size, program->wave_size) / program->wave_size; in calc_waves_per_workgroup()
355 unsigned max_waves_per_simd = program->dev.max_wave64_per_simd * (64 / program->wave_size); in update_vgpr_sgpr_demand()
Daco_ir.cpp88 program->wave_size = info->wave_size; in init_program()
89 program->lane_mask = program->wave_size == 32 ? s1 : s2; in init_program()
104 program->dev.physical_vgprs = program->wave_size == 32 ? 1024 : 512; in init_program()
109 program->dev.vgpr_alloc_granule = program->wave_size == 32 ? 16 : 8; in init_program()
111 program->dev.vgpr_alloc_granule = program->wave_size == 32 ? 8 : 4; in init_program()
Daco_statistics.cpp360 bool dual_issue = program->chip_class >= GFX10 && program->wave_size == 64 && in add()
539 double wave64_per_cycle = waves_per_cycle * (program->wave_size / 64.0); in collect_preasm_stats()
544 program->workgroup_size / (double)align(program->workgroup_size, program->wave_size); in collect_preasm_stats()
/third_party/mesa3d/src/amd/compiler/tests/
Dhelpers.h74 unsigned wave_size=64, enum radeon_family family=CHIP_UNKNOWN);
77 unsigned wave_size=64);
Dhelpers.cpp75 void create_program(enum chip_class chip_class, Stage stage, unsigned wave_size, enum radeon_family… in create_program() argument
78 info.wave_size = wave_size; in create_program()
103 unsigned wave_size) in setup_cs() argument
113 create_program(chip_class, compute_cs, wave_size, family); in setup_cs()
180 program->workgroup_size = program->wave_size; in finish_ra_test()
/third_party/mesa3d/src/gallium/drivers/radeonsi/
Dsi_debug.c854 gl_shader_stage stage, unsigned wave_size) in si_add_split_disasm() argument
859 .wave_size = wave_size, in si_add_split_disasm()
928 unsigned wave_size = si_get_shader_wave_size(shader); in si_print_annotated_shader() local
935 instructions, stage, wave_size); in si_print_annotated_shader()
939 &num_inst, instructions, stage, wave_size); in si_print_annotated_shader()
943 &num_inst, instructions, stage, wave_size); in si_print_annotated_shader()
946 instructions, stage, wave_size); in si_print_annotated_shader()
949 instructions, stage, wave_size); in si_print_annotated_shader()
Dsi_compute_blit.c175 unsigned wave_size = sctx->screen->compute_wave_size; in si_compute_clear_buffer_rmw() local
176 unsigned dwords_per_wave = dwords_per_instruction * wave_size; in si_compute_clear_buffer_rmw()
182 info.block[0] = MIN2(wave_size, num_instructions); in si_compute_clear_buffer_rmw()
260 unsigned wave_size = sctx->screen->compute_wave_size; in si_compute_do_clear_or_copy() local
261 unsigned dwords_per_wave = dwords_per_thread * wave_size; in si_compute_do_clear_or_copy()
267 info.block[0] = MIN2(wave_size, num_instructions); in si_compute_do_clear_or_copy()
Dsi_shader.c810 .wave_size = si_get_shader_wave_size(shader), in si_shader_binary_open()
896 gl_shader_stage stage, unsigned wave_size, in si_shader_dump_disassembly() argument
905 .wave_size = wave_size, in si_shader_dump_disassembly()
1114 unsigned wave_size = si_get_shader_wave_size(shader); in si_shader_dump() local
1119 si_shader_dump_disassembly(sscreen, &shader->prolog->binary, stage, wave_size, debug, in si_shader_dump()
1123 wave_size, debug, "previous stage", file); in si_shader_dump()
1125 si_shader_dump_disassembly(sscreen, &shader->prolog2->binary, stage, wave_size, in si_shader_dump()
1128 si_shader_dump_disassembly(sscreen, &shader->binary, stage, wave_size, debug, "main", in si_shader_dump()
1132 si_shader_dump_disassembly(sscreen, &shader->epilog->binary, stage, wave_size, debug, in si_shader_dump()
1479 unsigned wave_size = sscreen->compute_wave_size; in si_compile_shader() local
[all …]
Dgfx10_shader_ngg.c46 LLVMConstInt(ctx->ac.i32, ctx->ac.wave_size, false), ""); in get_thread_id_in_tg()
481 primemit_scan[stream].maxwaves = DIV_ROUND_UP(256, ctx->ac.wave_size); in build_streamout()
484 ctx->ac.wave_size); in build_streamout()
777 LLVMConstInt(ctx->ac.i32, ctx->ac.wave_size, 0), ""); in update_thread_counts()
781 ac_build_imin(&ctx->ac, *new_num_threads, LLVMConstInt(ctx->ac.i32, ctx->ac.wave_size, 0)); in update_thread_counts()
824 unsigned max_waves = DIV_ROUND_UP(ctx->screen->ngg_subgroup_size, ctx->ac.wave_size); in gfx10_emit_ngg_culling_epilogue()
1129 LLVMConstInt(ctx->ac.i32, ctx->ac.wave_size, 0), ""), in gfx10_emit_ngg_culling_epilogue()
1763 numprims = ac_build_reduce(&ctx->ac, numprims, nir_op_iadd, ctx->ac.wave_size); in gfx10_ngg_gs_emit_epilogue()
1883 vertlive_scan.maxwaves = DIV_ROUND_UP(256, ctx->ac.wave_size); in gfx10_ngg_gs_emit_epilogue()
Dsi_shader_llvm_gs.c152 LLVMConstInt(ctx->ac.i32, ctx->ac.wave_size, false), ""), in si_llvm_emit_es_epilogue()
369 num_records = ctx->ac.wave_size; in si_preload_gs_rings()
374 stream_offset += stride * ctx->ac.wave_size; in si_preload_gs_rings()
/third_party/mesa3d/src/amd/llvm/
Dac_llvm_util.c324 ctx->chip_class >= GFX10 && ctx->wave_size == 64 ? in ac_llvm_set_target_features()

12