radv: allow to emit a vertex to a specified stream
[mesa.git] / src / amd / vulkan / radv_nir_to_llvm.c
index 15c1049302282d750f72b7be6b53114a5b153864..07dc6a2301dca0b439ce1d7aee2b5c85a0e70074 100644 (file)
@@ -27,6 +27,7 @@
 
 #include "radv_private.h"
 #include "radv_shader.h"
+#include "radv_shader_helper.h"
 #include "nir/nir.h"
 
 #include <llvm-c/Core.h>
@@ -77,7 +78,7 @@ struct radv_shader_context {
        LLVMValueRef gs_vtx_offset[6];
 
        LLVMValueRef esgs_ring;
-       LLVMValueRef gsvs_ring;
+       LLVMValueRef gsvs_ring[4];
        LLVMValueRef hs_ring_tess_offchip;
        LLVMValueRef hs_ring_tess_factor;
 
@@ -90,11 +91,9 @@ struct radv_shader_context {
 
        uint64_t input_mask;
        uint64_t output_mask;
-       uint8_t num_output_clips;
-       uint8_t num_output_culls;
 
        bool is_gs_copy_shader;
-       LLVMValueRef gs_next_vertex;
+       LLVMValueRef gs_next_vertex[4];
        unsigned gs_max_out_vertices;
 
        unsigned tes_primitive_mode;
@@ -401,10 +400,8 @@ get_tcs_out_current_patch_offset(struct radv_shader_context *ctx)
        LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx);
        LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
 
-       return LLVMBuildAdd(ctx->ac.builder, patch0_offset,
-                           LLVMBuildMul(ctx->ac.builder, patch_stride,
-                                        rel_patch_id, ""),
-                           "");
+       return ac_build_imad(&ctx->ac, patch_stride, rel_patch_id,
+                            patch0_offset);
 }
 
 static LLVMValueRef
@@ -415,13 +412,11 @@ get_tcs_out_current_patch_data_offset(struct radv_shader_context *ctx)
        LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx);
        LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
 
-       return LLVMBuildAdd(ctx->ac.builder, patch0_patch_data_offset,
-                           LLVMBuildMul(ctx->ac.builder, patch_stride,
-                                        rel_patch_id, ""),
-                           "");
+       return ac_build_imad(&ctx->ac, patch_stride, rel_patch_id,
+                            patch0_patch_data_offset);
 }
 
