nir: change "user_data_amd" sysval from 4 to 8 components

so that we can pass more fast constants to compute shaders (without
reading memory in the shader).

Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28606>

src/amd/common/ac_shader_args.c

@@ -12,6 +12,7 @@ void ac_add_arg(struct ac_shader_args *info, enum ac_arg_regfile regfile, unsign
                 enum ac_arg_type type, struct ac_arg *arg)
 {
    assert(info->arg_count < AC_MAX_ARGS);
+   assert(nir_num_components_valid(size));
 
    unsigned offset;
    if (regfile == AC_ARG_SGPR) {

src/compiler/nir/nir_intrinsics.py

@@ -1012,7 +1012,7 @@ system_value("color0", 4)
 system_value("color1", 4)
 
 # System value for internal compute shaders in radeonsi.
-system_value("user_data_amd", 4)
+system_value("user_data_amd", 8)
 
 # In a fragment shader, the current sample mask. At the beginning of the shader,
 # this is the same as load_sample_mask_in, but as the shader is executed, it may

src/compiler/shader_info.h

@@ -532,7 +532,7 @@ typedef struct shader_info {
       struct {
          uint16_t workgroup_size_hint[3];
 
-         uint8_t user_data_components_amd:3;
+         uint8_t user_data_components_amd:4;
 
          /*
           * Arrangement of invocations used to calculate derivatives in a compute

src/gallium/drivers/radeonsi/si_nir_lower_abi.c

@@ -746,7 +746,7 @@ static bool lower_intrinsic(nir_builder *b, nir_instr *instr, struct lower_abi_s
    }
    case nir_intrinsic_load_user_data_amd:
       replacement = ac_nir_load_arg(b, &args->ac, args->cs_user_data);
-      replacement = nir_pad_vec4(b, replacement);
+      replacement = nir_pad_vector(b, replacement, 8);
       break;
    default:
       return false;

src/gallium/drivers/radeonsi/si_pipe.h

@@ -1159,7 +1159,7 @@ struct si_context {
    unsigned border_color_count;
    unsigned num_vs_blit_sgprs;
    uint32_t vs_blit_sh_data[MAX_SI_VS_BLIT_SGPRS];
-   uint32_t cs_user_data[4];
+   uint32_t cs_user_data[8];
 
    /* Vertex buffers. */
    bool vertex_buffers_dirty;