radv: allow to emit a vertex to a specified stream
[mesa.git] / src / amd / vulkan / radv_nir_to_llvm.c
index 17c76332e18268afd1da0282b4f2f09e7c7ddda8..07dc6a2301dca0b439ce1d7aee2b5c85a0e70074 100644 (file)
@@ -78,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;
 
@@ -91,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;
@@ -418,7 +416,7 @@ get_tcs_out_current_patch_data_offset(struct radv_shader_context *ctx)
                             patch0_patch_data_offset);
 }
 
-#define MAX_ARGS 23
+#define MAX_ARGS 64
 struct arg_info {
        LLVMTypeRef types[MAX_ARGS];
        LLVMValueRef *assign[MAX_ARGS];
@@ -541,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;
 }
 
@@ -559,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
@@ -573,15 +570,16 @@ 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_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);
-       if (indirect_offset == 0)
+       set_loc(ud_info, sgpr_idx, HAVE_32BIT_POINTERS ? 1 : 2, indirect);
+
+       if (!indirect)
                locs->descriptor_sets_enabled |= 1 << idx;
 }
 
@@ -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), "");
@@ -1704,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
@@ -1723,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;
-               }
-
                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 = 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
@@ -2243,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;
                }
        }
@@ -2296,7 +2286,7 @@ si_llvm_init_export_args(struct radv_shader_context *ctx,
                return;
 
        bool is_16bit = ac_get_type_size(LLVMTypeOf(values[0])) == 2;
-       if (ctx->stage == MESA_SHADER_FRAGMENT && target >= V_008DFC_SQ_EXP_MRT) {
+       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;
@@ -2477,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],
@@ -2505,7 +2508,7 @@ handle_vs_outputs_post(struct radv_shader_context *ctx,
                /* 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 (ctx->num_output_clips + ctx->num_output_culls > 4) {
+               if (length > 4) {
                        radv_export_param(ctx, param_count, &slots[4], 0xf);
                        outinfo->vs_output_param_offset[VARYING_SLOT_CLIP_DIST1] = param_count++;
                }
@@ -2662,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);
 
@@ -2681,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), ""), "");
@@ -2711,7 +2722,7 @@ handle_es_outputs_post(struct radv_shader_context *ctx,
                }
 
                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);
 
@@ -2758,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;
 
@@ -2765,7 +2778,8 @@ 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),
@@ -3124,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, "");
 
-               ctx->gsvs_ring = LLVMBuildBitCast(ctx->ac.builder, ctx->gsvs_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, "");
 
-               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, "");
+                       ring = LLVMBuildInsertElement(ctx->ac.builder, ring,
+                                                     LLVMConstInt(ctx->ac.i32,
+                                                                  num_records, false),
+                                                     LLVMConstInt(ctx->ac.i32, 2, false), "");
+
+                       ctx->gsvs_ring[stream] = ring;
+               }
        }
 
        if (ctx->stage == MESA_SHADER_TESS_CTRL ||
@@ -3175,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, "");
@@ -3188,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);
 }
 
 
@@ -3261,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;
@@ -3329,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, "");
@@ -3579,30 +3639,30 @@ 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);
@@ -3616,7 +3676,6 @@ ac_gs_copy_shader_emit(struct radv_shader_context *ctx)
                        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);
 }
@@ -3653,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,