| /xsrc/external/mit/MesaLib/dist/src/amd/vulkan/ |
| radv_llvm_helper.h | 33 enum ac_target_machine_options tm_options, unsigned wave_size);
|
| radv_llvm_helper.cpp | 31 : family(arg_family), tm_options(arg_tm_options), wave_size(arg_wave_size), passes(NULL)
60 if (arg_family == family && arg_tm_options == tm_options && arg_wave_size == wave_size)
69 unsigned wave_size; member in class:radv_llvm_per_thread_info
101 enum ac_target_machine_options tm_options, unsigned wave_size)
104 if (I.is_same(family, tm_options, wave_size)) {
110 radv_llvm_per_thread_list.emplace_back(family, tm_options, wave_size);
|
| radv_shader.c | 1009 info->wave_size,
1020 nir, info->wave_size, info->workgroup_size,
1367 config_out->rsrc1 = S_00B848_VGPRS((num_vgprs - 1) / (info->wave_size == 32 ? 8 : 4)) |
1601 .wave_size = binary->info.wave_size,
1906 info.wave_size = 64;
1918 upload_vs_prolog(struct radv_device *device, struct radv_prolog_binary *bin, unsigned wave_size)
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...] |
| radv_pipeline.c | 1942 pipeline->device->physical_device->rad_info.chip_class, infos[es_stage].wave_size,
2116 wavesize = gs_info->wave_size;
2118 wavesize = nir[MESA_SHADER_TESS_CTRL] ? infos[MESA_SHADER_TESS_EVAL].wave_size
2119 : infos[MESA_SHADER_VERTEX].wave_size;
2223 unsigned wave_size = 64; local
2237 align(gs->vgt_esgs_ring_itemsize * 4 * gs_vertex_reuse * wave_size, alignment);
2240 max_gs_waves * 2 * wave_size * gs->vgt_esgs_ring_itemsize * 4 * gs_info->gs.vertices_in;
2241 unsigned gsvs_ring_size = max_gs_waves * 2 * wave_size * gs_info->gs.max_gsvs_emit_size;
2974 infos[i].wave_size = radv_get_wave_size(pipeline->device, pStages[i], i, &infos[i]);
2982 infos[MESA_SHADER_FRAGMENT].workgroup_size = infos[MESA_SHADER_FRAGMENT].wave_size;
[all...] |
| /xsrc/external/mit/MesaLib/dist/src/amd/common/ |
| ac_binary.h | 53 void ac_parse_shader_binary_config(const char *data, size_t nbytes, unsigned wave_size,
|
| ac_nir.h | 99 unsigned wave_size,
110 unsigned wave_size,
|
| ac_rtld.h | 59 unsigned wave_size; member in struct:ac_rtld_binary
96 unsigned wave_size; member in struct:ac_rtld_open_info
|
| ac_binary.c | 39 void ac_parse_shader_binary_config(const char *data, size_t nbytes, unsigned wave_size,
54 if (wave_size == 32 || info->wave64_vgpr_alloc_granularity == 8)
136 conf->num_vgprs = align(conf->num_vgprs, wave_size == 32 ? 16 : 8);
|
| ac_shader_util.h | 115 unsigned ac_compute_esgs_workgroup_size(enum chip_class chip_class, unsigned wave_size,
|
| ac_nir_lower_ngg.c | 55 unsigned wave_size; member in struct:__anon395
84 unsigned wave_size; member in struct:__anon397
210 unsigned wave_size)
220 nir_ssa_def *input_mask = nir_build_ballot(b, 1, wave_size, input_bool);
1196 nogs_state->max_num_waves, nogs_state->wave_size);
1268 unsigned wave_size,
1279 assert(max_num_es_vertices && max_workgroup_size && wave_size);
1298 .max_num_waves = DIV_ROUND_UP(max_workgroup_size, wave_size),
1300 .wave_size = wave_size,
[all...] |
| ac_nir_lower_esgs_io_to_mem.c | 219 unsigned wave_size = 64u; /* GFX6-8 only support wave64 */ local
221 return emit_split_buffer_load(b, ring, off, nir_imm_zero(b, 1, 32), 4u * wave_size,
|
| ac_shader_util.c | 547 unsigned ac_compute_esgs_workgroup_size(enum chip_class chip_class, unsigned wave_size,
558 return wave_size;
|
| ac_rtld.c | 263 binary->wave_size = i.wave_size;
533 ac_parse_shader_binary_config(config_data, config_nbytes, binary->wave_size, true, info, &c);
|
| /xsrc/external/mit/MesaLib/dist/src/amd/compiler/tests/ |
| helpers.h | 74 unsigned wave_size=64, enum radeon_family family=CHIP_UNKNOWN);
77 unsigned wave_size=64);
|
| helpers.cpp | 75 void create_program(enum chip_class chip_class, Stage stage, unsigned wave_size, enum radeon_family family)
78 info.wave_size = wave_size;
103 unsigned wave_size)
113 create_program(chip_class, compute_cs, wave_size, family);
180 program->workgroup_size = program->wave_size;
|
| /xsrc/external/mit/MesaLib/dist/src/amd/compiler/ |
| aco_spill.cpp | 83 unsigned wave_size; member in struct:aco::__anon500::spill_ctx
90 processed(program->blocks.size(), false), wave_size(program->wave_size)
1414 S_008F0C_ADD_TID_ENABLE(1) | S_008F0C_INDEX_STRIDE(ctx.program->wave_size == 64 ? 3 : 2);
1447 find_available_slot(std::vector<bool>& used, unsigned wave_size, unsigned size, bool is_sgpr,
1450 unsigned wave_size_minus_one = wave_size - 1;
1466 if (is_sgpr && ((slot & wave_size_minus_one) > wave_size - size)) {
1467 slot = align(slot, wave_size);
1499 find_available_slot(slots_used, ctx.wave_size, ctx.interferences[vec[0]].first.size(),
1520 find_available_slot(slots_used, ctx.wave_size, ctx.interferences[id].first.size()
[all...] |
| aco_live_var_analysis.cpp | 269 /* When workgroup size is not known, just go with wave_size */
271 program->workgroup_size == UINT_MAX ? program->wave_size : program->workgroup_size;
273 return align(workgroup_size, program->wave_size) / program->wave_size;
355 unsigned max_waves_per_simd = program->dev.max_wave64_per_simd * (64 / program->wave_size);
|
| aco_ir.cpp | 88 program->wave_size = info->wave_size;
89 program->lane_mask = program->wave_size == 32 ? s1 : s2;
104 program->dev.physical_vgprs = program->wave_size == 32 ? 1024 : 512;
109 program->dev.vgpr_alloc_granule = program->wave_size == 32 ? 16 : 8;
111 program->dev.vgpr_alloc_granule = program->wave_size == 32 ? 8 : 4;
|
| aco_statistics.cpp | 360 bool dual_issue = program->chip_class >= GFX10 && program->wave_size == 64 &&
539 double wave64_per_cycle = waves_per_cycle * (program->wave_size / 64.0);
544 program->workgroup_size / (double)align(program->workgroup_size, program->wave_size);
|
| aco_print_asm.cpp | 216 if (program->chip_class >= GFX10 && program->wave_size == 64) {
|
| /xsrc/external/mit/MesaLib/dist/src/gallium/drivers/radeonsi/ |
| si_debug.c | 854 gl_shader_stage stage, unsigned wave_size)
859 .wave_size = wave_size,
928 unsigned wave_size = si_get_shader_wave_size(shader); local
935 instructions, stage, wave_size);
939 &num_inst, instructions, stage, wave_size);
943 &num_inst, instructions, stage, wave_size);
946 instructions, stage, wave_size);
949 instructions, stage, wave_size);
|
| si_compute_blit.c | 175 unsigned wave_size = sctx->screen->compute_wave_size; local
176 unsigned dwords_per_wave = dwords_per_instruction * wave_size;
182 info.block[0] = MIN2(wave_size, num_instructions);
260 unsigned wave_size = sctx->screen->compute_wave_size; local
261 unsigned dwords_per_wave = dwords_per_thread * wave_size;
267 info.block[0] = MIN2(wave_size, num_instructions);
|
| si_shader.c | 810 .wave_size = si_get_shader_wave_size(shader),
896 gl_shader_stage stage, unsigned wave_size,
905 .wave_size = wave_size,
1114 unsigned wave_size = si_get_shader_wave_size(shader); local
1119 si_shader_dump_disassembly(sscreen, &shader->prolog->binary, stage, wave_size, debug,
1123 wave_size, debug, "previous stage", file);
1125 si_shader_dump_disassembly(sscreen, &shader->prolog2->binary, stage, wave_size,
1128 si_shader_dump_disassembly(sscreen, &shader->binary, stage, wave_size, debug, "main",
1132 si_shader_dump_disassembly(sscreen, &shader->epilog->binary, stage, wave_size, debug
1479 unsigned wave_size = sscreen->compute_wave_size; local
[all...] |
| si_shader_internal.h | 203 struct ac_llvm_compiler *compiler, unsigned wave_size);
|
| /xsrc/external/mit/MesaLib/dist/src/amd/llvm/ |
| ac_llvm_util.c | 324 ctx->chip_class >= GFX10 && ctx->wave_size == 64 ?
|