-#define MAX_ARGS 23
+#define MAX_ARGS 64
 struct arg_info {
        LLVMTypeRef types[MAX_ARGS];
        LLVMValueRef *assign[MAX_ARGS];
@@ -544,13 +539,12 @@ create_llvm_function(LLVMContextRef ctx, LLVMModuleRef module,
 
 
 static void
-set_loc(struct radv_userdata_info *ud_info, uint8_t *sgpr_idx, uint8_t num_sgprs,
-       uint32_t indirect_offset)
+set_loc(struct radv_userdata_info *ud_info, uint8_t *sgpr_idx,
+       uint8_t num_sgprs, bool indirect)
 {
        ud_info->sgpr_idx = *sgpr_idx;
        ud_info->num_sgprs = num_sgprs;
-       ud_info->indirect = indirect_offset > 0;
-       ud_info->indirect_offset = indirect_offset;
+       ud_info->indirect = indirect;
        *sgpr_idx += num_sgprs;
 }
 
@@ -562,7 +556,7 @@ set_loc_shader(struct radv_shader_context *ctx, int idx, uint8_t *sgpr_idx,
                &ctx->shader_info->user_sgprs_locs.shader_data[idx];
        assert(ud_info);
 
-       set_loc(ud_info, sgpr_idx, num_sgprs, 0);
+       set_loc(ud_info, sgpr_idx, num_sgprs, false);
 }
 
 static void
@@ -576,13 +570,17 @@ set_loc_shader_ptr(struct radv_shader_context *ctx, int idx, uint8_t *sgpr_idx)
 
 static void
 set_loc_desc(struct radv_shader_context *ctx, int idx,  uint8_t *sgpr_idx,
-            uint32_t indirect_offset)
+            bool indirect)
 {
-       struct radv_userdata_info *ud_info =
-               &ctx->shader_info->user_sgprs_locs.descriptor_sets[idx];
+       struct radv_userdata_locations *locs =
+               &ctx->shader_info->user_sgprs_locs;
+       struct radv_userdata_info *ud_info = &locs->descriptor_sets[idx];
        assert(ud_info);
 
-       set_loc(ud_info, sgpr_idx, HAVE_32BIT_POINTERS ? 1 : 2, indirect_offset);
+       set_loc(ud_info, sgpr_idx, HAVE_32BIT_POINTERS ? 1 : 2, indirect);
+
+       if (!indirect)
+               locs->descriptor_sets_enabled |= 1 << idx;
 }
 
 struct user_sgpr_info {
@@ -691,7 +689,7 @@ static void allocate_user_sgprs(struct radv_shader_context *ctx,
        if (ctx->shader_info->info.loads_push_constants)
                user_sgpr_count += HAVE_32BIT_POINTERS ? 1 : 2;
 
-       uint32_t available_sgprs = ctx->options->chip_class >= GFX9 ? 32 : 16;
+       uint32_t available_sgprs = ctx->options->chip_class >= GFX9 && stage != MESA_SHADER_COMPUTE ? 32 : 16;
        uint32_t remaining_sgprs = available_sgprs - user_sgpr_count;
        uint32_t num_desc_set =
                util_bitcount(ctx->shader_info->info.desc_set_used_mask);
@@ -802,7 +800,7 @@ set_global_input_locs(struct radv_shader_context *ctx, gl_shader_stage stage,
                for (unsigned i = 0; i < num_sets; ++i) {
                        if ((ctx->shader_info->info.desc_set_used_mask & (1 << i)) &&
                            ctx->options->layout->set[i].layout->shader_stages & stage_mask) {
-                               set_loc_desc(ctx, i, user_sgpr_idx, 0);
+                               set_loc_desc(ctx, i, user_sgpr_idx, false);
                        } else
                                ctx->descriptor_sets[i] = NULL;
                }
@@ -813,7 +811,6 @@ set_global_input_locs(struct radv_shader_context *ctx, gl_shader_stage stage,
                for (unsigned i = 0; i < num_sets; ++i) {
                        if ((ctx->shader_info->info.desc_set_used_mask & (1 << i)) &&
                            ctx->options->layout->set[i].layout->shader_stages & stage_mask) {
-                               set_loc_desc(ctx, i, user_sgpr_idx, i * 8);
                                ctx->descriptor_sets[i] =
                                        ac_build_load_to_sgpr(&ctx->ac,
                                                              desc_sets,
@@ -1139,7 +1136,7 @@ static void create_function(struct radv_shader_context *ctx,
                                   &user_sgpr_idx);
                if (ctx->options->supports_spill) {
                        ctx->ring_offsets = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.implicit.buffer.ptr",
-                                                              LLVMPointerType(ctx->ac.i8, AC_CONST_ADDR_SPACE),
+                                                              LLVMPointerType(ctx->ac.i8, AC_ADDR_SPACE_CONST),
                                                               NULL, 0, AC_FUNC_ATTR_READNONE);
                        ctx->ring_offsets = LLVMBuildBitCast(ctx->ac.builder, ctx->ring_offsets,
                                                             ac_array_in_const_addr_space(ctx->ac.v4i32), "");
@@ -1226,9 +1223,8 @@ radv_load_resource(struct ac_shader_abi *abi, LLVMValueRef index,
        } else
                stride = LLVMConstInt(ctx->ac.i32, layout->binding[binding].size, false);
 
-       offset = LLVMConstInt(ctx->ac.i32, base_offset, false);
-       index = LLVMBuildMul(ctx->ac.builder, index, stride, "");
-       offset = LLVMBuildAdd(ctx->ac.builder, offset, index, "");
+       offset = ac_build_imad(&ctx->ac, index, stride,
+                              LLVMConstInt(ctx->ac.i32, base_offset, false));
 
        desc_ptr = ac_build_gep0(&ctx->ac, desc_ptr, offset);
        desc_ptr = ac_cast_ptr(&ctx->ac, desc_ptr, ctx->ac.v4i32);
@@ -1293,11 +1289,8 @@ static LLVMValueRef get_tcs_tes_buffer_address(struct radv_shader_context *ctx,
        constant16 = LLVMConstInt(ctx->ac.i32, 16, false);
        param_stride = calc_param_stride(ctx, vertex_index);
        if (vertex_index) {
-               base_addr = LLVMBuildMul(ctx->ac.builder, rel_patch_id,
-                                        vertices_per_patch, "");
-
-               base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
-                                        vertex_index, "");
+               base_addr = ac_build_imad(&ctx->ac, rel_patch_id,
+                                         vertices_per_patch, vertex_index);
        } else {
                base_addr = rel_patch_id;
        }
@@ -1475,6 +1468,8 @@ store_tcs_output(struct ac_shader_abi *abi,
                if (!(writemask & (1 << chan)))
                        continue;
                LLVMValueRef value = ac_llvm_extract_elem(&ctx->ac, src, chan - component);
+               value = ac_to_integer(&ctx->ac, value);
+               value = LLVMBuildZExtOrBitCast(ctx->ac.builder, value, ctx->ac.i32, "");
 
                if (store_lds || is_tess_factor) {
                        LLVMValueRef dw_addr_chan =
@@ -1571,10 +1566,13 @@ load_gs_input(struct ac_shader_abi *abi,
                                                        ctx->ac.i32_0,
                                                        vtx_offset, soffset,
                                                        0, 1, 0, true, false);
+               }
 
-                       value[i] = LLVMBuildBitCast(ctx->ac.builder, value[i],
-                                                   type, "");
+               if (ac_get_type_size(type) == 2) {
+                       value[i] = LLVMBuildBitCast(ctx->ac.builder, value[i], ctx->ac.i32, "");
+                       value[i] = LLVMBuildTrunc(ctx->ac.builder, value[i], ctx->ac.i16, "");
                }
+               value[i] = LLVMBuildBitCast(ctx->ac.builder, value[i], type, "");
        }
        result = ac_build_varying_gather_values(&ctx->ac, value, num_components, component);
        result = ac_to_integer(&ctx->ac, result);
@@ -1703,14 +1701,12 @@ visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addr
 {
        LLVMValueRef gs_next_vertex;
        LLVMValueRef can_emit;
-       int idx;
+       unsigned offset = 0;
        struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
 
-       assert(stream == 0);
-
        /* Write vertex attribute values to GSVS ring */
        gs_next_vertex = LLVMBuildLoad(ctx->ac.builder,
-                                      ctx->gs_next_vertex,
+                                      ctx->gs_next_vertex[stream],
                                       "");
 
        /* If this thread has already emitted the declared maximum number of
@@ -1722,52 +1718,51 @@ visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addr
                                 LLVMConstInt(ctx->ac.i32, ctx->gs_max_out_vertices, false), "");
        ac_build_kill_if_false(&ctx->ac, can_emit);
 
-       /* loop num outputs */
-       idx = 0;
        for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
                unsigned output_usage_mask =
                        ctx->shader_info->info.gs.output_usage_mask[i];
+               uint8_t output_stream =
+                       ctx->shader_info->info.gs.output_streams[i];
                LLVMValueRef *out_ptr = &addrs[i * 4];
-               int length = 4;
-               int slot = idx;
-               int slot_inc = 1;
+               int length = util_last_bit(output_usage_mask);
 
-               if (!(ctx->output_mask & (1ull << i)))
+               if (!(ctx->output_mask & (1ull << i)) ||
+                   output_stream != stream)
                        continue;
 
-               if (i == VARYING_SLOT_CLIP_DIST0) {
-                       /* pack clip and cull into a single set of slots */
-                       length = ctx->num_output_clips + ctx->num_output_culls;
-                       if (length > 4)
-                               slot_inc = 2;
-                       output_usage_mask = (1 << length) - 1;
-               }
-
                for (unsigned j = 0; j < length; j++) {
                        if (!(output_usage_mask & (1 << j)))
                                continue;
 
                        LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder,
                                                             out_ptr[j], "");
-                       LLVMValueRef voffset = LLVMConstInt(ctx->ac.i32, (slot * 4 + j) * ctx->gs_max_out_vertices, false);
+                       LLVMValueRef voffset =
+                               LLVMConstInt(ctx->ac.i32, offset *
+                                            ctx->gs_max_out_vertices, false);
+
+                       offset++;
+
                        voffset = LLVMBuildAdd(ctx->ac.builder, voffset, gs_next_vertex, "");
                        voffset = LLVMBuildMul(ctx->ac.builder, voffset, LLVMConstInt(ctx->ac.i32, 4, false), "");
 
-                       out_val = LLVMBuildBitCast(ctx->ac.builder, out_val, ctx->ac.i32, "");
+                       out_val = ac_to_integer(&ctx->ac, out_val);
+                       out_val = LLVMBuildZExtOrBitCast(ctx->ac.builder, out_val, ctx->ac.i32, "");
 
-                       ac_build_buffer_store_dword(&ctx->ac, ctx->gsvs_ring,
+                       ac_build_buffer_store_dword(&ctx->ac,
+                                                   ctx->gsvs_ring[stream],
                                                    out_val, 1,
                                                    voffset, ctx->gs2vs_offset, 0,
                                                    1, 1, true, true);
                }
-               idx += slot_inc;
        }
 
        gs_next_vertex = LLVMBuildAdd(ctx->ac.builder, gs_next_vertex,
                                      ctx->ac.i32_1, "");
-       LLVMBuildStore(ctx->ac.builder, gs_next_vertex, ctx->gs_next_vertex);
+       LLVMBuildStore(ctx->ac.builder, gs_next_vertex, ctx->gs_next_vertex[stream]);
 
-       ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (0 << 8), ctx->gs_wave_id);
+       ac_build_sendmsg(&ctx->ac,
+                        AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (stream << 8),
+                        ctx->gs_wave_id);
 }
 
 static void
@@ -1972,6 +1967,7 @@ handle_vs_input_decl(struct radv_shader_context *ctx,
 
        variable->data.driver_location = variable->data.location * 4;
 
+       enum glsl_base_type type = glsl_get_base_type(variable->type);
        for (unsigned i = 0; i < attrib_count; ++i) {
                LLVMValueRef output[4];
                unsigned attrib_index = variable->data.location + i - VERT_ATTRIB_GENERIC0;
@@ -1980,8 +1976,7 @@ handle_vs_input_decl(struct radv_shader_context *ctx,
                        uint32_t divisor = ctx->options->key.vs.instance_rate_divisors[attrib_index];
 
                        if (divisor) {
-                               buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.instance_id,
-                                                           ctx->abi.start_instance, "");
+                               buffer_index = ctx->abi.instance_id;
 
                                if (divisor != 1) {
                                        buffer_index = LLVMBuildUDiv(ctx->ac.builder, buffer_index,
@@ -1998,6 +1993,8 @@ handle_vs_input_decl(struct radv_shader_context *ctx,
                        } else {
                                buffer_index = ctx->ac.i32_0;
                        }
+
+                       buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.start_instance, buffer_index, "");
                } else
                        buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.vertex_id,
                                                    ctx->abi.base_vertex, "");
@@ -2015,14 +2012,21 @@ handle_vs_input_decl(struct radv_shader_context *ctx,
                for (unsigned chan = 0; chan < 4; chan++) {
                        LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, chan, false);
                        output[chan] = LLVMBuildExtractElement(ctx->ac.builder, input, llvm_chan, "");
+                       if (type == GLSL_TYPE_FLOAT16) {
+                               output[chan] = LLVMBuildBitCast(ctx->ac.builder, output[chan], ctx->ac.f32, "");
+                               output[chan] = LLVMBuildFPTrunc(ctx->ac.builder, output[chan], ctx->ac.f16, "");
+                       }
                }
 
                unsigned alpha_adjust = (ctx->options->key.vs.alpha_adjust >> (attrib_index * 2)) & 3;
                output[3] = adjust_vertex_fetch_alpha(ctx, alpha_adjust, output[3]);
 
                for (unsigned chan = 0; chan < 4; chan++) {
-                       ctx->inputs[ac_llvm_reg_index_soa(variable->data.location + i, chan)] =
-                               ac_to_integer(&ctx->ac, output[chan]);
+                       output[chan] = ac_to_integer(&ctx->ac, output[chan]);
+                       if (type == GLSL_TYPE_UINT16 || type == GLSL_TYPE_INT16)
+                               output[chan] = LLVMBuildTrunc(ctx->ac.builder, output[chan], ctx->ac.i16, "");
+
+                       ctx->inputs[ac_llvm_reg_index_soa(variable->data.location + i, chan)] = output[chan];
                }
        }
 }
@@ -2036,7 +2040,7 @@ static void interp_fs_input(struct radv_shader_context *ctx,
        LLVMValueRef attr_number;
        unsigned chan;
        LLVMValueRef i, j;
-       bool interp = interp_param != NULL;
+       bool interp = !LLVMIsUndef(interp_param);
 
        attr_number = LLVMConstInt(ctx->ac.i32, attr, false);
 
@@ -2074,6 +2078,8 @@ static void interp_fs_input(struct radv_shader_context *ctx,
                                                              llvm_chan,
                                                              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), "");
                }
        }
 }
@@ -2084,10 +2090,11 @@ handle_fs_input_decl(struct radv_shader_context *ctx,
 {
        int idx = variable->data.location;
        unsigned attrib_count = glsl_count_attribute_slots(variable->type, false);
-       LLVMValueRef interp;
+       LLVMValueRef interp = NULL;
+       uint64_t mask;
 
        variable->data.driver_location = idx * 4;
-       ctx->input_mask |= ((1ull << attrib_count) - 1) << variable->data.location;
+       mask = ((1ull << attrib_count) - 1) << variable->data.location;
 
        if (glsl_get_base_type(glsl_without_array(variable->type)) == GLSL_TYPE_FLOAT) {
                unsigned interp_type;
@@ -2099,12 +2106,24 @@ handle_fs_input_decl(struct radv_shader_context *ctx,
                        interp_type = INTERP_CENTER;
 
                interp = lookup_interp_param(&ctx->abi, variable->data.interpolation, interp_type);
-       } else
-               interp = NULL;
+       }
+       bool is_16bit = glsl_type_is_16bit(variable->type);
+       LLVMTypeRef type = is_16bit ? ctx->ac.i16 : ctx->ac.i32;
+       if (interp == NULL)
+               interp = LLVMGetUndef(type);
 
        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;
 }
 
 static void
@@ -2150,8 +2169,10 @@ handle_fs_inputs(struct radv_shader_context *ctx,
        unsigned index = 0;
 
        if (ctx->shader_info->info.ps.uses_input_attachments ||
-           ctx->shader_info->info.needs_multiview_view_index)
+           ctx->shader_info->info.needs_multiview_view_index) {
                ctx->input_mask |= 1ull << VARYING_SLOT_LAYER;
+               ctx->inputs[ac_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)] = LLVMGetUndef(ctx->ac.i32);
+       }
 
        for (unsigned i = 0; i < RADEON_LLVM_MAX_INPUTS; ++i) {
                LLVMValueRef interp_param;
@@ -2166,9 +2187,20 @@ handle_fs_inputs(struct radv_shader_context *ctx,
                        interp_fs_input(ctx, index, interp_param, ctx->abi.prim_mask,
                                        inputs);
 
-                       if (!interp_param)
+                       if (LLVMIsUndef(interp_param))
                                ctx->shader_info->fs.flat_shaded_mask |= 1u << index;
                        ++index;
+               } else if (i == VARYING_SLOT_CLIP_DIST0) {
+                       int length = ctx->shader_info->info.ps.num_input_clips_culls;
+
+                       for (unsigned j = 0; j < length; j += 4) {
+                               inputs = ctx->inputs + ac_llvm_reg_index_soa(i, j);
+
+                               interp_param = *inputs;
+                               interp_fs_input(ctx, index, interp_param,
+                                               ctx->abi.prim_mask, inputs);
+                               ++index;
+                       }
                } else if (i == VARYING_SLOT_POS) {
                        for(int i = 0; i < 3; ++i)
                                inputs[i] = ctx->abi.frag_pos[i];
@@ -2205,21 +2237,17 @@ scan_shader_output_decl(struct radv_shader_context *ctx,
            stage == MESA_SHADER_TESS_EVAL ||
            stage == MESA_SHADER_GEOMETRY) {
                if (idx == VARYING_SLOT_CLIP_DIST0) {
-                       int length = shader->info.clip_distance_array_size +
-                                    shader->info.cull_distance_array_size;
                        if (stage == MESA_SHADER_VERTEX) {
                                ctx->shader_info->vs.outinfo.clip_dist_mask = (1 << shader->info.clip_distance_array_size) - 1;
                                ctx->shader_info->vs.outinfo.cull_dist_mask = (1 << shader->info.cull_distance_array_size) - 1;
+                               ctx->shader_info->vs.outinfo.cull_dist_mask <<= shader->info.clip_distance_array_size;
                        }
                        if (stage == MESA_SHADER_TESS_EVAL) {
                                ctx->shader_info->tes.outinfo.clip_dist_mask = (1 << shader->info.clip_distance_array_size) - 1;
                                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;
                        }
 
-                       if (length > 4)
-                               attrib_count = 2;
-                       else
-                               attrib_count = 1;
                        mask_attribs = 1ull << idx;
                }
        }
@@ -2254,7 +2282,11 @@ si_llvm_init_export_args(struct radv_shader_context *ctx,
        args->out[2] = LLVMGetUndef(ctx->ac.f32);
        args->out[3] = LLVMGetUndef(ctx->ac.f32);
 
-       if (ctx->stage == MESA_SHADER_FRAGMENT && target >= V_008DFC_SQ_EXP_MRT) {
+       if (!values)
+               return;
+
+       bool is_16bit = ac_get_type_size(LLVMTypeOf(values[0])) == 2;
+       if (ctx->stage == MESA_SHADER_FRAGMENT) {
                unsigned index = target - V_008DFC_SQ_EXP_MRT;
                unsigned col_format = (ctx->options->key.fs.col_format >> (4 * index)) & 0xf;
                bool is_int8 = (ctx->options->key.fs.is_int8 >> index) & 1;
@@ -2291,6 +2323,12 @@ si_llvm_init_export_args(struct radv_shader_context *ctx,
                case V_028714_SPI_SHADER_FP16_ABGR:
                        args->enabled_channels = 0x5;
                        packf = ac_build_cvt_pkrtz_f16;
+                       if (is_16bit) {
+                               for (unsigned chan = 0; chan < 4; chan++)
+                                       values[chan] = LLVMBuildFPExt(ctx->ac.builder,
+                                                                     values[chan],
+                                                                     ctx->ac.f32, "");
+                       }
                        break;
 
                case V_028714_SPI_SHADER_UNORM16_ABGR:
@@ -2306,11 +2344,23 @@ si_llvm_init_export_args(struct radv_shader_context *ctx,
                case V_028714_SPI_SHADER_UINT16_ABGR:
                        args->enabled_channels = 0x5;
                        packi = ac_build_cvt_pk_u16;
+                       if (is_16bit) {
+                               for (unsigned chan = 0; chan < 4; chan++)
+                                       values[chan] = LLVMBuildZExt(ctx->ac.builder,
+                                                                     values[chan],
+                                                                     ctx->ac.i32, "");
+                       }
                        break;
 
                case V_028714_SPI_SHADER_SINT16_ABGR:
                        args->enabled_channels = 0x5;
                        packi = ac_build_cvt_pk_i16;
+                       if (is_16bit) {
+                               for (unsigned chan = 0; chan < 4; chan++)
+                                       values[chan] = LLVMBuildSExt(ctx->ac.builder,
+                                                                     values[chan],
+                                                                     ctx->ac.i32, "");
+                       }
                        break;
 
                default:
@@ -2353,7 +2403,13 @@ si_llvm_init_export_args(struct radv_shader_context *ctx,
                return;
        }
 
-       memcpy(&args->out[0], values, sizeof(values[0]) * 4);
+       if (is_16bit) {
+               for (unsigned chan = 0; chan < 4; chan++) {
+                       values[chan] = LLVMBuildBitCast(ctx->ac.builder, values[chan], ctx->ac.i16, "");
+                       args->out[chan] = LLVMBuildZExt(ctx->ac.builder, values[chan], ctx->ac.i32, "");
+               }
+       } else
+               memcpy(&args->out[0], values, sizeof(values[0]) * 4);
 
        for (unsigned i = 0; i < 4; ++i) {
                if (!(args->enabled_channels & (1 << i)))
@@ -2411,20 +2467,33 @@ handle_vs_outputs_post(struct radv_shader_context *ctx,
               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 (outinfo->cull_dist_mask)
-                       outinfo->cull_dist_mask <<= ctx->num_output_clips;
+               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];
+               }
+
+               length = util_last_bit(output_usage_mask);
 
                i = VARYING_SLOT_CLIP_DIST0;
-               for (j = 0; j < ctx->num_output_clips + ctx->num_output_culls; j++)
+               for (j = 0; j < length; j++)
                        slots[j] = ac_to_float(&ctx->ac, radv_load_output(ctx, i, j));
 
-               for (i = ctx->num_output_clips + ctx->num_output_culls; i < 8; i++)
+               for (i = length; i < 8; i++)
                        slots[i] = LLVMGetUndef(ctx->ac.f32);
 
-               if (ctx->num_output_clips + ctx->num_output_culls > 4) {
+               if (length > 4) {
                        target = V_008DFC_SQ_EXP_POS + 3;
                        si_llvm_init_export_args(ctx, &slots[4], 0xf, target, &args);
                        memcpy(&pos_args[target - V_008DFC_SQ_EXP_POS],
@@ -2436,6 +2505,13 @@ 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[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++;
+               }
        }
 
        LLVMValueRef pos_values[4] = {ctx->ac.f32_0, ctx->ac.f32_0, ctx->ac.f32_0, ctx->ac.f32_1};
@@ -2589,14 +2665,24 @@ 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 = ctx->num_output_clips + ctx->num_output_culls;
+                       length = util_last_bit(output_usage_mask);
 
                param_index = shader_io_get_unique_index(i);
 
@@ -2608,9 +2694,7 @@ handle_es_outputs_post(struct radv_shader_context *ctx,
        if (ctx->ac.chip_class  >= GFX9) {
                unsigned itemsize_dw = outinfo->esgs_itemsize / 4;
                LLVMValueRef vertex_idx = ac_get_thread_id(&ctx->ac);
-               LLVMValueRef wave_idx = ac_build_bfe(&ctx->ac, ctx->merged_wave_info,
-                                                    LLVMConstInt(ctx->ac.i32, 24, false),
-                                                    LLVMConstInt(ctx->ac.i32, 4, false), false);
+               LLVMValueRef wave_idx = ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 24, 4);
                vertex_idx = LLVMBuildOr(ctx->ac.builder, vertex_idx,
                                         LLVMBuildMul(ctx->ac.builder, wave_idx,
                                                      LLVMConstInt(ctx->ac.i32, 64, false), ""), "");
@@ -2637,10 +2721,8 @@ 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 = ctx->num_output_clips + ctx->num_output_culls;
-                       output_usage_mask = (1 << length) - 1;
-               }
+               if (i == VARYING_SLOT_CLIP_DIST0)
+                       length = util_last_bit(output_usage_mask);
 
                param_index = shader_io_get_unique_index(i);
 
@@ -2655,7 +2737,8 @@ handle_es_outputs_post(struct radv_shader_context *ctx,
                                continue;
 
                        LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, out_ptr[j], "");
-                       out_val = LLVMBuildBitCast(ctx->ac.builder, out_val, ctx->ac.i32, "");
+                       out_val = ac_to_integer(&ctx->ac, out_val);
+                       out_val = LLVMBuildZExtOrBitCast(ctx->ac.builder, out_val, ctx->ac.i32, "");
 
                        if (ctx->ac.chip_class  >= GFX9) {
                                LLVMValueRef dw_addr_offset =
@@ -2663,8 +2746,7 @@ handle_es_outputs_post(struct radv_shader_context *ctx,
                                                     LLVMConstInt(ctx->ac.i32,
                                                                  j, false), "");
 
-                               ac_lds_store(&ctx->ac, dw_addr_offset,
-                                            LLVMBuildLoad(ctx->ac.builder, out_ptr[j], ""));
+                               ac_lds_store(&ctx->ac, dw_addr_offset, out_val);
                        } else {
                                ac_build_buffer_store_dword(&ctx->ac,
                                                            ctx->esgs_ring,
@@ -2687,6 +2769,8 @@ 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;
 
@@ -2694,14 +2778,17 @@ handle_ls_outputs_post(struct radv_shader_context *ctx)
                        continue;
 
                if (i == VARYING_SLOT_CLIP_DIST0)
-                       length = ctx->num_output_clips + ctx->num_output_culls;
+                       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++) {
-                       ac_lds_store(&ctx->ac, dw_addr,
-                                    LLVMBuildLoad(ctx->ac.builder, out_ptr[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, "");
+                       ac_lds_store(&ctx->ac, dw_addr, value);
                        dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr, ctx->ac.i32_1, "");
                }
        }
@@ -3051,24 +3138,76 @@ ac_setup_rings(struct radv_shader_context *ctx)
        }
 
        if (ctx->is_gs_copy_shader) {
-               ctx->gsvs_ring = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_GSVS_VS, false));
+               ctx->gsvs_ring[0] =
+                       ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets,
+                                             LLVMConstInt(ctx->ac.i32,
+                                                          RING_GSVS_VS, false));
        }
+
        if (ctx->stage == MESA_SHADER_GEOMETRY) {
-               LLVMValueRef tmp;
-               uint32_t num_entries = 64;
-               LLVMValueRef gsvs_ring_stride = LLVMConstInt(ctx->ac.i32, ctx->max_gsvs_emit_size, false);
-               LLVMValueRef gsvs_ring_desc = LLVMConstInt(ctx->ac.i32, ctx->max_gsvs_emit_size << 16, false);
-               ctx->gsvs_ring = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_GSVS_GS, false));
+               /* The conceptual layout of the GSVS ring is
+                *   v0c0 .. vLv0 v0c1 .. vLc1 ..
+                * but the real memory layout is swizzled across
+                * threads:
+                *   t0v0c0 .. t15v0c0 t0v1c0 .. t15v1c0 ... t15vLcL
+                *   t16v0c0 ..
+                * Override the buffer descriptor accordingly.
+                */
+               LLVMTypeRef v2i64 = LLVMVectorType(ctx->ac.i64, 2);
+               uint64_t stream_offset = 0;
+               unsigned num_records = 64;
+               LLVMValueRef base_ring;
+
+               base_ring =
+                       ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets,
+                                             LLVMConstInt(ctx->ac.i32,
+                                                          RING_GSVS_GS, false));
+
+               for (unsigned stream = 0; stream < 4; stream++) {
+                       unsigned num_components, stride;
+                       LLVMValueRef ring, tmp;
+
+                       num_components =
+                               ctx->shader_info->info.gs.num_stream_output_components[stream];
+
+                       if (!num_components)
+                               continue;
+
+                       stride = 4 * num_components * ctx->gs_max_out_vertices;
+
+                       /* Limit on the stride field for <= CIK. */
+                       assert(stride < (1 << 14));
+
+                       ring = LLVMBuildBitCast(ctx->ac.builder,
+                                               base_ring, v2i64, "");
+                       tmp = LLVMBuildExtractElement(ctx->ac.builder,
+                                                     ring, ctx->ac.i32_0, "");
+                       tmp = LLVMBuildAdd(ctx->ac.builder, tmp,
+                                          LLVMConstInt(ctx->ac.i64,
+                                                       stream_offset, 0), "");
+                       ring = LLVMBuildInsertElement(ctx->ac.builder,
+                                                     ring, tmp, ctx->ac.i32_0, "");
+
+                       stream_offset += stride * 64;
+
+                       ring = LLVMBuildBitCast(ctx->ac.builder, ring,
+                                               ctx->ac.v4i32, "");
+
+                       tmp = LLVMBuildExtractElement(ctx->ac.builder, ring,
+                                                     ctx->ac.i32_1, "");
+                       tmp = LLVMBuildOr(ctx->ac.builder, tmp,
+                                         LLVMConstInt(ctx->ac.i32,
+                                                      S_008F04_STRIDE(stride), false), "");
+                       ring = LLVMBuildInsertElement(ctx->ac.builder, ring, tmp,
+                                                     ctx->ac.i32_1, "");
 
-               ctx->gsvs_ring = LLVMBuildBitCast(ctx->ac.builder, ctx->gsvs_ring, ctx->ac.v4i32, "");
+                       ring = LLVMBuildInsertElement(ctx->ac.builder, ring,
+                                                     LLVMConstInt(ctx->ac.i32,
+                                                                  num_records, false),
+                                                     LLVMConstInt(ctx->ac.i32, 2, false), "");
 
-               tmp = LLVMConstInt(ctx->ac.i32, num_entries, false);
-               if (ctx->options->chip_class >= VI)
-                       tmp = LLVMBuildMul(ctx->ac.builder, gsvs_ring_stride, tmp, "");
-               ctx->gsvs_ring = LLVMBuildInsertElement(ctx->ac.builder, ctx->gsvs_ring, tmp, LLVMConstInt(ctx->ac.i32, 2, false), "");
-               tmp = LLVMBuildExtractElement(ctx->ac.builder, ctx->gsvs_ring, ctx->ac.i32_1, "");
-               tmp = LLVMBuildOr(ctx->ac.builder, tmp, gsvs_ring_desc, "");
-               ctx->gsvs_ring = LLVMBuildInsertElement(ctx->ac.builder, ctx->gsvs_ring, tmp, ctx->ac.i32_1, "");
+                       ctx->gsvs_ring[stream] = ring;
+               }
        }
 
        if (ctx->stage == MESA_SHADER_TESS_CTRL ||
@@ -3102,9 +3241,7 @@ ac_nir_get_max_workgroup_size(enum chip_class chip_class,
 /* Fixup the HW not emitting the TCS regs if there are no HS threads. */
 static void ac_nir_fixup_ls_hs_input_vgprs(struct radv_shader_context *ctx)
 {
-       LLVMValueRef count = ac_build_bfe(&ctx->ac, ctx->merged_wave_info,
-                                         LLVMConstInt(ctx->ac.i32, 8, false),
-                                         LLVMConstInt(ctx->ac.i32, 8, false), false);
+       LLVMValueRef count = ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 8, 8);
        LLVMValueRef hs_empty = LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ, count,
                                              ctx->ac.i32_0, "");
        ctx->abi.instance_id = LLVMBuildSelect(ctx->ac.builder, hs_empty, ctx->rel_auto_id, ctx->abi.instance_id, "");
@@ -3115,14 +3252,11 @@ static void ac_nir_fixup_ls_hs_input_vgprs(struct radv_shader_context *ctx)
 static void prepare_gs_input_vgprs(struct radv_shader_context *ctx)
 {
        for(int i = 5; i >= 0; --i) {
-               ctx->gs_vtx_offset[i] = ac_build_bfe(&ctx->ac, ctx->gs_vtx_offset[i & ~1],
-                                                    LLVMConstInt(ctx->ac.i32, (i & 1) * 16, false),
-                                                    LLVMConstInt(ctx->ac.i32, 16, false), false);
+               ctx->gs_vtx_offset[i] = ac_unpack_param(&ctx->ac, ctx->gs_vtx_offset[i & ~1],
+                                                       (i & 1) * 16, 16);
        }
 
-       ctx->gs_wave_id = ac_build_bfe(&ctx->ac, ctx->merged_wave_info,
-                                      LLVMConstInt(ctx->ac.i32, 16, false),
-                                      LLVMConstInt(ctx->ac.i32, 8, false), false);
+       ctx->gs_wave_id = ac_unpack_param(&ctx->ac, ctx->merged_wave_info, 16, 8);
 }
 
 
@@ -3188,11 +3322,12 @@ LLVMModuleRef ac_translate_nir_to_llvm(struct ac_llvm_compiler *ac_llvm,
        for(int i = 0; i < shader_count; ++i) {
                ctx.stage = shaders[i]->info.stage;
                ctx.output_mask = 0;
-               ctx.num_output_clips = shaders[i]->info.clip_distance_array_size;
-               ctx.num_output_culls = shaders[i]->info.cull_distance_array_size;
 
                if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY) {
-                       ctx.gs_next_vertex = ac_build_alloca(&ctx.ac, ctx.ac.i32, "gs_next_vertex");
+                       for (int i = 0; i < 4; i++) {
+                               ctx.gs_next_vertex[i] =
+                                       ac_build_alloca(&ctx.ac, ctx.ac.i32, "");
+                       }
                        ctx.gs_max_out_vertices = shaders[i]->info.gs.vertices_out;
                        ctx.abi.load_inputs = load_gs_input;
                        ctx.abi.emit_primitive = visit_end_primitive;
@@ -3256,9 +3391,7 @@ LLVMModuleRef ac_translate_nir_to_llvm(struct ac_llvm_compiler *ac_llvm,
                        LLVMBasicBlockRef then_block = LLVMAppendBasicBlockInContext(ctx.ac.context, fn, "");
                        merge_block = LLVMAppendBasicBlockInContext(ctx.ac.context, fn, "");
 
-                       LLVMValueRef count = ac_build_bfe(&ctx.ac, ctx.merged_wave_info,
-                                                         LLVMConstInt(ctx.ac.i32, 8 * i, false),
-                                                         LLVMConstInt(ctx.ac.i32, 8, false), false);
+                       LLVMValueRef count = ac_unpack_param(&ctx.ac, ctx.merged_wave_info, 8 * i, 8);
                        LLVMValueRef thread_id = ac_get_thread_id(&ctx.ac);
                        LLVMValueRef cond = LLVMBuildICmp(ctx.ac.builder, LLVMIntULT,
                                                          thread_id, count, "");
@@ -3328,12 +3461,7 @@ static unsigned ac_llvm_compile(LLVMModuleRef M,
                                 struct ac_llvm_compiler *ac_llvm)
 {
        unsigned retval = 0;
-       char *err;
        LLVMContextRef llvm_ctx;
-       LLVMMemoryBufferRef out_buffer;
-       unsigned buffer_size;
-       const char *buffer_data;
-       LLVMBool mem_err;
 
        /* Setup Diagnostic Handler*/
        llvm_ctx = LLVMGetModuleContext(M);
@@ -3342,27 +3470,8 @@ static unsigned ac_llvm_compile(LLVMModuleRef M,
                                        &retval);
 
        /* Compile IR*/
-       mem_err = LLVMTargetMachineEmitToMemoryBuffer(ac_llvm->tm, M, LLVMObjectFile,
-                                                     &err, &out_buffer);
-
-       /* Process Errors/Warnings */
-       if (mem_err) {
-               fprintf(stderr, "%s: %s", __FUNCTION__, err);
-               free(err);
+       if (!radv_compile_to_binary(ac_llvm, M, binary))
                retval = 1;
-               goto out;
-       }
-
-       /* Extract Shader Code*/
-       buffer_size = LLVMGetBufferSize(out_buffer);
-       buffer_data = LLVMGetBufferStart(out_buffer);
-
-       ac_elf_read(buffer_data, buffer_size, binary);
-
-       /* Clean up */
-       LLVMDisposeMemoryBuffer(out_buffer);
-
-out:
        return retval;
 }
 
@@ -3530,38 +3639,43 @@ ac_gs_copy_shader_emit(struct radv_shader_context *ctx)
        LLVMValueRef vtx_offset =
                LLVMBuildMul(ctx->ac.builder, ctx->abi.vertex_id,
                             LLVMConstInt(ctx->ac.i32, 4, false), "");
-       int idx = 0;
+       unsigned offset = 0;
 
        for (unsigned i = 0; i < AC_LLVM_MAX_OUTPUTS; ++i) {
-               int length = 4;
-               int slot = idx;
-               int slot_inc = 1;
+               unsigned output_usage_mask =
+                       ctx->shader_info->info.gs.output_usage_mask[i];
+               int length = util_last_bit(output_usage_mask);
+
                if (!(ctx->output_mask & (1ull << i)))
                        continue;
 
-               if (i == VARYING_SLOT_CLIP_DIST0) {
-                       /* unpack clip and cull from a single set of slots */
-                       length = ctx->num_output_clips + ctx->num_output_culls;
-                       if (length > 4)
-                               slot_inc = 2;
-               }
-
                for (unsigned j = 0; j < length; j++) {
                        LLVMValueRef value, soffset;
 
+                       if (!(output_usage_mask & (1 << j)))
+                               continue;
+
                        soffset = LLVMConstInt(ctx->ac.i32,
-                                              (slot * 4 + j) *
+                                              offset *
                                               ctx->gs_max_out_vertices * 16 * 4, false);
 
-                       value = ac_build_buffer_load(&ctx->ac, ctx->gsvs_ring,
+                       offset++;
+
+                       value = ac_build_buffer_load(&ctx->ac,
+                                                    ctx->gsvs_ring[0],
                                                     1, ctx->ac.i32_0,
                                                     vtx_offset, soffset,
                                                     0, 1, 1, true, false);
 
+                       LLVMTypeRef type = LLVMGetAllocatedType(ctx->abi.outputs[ac_llvm_reg_index_soa(i, j)]);
+                       if (ac_get_type_size(type) == 2) {
+                               value = LLVMBuildBitCast(ctx->ac.builder, value, ctx->ac.i32, "");
+                               value = LLVMBuildTrunc(ctx->ac.builder, value, ctx->ac.i16, "");
+                       }
+
                        LLVMBuildStore(ctx->ac.builder,
                                       ac_to_float(&ctx->ac, value), ctx->abi.outputs[ac_llvm_reg_index_soa(i, j)]);
                }
-               idx += slot_inc;
        }
        handle_vs_outputs_post(ctx, false, false, &ctx->shader_info->vs.outinfo);
 }
@@ -3598,9 +3712,6 @@ radv_compile_gs_copy_shader(struct ac_llvm_compiler *ac_llvm,
        ctx.gs_max_out_vertices = geom_shader->info.gs.vertices_out;
        ac_setup_rings(&ctx);
 
-       ctx.num_output_clips = geom_shader->info.clip_distance_array_size;
-       ctx.num_output_culls = geom_shader->info.cull_distance_array_size;
-
        nir_foreach_variable(variable, &geom_shader->outputs) {
                scan_shader_output_decl(&ctx, variable, geom_shader, MESA_SHADER_VERTEX);
                ac_handle_shader_output_decl(&ctx.ac, &ctx.abi, geom_shader,