diff --git a/src/amd/llvm/ac_llvm_build.c b/src/amd/llvm/ac_llvm_build.c
index 12c5e2fb296..312bbd20b8a 100644
--- a/src/amd/llvm/ac_llvm_build.c
+++ b/src/amd/llvm/ac_llvm_build.c
@@ -2243,17 +2243,6 @@ void ac_init_exec_full_mask(struct ac_llvm_context *ctx)
    ac_build_intrinsic(ctx, "llvm.amdgcn.init.exec", ctx->voidt, &full_mask, 1, 0);
 }
 
-void ac_declare_lds_as_pointer(struct ac_llvm_context *ctx)
-{
-   unsigned lds_size = ctx->gfx_level >= GFX7 ? 65536 : 32768;
-   LLVMTypeRef type = LLVMArrayType(ctx->i32, lds_size / 4);
-   ctx->lds = (struct ac_llvm_pointer) {
-      .value = LLVMBuildIntToPtr(ctx->builder, ctx->i32_0,
-                  LLVMPointerType(type, AC_ADDR_SPACE_LDS), "lds"),
-      .pointee_type = type
-   };
-}
-
 LLVMValueRef ac_find_lsb(struct ac_llvm_context *ctx, LLVMTypeRef dst_type, LLVMValueRef src0)
 {
    unsigned src0_bitsize = ac_get_elem_bits(ctx, LLVMTypeOf(src0));
diff --git a/src/amd/llvm/ac_llvm_build.h b/src/amd/llvm/ac_llvm_build.h
index 1e75e29c40a..fa354e8b0b2 100644
--- a/src/amd/llvm/ac_llvm_build.h
+++ b/src/amd/llvm/ac_llvm_build.h
@@ -142,8 +142,6 @@ struct ac_llvm_context {
    bool exports_color_null;
    bool exports_mrtz;
 
-   struct ac_llvm_pointer lds;
-
    LLVMValueRef ring_offsets;
    int ring_offsets_index;
 };
@@ -408,8 +406,6 @@ LLVMValueRef ac_build_sudot_4x8(struct ac_llvm_context *ctx, LLVMValueRef s0, LL
 
 void ac_init_exec_full_mask(struct ac_llvm_context *ctx);
 
-void ac_declare_lds_as_pointer(struct ac_llvm_context *ac);
-
 LLVMValueRef ac_find_lsb(struct ac_llvm_context *ctx, LLVMTypeRef dst_type, LLVMValueRef src0);
 
 LLVMTypeRef ac_array_in_const_addr_space(LLVMTypeRef elem_type);
diff --git a/src/amd/llvm/ac_nir_to_llvm.c b/src/amd/llvm/ac_nir_to_llvm.c
index 1dfb37fdf34..af3c345b1d6 100644
--- a/src/amd/llvm/ac_nir_to_llvm.c
+++ b/src/amd/llvm/ac_nir_to_llvm.c
@@ -31,6 +31,7 @@ struct ac_nir_context {
 
    struct ac_llvm_pointer scratch;
    struct ac_llvm_pointer constant_data;
+   struct ac_llvm_pointer lds;
 
    struct hash_table *defs;
    struct hash_table *phis;
@@ -57,10 +58,20 @@ static LLVMValueRef get_src(struct ac_nir_context *nir, nir_src src)
 
 static LLVMValueRef get_memory_ptr(struct ac_nir_context *ctx, nir_src src, unsigned c_off)
 {
+   if (!ctx->lds.value) {
+      unsigned lds_size = ctx->ac.gfx_level >= GFX7 ? 65536 : 32768;
+      LLVMTypeRef type = LLVMArrayType(ctx->ac.i32, lds_size / 4);
+      ctx->lds = (struct ac_llvm_pointer) {
+         .value = LLVMBuildIntToPtr(ctx->ac.builder, ctx->ac.i32_0,
+                     LLVMPointerType(type, AC_ADDR_SPACE_LDS), "lds"),
+         .pointee_type = type,
+      };
+   }
+
    LLVMValueRef ptr = get_src(ctx, src);
    ptr = LLVMBuildAdd(ctx->ac.builder, ptr, LLVMConstInt(ctx->ac.i32, c_off, 0), "");
    /* LDS is used here as a i8 pointer. */
-   return LLVMBuildGEP2(ctx->ac.builder, ctx->ac.i8, ctx->ac.lds.value, &ptr, 1, "");
+   return LLVMBuildGEP2(ctx->ac.builder, ctx->ac.i8, ctx->lds.value, &ptr, 1, "");
 }
 
 static LLVMBasicBlockRef get_block(struct ac_nir_context *nir, const struct nir_block *b)
@@ -4024,23 +4035,6 @@ static void setup_constant_data(struct ac_nir_context *ctx, struct nir_shader *s
    };
 }
 
-static void setup_shared(struct ac_nir_context *ctx, struct nir_shader *nir)
-{
-   if (ctx->ac.lds.value)
-      return;
-
-   LLVMTypeRef type = LLVMArrayType(ctx->ac.i8, nir->info.shared_size);
-
-   LLVMValueRef lds =
-      LLVMAddGlobalInAddressSpace(ctx->ac.module, type, "compute_lds", AC_ADDR_SPACE_LDS);
-   LLVMSetAlignment(lds, 64 * 1024);
-
-   ctx->ac.lds = (struct ac_llvm_pointer) {
-      .value = lds,
-      .pointee_type = type
-   };
-}
-
 static void setup_gds(struct ac_nir_context *ctx, nir_function_impl *impl)
 {
    bool has_gds_atomic = false;
@@ -4099,9 +4093,6 @@ bool ac_nir_translate(struct ac_llvm_context *ac, struct ac_shader_abi *abi,
    setup_constant_data(&ctx, nir);
    setup_gds(&ctx, func->impl);
 
-   if (gl_shader_stage_is_compute(nir->info.stage))
-      setup_shared(&ctx, nir);
-
    if ((ret = visit_cf_list(&ctx, &func->impl->body)))
       phi_post_pass(&ctx);
 
diff --git a/src/amd/vulkan/radv_nir_to_llvm.c b/src/amd/vulkan/radv_nir_to_llvm.c
index faf68746fd6..ac878a71311 100644
--- a/src/amd/vulkan/radv_nir_to_llvm.c
+++ b/src/amd/vulkan/radv_nir_to_llvm.c
@@ -108,13 +108,6 @@ create_function(struct radv_shader_context *ctx, gl_shader_stage stage, bool has
    ctx->main_function = create_llvm_function(&ctx->ac, ctx->ac.module, ctx->ac.builder, &ctx->args->ac,
                                              get_llvm_calling_convention(ctx->main_function.value, stage),
                                              ctx->max_workgroup_size, ctx->options);
-
-   if (stage == MESA_SHADER_TESS_CTRL || (stage == MESA_SHADER_VERTEX && ctx->shader_info->vs.as_ls) ||
-       ctx->shader_info->is_ngg ||
-       /* GFX9 has the ESGS ring buffer in LDS. */
-       (stage == MESA_SHADER_GEOMETRY && has_previous_stage)) {
-      ac_declare_lds_as_pointer(&ctx->ac);
-   }
 }
 
 static LLVMValueRef
diff --git a/src/gallium/drivers/radeonsi/si_shader_llvm.c b/src/gallium/drivers/radeonsi/si_shader_llvm.c
index 23f9f3bfd48..b9cfe5f8151 100644
--- a/src/gallium/drivers/radeonsi/si_shader_llvm.c
+++ b/src/gallium/drivers/radeonsi/si_shader_llvm.c
@@ -196,23 +196,6 @@ static void si_llvm_create_main_func(struct si_shader_context *ctx)
       ac_llvm_add_target_dep_function_attr(
          ctx->main_fn.value, "InitialPSInputAddr", SI_SPI_PS_INPUT_ADDR_FOR_PROLOG);
    }
-
-
-   if (ctx->stage <= MESA_SHADER_GEOMETRY &&
-       (shader->key.ge.as_ls || ctx->stage == MESA_SHADER_TESS_CTRL ||
-        shader->key.ge.as_es || shader->key.ge.as_ngg)) {
-      /* The LSHS size is not known until draw time, so we append it
-       * at the end of whatever LDS use there may be in the rest of
-       * the shader (currently none, unless LLVM decides to do its
-       * own LDS-based lowering).
-       */
-      ctx->ac.lds = (struct ac_llvm_pointer) {
-         .value = LLVMAddGlobalInAddressSpace(ctx->ac.module, LLVMArrayType(ctx->ac.i32, 0),
-                                                "__lds_end", AC_ADDR_SPACE_LDS),
-         .pointee_type = LLVMArrayType(ctx->ac.i32, 0)
-      };
-      LLVMSetAlignment(ctx->ac.lds.value, 256);
-   }
 }
 
 static void si_llvm_optimize_module(struct si_shader_context *ctx)
@@ -314,27 +297,6 @@ LLVMValueRef si_unpack_param(struct si_shader_context *ctx, struct ac_arg param,
    return unpack_llvm_param(ctx, value, rshift, bitwidth);
 }
 
-static void si_llvm_declare_compute_memory(struct si_shader_context *ctx)
-{
-   struct si_shader_selector *sel = ctx->shader->selector;
-   unsigned lds_size = sel->info.base.shared_size;
-
-   LLVMTypeRef i8p = LLVMPointerType(ctx->ac.i8, AC_ADDR_SPACE_LDS);
-   LLVMValueRef var;
-
-   assert(!ctx->ac.lds.value);
-
-   LLVMTypeRef type = LLVMArrayType(ctx->ac.i8, lds_size);
-   var = LLVMAddGlobalInAddressSpace(ctx->ac.module, type,
-                                     "compute_lds", AC_ADDR_SPACE_LDS);
-   LLVMSetAlignment(var, 64 * 1024);
-
-   ctx->ac.lds = (struct ac_llvm_pointer) {
-      .value = LLVMBuildBitCast(ctx->ac.builder, var, i8p, ""),
-      .pointee_type = type,
-   };
-}
-
 /**
  * Given two parts (LS/HS or ES/GS) of a merged shader, build a wrapper function that
  * runs them in sequence to form a monolithic shader.
@@ -481,16 +443,10 @@ static bool si_llvm_translate_nir(struct si_shader_context *ctx, struct si_shade
    si_llvm_create_main_func(ctx);
 
    switch (ctx->stage) {
-   case MESA_SHADER_VERTEX:
-      break;
-
    case MESA_SHADER_TESS_CTRL:
       si_llvm_init_tcs_callbacks(ctx);
       break;
 
-   case MESA_SHADER_GEOMETRY:
-      break;
-
    case MESA_SHADER_FRAGMENT: {
       ctx->abi.kill_ps_if_inf_interp =
          ctx->screen->options.no_infinite_interp &&
@@ -500,12 +456,6 @@ static bool si_llvm_translate_nir(struct si_shader_context *ctx, struct si_shade
       break;
    }
 
-   case MESA_SHADER_COMPUTE:
-   case MESA_SHADER_KERNEL:
-      if (ctx->shader->selector->info.base.shared_size)
-         si_llvm_declare_compute_memory(ctx);
-      break;
-
    default:
       break;
    }