diff --git a/src/amd/vulkan/meson.build b/src/amd/vulkan/meson.build
index 7715bb4044b..c3a1569c299 100644
--- a/src/amd/vulkan/meson.build
+++ b/src/amd/vulkan/meson.build
@@ -62,6 +62,7 @@ libradv_files = files(
   'meta/radv_meta_resolve.c',
   'meta/radv_meta_resolve_cs.c',
   'meta/radv_meta_resolve_fs.c',
+  'nir/radv_meta_nir.c',
   'nir/radv_nir.h',
   'nir/radv_nir_apply_pipeline_layout.c',
   'nir/radv_nir_export_multiview.c',
diff --git a/src/amd/vulkan/meta/radv_meta.h b/src/amd/vulkan/meta/radv_meta.h
index 27f695927de..fa3eff0bc72 100644
--- a/src/amd/vulkan/meta/radv_meta.h
+++ b/src/amd/vulkan/meta/radv_meta.h
@@ -275,6 +275,9 @@ nir_def *get_global_ids(nir_builder *b, unsigned num_components);
 
 void radv_break_on_count(nir_builder *b, nir_variable *var, nir_def *count);
 
+nir_shader *radv_meta_nir_build_buffer_fill_shader(struct radv_device *dev);
+nir_shader *radv_meta_nir_build_buffer_copy_shader(struct radv_device *dev);
+
 uint32_t radv_fill_buffer(struct radv_cmd_buffer *cmd_buffer, const struct radv_image *image,
                           struct radeon_winsys_bo *bo, uint64_t va, uint64_t size, uint32_t value);
 
diff --git a/src/amd/vulkan/meta/radv_meta_buffer.c b/src/amd/vulkan/meta/radv_meta_buffer.c
index e9edef6693a..a4bb806b061 100644
--- a/src/amd/vulkan/meta/radv_meta_buffer.c
+++ b/src/amd/vulkan/meta/radv_meta_buffer.c
@@ -1,35 +1,11 @@
-#include "nir/nir_builder.h"
 #include "radv_cp_dma.h"
 #include "radv_debug.h"
 #include "radv_meta.h"
 #include "radv_sdma.h"
 
 #include "radv_cs.h"
-#include "sid.h"
 #include "vk_common_entrypoints.h"
 
-static nir_shader *
-build_buffer_fill_shader(struct radv_device *dev)
-{
-   nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_buffer_fill");
-   b.shader->info.workgroup_size[0] = 64;
-
-   nir_def *pconst = nir_load_push_constant(&b, 4, 32, nir_imm_int(&b, 0), .range = 16);
-   nir_def *buffer_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst, 0b0011));
-   nir_def *max_offset = nir_channel(&b, pconst, 2);
-   nir_def *data = nir_swizzle(&b, nir_channel(&b, pconst, 3), (unsigned[]){0, 0, 0, 0}, 4);
-
-   nir_def *global_id =
-      nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b), 0), b.shader->info.workgroup_size[0]),
-               nir_load_local_invocation_index(&b));
-
-   nir_def *offset = nir_imin(&b, nir_imul_imm(&b, global_id, 16), max_offset);
-   nir_def *dst_addr = nir_iadd(&b, buffer_addr, nir_u2u64(&b, offset));
-   nir_build_store_global(&b, data, dst_addr, .align_mul = 4);
-
-   return b.shader;
-}
-
 struct fill_constants {
    uint64_t addr;
    uint32_t max_offset;
@@ -58,7 +34,7 @@ get_fill_pipeline(struct radv_device *device, VkPipeline *pipeline_out, VkPipeli
       return VK_SUCCESS;
    }
 
-   nir_shader *cs = build_buffer_fill_shader(device);
+   nir_shader *cs = radv_meta_nir_build_buffer_fill_shader(device);
 
    const VkPipelineShaderStageCreateInfo stage_info = {
       .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
@@ -82,29 +58,6 @@ get_fill_pipeline(struct radv_device *device, VkPipeline *pipeline_out, VkPipeli
    return result;
 }
 
-static nir_shader *
-build_buffer_copy_shader(struct radv_device *dev)
-{
-   nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_buffer_copy");
-   b.shader->info.workgroup_size[0] = 64;
-
-   nir_def *pconst = nir_load_push_constant(&b, 4, 32, nir_imm_int(&b, 0), .range = 16);
-   nir_def *max_offset = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .base = 16, .range = 4);
-   nir_def *src_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst, 0b0011));
-   nir_def *dst_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst, 0b1100));
-
-   nir_def *global_id =
-      nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b), 0), b.shader->info.workgroup_size[0]),
-               nir_load_local_invocation_index(&b));
-
-   nir_def *offset = nir_u2u64(&b, nir_imin(&b, nir_imul_imm(&b, global_id, 16), max_offset));
-
-   nir_def *data = nir_build_load_global(&b, 4, 32, nir_iadd(&b, src_addr, offset), .align_mul = 4);
-   nir_build_store_global(&b, data, nir_iadd(&b, dst_addr, offset), .align_mul = 4);
-
-   return b.shader;
-}
-
 struct copy_constants {
    uint64_t src_addr;
    uint64_t dst_addr;
@@ -133,7 +86,7 @@ get_copy_pipeline(struct radv_device *device, VkPipeline *pipeline_out, VkPipeli
       return VK_SUCCESS;
    }
 
-   nir_shader *cs = build_buffer_copy_shader(device);
+   nir_shader *cs = radv_meta_nir_build_buffer_copy_shader(device);
 
    const VkPipelineShaderStageCreateInfo stage_info = {
       .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
diff --git a/src/amd/vulkan/nir/radv_meta_nir.c b/src/amd/vulkan/nir/radv_meta_nir.c
new file mode 100644
index 00000000000..1807997ebda
--- /dev/null
+++ b/src/amd/vulkan/nir/radv_meta_nir.c
@@ -0,0 +1,56 @@
+/* Based on anv:
+ * Copyright © 2015 Intel Corporation
+ *
+ * Copyright © 2016 Red Hat Inc.
+ * Copyright © 2018 Valve Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ */
+
+#include "../meta/radv_meta.h"
+#include "nir_builder.h"
+
+nir_shader *
+radv_meta_nir_build_buffer_fill_shader(struct radv_device *dev)
+{
+   nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_buffer_fill");
+   b.shader->info.workgroup_size[0] = 64;
+
+   nir_def *pconst = nir_load_push_constant(&b, 4, 32, nir_imm_int(&b, 0), .range = 16);
+   nir_def *buffer_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst, 0b0011));
+   nir_def *max_offset = nir_channel(&b, pconst, 2);
+   nir_def *data = nir_swizzle(&b, nir_channel(&b, pconst, 3), (unsigned[]){0, 0, 0, 0}, 4);
+
+   nir_def *global_id =
+      nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b), 0), b.shader->info.workgroup_size[0]),
+               nir_load_local_invocation_index(&b));
+
+   nir_def *offset = nir_imin(&b, nir_imul_imm(&b, global_id, 16), max_offset);
+   nir_def *dst_addr = nir_iadd(&b, buffer_addr, nir_u2u64(&b, offset));
+   nir_build_store_global(&b, data, dst_addr, .align_mul = 4);
+
+   return b.shader;
+}
+
+nir_shader *
+radv_meta_nir_build_buffer_copy_shader(struct radv_device *dev)
+{
+   nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_buffer_copy");
+   b.shader->info.workgroup_size[0] = 64;
+
+   nir_def *pconst = nir_load_push_constant(&b, 4, 32, nir_imm_int(&b, 0), .range = 16);
+   nir_def *max_offset = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .base = 16, .range = 4);
+   nir_def *src_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst, 0b0011));
+   nir_def *dst_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst, 0b1100));
+
+   nir_def *global_id =
+      nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b), 0), b.shader->info.workgroup_size[0]),
+               nir_load_local_invocation_index(&b));
+
+   nir_def *offset = nir_u2u64(&b, nir_imin(&b, nir_imul_imm(&b, global_id, 16), max_offset));
+
+   nir_def *data = nir_build_load_global(&b, 4, 32, nir_iadd(&b, src_addr, offset), .align_mul = 4);
+   nir_build_store_global(&b, data, nir_iadd(&b, dst_addr, offset), .align_mul = 4);
+
+   return b.shader;
+}