diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c
index b4228a9b3e6..b4524aaf453 100644
--- a/src/gallium/drivers/radeonsi/si_shader.c
+++ b/src/gallium/drivers/radeonsi/si_shader.c
@@ -167,7 +167,10 @@ static void declare_streamout_params(struct si_shader_args *args, struct si_shad
 
 unsigned si_get_max_workgroup_size(const struct si_shader *shader)
 {
-   switch (shader->selector->stage) {
+   gl_shader_stage stage = shader->is_gs_copy_shader ?
+      MESA_SHADER_VERTEX : shader->selector->stage;
+
+   switch (stage) {
    case MESA_SHADER_VERTEX:
    case MESA_SHADER_TESS_EVAL:
       /* Use the largest workgroup size for streamout */
@@ -2329,7 +2332,7 @@ static void si_determine_use_aco(struct si_shader *shader)
       break;
    case MESA_SHADER_TESS_EVAL:
    case MESA_SHADER_GEOMETRY:
-      shader->use_aco = !si_is_multi_part_shader(shader);
+      shader->use_aco = !si_is_multi_part_shader(shader) || shader->is_gs_copy_shader;
       break;
    case MESA_SHADER_FRAGMENT:
       shader->use_aco = shader->is_monolithic;
@@ -2399,6 +2402,8 @@ si_nir_generate_gs_copy_shader(struct si_screen *sscreen,
                                    sscreen->options.vrs2x2,
                                    output_info);
 
+   si_determine_use_aco(shader);
+
    struct si_shader_args args;
    si_init_shader_args(shader, &args);
 
@@ -2406,13 +2411,20 @@ si_nir_generate_gs_copy_shader(struct si_screen *sscreen,
 
    si_nir_opts(gs_selector->screen, nir, false);
 
+   /* aco only accept scalar const */
+   if (shader->use_aco)
+      NIR_PASS_V(nir, nir_lower_load_const_to_scalar);
+
    if (si_can_dump_shader(sscreen, MESA_SHADER_GEOMETRY, SI_DUMP_NIR)) {
       fprintf(stderr, "GS Copy Shader:\n");
       nir_print_shader(nir, stderr);
    }
 
-   bool ok = false;
-   if (si_llvm_compile_shader(sscreen, compiler, shader, &args, debug, nir)) {
+   bool ok = shader->use_aco ?
+      si_aco_compile_shader(shader, &args, nir, debug) :
+      si_llvm_compile_shader(sscreen, compiler, shader, &args, debug, nir);
+
+   if (ok) {
       assert(!shader->config.scratch_bytes_per_wave);
       ok = si_shader_binary_upload(sscreen, shader, 0);
       si_shader_dump(sscreen, shader, debug, stderr, true);
diff --git a/src/gallium/drivers/radeonsi/si_shader_aco.c b/src/gallium/drivers/radeonsi/si_shader_aco.c
index ea9878ec1da..0836ed36840 100644
--- a/src/gallium/drivers/radeonsi/si_shader_aco.c
+++ b/src/gallium/drivers/radeonsi/si_shader_aco.c
@@ -62,6 +62,7 @@ si_fill_aco_shader_info(struct si_shader *shader, struct aco_shader_info *info)
 {
    const struct si_shader_selector *sel = shader->selector;
    const union si_shader_key *key = &shader->key;
+   gl_shader_stage stage = shader->is_gs_copy_shader ? MESA_SHADER_VERTEX : sel->stage;
 
    info->wave_size = shader->wave_size;
    info->workgroup_size = si_get_max_workgroup_size(shader);
@@ -71,13 +72,13 @@ si_fill_aco_shader_info(struct si_shader *shader, struct aco_shader_info *info)
 
    info->image_2d_view_of_3d = sel->screen->info.gfx_level == GFX9;
 
-   if (sel->stage <= MESA_SHADER_GEOMETRY && key->ge.as_ngg && !key->ge.as_es) {
+   if (stage <= MESA_SHADER_GEOMETRY && key->ge.as_ngg && !key->ge.as_es) {
       info->is_ngg = true;
       info->has_ngg_culling = key->ge.opt.ngg_culling;
       info->has_ngg_early_prim_export = gfx10_ngg_export_prim_early(shader);
    }
 
-   switch (sel->stage) {
+   switch (stage) {
    case MESA_SHADER_VERTEX:
       info->vs.as_es = key->ge.as_es;
       info->vs.as_ls = key->ge.as_ls;