X-Git-Url: https://git.libre-soc.org/?a=blobdiff_plain;f=src%2Famd%2Fvulkan%2Fradv_nir_to_llvm.c;h=84ab586caacf41b8112c25c6a072142e5b25a174;hb=2a83154b4aa4ae9a86134bfd66f315c91e91302f;hp=a589ff41559911913c8ec88c41a1e816d86437c9;hpb=c0110477b5503978adf37f64c65c30c9304e3fe4;p=mesa.git diff --git a/src/amd/vulkan/radv_nir_to_llvm.c b/src/amd/vulkan/radv_nir_to_llvm.c index a589ff41559..84ab586caac 100644 --- a/src/amd/vulkan/radv_nir_to_llvm.c +++ b/src/amd/vulkan/radv_nir_to_llvm.c @@ -36,7 +36,6 @@ #include #include "sid.h" -#include "gfx9d.h" #include "ac_binary.h" #include "ac_llvm_util.h" #include "ac_llvm_build.h" @@ -262,7 +261,7 @@ get_tcs_num_patches(struct radv_shader_context *ctx) * * Test: dEQP-VK.tessellation.shader_input_output.barrier */ - if (ctx->options->chip_class >= CIK && ctx->options->family != CHIP_STONEY) + if (ctx->options->chip_class >= GFX7 && ctx->options->family != CHIP_STONEY) hardware_lds_size = 65536; num_patches = MIN2(num_patches, hardware_lds_size / (input_patch_size + output_patch_size)); @@ -273,8 +272,8 @@ get_tcs_num_patches(struct radv_shader_context *ctx) */ num_patches = MIN2(num_patches, 40); - /* SI bug workaround - limit LS-HS threadgroups to only one wave. */ - if (ctx->options->chip_class == SI) { + /* GFX6 bug workaround - limit LS-HS threadgroups to only one wave. */ + if (ctx->options->chip_class == GFX6) { unsigned one_wave = 64 / MAX2(num_tcs_input_cp, num_tcs_output_cp); num_patches = MIN2(num_patches, one_wave); } @@ -518,11 +517,8 @@ create_llvm_function(LLVMContextRef ctx, LLVMModuleRef module, options->address32_hi); } - if (max_workgroup_size) { - ac_llvm_add_target_dep_function_attr(main_function, - "amdgpu-max-work-group-size", - max_workgroup_size); - } + ac_llvm_set_workgroup_size(main_function, max_workgroup_size); + if (options->unsafe_math) { /* These were copied from some LLVM test. */ LLVMAddTargetDependentFunctionAttr(main_function, @@ -1305,13 +1301,35 @@ radv_load_resource(struct ac_shader_abi *abi, LLVMValueRef index, } else stride = LLVMConstInt(ctx->ac.i32, layout->binding[binding].size, false); - offset = ac_build_imad(&ctx->ac, index, stride, - LLVMConstInt(ctx->ac.i32, base_offset, false)); + offset = LLVMConstInt(ctx->ac.i32, base_offset, false); - desc_ptr = ac_build_gep0(&ctx->ac, desc_ptr, offset); + if (layout->binding[binding].type != VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK_EXT) { + offset = ac_build_imad(&ctx->ac, index, stride, offset); + } + + desc_ptr = LLVMBuildGEP(ctx->ac.builder, desc_ptr, &offset, 1, ""); desc_ptr = ac_cast_ptr(&ctx->ac, desc_ptr, ctx->ac.v4i32); LLVMSetMetadata(desc_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md); + if (layout->binding[binding].type == VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK_EXT) { + uint32_t desc_type = S_008F0C_DST_SEL_X(V_008F0C_SQ_SEL_X) | + S_008F0C_DST_SEL_Y(V_008F0C_SQ_SEL_Y) | + S_008F0C_DST_SEL_Z(V_008F0C_SQ_SEL_Z) | + S_008F0C_DST_SEL_W(V_008F0C_SQ_SEL_W) | + S_008F0C_NUM_FORMAT(V_008F0C_BUF_NUM_FORMAT_FLOAT) | + S_008F0C_DATA_FORMAT(V_008F0C_BUF_DATA_FORMAT_32); + + LLVMValueRef desc_components[4] = { + LLVMBuildPtrToInt(ctx->ac.builder, desc_ptr, ctx->ac.intptr, ""), + LLVMConstInt(ctx->ac.i32, S_008F04_BASE_ADDRESS_HI(ctx->options->address32_hi), false), + /* High limit to support variable sizes. */ + LLVMConstInt(ctx->ac.i32, 0xffffffff, false), + LLVMConstInt(ctx->ac.i32, desc_type, false), + }; + + return ac_build_gather_values(&ctx->ac, desc_components, 4); + } + return desc_ptr; } @@ -1567,13 +1585,13 @@ store_tcs_output(struct ac_shader_abi *abi, if (!is_tess_factor && writemask != 0xF) ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, value, 1, buf_addr, ctx->oc_lds, - 4 * (base + chan), 1, 0, true, false); + 4 * (base + chan), ac_glc, false); } if (writemask == 0xF) { ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, src, 4, buf_addr, ctx->oc_lds, - (base * 4), 1, 0, true, false); + (base * 4), ac_glc, false); } } @@ -1612,7 +1630,7 @@ load_tes_input(struct ac_shader_abi *abi, buf_addr = LLVMBuildAdd(ctx->ac.builder, buf_addr, comp_offset, ""); result = ac_build_buffer_load(&ctx->ac, ctx->hs_ring_tess_offchip, num_components, NULL, - buf_addr, ctx->oc_lds, is_compact ? (4 * const_index) : 0, 1, 0, true, false); + buf_addr, ctx->oc_lds, is_compact ? (4 * const_index) : 0, ac_glc, true, false); result = ac_trim_vector(&ctx->ac, result, num_components); return result; } @@ -1655,7 +1673,7 @@ load_gs_input(struct ac_shader_abi *abi, ctx->esgs_ring, 1, ctx->ac.i32_0, vtx_offset, soffset, - 0, 1, 0, true, false); + 0, ac_glc, true, false); } if (ac_get_type_size(type) == 2) { @@ -1733,7 +1751,8 @@ static LLVMValueRef load_sample_position(struct ac_shader_abi *abi, struct radv_shader_context *ctx = radv_shader_context_from_abi(abi); LLVMValueRef result; - LLVMValueRef ptr = ac_build_gep0(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_PS_SAMPLE_POSITIONS, false)); + LLVMValueRef index = LLVMConstInt(ctx->ac.i32, RING_PS_SAMPLE_POSITIONS, false); + LLVMValueRef ptr = LLVMBuildGEP(ctx->ac.builder, ctx->ring_offsets, &index, 1, ""); ptr = LLVMBuildBitCast(ctx->ac.builder, ptr, ac_array_in_const_addr_space(ctx->ac.v2f32), ""); @@ -1839,7 +1858,7 @@ visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addr ctx->gsvs_ring[stream], out_val, 1, voffset, ctx->gs2vs_offset, 0, - 1, 1, true, true); + ac_glc | ac_slc, true); } } @@ -1910,6 +1929,11 @@ static LLVMValueRef radv_load_ubo(struct ac_shader_abi *abi, LLVMValueRef buffer struct radv_shader_context *ctx = radv_shader_context_from_abi(abi); LLVMValueRef result; + if (LLVMGetTypeKind(LLVMTypeOf(buffer_ptr)) != LLVMPointerTypeKind) { + /* Do not load the descriptor for inlined uniform blocks. */ + return buffer_ptr; + } + LLVMSetMetadata(buffer_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md); result = LLVMBuildLoad(ctx->ac.builder, buffer_ptr, ""); @@ -1951,8 +1975,9 @@ static LLVMValueRef radv_get_sampler_desc(struct ac_shader_abi *abi, break; case AC_DESC_SAMPLER: type = ctx->ac.v4i32; - if (binding->type == VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER) - offset += 64; + if (binding->type == VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER) { + offset += radv_combined_image_descriptor_sampler_offset(binding); + } type_size = 16; break; @@ -1960,6 +1985,13 @@ static LLVMValueRef radv_get_sampler_desc(struct ac_shader_abi *abi, type = ctx->ac.v4i32; type_size = 16; break; + case AC_DESC_PLANE_0: + case AC_DESC_PLANE_1: + case AC_DESC_PLANE_2: + type = ctx->ac.v8i32; + type_size = 32; + offset += 32 * (desc_type - AC_DESC_PLANE_0); + break; default: unreachable("invalid desc_type\n"); } @@ -1984,16 +2016,35 @@ static LLVMValueRef radv_get_sampler_desc(struct ac_shader_abi *abi, assert(stride % type_size == 0); - if (!index) - index = ctx->ac.i32_0; + LLVMValueRef adjusted_index = index; + if (!adjusted_index) + adjusted_index = ctx->ac.i32_0; - index = LLVMBuildMul(builder, index, LLVMConstInt(ctx->ac.i32, stride / type_size, 0), ""); + adjusted_index = LLVMBuildMul(builder, adjusted_index, LLVMConstInt(ctx->ac.i32, stride / type_size, 0), ""); - list = ac_build_gep0(&ctx->ac, list, LLVMConstInt(ctx->ac.i32, offset, 0)); + LLVMValueRef val_offset = LLVMConstInt(ctx->ac.i32, offset, 0); + list = LLVMBuildGEP(builder, list, &val_offset, 1, ""); list = LLVMBuildPointerCast(builder, list, ac_array_in_const32_addr_space(type), ""); - return ac_build_load_to_sgpr(&ctx->ac, list, index); + LLVMValueRef descriptor = ac_build_load_to_sgpr(&ctx->ac, list, adjusted_index); + + /* 3 plane formats always have same size and format for plane 1 & 2, so + * use the tail from plane 1 so that we can store only the first 16 bytes + * of the last plane. */ + if (desc_type == AC_DESC_PLANE_2) { + LLVMValueRef descriptor2 = radv_get_sampler_desc(abi, descriptor_set, base_index, constant_index, index, AC_DESC_PLANE_1,image, write, bindless); + + LLVMValueRef components[8]; + for (unsigned i = 0; i < 4; ++i) + components[i] = ac_llvm_extract_elem(&ctx->ac, descriptor, i); + + for (unsigned i = 4; i < 8; ++i) + components[i] = ac_llvm_extract_elem(&ctx->ac, descriptor2, i); + descriptor = ac_build_gather_values(&ctx->ac, components, 8); + } + + return descriptor; } /* For 2_10_10_10 formats the alpha is handled as unsigned by pre-vega HW. @@ -2008,6 +2059,8 @@ adjust_vertex_fetch_alpha(struct radv_shader_context *ctx, LLVMValueRef c30 = LLVMConstInt(ctx->ac.i32, 30, 0); + alpha = LLVMBuildBitCast(ctx->ac.builder, alpha, ctx->ac.f32, ""); + if (adjustment == RADV_ALPHA_ADJUST_SSCALED) alpha = LLVMBuildFPToUI(ctx->ac.builder, alpha, ctx->ac.i32, ""); else @@ -2035,7 +2088,7 @@ adjust_vertex_fetch_alpha(struct radv_shader_context *ctx, alpha = LLVMBuildSIToFP(ctx->ac.builder, alpha, ctx->ac.f32, ""); } - return alpha; + return LLVMBuildBitCast(ctx->ac.builder, alpha, ctx->ac.i32, ""); } static unsigned @@ -2096,7 +2149,7 @@ radv_fixup_vertex_input_fetches(struct radv_shader_context *ctx, for (unsigned i = num_channels; i < 4; i++) { chan[i] = i == 3 ? one : zero; - chan[i] = ac_to_float(&ctx->ac, chan[i]); + chan[i] = ac_to_integer(&ctx->ac, chan[i]); } return ac_build_gather_values(&ctx->ac, chan, 4); @@ -2138,14 +2191,6 @@ handle_vs_input_decl(struct radv_shader_context *ctx, buffer_index = LLVMBuildUDiv(ctx->ac.builder, buffer_index, LLVMConstInt(ctx->ac.i32, divisor, 0), ""); } - - if (ctx->options->key.vs.as_ls) { - ctx->shader_info->vs.vgpr_comp_cnt = - MAX2(2, ctx->shader_info->vs.vgpr_comp_cnt); - } else { - ctx->shader_info->vs.vgpr_comp_cnt = - MAX2(1, ctx->shader_info->vs.vgpr_comp_cnt); - } } else { buffer_index = ctx->ac.i32_0; } @@ -2154,20 +2199,54 @@ handle_vs_input_decl(struct radv_shader_context *ctx, } else buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.vertex_id, ctx->abi.base_vertex, ""); - t_offset = LLVMConstInt(ctx->ac.i32, attrib_index, false); - - t_list = ac_build_load_to_sgpr(&ctx->ac, t_list_ptr, t_offset); /* Adjust the number of channels to load based on the vertex * attribute format. */ unsigned num_format_channels = get_num_channels_from_data_format(data_format); unsigned num_channels = MIN2(num_input_channels, num_format_channels); + unsigned attrib_binding = ctx->options->key.vs.vertex_attribute_bindings[attrib_index]; + unsigned attrib_offset = ctx->options->key.vs.vertex_attribute_offsets[attrib_index]; + unsigned attrib_stride = ctx->options->key.vs.vertex_attribute_strides[attrib_index]; + + if (ctx->options->key.vs.post_shuffle & (1 << attrib_index)) { + /* Always load, at least, 3 channels for formats that + * need to be shuffled because X<->Z. + */ + num_channels = MAX2(num_channels, 3); + } - input = ac_build_buffer_load_format(&ctx->ac, t_list, + if (attrib_stride != 0 && attrib_offset > attrib_stride) { + LLVMValueRef buffer_offset = + LLVMConstInt(ctx->ac.i32, + attrib_offset / attrib_stride, false); + + buffer_index = LLVMBuildAdd(ctx->ac.builder, buffer_index, - ctx->ac.i32_0, - num_channels, false, true); + buffer_offset, ""); + + attrib_offset = attrib_offset % attrib_stride; + } + + t_offset = LLVMConstInt(ctx->ac.i32, attrib_binding, false); + t_list = ac_build_load_to_sgpr(&ctx->ac, t_list_ptr, t_offset); + + input = ac_build_struct_tbuffer_load(&ctx->ac, t_list, + buffer_index, + LLVMConstInt(ctx->ac.i32, attrib_offset, false), + ctx->ac.i32_0, ctx->ac.i32_0, + num_channels, + data_format, num_format, 0, true); + + if (ctx->options->key.vs.post_shuffle & (1 << attrib_index)) { + LLVMValueRef c[4]; + c[0] = ac_llvm_extract_elem(&ctx->ac, input, 2); + c[1] = ac_llvm_extract_elem(&ctx->ac, input, 1); + c[2] = ac_llvm_extract_elem(&ctx->ac, input, 0); + c[3] = ac_llvm_extract_elem(&ctx->ac, input, 3); + + input = ac_build_gather_values(&ctx->ac, c, 4); + } input = radv_fixup_vertex_input_fetches(ctx, input, num_channels, is_float); @@ -2268,7 +2347,7 @@ static void mark_16bit_fs_input(struct radv_shader_context *ctx, mark_16bit_fs_input(ctx, glsl_get_array_element(type), location + i * stride); } } else { - assert(glsl_type_is_struct(type)); + assert(glsl_type_is_struct_or_ifc(type)); for (unsigned i = 0; i < glsl_get_length(type); i++) { mark_16bit_fs_input(ctx, glsl_get_struct_field(type, i), location); location += glsl_count_attribute_slots(glsl_get_struct_field(type, i), false); @@ -2517,9 +2596,15 @@ si_llvm_init_export_args(struct radv_shader_context *ctx, break; case V_028714_SPI_SHADER_32_AR: - args->enabled_channels = 0x9; - args->out[0] = values[0]; - args->out[3] = values[3]; + if (ctx->ac.chip_class >= GFX10) { + args->enabled_channels = 0x3; + args->out[0] = values[0]; + args->out[1] = values[3]; + } else { + args->enabled_channels = 0x9; + args->out[0] = values[0]; + args->out[3] = values[3]; + } break; case V_028714_SPI_SHADER_FP16_ABGR: @@ -2676,14 +2761,16 @@ radv_emit_stream_output(struct radv_shader_context *ctx, /* fall through */ case 4: /* as v4i32 */ vdata = ac_build_gather_values(&ctx->ac, out, - util_next_power_of_two(num_comps)); + !ac_has_vec3_support(ctx->ac.chip_class, false) ? + util_next_power_of_two(num_comps) : + num_comps); break; } ac_build_buffer_store_dword(&ctx->ac, so_buffers[buf], vdata, num_comps, so_write_offsets[buf], ctx->ac.i32_0, offset, - 1, 1, true, false); + ac_glc | ac_slc, false); } static void @@ -2771,6 +2858,7 @@ radv_emit_streamout(struct radv_shader_context *ctx, unsigned stream) static void handle_vs_outputs_post(struct radv_shader_context *ctx, bool export_prim_id, bool export_layer_id, + bool export_clip_dists, struct radv_vs_output_info *outinfo) { uint32_t param_count = 0; @@ -2827,9 +2915,11 @@ handle_vs_outputs_post(struct radv_shader_context *ctx, memcpy(&pos_args[target - V_008DFC_SQ_EXP_POS], &args, sizeof(args)); - /* Export the clip/cull distances values to the next stage. */ - radv_export_param(ctx, param_count, &slots[0], 0xf); - outinfo->vs_output_param_offset[location] = param_count++; + if (export_clip_dists) { + /* Export the clip/cull distances values to the next stage. */ + radv_export_param(ctx, param_count, &slots[0], 0xf); + outinfo->vs_output_param_offset[location] = param_count++; + } } } @@ -2905,6 +2995,14 @@ handle_vs_outputs_post(struct radv_shader_context *ctx, num_pos_exports++; } + /* Navi10-14 skip POS0 exports if EXEC=0 and DONE=0, causing a hang. + * Setting valid_mask=1 prevents it and has no other effect. + */ + if (ctx->ac.family == CHIP_NAVI10 || + ctx->ac.family == CHIP_NAVI12 || + ctx->ac.family == CHIP_NAVI14) + pos_args[0].valid_mask = 1; + pos_idx = 0; for (i = 0; i < 4; i++) { if (!pos_args[i].out[0]) @@ -2954,8 +3052,6 @@ handle_vs_outputs_post(struct radv_shader_context *ctx, LLVMValueRef values[4]; values[0] = ctx->vs_prim_id; - ctx->shader_info->vs.vgpr_comp_cnt = MAX2(2, - ctx->shader_info->vs.vgpr_comp_cnt); for (unsigned j = 1; j < 4; j++) values[j] = ctx->ac.f32_0; @@ -3060,7 +3156,7 @@ handle_es_outputs_post(struct radv_shader_context *ctx, out_val, 1, NULL, ctx->es2gs_offset, (4 * param_index + j) * 4, - 1, 1, true, true); + ac_glc | ac_slc, true); } } } @@ -3184,7 +3280,7 @@ write_tess_factors(struct radv_shader_context *ctx) LLVMConstInt(ctx->ac.i32, 4 * stride, false), ""); unsigned tf_offset = 0; - if (ctx->options->chip_class <= VI) { + if (ctx->options->chip_class <= GFX8) { ac_nir_build_if(&inner_if_ctx, ctx, LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ, rel_patch_id, ctx->ac.i32_0, "")); @@ -3193,7 +3289,7 @@ write_tess_factors(struct radv_shader_context *ctx) ac_build_buffer_store_dword(&ctx->ac, buffer, LLVMConstInt(ctx->ac.i32, 0x80000000, false), 1, ctx->ac.i32_0, tf_base, - 0, 1, 0, true, false); + 0, ac_glc, false); tf_offset += 4; ac_nir_build_endif(&inner_if_ctx); @@ -3202,11 +3298,11 @@ write_tess_factors(struct radv_shader_context *ctx) /* Store the tessellation factors. */ ac_build_buffer_store_dword(&ctx->ac, buffer, vec0, MIN2(stride, 4), byteoffset, tf_base, - tf_offset, 1, 0, true, false); + tf_offset, ac_glc, false); if (vec1) ac_build_buffer_store_dword(&ctx->ac, buffer, vec1, stride - 4, byteoffset, tf_base, - 16 + tf_offset, 1, 0, true, false); + 16 + tf_offset, ac_glc, false); //store to offchip for TES to read - only if TES reads them if (ctx->options->key.tcs.tes_reads_tess_factors) { @@ -3223,7 +3319,7 @@ write_tess_factors(struct radv_shader_context *ctx) ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, outer_vec, outer_comps, tf_outer_offset, - ctx->oc_lds, 0, 1, 0, true, false); + ctx->oc_lds, 0, ac_glc, false); if (inner_comps) { param_inner = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_INNER); tf_inner_offset = get_tcs_tes_buffer_address(ctx, NULL, @@ -3233,7 +3329,7 @@ write_tess_factors(struct radv_shader_context *ctx) ac_build_gather_values(&ctx->ac, inner, inner_comps); ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, inner_vec, inner_comps, tf_inner_offset, - ctx->oc_lds, 0, 1, 0, true, false); + ctx->oc_lds, 0, ac_glc, false); } } ac_nir_build_endif(&if_ctx); @@ -3356,6 +3452,7 @@ handle_shader_outputs_post(struct ac_shader_abi *abi, unsigned max_outputs, else handle_vs_outputs_post(ctx, ctx->options->key.vs.export_prim_id, ctx->options->key.vs.export_layer_id, + ctx->options->key.vs.export_clip_dists, &ctx->shader_info->vs.outinfo); break; case MESA_SHADER_FRAGMENT: @@ -3373,6 +3470,7 @@ handle_shader_outputs_post(struct ac_shader_abi *abi, unsigned max_outputs, else handle_vs_outputs_post(ctx, ctx->options->key.tes.export_prim_id, ctx->options->key.tes.export_layer_id, + ctx->options->key.tes.export_clip_dists, &ctx->shader_info->tes.outinfo); break; default: @@ -3426,7 +3524,7 @@ ac_nir_eliminate_const_vs_outputs(struct radv_shader_context *ctx) static void ac_setup_rings(struct radv_shader_context *ctx) { - if (ctx->options->chip_class <= VI && + if (ctx->options->chip_class <= GFX8 && (ctx->stage == MESA_SHADER_GEOMETRY || ctx->options->key.vs.as_es || ctx->options->key.tes.as_es)) { unsigned ring = ctx->stage == MESA_SHADER_GEOMETRY ? RING_ESGS_GS @@ -3476,7 +3574,7 @@ ac_setup_rings(struct radv_shader_context *ctx) stride = 4 * num_components * ctx->gs_max_out_vertices; - /* Limit on the stride field for <= CIK. */ + /* Limit on the stride field for <= GFX7. */ assert(stride < (1 << 14)); ring = LLVMBuildBitCast(ctx->ac.builder, @@ -3524,7 +3622,7 @@ radv_nir_get_max_workgroup_size(enum chip_class chip_class, { switch (nir->info.stage) { case MESA_SHADER_TESS_CTRL: - return chip_class >= CIK ? 128 : 64; + return chip_class >= GFX7 ? 128 : 64; case MESA_SHADER_GEOMETRY: return chip_class >= GFX9 ? 128 : 64; case MESA_SHADER_COMPUTE: @@ -3615,10 +3713,17 @@ LLVMModuleRef ac_translate_nir_to_llvm(struct ac_llvm_compiler *ac_llvm, ctx.abi.clamp_shadow_reference = false; ctx.abi.gfx9_stride_size_workaround = ctx.ac.chip_class == GFX9 && HAVE_LLVM < 0x800; + /* Because the new raw/struct atomic intrinsics are buggy with LLVM 8, + * we fallback to the old intrinsics for atomic buffer image operations + * and thus we need to apply the indexing workaround... + */ + ctx.abi.gfx9_stride_size_workaround_for_atomic = ctx.ac.chip_class == GFX9 && HAVE_LLVM < 0x900; + if (shader_count >= 2) ac_init_exec_full_mask(&ctx.ac); - if (ctx.ac.chip_class == GFX9 && + if ((ctx.ac.family == CHIP_VEGA10 || + ctx.ac.family == CHIP_RAVEN) && shaders[shader_count - 1]->info.stage == MESA_SHADER_TESS_CTRL) ac_nir_fixup_ls_hs_input_vgprs(&ctx); @@ -3654,15 +3759,6 @@ LLVMModuleRef ac_translate_nir_to_llvm(struct ac_llvm_compiler *ac_llvm, ctx.tcs_vertices_per_patch = shaders[i]->info.tess.tcs_vertices_out; ctx.tcs_num_patches = ctx.options->key.tes.num_patches; } else if (shaders[i]->info.stage == MESA_SHADER_VERTEX) { - if (shader_info->info.vs.needs_instance_id) { - if (ctx.options->key.vs.as_ls) { - ctx.shader_info->vs.vgpr_comp_cnt = - MAX2(2, ctx.shader_info->vs.vgpr_comp_cnt); - } else { - ctx.shader_info->vs.vgpr_comp_cnt = - MAX2(1, ctx.shader_info->vs.vgpr_comp_cnt); - } - } ctx.abi.load_base_vertex = radv_load_base_vertex; } else if (shaders[i]->info.stage == MESA_SHADER_FRAGMENT) { shader_info->fs.can_discard = shaders[i]->info.fs.uses_discard; @@ -3759,9 +3855,9 @@ static void ac_diagnostic_handler(LLVMDiagnosticInfoRef di, void *context) LLVMDisposeMessage(description); } -static unsigned ac_llvm_compile(LLVMModuleRef M, - struct ac_shader_binary *binary, - struct ac_llvm_compiler *ac_llvm) +static unsigned radv_llvm_compile(LLVMModuleRef M, + char **pelf_buffer, size_t *pelf_size, + struct ac_llvm_compiler *ac_llvm) { unsigned retval = 0; LLVMContextRef llvm_ctx; @@ -3773,98 +3869,55 @@ static unsigned ac_llvm_compile(LLVMModuleRef M, &retval); /* Compile IR*/ - if (!radv_compile_to_binary(ac_llvm, M, binary)) + if (!radv_compile_to_elf(ac_llvm, M, pelf_buffer, pelf_size)) retval = 1; return retval; } static void ac_compile_llvm_module(struct ac_llvm_compiler *ac_llvm, LLVMModuleRef llvm_module, - struct ac_shader_binary *binary, - struct ac_shader_config *config, + struct radv_shader_binary **rbinary, struct radv_shader_variant_info *shader_info, gl_shader_stage stage, const struct radv_nir_compiler_options *options) { + char *elf_buffer = NULL; + size_t elf_size = 0; + char *llvm_ir_string = NULL; if (options->dump_shader) ac_dump_module(llvm_module); - memset(binary, 0, sizeof(*binary)); - if (options->record_llvm_ir) { char *llvm_ir = LLVMPrintModuleToString(llvm_module); - binary->llvm_ir_string = strdup(llvm_ir); + llvm_ir_string = strdup(llvm_ir); LLVMDisposeMessage(llvm_ir); } - int v = ac_llvm_compile(llvm_module, binary, ac_llvm); + int v = radv_llvm_compile(llvm_module, &elf_buffer, &elf_size, ac_llvm); if (v) { fprintf(stderr, "compile failed\n"); } - if (options->dump_shader) - fprintf(stderr, "disasm:\n%s\n", binary->disasm_string); - - ac_shader_binary_read_config(binary, config, 0, options->supports_spill); - LLVMContextRef ctx = LLVMGetModuleContext(llvm_module); LLVMDisposeModule(llvm_module); LLVMContextDispose(ctx); - if (stage == MESA_SHADER_FRAGMENT) { - shader_info->num_input_vgprs = 0; - if (G_0286CC_PERSP_SAMPLE_ENA(config->spi_ps_input_addr)) - shader_info->num_input_vgprs += 2; - if (G_0286CC_PERSP_CENTER_ENA(config->spi_ps_input_addr)) - shader_info->num_input_vgprs += 2; - if (G_0286CC_PERSP_CENTROID_ENA(config->spi_ps_input_addr)) - shader_info->num_input_vgprs += 2; - if (G_0286CC_PERSP_PULL_MODEL_ENA(config->spi_ps_input_addr)) - shader_info->num_input_vgprs += 3; - if (G_0286CC_LINEAR_SAMPLE_ENA(config->spi_ps_input_addr)) - shader_info->num_input_vgprs += 2; - if (G_0286CC_LINEAR_CENTER_ENA(config->spi_ps_input_addr)) - shader_info->num_input_vgprs += 2; - if (G_0286CC_LINEAR_CENTROID_ENA(config->spi_ps_input_addr)) - shader_info->num_input_vgprs += 2; - if (G_0286CC_LINE_STIPPLE_TEX_ENA(config->spi_ps_input_addr)) - shader_info->num_input_vgprs += 1; - if (G_0286CC_POS_X_FLOAT_ENA(config->spi_ps_input_addr)) - shader_info->num_input_vgprs += 1; - if (G_0286CC_POS_Y_FLOAT_ENA(config->spi_ps_input_addr)) - shader_info->num_input_vgprs += 1; - if (G_0286CC_POS_Z_FLOAT_ENA(config->spi_ps_input_addr)) - shader_info->num_input_vgprs += 1; - if (G_0286CC_POS_W_FLOAT_ENA(config->spi_ps_input_addr)) - shader_info->num_input_vgprs += 1; - if (G_0286CC_FRONT_FACE_ENA(config->spi_ps_input_addr)) - shader_info->num_input_vgprs += 1; - if (G_0286CC_ANCILLARY_ENA(config->spi_ps_input_addr)) - shader_info->num_input_vgprs += 1; - if (G_0286CC_SAMPLE_COVERAGE_ENA(config->spi_ps_input_addr)) - shader_info->num_input_vgprs += 1; - if (G_0286CC_POS_FIXED_PT_ENA(config->spi_ps_input_addr)) - shader_info->num_input_vgprs += 1; - } - config->num_vgprs = MAX2(config->num_vgprs, shader_info->num_input_vgprs); - - /* +3 for scratch wave offset and VCC */ - config->num_sgprs = MAX2(config->num_sgprs, - shader_info->num_input_sgprs + 3); - - /* Enable 64-bit and 16-bit denormals, because there is no performance - * cost. - * - * If denormals are enabled, all floating-point output modifiers are - * ignored. - * - * Don't enable denormals for 32-bit floats, because: - * - Floating-point output modifiers would be ignored by the hw. - * - Some opcodes don't support denormals, such as v_mad_f32. We would - * have to stop using those. - * - SI & CI would be very slow. - */ - config->float_mode |= V_00B028_FP_64_DENORMS; + size_t llvm_ir_size = llvm_ir_string ? strlen(llvm_ir_string) : 0; + size_t alloc_size = sizeof(struct radv_shader_binary_rtld) + elf_size + llvm_ir_size + 1; + struct radv_shader_binary_rtld *rbin = calloc(1, alloc_size); + memcpy(rbin->data, elf_buffer, elf_size); + if (llvm_ir_string) + memcpy(rbin->data + elf_size, llvm_ir_string, llvm_ir_size + 1); + + rbin->base.type = RADV_BINARY_TYPE_RTLD; + rbin->base.stage = stage; + rbin->base.total_size = alloc_size; + rbin->elf_size = elf_size; + rbin->llvm_ir_size = llvm_ir_size; + *rbinary = &rbin->base; + + free(llvm_ir_string); + free(elf_buffer); } static void @@ -3890,6 +3943,7 @@ ac_fill_shader_info(struct radv_shader_variant_info *shader_info, struct nir_sha shader_info->tes.ccw = nir->info.tess.ccw; shader_info->tes.point_mode = nir->info.tess.point_mode; shader_info->tes.as_es = options->key.tes.as_es; + shader_info->tes.export_prim_id = options->key.tes.export_prim_id; break; case MESA_SHADER_TESS_CTRL: shader_info->tcs.tcs_vertices_out = nir->info.tess.tcs_vertices_out; @@ -3897,9 +3951,7 @@ ac_fill_shader_info(struct radv_shader_variant_info *shader_info, struct nir_sha case MESA_SHADER_VERTEX: shader_info->vs.as_es = options->key.vs.as_es; shader_info->vs.as_ls = options->key.vs.as_ls; - /* in LS mode we need at least 1, invocation id needs 2, handled elsewhere */ - if (options->key.vs.as_ls) - shader_info->vs.vgpr_comp_cnt = MAX2(1, shader_info->vs.vgpr_comp_cnt); + shader_info->vs.export_prim_id = options->key.vs.export_prim_id; break; default: break; @@ -3908,8 +3960,7 @@ ac_fill_shader_info(struct radv_shader_variant_info *shader_info, struct nir_sha void radv_compile_nir_shader(struct ac_llvm_compiler *ac_llvm, - struct ac_shader_binary *binary, - struct ac_shader_config *config, + struct radv_shader_binary **rbinary, struct radv_shader_variant_info *shader_info, struct nir_shader *const *nir, int nir_count, @@ -3921,8 +3972,8 @@ radv_compile_nir_shader(struct ac_llvm_compiler *ac_llvm, llvm_module = ac_translate_nir_to_llvm(ac_llvm, nir, nir_count, shader_info, options); - ac_compile_llvm_module(ac_llvm, llvm_module, binary, config, shader_info, - nir[0]->info.stage, options); + ac_compile_llvm_module(ac_llvm, llvm_module, rbinary, shader_info, + nir[nir_count - 1]->info.stage, options); for (int i = 0; i < nir_count; ++i) ac_fill_shader_info(shader_info, nir[i], options); @@ -4003,7 +4054,7 @@ ac_gs_copy_shader_emit(struct radv_shader_context *ctx) ctx->gsvs_ring[0], 1, ctx->ac.i32_0, vtx_offset, soffset, - 0, 1, 1, true, false); + 0, ac_glc | ac_slc, true, false); LLVMTypeRef type = LLVMGetAllocatedType(ctx->abi.outputs[ac_llvm_reg_index_soa(i, j)]); if (ac_get_type_size(type) == 2) { @@ -4020,7 +4071,7 @@ ac_gs_copy_shader_emit(struct radv_shader_context *ctx) radv_emit_streamout(ctx, stream); if (stream == 0) { - handle_vs_outputs_post(ctx, false, false, + handle_vs_outputs_post(ctx, false, false, true, &ctx->shader_info->vs.outinfo); } @@ -4033,8 +4084,7 @@ ac_gs_copy_shader_emit(struct radv_shader_context *ctx) void radv_compile_gs_copy_shader(struct ac_llvm_compiler *ac_llvm, struct nir_shader *geom_shader, - struct ac_shader_binary *binary, - struct ac_shader_config *config, + struct radv_shader_binary **rbinary, struct radv_shader_variant_info *shader_info, const struct radv_nir_compiler_options *options) { @@ -4074,6 +4124,8 @@ radv_compile_gs_copy_shader(struct ac_llvm_compiler *ac_llvm, ac_llvm_finalize_module(&ctx, ac_llvm->passmgr, options); - ac_compile_llvm_module(ac_llvm, ctx.ac.module, binary, config, shader_info, + ac_compile_llvm_module(ac_llvm, ctx.ac.module, rbinary, shader_info, MESA_SHADER_VERTEX, options); + (*rbinary)->is_gs_copy_shader = true; + }