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=e80938527e5a41d28a7237927409f2ff82e61a40;hpb=afeef3cacf35376a0c31f0871b57628cafd90ccb;p=mesa.git diff --git a/src/amd/vulkan/radv_nir_to_llvm.c b/src/amd/vulkan/radv_nir_to_llvm.c index e80938527e5..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" @@ -92,6 +91,7 @@ struct radv_shader_context { gl_shader_stage stage; LLVMValueRef inputs[RADEON_LLVM_MAX_INPUTS * 4]; + uint64_t float16_shaded_mask; uint64_t input_mask; uint64_t output_mask; @@ -261,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)); @@ -272,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); } @@ -517,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, @@ -588,6 +585,7 @@ set_loc_desc(struct radv_shader_context *ctx, int idx, uint8_t *sgpr_idx) struct user_sgpr_info { bool need_ring_offsets; bool indirect_all_descriptor_sets; + uint8_t remaining_sgprs; }; static bool needs_view_index_sgpr(struct radv_shader_context *ctx, @@ -626,6 +624,50 @@ count_vs_user_sgprs(struct radv_shader_context *ctx) return count; } +static void allocate_inline_push_consts(struct radv_shader_context *ctx, + struct user_sgpr_info *user_sgpr_info) +{ + uint8_t remaining_sgprs = user_sgpr_info->remaining_sgprs; + + /* Only supported if shaders use push constants. */ + if (ctx->shader_info->info.min_push_constant_used == UINT8_MAX) + return; + + /* Only supported if shaders don't have indirect push constants. */ + if (ctx->shader_info->info.has_indirect_push_constants) + return; + + /* Only supported for 32-bit push constants. */ + if (!ctx->shader_info->info.has_only_32bit_push_constants) + return; + + uint8_t num_push_consts = + (ctx->shader_info->info.max_push_constant_used - + ctx->shader_info->info.min_push_constant_used) / 4; + + /* Check if the number of user SGPRs is large enough. */ + if (num_push_consts < remaining_sgprs) { + ctx->shader_info->info.num_inline_push_consts = num_push_consts; + } else { + ctx->shader_info->info.num_inline_push_consts = remaining_sgprs; + } + + /* Clamp to the maximum number of allowed inlined push constants. */ + if (ctx->shader_info->info.num_inline_push_consts > AC_MAX_INLINE_PUSH_CONSTS) + ctx->shader_info->info.num_inline_push_consts = AC_MAX_INLINE_PUSH_CONSTS; + + if (ctx->shader_info->info.num_inline_push_consts == num_push_consts && + !ctx->shader_info->info.loads_dynamic_offsets) { + /* Disable the default push constants path if all constants are + * inlined and if shaders don't use dynamic descriptors. + */ + ctx->shader_info->info.loads_push_constants = false; + } + + ctx->shader_info->info.base_inline_push_consts = + ctx->shader_info->info.min_push_constant_used / 4; +} + static void allocate_user_sgprs(struct radv_shader_context *ctx, gl_shader_stage stage, bool has_previous_stage, @@ -701,7 +743,12 @@ static void allocate_user_sgprs(struct radv_shader_context *ctx, if (remaining_sgprs < num_desc_set) { user_sgpr_info->indirect_all_descriptor_sets = true; + user_sgpr_info->remaining_sgprs = remaining_sgprs - 1; + } else { + user_sgpr_info->remaining_sgprs = remaining_sgprs - num_desc_set; } + + allocate_inline_push_consts(ctx, user_sgpr_info); } static void @@ -731,6 +778,13 @@ declare_global_input_sgprs(struct radv_shader_context *ctx, add_arg(args, ARG_SGPR, type, &ctx->abi.push_constants); } + for (unsigned i = 0; i < ctx->shader_info->info.num_inline_push_consts; i++) { + add_arg(args, ARG_SGPR, ctx->ac.i32, + &ctx->abi.inline_push_consts[i]); + } + ctx->abi.num_inline_push_consts = ctx->shader_info->info.num_inline_push_consts; + ctx->abi.base_inline_push_consts = ctx->shader_info->info.base_inline_push_consts; + if (ctx->shader_info->info.so.num_outputs) { add_arg(args, ARG_SGPR, ac_array_in_const32_addr_space(ctx->ac.v4i32), @@ -849,6 +903,11 @@ set_global_input_locs(struct radv_shader_context *ctx, set_loc_shader_ptr(ctx, AC_UD_PUSH_CONSTANTS, user_sgpr_idx); } + if (ctx->shader_info->info.num_inline_push_consts) { + set_loc_shader(ctx, AC_UD_INLINE_PUSH_CONSTANTS, user_sgpr_idx, + ctx->shader_info->info.num_inline_push_consts); + } + if (ctx->streamout_buffers) { set_loc_shader_ptr(ctx, AC_UD_STREAMOUT_BUFFERS, user_sgpr_idx); @@ -1242,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; } @@ -1441,7 +1522,7 @@ store_tcs_output(struct ac_shader_abi *abi, { struct radv_shader_context *ctx = radv_shader_context_from_abi(abi); const unsigned location = var->data.location; - const unsigned component = var->data.location_frac; + unsigned component = var->data.location_frac; const bool is_patch = var->data.patch; const bool is_compact = var->data.compact; LLVMValueRef dw_addr; @@ -1459,10 +1540,14 @@ store_tcs_output(struct ac_shader_abi *abi, } param = shader_io_get_unique_index(location); - if (location == VARYING_SLOT_CLIP_DIST0 && - is_compact && const_index > 3) { - const_index -= 3; - param++; + if ((location == VARYING_SLOT_CLIP_DIST0 || location == VARYING_SLOT_CLIP_DIST1) && is_compact) { + const_index += component; + component = 0; + + if (const_index >= 4) { + const_index -= 4; + param++; + } } if (!is_patch) { @@ -1500,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); } } @@ -1529,9 +1614,13 @@ load_tes_input(struct ac_shader_abi *abi, LLVMValueRef result; unsigned param = shader_io_get_unique_index(location); - if (location == VARYING_SLOT_CLIP_DIST0 && is_compact && const_index > 3) { - const_index -= 3; - param++; + if ((location == VARYING_SLOT_CLIP_DIST0 || location == VARYING_SLOT_CLIP_DIST1) && is_compact) { + const_index += component; + component = 0; + if (const_index >= 4) { + const_index -= 4; + param++; + } } buf_addr = get_tcs_tes_buffer_address_params(ctx, param, const_index, @@ -1541,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; } @@ -1584,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) { @@ -1662,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), ""); @@ -1768,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); } } @@ -1839,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, ""); @@ -1880,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; @@ -1889,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"); } @@ -1913,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. @@ -1937,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 @@ -1964,7 +2088,71 @@ 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 +get_num_channels_from_data_format(unsigned data_format) +{ + switch (data_format) { + case V_008F0C_BUF_DATA_FORMAT_8: + case V_008F0C_BUF_DATA_FORMAT_16: + case V_008F0C_BUF_DATA_FORMAT_32: + return 1; + case V_008F0C_BUF_DATA_FORMAT_8_8: + case V_008F0C_BUF_DATA_FORMAT_16_16: + case V_008F0C_BUF_DATA_FORMAT_32_32: + return 2; + case V_008F0C_BUF_DATA_FORMAT_10_11_11: + case V_008F0C_BUF_DATA_FORMAT_11_11_10: + case V_008F0C_BUF_DATA_FORMAT_32_32_32: + return 3; + case V_008F0C_BUF_DATA_FORMAT_8_8_8_8: + case V_008F0C_BUF_DATA_FORMAT_10_10_10_2: + case V_008F0C_BUF_DATA_FORMAT_2_10_10_10: + case V_008F0C_BUF_DATA_FORMAT_16_16_16_16: + case V_008F0C_BUF_DATA_FORMAT_32_32_32_32: + return 4; + default: + break; + } + + return 4; +} + +static LLVMValueRef +radv_fixup_vertex_input_fetches(struct radv_shader_context *ctx, + LLVMValueRef value, + unsigned num_channels, + bool is_float) +{ + LLVMValueRef zero = is_float ? ctx->ac.f32_0 : ctx->ac.i32_0; + LLVMValueRef one = is_float ? ctx->ac.f32_1 : ctx->ac.i32_1; + LLVMValueRef chan[4]; + + if (LLVMGetTypeKind(LLVMTypeOf(value)) == LLVMVectorTypeKind) { + unsigned vec_size = LLVMGetVectorSize(LLVMTypeOf(value)); + + if (num_channels == 4 && num_channels == vec_size) + return value; + + num_channels = MIN2(num_channels, vec_size); + + for (unsigned i = 0; i < num_channels; i++) + chan[i] = ac_llvm_extract_elem(&ctx->ac, value, i); + } else { + if (num_channels) { + assert(num_channels == 1); + chan[0] = value; + } + } + + for (unsigned i = num_channels; i < 4; i++) { + chan[i] = i == 3 ? one : zero; + chan[i] = ac_to_integer(&ctx->ac, chan[i]); + } + + return ac_build_gather_values(&ctx->ac, chan, 4); } static void @@ -1979,7 +2167,7 @@ handle_vs_input_decl(struct radv_shader_context *ctx, unsigned attrib_count = glsl_count_attribute_slots(variable->type, true); uint8_t input_usage_mask = ctx->shader_info->info.vs.input_usage_mask[variable->data.location]; - unsigned num_channels = util_last_bit(input_usage_mask); + unsigned num_input_channels = util_last_bit(input_usage_mask); variable->data.driver_location = variable->data.location * 4; @@ -1987,6 +2175,11 @@ handle_vs_input_decl(struct radv_shader_context *ctx, for (unsigned i = 0; i < attrib_count; ++i) { LLVMValueRef output[4]; unsigned attrib_index = variable->data.location + i - VERT_ATTRIB_GENERIC0; + unsigned attrib_format = ctx->options->key.vs.vertex_attribute_formats[attrib_index]; + unsigned data_format = attrib_format & 0x0f; + unsigned num_format = (attrib_format >> 4) & 0x07; + bool is_float = num_format != V_008F0C_BUF_NUM_FORMAT_UINT && + num_format != V_008F0C_BUF_NUM_FORMAT_SINT; if (ctx->options->key.vs.instance_rate_inputs & (1u << attrib_index)) { uint32_t divisor = ctx->options->key.vs.instance_rate_divisors[attrib_index]; @@ -1998,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; } @@ -2014,16 +2199,57 @@ 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 = ac_build_expand_to_vec4(&ctx->ac, input, num_channels); + input = radv_fixup_vertex_input_fetches(ctx, input, num_channels, + is_float); for (unsigned chan = 0; chan < 4; chan++) { LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, chan, false); @@ -2051,6 +2277,7 @@ static void interp_fs_input(struct radv_shader_context *ctx, unsigned attr, LLVMValueRef interp_param, LLVMValueRef prim_mask, + bool float16, LLVMValueRef result[4]) { LLVMValueRef attr_number; @@ -2083,7 +2310,12 @@ static void interp_fs_input(struct radv_shader_context *ctx, for (chan = 0; chan < 4; chan++) { LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, chan, false); - if (interp) { + if (interp && float16) { + result[chan] = ac_build_fs_interp_f16(&ctx->ac, + llvm_chan, + attr_number, + prim_mask, i, j); + } else if (interp) { result[chan] = ac_build_fs_interp(&ctx->ac, llvm_chan, attr_number, @@ -2095,7 +2327,30 @@ static void interp_fs_input(struct radv_shader_context *ctx, attr_number, prim_mask); result[chan] = LLVMBuildBitCast(ctx->ac.builder, result[chan], ctx->ac.i32, ""); - result[chan] = LLVMBuildTruncOrBitCast(ctx->ac.builder, result[chan], LLVMTypeOf(interp_param), ""); + result[chan] = LLVMBuildTruncOrBitCast(ctx->ac.builder, result[chan], float16 ? ctx->ac.i16 : ctx->ac.i32, ""); + } + } +} + +static void mark_16bit_fs_input(struct radv_shader_context *ctx, + const struct glsl_type *type, + int location) +{ + if (glsl_type_is_scalar(type) || glsl_type_is_vector(type) || glsl_type_is_matrix(type)) { + unsigned attrib_count = glsl_count_attribute_slots(type, false); + if (glsl_type_is_16bit(type)) { + ctx->float16_shaded_mask |= ((1ull << attrib_count) - 1) << location; + } + } else if (glsl_type_is_array(type)) { + unsigned stride = glsl_count_attribute_slots(glsl_get_array_element(type), false); + for (unsigned i = 0; i < glsl_get_length(type); ++i) { + mark_16bit_fs_input(ctx, glsl_get_array_element(type), location + i * stride); + } + } else { + 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); } } } @@ -2110,9 +2365,20 @@ handle_fs_input_decl(struct radv_shader_context *ctx, uint64_t mask; variable->data.driver_location = idx * 4; + + + if (variable->data.compact) { + unsigned component_count = variable->data.location_frac + + glsl_get_length(variable->type); + attrib_count = (component_count + 3) / 4; + } else + mark_16bit_fs_input(ctx, variable->type, idx); + mask = ((1ull << attrib_count) - 1) << variable->data.location; - if (glsl_get_base_type(glsl_without_array(variable->type)) == GLSL_TYPE_FLOAT) { + if (glsl_get_base_type(glsl_without_array(variable->type)) == GLSL_TYPE_FLOAT || + glsl_get_base_type(glsl_without_array(variable->type)) == GLSL_TYPE_FLOAT16 || + glsl_get_base_type(glsl_without_array(variable->type)) == GLSL_TYPE_STRUCT) { unsigned interp_type; if (variable->data.sample) interp_type = INTERP_SAMPLE; @@ -2123,22 +2389,12 @@ handle_fs_input_decl(struct radv_shader_context *ctx, interp = lookup_interp_param(&ctx->abi, variable->data.interpolation, interp_type); } - bool is_16bit = glsl_type_is_16bit(glsl_without_array(variable->type)); - LLVMTypeRef type = is_16bit ? ctx->ac.i16 : ctx->ac.i32; if (interp == NULL) - interp = LLVMGetUndef(type); + interp = LLVMGetUndef(ctx->ac.i32); for (unsigned i = 0; i < attrib_count; ++i) ctx->inputs[ac_llvm_reg_index_soa(idx + i, 0)] = interp; - if (idx == VARYING_SLOT_CLIP_DIST0) { - /* Do not account for the number of components inside the array - * of clip/cull distances because this might wrongly set other - * bits like primitive ID or layer. - */ - mask = 1ull << VARYING_SLOT_CLIP_DIST0; - } - ctx->input_mask |= mask; } @@ -2200,11 +2456,14 @@ handle_fs_inputs(struct radv_shader_context *ctx, if (i >= VARYING_SLOT_VAR0 || i == VARYING_SLOT_PNTC || i == VARYING_SLOT_PRIMITIVE_ID || i == VARYING_SLOT_LAYER) { interp_param = *inputs; - interp_fs_input(ctx, index, interp_param, ctx->abi.prim_mask, + bool float16 = (ctx->float16_shaded_mask >> i) & 1; + interp_fs_input(ctx, index, interp_param, ctx->abi.prim_mask, float16, inputs); if (LLVMIsUndef(interp_param)) ctx->shader_info->fs.flat_shaded_mask |= 1u << index; + if (float16) + ctx->shader_info->fs.float16_shaded_mask |= 1u << index; if (i >= VARYING_SLOT_VAR0) ctx->abi.fs_input_attr_indices[i - VARYING_SLOT_VAR0] = index; ++index; @@ -2216,7 +2475,7 @@ handle_fs_inputs(struct radv_shader_context *ctx, interp_param = *inputs; interp_fs_input(ctx, index, interp_param, - ctx->abi.prim_mask, inputs); + ctx->abi.prim_mask, false, inputs); ++index; } } else if (i == VARYING_SLOT_POS) { @@ -2250,6 +2509,12 @@ scan_shader_output_decl(struct radv_shader_context *ctx, if (stage == MESA_SHADER_TESS_CTRL) return; + if (variable->data.compact) { + unsigned component_count = variable->data.location_frac + + glsl_get_length(variable->type); + attrib_count = (component_count + 3) / 4; + } + mask_attribs = ((1ull << attrib_count) - 1) << idx; if (stage == MESA_SHADER_VERTEX || stage == MESA_SHADER_TESS_EVAL || @@ -2265,8 +2530,6 @@ scan_shader_output_decl(struct radv_shader_context *ctx, ctx->shader_info->tes.outinfo.cull_dist_mask = (1 << shader->info.cull_distance_array_size) - 1; ctx->shader_info->tes.outinfo.cull_dist_mask <<= shader->info.clip_distance_array_size; } - - mask_attribs = 1ull << idx; } } @@ -2333,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: @@ -2365,7 +2634,7 @@ si_llvm_init_export_args(struct radv_shader_context *ctx, if (is_16bit) { for (unsigned chan = 0; chan < 4; chan++) values[chan] = LLVMBuildZExt(ctx->ac.builder, - values[chan], + ac_to_integer(&ctx->ac, values[chan]), ctx->ac.i32, ""); } break; @@ -2376,7 +2645,7 @@ si_llvm_init_export_args(struct radv_shader_context *ctx, if (is_16bit) { for (unsigned chan = 0; chan < 4; chan++) values[chan] = LLVMBuildSExt(ctx->ac.builder, - values[chan], + ac_to_integer(&ctx->ac, values[chan]), ctx->ac.i32, ""); } break; @@ -2429,12 +2698,8 @@ si_llvm_init_export_args(struct radv_shader_context *ctx, } else memcpy(&args->out[0], values, sizeof(values[0]) * 4); - for (unsigned i = 0; i < 4; ++i) { - if (!(args->enabled_channels & (1 << i))) - continue; - + for (unsigned i = 0; i < 4; ++i) args->out[i] = ac_to_float(&ctx->ac, args->out[i]); - } } static void @@ -2496,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 @@ -2591,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; @@ -2615,51 +2883,43 @@ handle_vs_outputs_post(struct radv_shader_context *ctx, memset(outinfo->vs_output_param_offset, AC_EXP_PARAM_UNDEFINED, sizeof(outinfo->vs_output_param_offset)); - if (ctx->output_mask & (1ull << VARYING_SLOT_CLIP_DIST0)) { - unsigned output_usage_mask, length; - LLVMValueRef slots[8]; - unsigned j; - - if (ctx->stage == MESA_SHADER_VERTEX && - !ctx->is_gs_copy_shader) { - output_usage_mask = - ctx->shader_info->info.vs.output_usage_mask[VARYING_SLOT_CLIP_DIST0]; - } else if (ctx->stage == MESA_SHADER_TESS_EVAL) { - output_usage_mask = - ctx->shader_info->info.tes.output_usage_mask[VARYING_SLOT_CLIP_DIST0]; - } else { - assert(ctx->is_gs_copy_shader); - output_usage_mask = - ctx->shader_info->info.gs.output_usage_mask[VARYING_SLOT_CLIP_DIST0]; - } + for(unsigned location = VARYING_SLOT_CLIP_DIST0; location <= VARYING_SLOT_CLIP_DIST1; ++location) { + if (ctx->output_mask & (1ull << location)) { + unsigned output_usage_mask, length; + LLVMValueRef slots[4]; + unsigned j; + + if (ctx->stage == MESA_SHADER_VERTEX && + !ctx->is_gs_copy_shader) { + output_usage_mask = + ctx->shader_info->info.vs.output_usage_mask[location]; + } else if (ctx->stage == MESA_SHADER_TESS_EVAL) { + output_usage_mask = + ctx->shader_info->info.tes.output_usage_mask[location]; + } else { + assert(ctx->is_gs_copy_shader); + output_usage_mask = + ctx->shader_info->info.gs.output_usage_mask[location]; + } - length = util_last_bit(output_usage_mask); + length = util_last_bit(output_usage_mask); - i = VARYING_SLOT_CLIP_DIST0; - for (j = 0; j < length; j++) - slots[j] = ac_to_float(&ctx->ac, radv_load_output(ctx, i, j)); + for (j = 0; j < length; j++) + slots[j] = ac_to_float(&ctx->ac, radv_load_output(ctx, location, j)); - for (i = length; i < 8; i++) - slots[i] = LLVMGetUndef(ctx->ac.f32); + for (i = length; i < 4; i++) + slots[i] = LLVMGetUndef(ctx->ac.f32); - if (length > 4) { - target = V_008DFC_SQ_EXP_POS + 3; - si_llvm_init_export_args(ctx, &slots[4], 0xf, target, &args); + target = V_008DFC_SQ_EXP_POS + 2 + (location - VARYING_SLOT_CLIP_DIST0); + si_llvm_init_export_args(ctx, &slots[0], 0xf, target, &args); memcpy(&pos_args[target - V_008DFC_SQ_EXP_POS], - &args, sizeof(args)); - } + &args, sizeof(args)); - target = V_008DFC_SQ_EXP_POS + 2; - si_llvm_init_export_args(ctx, &slots[0], 0xf, target, &args); - 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[VARYING_SLOT_CLIP_DIST0] = param_count++; - if (length > 4) { - radv_export_param(ctx, param_count, &slots[4], 0xf); - outinfo->vs_output_param_offset[VARYING_SLOT_CLIP_DIST1] = 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++; + } } } @@ -2735,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]) @@ -2784,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; @@ -2820,28 +3086,14 @@ handle_es_outputs_post(struct radv_shader_context *ctx, LLVMValueRef lds_base = NULL; for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) { - unsigned output_usage_mask; int param_index; - int length = 4; if (!(ctx->output_mask & (1ull << i))) continue; - if (ctx->stage == MESA_SHADER_VERTEX) { - output_usage_mask = - ctx->shader_info->info.vs.output_usage_mask[i]; - } else { - assert(ctx->stage == MESA_SHADER_TESS_EVAL); - output_usage_mask = - ctx->shader_info->info.tes.output_usage_mask[i]; - } - - if (i == VARYING_SLOT_CLIP_DIST0) - length = util_last_bit(output_usage_mask); - param_index = shader_io_get_unique_index(i); - max_output_written = MAX2(param_index + (length > 4), max_output_written); + max_output_written = MAX2(param_index, max_output_written); } outinfo->esgs_itemsize = (max_output_written + 1) * 16; @@ -2862,7 +3114,6 @@ handle_es_outputs_post(struct radv_shader_context *ctx, LLVMValueRef *out_ptr = &ctx->abi.outputs[i * 4]; unsigned output_usage_mask; int param_index; - int length = 4; if (!(ctx->output_mask & (1ull << i))) continue; @@ -2876,9 +3127,6 @@ handle_es_outputs_post(struct radv_shader_context *ctx, ctx->shader_info->info.tes.output_usage_mask[i]; } - if (i == VARYING_SLOT_CLIP_DIST0) - length = util_last_bit(output_usage_mask); - param_index = shader_io_get_unique_index(i); if (lds_base) { @@ -2887,7 +3135,7 @@ handle_es_outputs_post(struct radv_shader_context *ctx, ""); } - for (j = 0; j < length; j++) { + for (j = 0; j < 4; j++) { if (!(output_usage_mask & (1 << j))) continue; @@ -2908,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); } } } @@ -2924,22 +3172,16 @@ handle_ls_outputs_post(struct radv_shader_context *ctx) vertex_dw_stride, ""); for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) { - unsigned output_usage_mask = - ctx->shader_info->info.vs.output_usage_mask[i]; LLVMValueRef *out_ptr = &ctx->abi.outputs[i * 4]; - int length = 4; if (!(ctx->output_mask & (1ull << i))) continue; - if (i == VARYING_SLOT_CLIP_DIST0) - length = util_last_bit(output_usage_mask); - int param = shader_io_get_unique_index(i); LLVMValueRef dw_addr = LLVMBuildAdd(ctx->ac.builder, base_dw_addr, LLVMConstInt(ctx->ac.i32, param * 4, false), ""); - for (unsigned j = 0; j < length; j++) { + for (unsigned j = 0; j < 4; j++) { LLVMValueRef value = LLVMBuildLoad(ctx->ac.builder, out_ptr[j], ""); value = ac_to_integer(&ctx->ac, value); value = LLVMBuildZExtOrBitCast(ctx->ac.builder, value, ctx->ac.i32, ""); @@ -3038,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, "")); @@ -3047,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); @@ -3056,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) { @@ -3077,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, @@ -3087,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); @@ -3210,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: @@ -3227,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: @@ -3280,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 @@ -3330,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, @@ -3372,13 +3616,13 @@ ac_setup_rings(struct radv_shader_context *ctx) } } -static unsigned -ac_nir_get_max_workgroup_size(enum chip_class chip_class, - const struct nir_shader *nir) +unsigned +radv_nir_get_max_workgroup_size(enum chip_class chip_class, + const struct nir_shader *nir) { 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: @@ -3439,6 +3683,8 @@ LLVMModuleRef ac_translate_nir_to_llvm(struct ac_llvm_compiler *ac_llvm, memset(shader_info, 0, sizeof(*shader_info)); + radv_nir_shader_info_init(&shader_info->info); + for(int i = 0; i < shader_count; ++i) radv_nir_shader_info_pass(shaders[i], options, &shader_info->info); @@ -3450,7 +3696,7 @@ LLVMModuleRef ac_translate_nir_to_llvm(struct ac_llvm_compiler *ac_llvm, ctx.max_workgroup_size = 0; for (int i = 0; i < shader_count; ++i) { ctx.max_workgroup_size = MAX2(ctx.max_workgroup_size, - ac_nir_get_max_workgroup_size(ctx.options->chip_class, + radv_nir_get_max_workgroup_size(ctx.options->chip_class, shaders[i])); } @@ -3467,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); @@ -3506,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; @@ -3611,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; @@ -3625,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 @@ -3742,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; @@ -3749,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; @@ -3760,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, @@ -3773,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); @@ -3855,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) { @@ -3872,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); } @@ -3885,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) { @@ -3926,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; + }