diff --git a/source/blender/draw/intern/draw_command.cc b/source/blender/draw/intern/draw_command.cc
index 33082a07c9c..0597ff2b344 100644
--- a/source/blender/draw/intern/draw_command.cc
+++ b/source/blender/draw/intern/draw_command.cc
@@ -179,10 +179,12 @@ void Draw::execute(RecordingState &state) const
 
   if (is_primitive_expansion()) {
     /* Expanded drawcall. */
-    IndexRange vert_range = GPU_batch_draw_expanded_parameter_get(
-        batch, GPUPrimType(expand_prim_type), vertex_len, vertex_first);
-    IndexRange expanded_range = {vert_range.start() * expand_prim_len,
-                                 vert_range.size() * expand_prim_len};
+    IndexRange expanded_range = GPU_batch_draw_expanded_parameter_get(
+        batch->prim_type,
+        GPUPrimType(expand_prim_type),
+        vertex_len,
+        vertex_first,
+        expand_prim_len);
 
     if (expanded_range.is_empty()) {
       /* Nothing to draw, and can lead to asserts in GPU_batch_bind_as_resources. */
@@ -803,13 +805,14 @@ void DrawMultiBuf::generate_commands(Vector<Header, 0> & /*headers*/,
     if (group.desc.expand_prim_type != GPU_PRIM_NONE) {
       /* Expanded drawcall. */
       IndexRange vert_range = GPU_batch_draw_expanded_parameter_get(
-          group.desc.gpu_batch,
+          group.desc.gpu_batch->prim_type,
           GPUPrimType(group.desc.expand_prim_type),
           group.vertex_len,
-          group.vertex_first);
+          group.vertex_first,
+          group.desc.expand_prim_len);
 
-      group.vertex_first = vert_range.start() * group.desc.expand_prim_len;
-      group.vertex_len = vert_range.size() * group.desc.expand_prim_len;
+      group.vertex_first = vert_range.start();
+      group.vertex_len = vert_range.size();
       /* Override base index to -1 as the generated drawcall will not use an index buffer and do
        * the indirection manually inside the shader. */
       group.base_index = -1;
diff --git a/source/blender/gpu/CMakeLists.txt b/source/blender/gpu/CMakeLists.txt
index 5edf87e4ef3..335db665910 100644
--- a/source/blender/gpu/CMakeLists.txt
+++ b/source/blender/gpu/CMakeLists.txt
@@ -504,9 +504,7 @@ set(GLSL_SRC
   shaders/gpu_shader_3D_flat_color_vert.glsl
   shaders/gpu_shader_3D_line_dashed_uniform_color_vert.glsl
   shaders/gpu_shader_3D_polyline_frag.glsl
-  shaders/gpu_shader_3D_polyline_geom.glsl
   shaders/gpu_shader_3D_polyline_vert.glsl
-  shaders/gpu_shader_3D_polyline_vert_no_geom.glsl
   shaders/gpu_shader_3D_smooth_color_vert.glsl
   shaders/gpu_shader_3D_smooth_color_frag.glsl
   shaders/gpu_shader_3D_clipped_uniform_color_vert.glsl
diff --git a/source/blender/gpu/GPU_batch.hh b/source/blender/gpu/GPU_batch.hh
index ce08fb4811b..b0f395cb0dd 100644
--- a/source/blender/gpu/GPU_batch.hh
+++ b/source/blender/gpu/GPU_batch.hh
@@ -425,10 +425,11 @@ void GPU_batch_draw_parameter_get(blender::gpu::Batch *batch,
 /**
  * Return vertex range for this #blender::gpu::Batch when using primitive expansions.
  */
-blender::IndexRange GPU_batch_draw_expanded_parameter_get(const blender::gpu::Batch *batch,
-                                                          GPUPrimType expanded_prim_type,
+blender::IndexRange GPU_batch_draw_expanded_parameter_get(GPUPrimType input_prim_type,
+                                                          GPUPrimType output_prim_type,
                                                           int vertex_count,
-                                                          int vertex_first);
+                                                          int vertex_first,
+                                                          int output_primitive_cout);
 
 /** \} */
 
diff --git a/source/blender/gpu/GPU_index_buffer.hh b/source/blender/gpu/GPU_index_buffer.hh
index 55f15b8f986..80bc899790a 100644
--- a/source/blender/gpu/GPU_index_buffer.hh
+++ b/source/blender/gpu/GPU_index_buffer.hh
@@ -134,6 +134,14 @@ inline int indices_per_primitive(GPUPrimType prim_type)
       return 4;
     case GPU_PRIM_TRIS_ADJ:
       return 6;
+    /** IMPORTANT: These last two expects no restart primitive.
+     * Asserting for this would be too slow. Just don't be stupid.
+     * This is needed for polylines but should be deprecated.
+     * See GPU_batch_draw_expanded_parameter_get */
+    case GPU_PRIM_LINE_STRIP:
+      return 1; /* Minus one for the whole length. */
+    case GPU_PRIM_LINE_LOOP:
+      return 1;
     default:
       return -1;
   }
diff --git a/source/blender/gpu/GPU_shader.hh b/source/blender/gpu/GPU_shader.hh
index cd5d65fe988..38248862d46 100644
--- a/source/blender/gpu/GPU_shader.hh
+++ b/source/blender/gpu/GPU_shader.hh
@@ -191,6 +191,7 @@ void GPU_shader_uniform_2fv(GPUShader *sh, const char *name, const float data[2]
 void GPU_shader_uniform_3fv(GPUShader *sh, const char *name, const float data[3]);
 void GPU_shader_uniform_4fv(GPUShader *sh, const char *name, const float data[4]);
 void GPU_shader_uniform_2iv(GPUShader *sh, const char *name, const int data[2]);
+void GPU_shader_uniform_3iv(GPUShader *sh, const char *name, const int data[3]);
 void GPU_shader_uniform_mat4(GPUShader *sh, const char *name, const float data[4][4]);
 void GPU_shader_uniform_mat3_as_mat4(GPUShader *sh, const char *name, const float data[3][3]);
 void GPU_shader_uniform_1f_array(GPUShader *sh, const char *name, int len, const float *val);
diff --git a/source/blender/gpu/intern/gpu_batch.cc b/source/blender/gpu/intern/gpu_batch.cc
index 16aa4279aef..622a28045ec 100644
--- a/source/blender/gpu/intern/gpu_batch.cc
+++ b/source/blender/gpu/intern/gpu_batch.cc
@@ -350,44 +350,83 @@ void GPU_batch_draw_parameter_get(Batch *gpu_batch,
   *r_instance_count = i_count;
 }
 
-blender::IndexRange GPU_batch_draw_expanded_parameter_get(const blender::gpu::Batch *batch,
-                                                          GPUPrimType expanded_prim_type,
+blender::IndexRange GPU_batch_draw_expanded_parameter_get(GPUPrimType input_prim_type,
+                                                          GPUPrimType output_prim_type,
                                                           int vertex_count,
-                                                          int vertex_first)
+                                                          int vertex_first,
+                                                          int output_primitive_cout)
 {
-  int vert_per_original_primitive = indices_per_primitive(batch->prim_type);
-  int vert_per_expanded_primitive = indices_per_primitive(expanded_prim_type);
-
-  BLI_assert_msg(vert_per_original_primitive != -1,
-                 "Primitive expansion only works for primitives with known amount of vertices");
+  int vert_per_original_primitive = indices_per_primitive(input_prim_type);
+  int vert_per_expanded_primitive = indices_per_primitive(output_prim_type);
 
   int prim_first = vertex_first / vert_per_original_primitive;
   int prim_len = vertex_count / vert_per_original_primitive;
 
-  int out_vertex_first = prim_first * vert_per_expanded_primitive;
-  int out_vertex_count = prim_len * vert_per_expanded_primitive;
+  BLI_assert_msg(vert_per_original_primitive != -1,
+                 "Primitive expansion only works for primitives with known amount of vertices");
+
+  /* WORKAROUND: Needed for polyline_draw_workaround. */
+  if (input_prim_type == GPU_PRIM_LINE_STRIP) {
+    prim_len = vertex_count - 1;
+  }
+
+  int out_vertex_first = prim_first * vert_per_expanded_primitive * output_primitive_cout;
+  int out_vertex_count = prim_len * vert_per_expanded_primitive * output_primitive_cout;
 
   return blender::IndexRange(out_vertex_first, out_vertex_count);
 }
 
+static void polyline_draw_workaround(
+    Batch *batch, int vertex_first, int vertex_count, int instance_first, int instance_count)
+{
+  /* Check compatible input primitive. */
+  BLI_assert(ELEM(batch->prim_type, GPU_PRIM_LINES, GPU_PRIM_LINE_STRIP, GPU_PRIM_LINE_LOOP));
+
+  GPU_batch_bind_as_resources(batch, batch->shader);
+  blender::IndexRange range = GPU_batch_draw_expanded_parameter_get(
+      batch->prim_type, GPU_PRIM_TRIS, vertex_count, vertex_first, 2);
+  Batch *tri_batch = Context::get()->polyline_batch_get();
+  GPU_batch_set_shader(tri_batch, batch->shader);
+
+  int vert_stride_count[3] = {(batch->prim_type == GPU_PRIM_LINES) ? 2 : 1, int(vertex_count), 0};
+  GPU_shader_uniform_3iv(batch->shader, "gpu_vert_stride_count_offset", vert_stride_count);
+  /* Assume GPU_FETCH_FLOAT for now. A bit cumbersome to assert for this or to find the correct
+   * attribute. */
+  GPU_shader_uniform_1b(batch->shader, "gpu_attr_0_fetch_int", false);
+
+  GPU_batch_draw_advanced(tri_batch, range.start(), range.size(), instance_first, instance_count);
+}
+
 void GPU_batch_draw(Batch *batch)
 {
   BLI_assert(batch != nullptr);
   GPU_shader_bind(batch->shader);
-  GPU_batch_draw_advanced(batch, 0, 0, 0, 0);
+  if (unwrap(batch->shader)->is_polyline) {
+    polyline_draw_workaround(batch, 0, batch->vertex_count_get(), 0, 0);
+  }
+  else {
+    GPU_batch_draw_advanced(batch, 0, 0, 0, 0);
+  }
 }
 
 void GPU_batch_draw_range(Batch *batch, int vertex_first, int vertex_count)
 {
   BLI_assert(batch != nullptr);
   GPU_shader_bind(batch->shader);
-  GPU_batch_draw_advanced(batch, vertex_first, vertex_count, 0, 0);
+  if (unwrap(batch->shader)->is_polyline) {
+    polyline_draw_workaround(batch, vertex_first, vertex_count, 0, 0);
+  }
+  else {
+    GPU_batch_draw_advanced(batch, vertex_first, vertex_count, 0, 0);
+  }
 }
 
 void GPU_batch_draw_instance_range(Batch *batch, int instance_first, int instance_count)
 {
   BLI_assert(batch != nullptr);
   BLI_assert(batch->inst[0] == nullptr);
+  /* Not polyline shaders support instancing. */
+  BLI_assert(unwrap(batch->shader)->is_polyline == false);
 
   GPU_shader_bind(batch->shader);
   GPU_batch_draw_advanced(batch, 0, 0, instance_first, instance_count);
diff --git a/source/blender/gpu/intern/gpu_context.cc b/source/blender/gpu/intern/gpu_context.cc
index 0840650ceba..f0447bbc42b 100644
--- a/source/blender/gpu/intern/gpu_context.cc
+++ b/source/blender/gpu/intern/gpu_context.cc
@@ -76,6 +76,7 @@ Context::Context()
 Context::~Context()
 {
   GPU_matrix_state_discard(matrix_state);
+  GPU_BATCH_DISCARD_SAFE(polyline_batch);
   delete state_manager;
   delete front_left;
   delete back_left;
@@ -94,6 +95,22 @@ Context *Context::get()
   return active_ctx;
 }
 
+Batch *Context::polyline_batch_get()
+{
+  if (polyline_batch) {
+    return polyline_batch;
+  }
+
+  /* TODO(fclem): get rid of this dummy VBO. */
+  GPUVertFormat format = {0};
+  GPU_vertformat_attr_add(&format, "dummy", GPU_COMP_F32, 1, GPU_FETCH_FLOAT);
+  blender::gpu::VertBuf *vbo = GPU_vertbuf_create_with_format(format);
+  GPU_vertbuf_data_alloc(*vbo, 1);
+
+  polyline_batch = GPU_batch_create_ex(GPU_PRIM_TRIS, vbo, nullptr, GPU_BATCH_OWNS_VBO);
+  return polyline_batch;
+}
+
 }  // namespace blender::gpu
 
 /** \} */
diff --git a/source/blender/gpu/intern/gpu_context_private.hh b/source/blender/gpu/intern/gpu_context_private.hh
index ec34cbb50aa..8362e7ee21a 100644
--- a/source/blender/gpu/intern/gpu_context_private.hh
+++ b/source/blender/gpu/intern/gpu_context_private.hh
@@ -12,6 +12,7 @@
 
 #include "MEM_guardedalloc.h"
 
+#include "GPU_batch.hh"
 #include "GPU_context.hh"
 
 #include "gpu_debug_private.hh"
@@ -63,6 +64,9 @@ class Context {
   /* Used as a stack. Each render_begin/end pair will push pop from the stack. */
   Vector<GPUStorageBuf *> printf_buf;
 
+  /** Dummy triangle batch for polyline workaround. */
+  Batch *polyline_batch = nullptr;
+
  protected:
   /** Thread on which this context is active. */
   pthread_t thread_;
@@ -104,6 +108,8 @@ class Context {
   virtual void debug_unbind_all_ssbo() = 0;
 
   bool is_active_on_thread();
+
+  Batch *polyline_batch_get();
 };
 
 /* Syntactic sugar. */
diff --git a/source/blender/gpu/intern/gpu_immediate.cc b/source/blender/gpu/intern/gpu_immediate.cc
index f2fefac87c3..04523680112 100644
--- a/source/blender/gpu/intern/gpu_immediate.cc
+++ b/source/blender/gpu/intern/gpu_immediate.cc
@@ -277,6 +277,71 @@ void immEnd()
   wide_line_workaround_end();
 }
 
+void Immediate::polyline_draw_workaround(uint64_t offset)
+{
+  /* Check compatible input primitive. */
+  BLI_assert(ELEM(imm->prim_type, GPU_PRIM_LINES, GPU_PRIM_LINE_STRIP, GPU_PRIM_LINE_LOOP));
+
+  Batch *tri_batch = Context::get()->polyline_batch_get();
+  GPU_batch_set_shader(tri_batch, imm->shader);
+
+  BLI_assert(offset % 4 == 0);
+
+  /* Setup primitive and index buffer. */
+  int stride = (imm->prim_type == GPU_PRIM_LINES) ? 2 : 1;
+  int data[3] = {stride, int(imm->vertex_idx), int(offset / 4)};
+  GPU_shader_uniform_3iv(imm->shader, "gpu_vert_stride_count_offset", data);
+  GPU_shader_uniform_1b(imm->shader, "gpu_index_no_buffer", true);
+
+  {
+    /* Setup attributes metadata uniforms. */
+    const GPUVertFormat &format = imm->vertex_format;
+    /* Only support 4byte aligned formats. */
+    BLI_assert((format.stride % 4) == 0);
+    BLI_assert(format.attr_len > 0);
+
+    int pos_attr_id = -1;
+    int col_attr_id = -1;
+
+    for (uint a_idx = 0; a_idx < format.attr_len; a_idx++) {
+      const GPUVertAttr *a = &format.attrs[a_idx];
+      const char *name = GPU_vertformat_attr_name_get(&format, a, 0);
+      if (pos_attr_id == -1 && blender::StringRefNull(name) == "pos") {
+        int descriptor[2] = {int(format.stride) / 4, int(a->offset) / 4};
+        BLI_assert(ELEM(a->comp_type, GPU_COMP_F32, GPU_COMP_I32));
+        BLI_assert(ELEM(a->fetch_mode, GPU_FETCH_FLOAT, GPU_FETCH_INT_TO_FLOAT));
+        BLI_assert_msg((a->offset % 4) == 0, "Only support 4byte aligned attributes");
+        const bool fetch_int = a->fetch_mode == GPU_FETCH_INT_TO_FLOAT;
+        GPU_shader_uniform_2iv(imm->shader, "gpu_attr_0", descriptor);
+        GPU_shader_uniform_1i(imm->shader, "gpu_attr_0_len", a->comp_len);
+        GPU_shader_uniform_1b(imm->shader, "gpu_attr_0_fetch_int", fetch_int);
+        pos_attr_id = a_idx;
+      }
+      else if (col_attr_id == -1 && blender::StringRefNull(name) == "color") {
+        int descriptor[2] = {int(format.stride) / 4, int(a->offset) / 4};
+        /* Maybe we can relax this if needed. */
+        BLI_assert_msg(a->comp_type == GPU_COMP_F32, "Only support float attributes");
+        BLI_assert_msg((a->offset % 4) == 0, "Only support 4byte aligned attributes");
+        GPU_shader_uniform_2iv(imm->shader, "gpu_attr_1", descriptor);
+        GPU_shader_uniform_1i(imm->shader, "gpu_attr_1_len", a->comp_len);
+        col_attr_id = a_idx;
+      }
+      if (pos_attr_id != -1 && col_attr_id != -1) {
+        break;
+      }
+    }
+
+    BLI_assert(pos_attr_id != -1);
+    /* Could check for color attribute but we need to know which variant of the polyline shader is
+     * the one we are rendering with. */
+    // BLI_assert(pos_attr_id != -1);
+  }
+
+  blender::IndexRange range = GPU_batch_draw_expanded_parameter_get(
+      imm->prim_type, GPU_PRIM_TRIS, imm->vertex_idx, 0, 2);
+  GPU_batch_draw_advanced(tri_batch, range.start(), range.size(), 0, 0);
+}
+
 static void setAttrValueBit(uint attr_id)
 {
   uint16_t mask = 1 << attr_id;
diff --git a/source/blender/gpu/intern/gpu_immediate_private.hh b/source/blender/gpu/intern/gpu_immediate_private.hh
index 7f83fc450af..cc1ab3ce11c 100644
--- a/source/blender/gpu/intern/gpu_immediate_private.hh
+++ b/source/blender/gpu/intern/gpu_immediate_private.hh
@@ -56,6 +56,9 @@ class Immediate {
 
   virtual uchar *begin() = 0;
   virtual void end() = 0;
+
+  /* To be called after polyline SSBO binding. */
+  void polyline_draw_workaround(uint64_t offset);
 };
 
 }  // namespace blender::gpu
diff --git a/source/blender/gpu/intern/gpu_shader.cc b/source/blender/gpu/intern/gpu_shader.cc
index d38c3792678..1d5c135e62b 100644
--- a/source/blender/gpu/intern/gpu_shader.cc
+++ b/source/blender/gpu/intern/gpu_shader.cc
@@ -786,6 +786,12 @@ void GPU_shader_uniform_2iv(GPUShader *sh, const char *name, const int data[2])
   GPU_shader_uniform_int_ex(sh, loc, 2, 1, data);
 }
 
+void GPU_shader_uniform_3iv(GPUShader *sh, const char *name, const int data[3])
+{
+  const int loc = GPU_shader_get_uniform(sh, name);
+  GPU_shader_uniform_int_ex(sh, loc, 3, 1, data);
+}
+
 void GPU_shader_uniform_mat4(GPUShader *sh, const char *name, const float data[4][4])
 {
   const int loc = GPU_shader_get_uniform(sh, name);
diff --git a/source/blender/gpu/intern/gpu_shader_builtin.cc b/source/blender/gpu/intern/gpu_shader_builtin.cc
index f047563f66a..63bf7fe0ba2 100644
--- a/source/blender/gpu/intern/gpu_shader_builtin.cc
+++ b/source/blender/gpu/intern/gpu_shader_builtin.cc
@@ -11,6 +11,8 @@
 #include "GPU_capabilities.hh"
 #include "GPU_shader.hh"
 
+#include "gpu_shader_private.hh"
+
 /* Cache of built-in shaders (each is created on first use). */
 static GPUShader *builtin_shaders[GPU_SHADER_CFG_LEN][GPU_SHADER_BUILTIN_LEN] = {{nullptr}};
 
@@ -157,6 +159,8 @@ GPUShader *GPU_shader_get_builtin_shader_with_config(eGPUBuiltinShader shader,
          * Ideally this value should be set by the caller. */
         GPU_shader_bind(*sh_p);
         GPU_shader_uniform_1i(*sh_p, "lineSmooth", 1);
+        /* WORKAROUND: See is_polyline declaration. */
+        blender::gpu::unwrap(*sh_p)->is_polyline = true;
       }
     }
     else if (sh_cfg == GPU_SHADER_CFG_CLIPPED) {
diff --git a/source/blender/gpu/intern/gpu_shader_create_info.cc b/source/blender/gpu/intern/gpu_shader_create_info.cc
index e6b17f4512f..4ee007f2528 100644
--- a/source/blender/gpu/intern/gpu_shader_create_info.cc
+++ b/source/blender/gpu/intern/gpu_shader_create_info.cc
@@ -494,13 +494,6 @@ void gpu_shader_create_info_init()
 #ifdef WITH_METAL_BACKEND
   /* Metal-specific alternatives for Geometry shaders. */
   if (GPU_type_matches_ex(GPU_DEVICE_ANY, GPU_OS_MAC, GPU_DRIVER_ANY, GPU_BACKEND_METAL)) {
-    /* 3D polyline. */
-    gpu_shader_3D_polyline_uniform_color = gpu_shader_3D_polyline_uniform_color_no_geom;
-    gpu_shader_3D_polyline_flat_color = gpu_shader_3D_polyline_flat_color_no_geom;
-    gpu_shader_3D_polyline_smooth_color = gpu_shader_3D_polyline_smooth_color_no_geom;
-    gpu_shader_3D_polyline_uniform_color_clipped =
-        gpu_shader_3D_polyline_uniform_color_clipped_no_geom;
-
     /* Overlay Edit Mesh. */
     overlay_edit_mesh_edge = overlay_edit_mesh_edge_no_geom;
     overlay_edit_mesh_edge_flat = overlay_edit_mesh_edge_flat_no_geom;
diff --git a/source/blender/gpu/intern/gpu_shader_create_info.hh b/source/blender/gpu/intern/gpu_shader_create_info.hh
index 057afa4b9be..8be382235c7 100644
--- a/source/blender/gpu/intern/gpu_shader_create_info.hh
+++ b/source/blender/gpu/intern/gpu_shader_create_info.hh
@@ -38,6 +38,9 @@
 
 /* Used for primitive expansion. */
 #define GPU_SSBO_INDEX_BUF_SLOT 7
+/* Used for polylines. */
+#define GPU_SSBO_POLYLINE_POS_BUF_SLOT 0
+#define GPU_SSBO_POLYLINE_COL_BUF_SLOT 1
 
 #if defined(GLSL_CPP_STUBS)
 #  define GPU_SHADER_NAMED_INTERFACE_INFO(_interface, _inst_name) \
diff --git a/source/blender/gpu/intern/gpu_shader_private.hh b/source/blender/gpu/intern/gpu_shader_private.hh
index b1e94b6dc36..ae31d129582 100644
--- a/source/blender/gpu/intern/gpu_shader_private.hh
+++ b/source/blender/gpu/intern/gpu_shader_private.hh
@@ -64,6 +64,10 @@ class Shader {
     bool is_dirty;
   } constants;
 
+  /* WORKAROUND: True if this shader is a polyline shader and needs an appropriate setup to render.
+   * Eventually, in the future, we should modify the user code instead of relying on such hacks. */
+  bool is_polyline = false;
+
  protected:
   /** For debugging purpose. */
   char name[64];
diff --git a/source/blender/gpu/metal/mtl_context.mm b/source/blender/gpu/metal/mtl_context.mm
index 6b134a66303..680ed01239a 100644
--- a/source/blender/gpu/metal/mtl_context.mm
+++ b/source/blender/gpu/metal/mtl_context.mm
@@ -246,15 +246,15 @@ MTLContext::MTLContext(void *ghost_window, void *ghost_context)
   MTLBackend::platform_init(this);
   MTLBackend::capabilities_init(this);
 
+  /* Ensure global memory manager is initialized. */
+  MTLContext::global_memory_manager_acquire_ref();
+  MTLContext::get_global_memory_manager()->init(this->device);
+
   /* Initialize Metal modules. */
   this->memory_manager.init();
   this->state_manager = new MTLStateManager(this);
   this->imm = new MTLImmediate(this);
 
-  /* Ensure global memory manager is initialized. */
-  MTLContext::global_memory_manager_acquire_ref();
-  MTLContext::get_global_memory_manager()->init(this->device);
-
   /* Initialize texture read/update structures. */
   this->get_texture_utils().init();
 
diff --git a/source/blender/gpu/metal/mtl_immediate.mm b/source/blender/gpu/metal/mtl_immediate.mm
index 75406dc715d..c414f29eb45 100644
--- a/source/blender/gpu/metal/mtl_immediate.mm
+++ b/source/blender/gpu/metal/mtl_immediate.mm
@@ -268,7 +268,14 @@ void MTLImmediate::end()
                                 (const int *)(&this->vertex_idx));
     }
 
+    if (unwrap(this->shader)->is_polyline) {
+      context_->get_scratchbuffer_manager().bind_as_ssbo(GPU_SSBO_POLYLINE_POS_BUF_SLOT);
+      context_->get_scratchbuffer_manager().bind_as_ssbo(GPU_SSBO_POLYLINE_COL_BUF_SLOT);
+      context_->get_scratchbuffer_manager().bind_as_ssbo(GPU_SSBO_INDEX_BUF_SLOT);
+    }
+
     MTLPrimitiveType mtl_prim_type = gpu_prim_type_to_metal(this->prim_type);
+
     if (context_->ensure_render_pipeline_state(mtl_prim_type)) {
 
       /* Issue draw call. */
@@ -364,7 +371,10 @@ void MTLImmediate::end()
         /* Set depth stencil state (requires knowledge of primitive type). */
         context_->ensure_depth_stencil_state(primitive_type);
 
-        if (active_mtl_shader->get_uses_ssbo_vertex_fetch()) {
+        if (unwrap(this->shader)->is_polyline) {
+          this->polyline_draw_workaround(current_allocation_.buffer_offset);
+        }
+        else if (active_mtl_shader->get_uses_ssbo_vertex_fetch()) {
 
           /* Bind Null Buffers for empty/missing bind slots. */
           id<MTLBuffer> null_buffer = context_->get_null_buffer();
@@ -412,6 +422,17 @@ void MTLImmediate::end()
     if (G.debug & G_DEBUG_GPU) {
       [rec popDebugGroup];
     }
+
+    if (unwrap(this->shader)->is_polyline) {
+      context_->get_scratchbuffer_manager().unbind_as_ssbo();
+
+      context_->pipeline_state.ssbo_bindings[GPU_SSBO_POLYLINE_POS_BUF_SLOT].ssbo = nil;
+      context_->pipeline_state.ssbo_bindings[GPU_SSBO_POLYLINE_COL_BUF_SLOT].ssbo = nil;
+      context_->pipeline_state.ssbo_bindings[GPU_SSBO_INDEX_BUF_SLOT].ssbo = nil;
+      context_->pipeline_state.ssbo_bindings[GPU_SSBO_POLYLINE_POS_BUF_SLOT].bound = false;
+      context_->pipeline_state.ssbo_bindings[GPU_SSBO_POLYLINE_COL_BUF_SLOT].bound = false;
+      context_->pipeline_state.ssbo_bindings[GPU_SSBO_INDEX_BUF_SLOT].bound = false;
+    }
   }
 
   /* Reset allocation after draw submission. */
diff --git a/source/blender/gpu/metal/mtl_memory.hh b/source/blender/gpu/metal/mtl_memory.hh
index f6d9f5e1e50..566bb7945aa 100644
--- a/source/blender/gpu/metal/mtl_memory.hh
+++ b/source/blender/gpu/metal/mtl_memory.hh
@@ -99,6 +99,7 @@ namespace blender::gpu {
 class MTLContext;
 class MTLCommandBufferManager;
 class MTLUniformBuf;
+class MTLStorageBuf;
 
 /* -------------------------------------------------------------------- */
 /** \name Memory Management.
@@ -211,6 +212,8 @@ class MTLCircularBuffer {
 
   /* Wrapped MTLBuffer allocation handled. */
   gpu::MTLBuffer *cbuffer_;
+  /* Allocated SSBO that serves as source for cbuffer. */
+  MTLStorageBuf *ssbo_source_ = nullptr;
 
   /* Current offset where next allocation will begin. */
   uint64_t current_offset_;
@@ -513,6 +516,10 @@ class MTLScratchBufferManager {
    * the last offset the data was flushed from, to the current offset. */
   void flush_active_scratch_buffer();
 
+  /* Bind the whole scratch buffer as a SSBO resource. */
+  void bind_as_ssbo(int slot);
+  void unbind_as_ssbo();
+
   MEM_CXX_CLASS_ALLOC_FUNCS("MTLBufferPool");
 };
 
diff --git a/source/blender/gpu/metal/mtl_memory.mm b/source/blender/gpu/metal/mtl_memory.mm
index af3cbba925e..ff0da6a6657 100644
--- a/source/blender/gpu/metal/mtl_memory.mm
+++ b/source/blender/gpu/metal/mtl_memory.mm
@@ -6,9 +6,12 @@
 
 #include "DNA_userdef_types.h"
 
+#include "BLI_math_base.h"
+
 #include "mtl_context.hh"
 #include "mtl_debug.hh"
 #include "mtl_memory.hh"
+#include "mtl_storage_buffer.hh"
 
 using namespace blender;
 using namespace blender::gpu;
@@ -646,7 +649,6 @@ bool MTLSafeFreeList::should_flush()
 /** \name MTLBuffer wrapper class implementation.
  * \{ */
 
-/* Construct a gpu::MTLBuffer wrapper around a newly created metal::MTLBuffer. */
 MTLBuffer::MTLBuffer(id<MTLDevice> mtl_device,
                      uint64_t size,
                      MTLResourceOptions options,
@@ -916,15 +918,29 @@ void MTLScratchBufferManager::flush_active_scratch_buffer()
   active_scratch_buf->flush();
 }
 
+void MTLScratchBufferManager::bind_as_ssbo(int slot)
+{
+  /* Fetch active scratch buffer and verify context. */
+  MTLCircularBuffer *active_scratch_buf = scratch_buffers_[current_scratch_buffer_];
+  BLI_assert(&active_scratch_buf->own_context_ == &context_);
+  active_scratch_buf->ssbo_source_->bind(slot);
+}
+
+void MTLScratchBufferManager::unbind_as_ssbo()
+{
+  /* Fetch active scratch buffer and verify context. */
+  MTLCircularBuffer *active_scratch_buf = scratch_buffers_[current_scratch_buffer_];
+  BLI_assert(&active_scratch_buf->own_context_ == &context_);
+  active_scratch_buf->ssbo_source_->unbind();
+}
+
 /* MTLCircularBuffer implementation. */
 MTLCircularBuffer::MTLCircularBuffer(MTLContext &ctx, uint64_t initial_size, bool allow_grow)
     : own_context_(ctx)
 {
   BLI_assert(this);
-  MTLResourceOptions options = ([own_context_.device hasUnifiedMemory]) ?
-                                   MTLResourceStorageModeShared :
-                                   MTLResourceStorageModeManaged;
-  cbuffer_ = new gpu::MTLBuffer(own_context_.device, initial_size, options, 256);
+  ssbo_source_ = new gpu::MTLStorageBuf(initial_size);
+  cbuffer_ = ssbo_source_->metal_buffer_;
   current_offset_ = 0;
   can_resize_ = allow_grow;
   cbuffer_->flag_in_use(true);
@@ -940,7 +956,7 @@ MTLCircularBuffer::MTLCircularBuffer(MTLContext &ctx, uint64_t initial_size, boo
 
 MTLCircularBuffer::~MTLCircularBuffer()
 {
-  delete cbuffer_;
+  delete ssbo_source_;
 }
 
 MTLTemporaryBuffer MTLCircularBuffer::allocate_range(uint64_t alloc_size)
@@ -1033,10 +1049,9 @@ MTLTemporaryBuffer MTLCircularBuffer::allocate_range_aligned(uint64_t alloc_size
 
     /* Discard old buffer and create a new one - Relying on Metal reference counting to track
      * in-use buffers */
-    MTLResourceOptions prev_options = cbuffer_->get_resource_options();
-    uint prev_alignment = cbuffer_->get_alignment();
-    delete cbuffer_;
-    cbuffer_ = new gpu::MTLBuffer(own_context_.device, new_size, prev_options, prev_alignment);
+    delete ssbo_source_;
+    ssbo_source_ = new gpu::MTLStorageBuf(new_size);
+    cbuffer_ = ssbo_source_->metal_buffer_;
     cbuffer_->flag_in_use(true);
     current_offset_ = 0;
     last_flush_base_offset_ = 0;
diff --git a/source/blender/gpu/metal/mtl_storage_buffer.hh b/source/blender/gpu/metal/mtl_storage_buffer.hh
index c2ad9d477d1..57f5d03ffaa 100644
--- a/source/blender/gpu/metal/mtl_storage_buffer.hh
+++ b/source/blender/gpu/metal/mtl_storage_buffer.hh
@@ -18,11 +18,14 @@ namespace blender::gpu {
 class MTLUniformBuf;
 class MTLVertBuf;
 class MTLIndexBuf;
+class MTLCircularBuffer;
 
 /**
  * Implementation of Storage Buffers using Metal.
  */
 class MTLStorageBuf : public StorageBuf {
+  friend MTLCircularBuffer;
+
  private:
   /** Allocation Handle or indirect wrapped instance.
    * MTLStorageBuf can wrap a MTLVertBuf, MTLIndexBuf or MTLUniformBuf for binding as a writeable
@@ -68,6 +71,9 @@ class MTLStorageBuf : public StorageBuf {
   MTLStorageBuf(MTLIndexBuf *index_buf, size_t size);
   MTLStorageBuf(MTLTexture *texture, size_t size);
 
+  /* Only used internally to create a bindable buffer for #Immediate. */
+  MTLStorageBuf(size_t size);
+
   void update(const void *data) override;
   void bind(int slot) override;
   void unbind() override;
diff --git a/source/blender/gpu/metal/mtl_storage_buffer.mm b/source/blender/gpu/metal/mtl_storage_buffer.mm
index 52cc53a5502..552101c831d 100644
--- a/source/blender/gpu/metal/mtl_storage_buffer.mm
+++ b/source/blender/gpu/metal/mtl_storage_buffer.mm
@@ -17,9 +17,11 @@
 #include "mtl_context.hh"
 #include "mtl_debug.hh"
 #include "mtl_index_buffer.hh"
+#include "mtl_memory.hh"
 #include "mtl_storage_buffer.hh"
 #include "mtl_uniform_buffer.hh"
 #include "mtl_vertex_buffer.hh"
+#include <cstdio>
 
 namespace blender::gpu {
 
@@ -27,6 +29,13 @@ namespace blender::gpu {
 /** \name Creation & Deletion
  * \{ */
 
+MTLStorageBuf::MTLStorageBuf(size_t size) : StorageBuf(size, "Immediate")
+{
+  usage_ = GPU_USAGE_STREAM;
+  storage_source_ = MTL_STORAGE_BUF_TYPE_DEFAULT;
+  metal_buffer_ = MTLContext::get_global_memory_manager()->allocate_aligned(size, 256, true);
+}
+
 MTLStorageBuf::MTLStorageBuf(size_t size, GPUUsageType usage, const char *name)
     : StorageBuf(size, name)
 {
diff --git a/source/blender/gpu/opengl/gl_immediate.cc b/source/blender/gpu/opengl/gl_immediate.cc
index 26a5cda778f..b0d36ea31d1 100644
--- a/source/blender/gpu/opengl/gl_immediate.cc
+++ b/source/blender/gpu/opengl/gl_immediate.cc
@@ -69,6 +69,15 @@ uchar *GLImmediate::begin()
   /* Does the current buffer have enough room? */
   const size_t available_bytes = buffer_size() - buffer_offset();
 
+#ifndef NDEBUG
+  if (unwrap(this->shader)->is_polyline) {
+    /* Silence error. These are bound inside `immEnd()`. */
+    GLContext::get()->bound_ssbo_slots |= 1 << GPU_SSBO_POLYLINE_POS_BUF_SLOT;
+    GLContext::get()->bound_ssbo_slots |= 1 << GPU_SSBO_POLYLINE_COL_BUF_SLOT;
+    GLContext::get()->bound_ssbo_slots |= 1 << GPU_SSBO_INDEX_BUF_SLOT;
+  }
+#endif
+
   GL_CHECK_RESOURCES("Immediate");
 
   glBindBuffer(GL_ARRAY_BUFFER, vbo_id());
@@ -135,7 +144,26 @@ void GLImmediate::end()
   }
   glUnmapBuffer(GL_ARRAY_BUFFER);
 
-  if (vertex_len > 0) {
+  if (vertex_len == 0) {
+    /* Noop. Nothing to draw. */
+  }
+  else if (unwrap(this->shader)->is_polyline) {
+    GLintptr offset = buffer_offset();
+    GLenum target = GL_SHADER_STORAGE_BUFFER;
+    glBindBufferRange(target, GPU_SSBO_POLYLINE_POS_BUF_SLOT, vbo_id(), offset, buffer_bytes_used);
+    glBindBufferRange(target, GPU_SSBO_POLYLINE_COL_BUF_SLOT, vbo_id(), offset, buffer_bytes_used);
+    /* Not used. Satisfy the binding. */
+    glBindBufferRange(target, GPU_SSBO_INDEX_BUF_SLOT, vbo_id(), offset, buffer_bytes_used);
+
+    this->polyline_draw_workaround(0);
+
+#ifndef NDEBUG
+    GLContext::get()->bound_ssbo_slots &= ~(1 << GPU_SSBO_POLYLINE_POS_BUF_SLOT);
+    GLContext::get()->bound_ssbo_slots &= ~(1 << GPU_SSBO_POLYLINE_COL_BUF_SLOT);
+    GLContext::get()->bound_ssbo_slots &= ~(1 << GPU_SSBO_INDEX_BUF_SLOT);
+#endif
+  }
+  else {
     GLContext::get()->state_manager->apply_state();
 
     /* We convert the offset in vertex offset from the buffer's start.
diff --git a/source/blender/gpu/shaders/gpu_shader_3D_polyline_frag.glsl b/source/blender/gpu/shaders/gpu_shader_3D_polyline_frag.glsl
index 065bdaaf836..331d51be275 100644
--- a/source/blender/gpu/shaders/gpu_shader_3D_polyline_frag.glsl
+++ b/source/blender/gpu/shaders/gpu_shader_3D_polyline_frag.glsl
@@ -11,14 +11,13 @@ FRAGMENT_SHADER_CREATE_INFO(gpu_shader_3D_polyline_uniform_color)
 void main()
 {
 #ifdef CLIP
-  if (interp.clip < 0.0) {
+  if (clip < 0.0) {
     discard;
   }
 #endif
-  fragColor = interp.final_color;
+  fragColor = final_color;
   if (lineSmooth) {
-    fragColor.a *= clamp(
-        (lineWidth + SMOOTH_WIDTH) * 0.5 - abs(interp_noperspective.smoothline), 0.0, 1.0);
+    fragColor.a *= clamp((lineWidth + SMOOTH_WIDTH) * 0.5 - abs(smoothline), 0.0, 1.0);
   }
   fragColor = blender_srgb_to_framebuffer_space(fragColor);
 }
diff --git a/source/blender/gpu/shaders/gpu_shader_3D_polyline_geom.glsl b/source/blender/gpu/shaders/gpu_shader_3D_polyline_geom.glsl
deleted file mode 100644
index 03670a630b2..00000000000
--- a/source/blender/gpu/shaders/gpu_shader_3D_polyline_geom.glsl
+++ /dev/null
@@ -1,68 +0,0 @@
-/* SPDX-FileCopyrightText: 2020-2023 Blender Authors
- *
- * SPDX-License-Identifier: GPL-2.0-or-later */
-
-/* Clips point to near clip plane before perspective divide. */
-vec4 clip_line_point_homogeneous_space(vec4 p, vec4 q)
-{
-  if (p.z < -p.w) {
-    /* Just solves p + (q - p) * A; for A when p.z / p.w = -1.0. */
-    float denom = q.z - p.z + q.w - p.w;
-    if (denom == 0.0) {
-      /* No solution. */
-      return p;
-    }
-    float A = (-p.z - p.w) / denom;
-    p = p + (q - p) * A;
-  }
-  return p;
-}
-
-void do_vertex(const int i, vec4 pos, vec2 ofs)
-{
-#if defined(UNIFORM)
-  interp_out.final_color = color;
-
-#elif defined(FLAT)
-  /* WATCH: Assuming last provoking vertex. */
-  interp_out.final_color = interp_in[1].final_color;
-
-#elif defined(SMOOTH)
-  interp_out.final_color = interp_in[i].final_color;
-#endif
-
-#ifdef CLIP
-  interp_out.clip = interp_in[i].clip;
-#endif
-
-  interp_noperspective_out.smoothline = (lineWidth + SMOOTH_WIDTH * float(lineSmooth)) * 0.5;
-  gl_Position = pos;
-  gl_Position.xy += ofs * pos.w;
-  gpu_EmitVertex();
-
-  interp_noperspective_out.smoothline = -(lineWidth + SMOOTH_WIDTH * float(lineSmooth)) * 0.5;
-  gl_Position = pos;
-  gl_Position.xy -= ofs * pos.w;
-  gpu_EmitVertex();
-}
-
-void main()
-{
-  vec4 p0 = clip_line_point_homogeneous_space(gl_in[0].gl_Position, gl_in[1].gl_Position);
-  vec4 p1 = clip_line_point_homogeneous_space(gl_in[1].gl_Position, gl_in[0].gl_Position);
-  vec2 e = normalize(((p1.xy / p1.w) - (p0.xy / p0.w)) * viewportSize.xy);
-
-#if 0 /* Hard turn when line direction changes quadrant. */
-  e = abs(e);
-  vec2 ofs = (e.x > e.y) ? vec2(0.0, 1.0 / e.x) : vec2(1.0 / e.y, 0.0);
-#else /* Use perpendicular direction. */
-  vec2 ofs = vec2(-e.y, e.x);
-#endif
-  ofs /= viewportSize.xy;
-  ofs *= lineWidth + SMOOTH_WIDTH * float(lineSmooth);
-
-  do_vertex(0, p0, ofs);
-  do_vertex(1, p1, ofs);
-
-  EndPrimitive();
-}
diff --git a/source/blender/gpu/shaders/gpu_shader_3D_polyline_vert.glsl b/source/blender/gpu/shaders/gpu_shader_3D_polyline_vert.glsl
index 7b1ad5e5fd6..2d842cb2659 100644
--- a/source/blender/gpu/shaders/gpu_shader_3D_polyline_vert.glsl
+++ b/source/blender/gpu/shaders/gpu_shader_3D_polyline_vert.glsl
@@ -2,13 +2,206 @@
  *
  * SPDX-License-Identifier: GPL-2.0-or-later */
 
-void main()
+#include "gpu_shader_attribute_load_lib.glsl"
+#include "gpu_shader_index_load_lib.glsl"
+#include "gpu_shader_math_base_lib.glsl"
+#include "gpu_shader_utildefines_lib.glsl"
+
+struct VertIn {
+  vec3 ls_P;
+  vec4 final_color;
+};
+
+VertIn input_assembly(uint in_vertex_id)
 {
-  gl_Position = ModelViewProjectionMatrix * vec4(pos, 1.0);
+  uint v_i = gpu_index_load(in_vertex_id);
+  uint ofs = uint(gpu_vert_stride_count_offset.z);
+
+  VertIn vert_in;
+  vert_in.ls_P = vec3(0.0, 0.0, 0.0);
+  /* Need to support 1, 2 and 3 dimensional input (sigh). */
+  vert_in.ls_P.x = pos[gpu_attr_load_index(v_i, gpu_attr_0) + 0 + ofs];
+  if (gpu_attr_0_len >= 2) {
+    vert_in.ls_P.y = pos[gpu_attr_load_index(v_i, gpu_attr_0) + 1 + ofs];
+  }
+  if (gpu_attr_0_len >= 3) {
+    vert_in.ls_P.z = pos[gpu_attr_load_index(v_i, gpu_attr_0) + 2 + ofs];
+  }
+
+  if (gpu_attr_0_fetch_int) {
+    vert_in.ls_P = vec3(floatBitsToInt(vert_in.ls_P));
+  }
 #ifndef UNIFORM
-  interp.final_color = color;
+  vert_in.final_color = vec4(0.0, 0.0, 0.0, 1.0);
+  /* Need to support 1, 2, 3 and 4 dimensional input (sigh). */
+  vert_in.final_color.x = color[gpu_attr_load_index(v_i, gpu_attr_1) + 0 + ofs];
+  if (gpu_attr_1_len >= 2) {
+    vert_in.final_color.y = color[gpu_attr_load_index(v_i, gpu_attr_1) + 1 + ofs];
+  }
+  if (gpu_attr_1_len >= 3) {
+    vert_in.final_color.z = color[gpu_attr_load_index(v_i, gpu_attr_1) + 2 + ofs];
+  }
+  if (gpu_attr_1_len >= 4) {
+    vert_in.final_color.w = color[gpu_attr_load_index(v_i, gpu_attr_1) + 3 + ofs];
+  }
+#endif
+  return vert_in;
+}
+
+struct VertOut {
+  vec4 gpu_position;
+  vec4 final_color;
+  float clip;
+};
+
+VertOut vertex_main(VertIn vert_in)
+{
+  VertOut vert_out;
+  vert_out.gpu_position = ModelViewProjectionMatrix * vec4(vert_in.ls_P, 1.0);
+#ifndef UNIFORM
+  vert_out.final_color = vert_in.final_color;
 #endif
 #ifdef CLIP
-  interp.clip = dot(ModelMatrix * vec4(pos, 1.0), ClipPlane);
+  vert_out.clip = dot(ModelMatrix * vec4(vert_in.ls_P, 1.0), ClipPlane);
 #endif
+  return vert_out;
+}
+
+/* Clips point to near clip plane before perspective divide. */
+vec4 clip_line_point_homogeneous_space(vec4 p, vec4 q)
+{
+  if (p.z < -p.w) {
+    /* Just solves p + (q - p) * A; for A when p.z / p.w = -1.0. */
+    float denom = q.z - p.z + q.w - p.w;
+    if (denom == 0.0) {
+      /* No solution. */
+      return p;
+    }
+    float A = (-p.z - p.w) / denom;
+    p = p + (q - p) * A;
+  }
+  return p;
+}
+
+struct GeomOut {
+  vec4 gpu_position;
+  vec4 final_color;
+  float clip;
+  float smoothline;
+};
+
+void export_vertex(GeomOut geom_out)
+{
+  gl_Position = geom_out.gpu_position;
+  final_color = geom_out.final_color;
+  smoothline = geom_out.smoothline;
+  clip = geom_out.clip;
+}
+
+void strip_EmitVertex(const uint strip_index,
+                      uint out_vertex_id,
+                      uint out_primitive_id,
+                      GeomOut geom_out)
+{
+  bool is_odd_primitive = (out_primitive_id & 1u) != 0u;
+  /* Maps triangle list primitives to triangle strip indices. */
+  uint out_strip_index = (is_odd_primitive ? (2u - out_vertex_id) : out_vertex_id) +
+                         out_primitive_id;
+
+  if (out_strip_index == strip_index) {
+    export_vertex(geom_out);
+  }
+}
+
+void do_vertex(const uint i,
+               uint out_vertex_id,
+               uint out_primitive_id,
+               VertOut geom_in[2],
+               vec4 position,
+               vec2 ofs)
+{
+  GeomOut geom_out;
+#if defined(UNIFORM)
+  geom_out.final_color = color;
+
+#elif defined(FLAT)
+  /* WATCH: Assuming last provoking vertex. */
+  geom_out.final_color = geom_in[1].final_color;
+
+#elif defined(SMOOTH)
+  geom_out.final_color = geom_in[i].final_color;
+#endif
+
+#ifdef CLIP
+  geom_out.clip = geom_in[i].clip;
+#endif
+
+  geom_out.smoothline = (lineWidth + SMOOTH_WIDTH * float(lineSmooth)) * 0.5;
+  geom_out.gpu_position = position;
+  geom_out.gpu_position.xy += ofs * position.w;
+  strip_EmitVertex(i * 2u + 0u, out_vertex_id, out_primitive_id, geom_out);
+
+  geom_out.smoothline = -(lineWidth + SMOOTH_WIDTH * float(lineSmooth)) * 0.5;
+  geom_out.gpu_position = position;
+  geom_out.gpu_position.xy -= ofs * position.w;
+  strip_EmitVertex(i * 2u + 1u, out_vertex_id, out_primitive_id, geom_out);
+}
+
+void geometry_main(VertOut geom_in[2],
+                   uint out_vertex_id,
+                   uint out_primitive_id,
+                   uint out_invocation_id)
+{
+  vec4 p0 = clip_line_point_homogeneous_space(geom_in[0].gpu_position, geom_in[1].gpu_position);
+  vec4 p1 = clip_line_point_homogeneous_space(geom_in[1].gpu_position, geom_in[0].gpu_position);
+  vec2 e = normalize(((p1.xy / p1.w) - (p0.xy / p0.w)) * viewportSize.xy);
+
+#if 0 /* Hard turn when line direction changes quadrant. */
+  e = abs(e);
+  vec2 ofs = (e.x > e.y) ? vec2(0.0, 1.0 / e.x) : vec2(1.0 / e.y, 0.0);
+#else /* Use perpendicular direction. */
+  vec2 ofs = vec2(-e.y, e.x);
+#endif
+  ofs /= viewportSize.xy;
+  ofs *= lineWidth + SMOOTH_WIDTH * float(lineSmooth);
+
+  do_vertex(0u, out_vertex_id, out_primitive_id, geom_in, p0, ofs);
+  do_vertex(1u, out_vertex_id, out_primitive_id, geom_in, p1, ofs);
+}
+
+void main()
+{
+  /* Line list primitive. */
+  uint input_primitive_vertex_count = uint(gpu_vert_stride_count_offset.x);
+  /* Triangle list primitive (emulating triangle strip). */
+  const uint ouput_primitive_vertex_count = 3u;
+  const uint ouput_primitive_count = 2u;
+  const uint ouput_invocation_count = 1u;
+  const uint output_vertex_count_per_invocation = ouput_primitive_count *
+                                                  ouput_primitive_vertex_count;
+  const uint output_vertex_count_per_input_primitive = output_vertex_count_per_invocation *
+                                                       ouput_invocation_count;
+
+  uint in_primitive_id = uint(gl_VertexID) / output_vertex_count_per_input_primitive;
+  uint in_primitive_first_vertex = in_primitive_id * input_primitive_vertex_count;
+
+  uint out_vertex_id = uint(gl_VertexID) % ouput_primitive_vertex_count;
+  uint out_primitive_id = (uint(gl_VertexID) / ouput_primitive_vertex_count) %
+                          ouput_primitive_count;
+  uint out_invocation_id = (uint(gl_VertexID) / output_vertex_count_per_invocation) %
+                           ouput_invocation_count;
+  /* Used to wrap around for the line loop case. */
+  uint input_total_vertex_count = uint(gpu_vert_stride_count_offset.y);
+
+  VertIn vert_in[2];
+  vert_in[0] = input_assembly(in_primitive_first_vertex + 0u);
+  vert_in[1] = input_assembly((in_primitive_first_vertex + 1u) % input_total_vertex_count);
+
+  VertOut vert_out[2];
+  vert_out[0] = vertex_main(vert_in[0]);
+  vert_out[1] = vertex_main(vert_in[1]);
+
+  /* Discard by default. */
+  gl_Position = vec4(NAN_FLT);
+  geometry_main(vert_out, out_vertex_id, out_primitive_id, out_invocation_id);
 }
diff --git a/source/blender/gpu/shaders/gpu_shader_3D_polyline_vert_no_geom.glsl b/source/blender/gpu/shaders/gpu_shader_3D_polyline_vert_no_geom.glsl
deleted file mode 100644
index 853c674be58..00000000000
--- a/source/blender/gpu/shaders/gpu_shader_3D_polyline_vert_no_geom.glsl
+++ /dev/null
@@ -1,165 +0,0 @@
-/* SPDX-FileCopyrightText: 2022-2023 Blender Authors
- *
- * SPDX-License-Identifier: GPL-2.0-or-later */
-
-#pragma USE_SSBO_VERTEX_FETCH(TriangleList, 6)
-
-/* Local vars to store results per input vertex. */
-#if !defined(UNIFORM)
-vec4 finalColor_g[2];
-#endif
-
-#ifdef CLIP
-float clip_g[2];
-#endif
-
-#define SMOOTH_WIDTH 1.0
-
-/* Clips point to near clip plane before perspective divide. */
-vec4 clip_line_point_homogeneous_space(vec4 p, vec4 q)
-{
-  if (p.z < -p.w) {
-    /* Just solves p + (q - p) * A; for A when p.z / p.w = -1.0. */
-    float denom = q.z - p.z + q.w - p.w;
-    if (denom == 0.0) {
-      /* No solution. */
-      return p;
-    }
-    float A = (-p.z - p.w) / denom;
-    p = p + (q - p) * A;
-  }
-  return p;
-}
-
-void do_vertex(int index, vec4 pos, vec2 ofs, float flip)
-{
-#if defined(UNIFORM)
-  interp.final_color = color;
-
-#elif defined(FLAT)
-  /* WATCH: Assuming last provoking vertex. */
-  interp.final_color = finalColor_g[index];
-
-#elif defined(SMOOTH)
-  interp.final_color = finalColor_g[index];
-#endif
-
-#ifdef CLIP
-  interp.clip = clip_g[index];
-#endif
-
-  interp_noperspective.smoothline = flip * (lineWidth + SMOOTH_WIDTH * float(lineSmooth)) * 0.5;
-  gl_Position = pos;
-  gl_Position.xy += flip * ofs * pos.w;
-}
-
-void main()
-{
-  /** Determine output quad primitive structure. */
-  /* Index of the quad primitive. Each quad corresponds to one line in the input primitive. */
-  int quad_id = gl_VertexID / 6;
-
-  /* Determine vertex within the quad (A, B, C)(A, C, D). */
-  int quad_vertex_id = gl_VertexID % 6;
-
-  uint src_index_a;
-  uint src_index_b;
-  if (vertex_fetch_get_input_prim_type() == GPU_PRIM_LINE_STRIP) {
-    src_index_a = quad_id;
-    src_index_b = quad_id + 1;
-  }
-  else if (vertex_fetch_get_input_prim_type() == GPU_PRIM_LINES) {
-    src_index_a = quad_id * 2;
-    src_index_b = quad_id * 2 + 1;
-  }
-  else if (vertex_fetch_get_input_prim_type() == GPU_PRIM_LINE_LOOP) {
-    src_index_a = quad_id;
-    src_index_b = quad_id + 1;
-    if (quad_id == vertex_fetch_get_input_vert_count() - 1) {
-      src_index_b = 0;
-    }
-  }
-  else {
-    src_index_a = 0;
-    src_index_b = 0;
-  }
-
-  /* Fetch input attributes for line prims -- either provided as vec2 or vec3 -- So we need to
-   * query the type. */
-  vec3 in_pos0, in_pos1;
-  in_pos0 = vec3(0.0);
-  in_pos1 = vec3(0.0);
-  if (vertex_fetch_get_attr_type(pos) == GPU_SHADER_ATTR_TYPE_VEC4) {
-    in_pos0 = vertex_fetch_attribute(src_index_a, pos, vec4).xyz;
-    in_pos1 = vertex_fetch_attribute(src_index_b, pos, vec4).xyz;
-  }
-  else if (vertex_fetch_get_attr_type(pos) == GPU_SHADER_ATTR_TYPE_VEC3) {
-    in_pos0 = vertex_fetch_attribute(src_index_a, pos, vec3);
-    in_pos1 = vertex_fetch_attribute(src_index_b, pos, vec3);
-  }
-  else if (vertex_fetch_get_attr_type(pos) == GPU_SHADER_ATTR_TYPE_VEC2) {
-    in_pos0 = vec3(vertex_fetch_attribute(src_index_a, pos, vec2), 0.0);
-    in_pos1 = vec3(vertex_fetch_attribute(src_index_b, pos, vec2), 0.0);
-  }
-#if !defined(UNIFORM)
-  vec4 in_color0 = vec4(0.0);
-  vec4 in_color1 = vec4(0.0);
-
-  if (vertex_fetch_get_attr_type(color) == GPU_SHADER_ATTR_TYPE_VEC4) {
-    in_color0 = vertex_fetch_attribute(src_index_a, color, vec4);
-    in_color1 = vertex_fetch_attribute(src_index_b, color, vec4);
-  }
-  else if (vertex_fetch_get_attr_type(color) == GPU_SHADER_ATTR_TYPE_VEC3) {
-    in_color0 = vec4(vertex_fetch_attribute(src_index_a, color, vec3), 1.0);
-    in_color1 = vec4(vertex_fetch_attribute(src_index_b, color, vec3), 1.0);
-  }
-  else if (vertex_fetch_get_attr_type(color) == GPU_SHADER_ATTR_TYPE_UCHAR4_NORM) {
-    in_color0 = vec4(vertex_fetch_attribute(src_index_a, color, uchar4)) / vec4(255.0);
-    in_color1 = vec4(vertex_fetch_attribute(src_index_b, color, uchar4)) / vec4(255.0);
-  }
-  else if (vertex_fetch_get_attr_type(color) == GPU_SHADER_ATTR_TYPE_UCHAR3_NORM) {
-    in_color0 = vec4(vec3(vertex_fetch_attribute(src_index_a, color, uchar3)) / vec3(255.0), 1.0);
-    in_color1 = vec4(vec3(vertex_fetch_attribute(src_index_b, color, uchar3)) / vec3(255.0), 1.0);
-  }
-#endif
-
-  /* Calculate Vertex shader for both points in Line. */
-  vec4 out_pos0 = ModelViewProjectionMatrix * vec4(in_pos0, 1.0);
-  vec4 out_pos1 = ModelViewProjectionMatrix * vec4(in_pos1, 1.0);
-#if !defined(UNIFORM)
-  finalColor_g[0] = in_color0;
-  finalColor_g[1] = in_color1;
-#endif
-#ifdef CLIP
-  clip_g[0] = dot(ModelMatrix * vec4(in_pos0, 1.0), ClipPlane);
-  clip_g[1] = dot(ModelMatrix * vec4(in_pos1, 1.0), ClipPlane);
-#endif
-
-  /** Geometry Shader Alternative. */
-  vec4 p0 = clip_line_point_homogeneous_space(out_pos0, out_pos1);
-  vec4 p1 = clip_line_point_homogeneous_space(out_pos1, out_pos0);
-  vec2 e = normalize(((p1.xy / p1.w) - (p0.xy / p0.w)) * viewportSize.xy);
-
-#if 0 /* Hard turn when line direction changes quadrant. */
-  e = abs(e);
-  vec2 ofs = (e.x > e.y) ? vec2(0.0, 1.0 / e.x) : vec2(1.0 / e.y, 0.0);
-#else /* Use perpendicular direction. */
-  vec2 ofs = vec2(-e.y, e.x);
-#endif
-
-  ofs /= viewportSize.xy;
-  ofs *= lineWidth + SMOOTH_WIDTH * float(lineSmooth);
-
-  if (quad_vertex_id == 0) {
-    do_vertex(0, p0, ofs, 1.0);
-  }
-  else if (quad_vertex_id == 1 || quad_vertex_id == 3) {
-    do_vertex(0, p0, ofs, -1.0);
-  }
-  else if (quad_vertex_id == 2 || quad_vertex_id == 5) {
-    do_vertex(1, p1, ofs, 1.0);
-  }
-  else if (quad_vertex_id == 4) {
-    do_vertex(1, p1, ofs, -1.0);
-  }
-}
diff --git a/source/blender/gpu/shaders/infos/gpu_shader_3D_polyline_info.hh b/source/blender/gpu/shaders/infos/gpu_shader_3D_polyline_info.hh
index 96daf65155f..785f51bd3e4 100644
--- a/source/blender/gpu/shaders/infos/gpu_shader_3D_polyline_info.hh
+++ b/source/blender/gpu/shaders/infos/gpu_shader_3D_polyline_info.hh
@@ -11,6 +11,7 @@
 #  include "gpu_glsl_cpp_stubs.hh"
 
 #  include "GPU_shader_shared.hh"
+#  include "gpu_index_load_info.hh"
 #  include "gpu_srgb_to_framebuffer_space_info.hh"
 #  define SMOOTH_WIDTH 1.0
 #endif
@@ -18,14 +19,11 @@
 #include "gpu_interface_info.hh"
 #include "gpu_shader_create_info.hh"
 
-GPU_SHADER_NAMED_INTERFACE_INFO(gpu_shader_3D_polyline_iface, interp)
+GPU_SHADER_INTERFACE_INFO(gpu_shader_3D_polyline_iface)
 SMOOTH(VEC4, final_color)
 SMOOTH(FLOAT, clip)
-GPU_SHADER_NAMED_INTERFACE_END(interp)
-
-GPU_SHADER_NAMED_INTERFACE_INFO(gpu_shader_3D_polyline_noperspective_iface, interp_noperspective)
 NO_PERSPECTIVE(FLOAT, smoothline)
-GPU_SHADER_NAMED_INTERFACE_END(interp_noperspective)
+GPU_SHADER_INTERFACE_END()
 
 GPU_SHADER_CREATE_INFO(gpu_shader_3D_polyline)
 DEFINE_VALUE("SMOOTH_WIDTH", "1.0")
@@ -33,32 +31,17 @@ PUSH_CONSTANT(MAT4, ModelViewProjectionMatrix)
 PUSH_CONSTANT(VEC2, viewportSize)
 PUSH_CONSTANT(FLOAT, lineWidth)
 PUSH_CONSTANT(BOOL, lineSmooth)
-VERTEX_IN(0, VEC3, pos)
+STORAGE_BUF_FREQ(GPU_SSBO_POLYLINE_POS_BUF_SLOT, READ, float, pos[], GEOMETRY)
+PUSH_CONSTANT(IVEC2, gpu_attr_0)
+PUSH_CONSTANT(IVEC3, gpu_vert_stride_count_offset)
+PUSH_CONSTANT(INT, gpu_attr_0_len)
+PUSH_CONSTANT(BOOL, gpu_attr_0_fetch_int)
 VERTEX_OUT(gpu_shader_3D_polyline_iface)
-VERTEX_OUT(gpu_shader_3D_polyline_noperspective_iface)
-GEOMETRY_LAYOUT(PrimitiveIn::LINES, PrimitiveOut::TRIANGLE_STRIP, 4)
-GEOMETRY_OUT(gpu_shader_3D_polyline_iface)
-GEOMETRY_OUT(gpu_shader_3D_polyline_noperspective_iface)
 FRAGMENT_OUT(0, VEC4, fragColor)
 VERTEX_SOURCE("gpu_shader_3D_polyline_vert.glsl")
-GEOMETRY_SOURCE("gpu_shader_3D_polyline_geom.glsl")
-FRAGMENT_SOURCE("gpu_shader_3D_polyline_frag.glsl")
-ADDITIONAL_INFO(gpu_srgb_to_framebuffer_space)
-GPU_SHADER_CREATE_END()
-
-GPU_SHADER_CREATE_INFO(gpu_shader_3D_polyline_no_geom)
-DEFINE_VALUE("SMOOTH_WIDTH", "1.0")
-PUSH_CONSTANT(MAT4, ModelViewProjectionMatrix)
-PUSH_CONSTANT(VEC2, viewportSize)
-PUSH_CONSTANT(FLOAT, lineWidth)
-PUSH_CONSTANT(BOOL, lineSmooth)
-VERTEX_IN(0, VEC3, pos)
-VERTEX_OUT(gpu_shader_3D_polyline_iface)
-VERTEX_OUT(gpu_shader_3D_polyline_noperspective_iface)
-FRAGMENT_OUT(0, VEC4, fragColor)
-VERTEX_SOURCE("gpu_shader_3D_polyline_vert_no_geom.glsl")
 FRAGMENT_SOURCE("gpu_shader_3D_polyline_frag.glsl")
 ADDITIONAL_INFO(gpu_srgb_to_framebuffer_space)
+ADDITIONAL_INFO(gpu_index_buffer_load)
 GPU_SHADER_CREATE_END()
 
 GPU_SHADER_CREATE_INFO(gpu_shader_3D_polyline_uniform_color)
@@ -68,14 +51,6 @@ PUSH_CONSTANT(VEC4, color)
 ADDITIONAL_INFO(gpu_shader_3D_polyline)
 GPU_SHADER_CREATE_END()
 
-GPU_SHADER_CREATE_INFO(gpu_shader_3D_polyline_uniform_color_no_geom)
-METAL_BACKEND_ONLY()
-DO_STATIC_COMPILATION()
-DEFINE("UNIFORM")
-PUSH_CONSTANT(VEC4, color)
-ADDITIONAL_INFO(gpu_shader_3D_polyline_no_geom)
-GPU_SHADER_CREATE_END()
-
 GPU_SHADER_CREATE_INFO(gpu_shader_3D_polyline_uniform_color_clipped)
 DO_STATIC_COMPILATION()
 /* TODO(fclem): Put in a UBO to fit the 128byte requirement. */
@@ -85,42 +60,20 @@ DEFINE("CLIP")
 ADDITIONAL_INFO(gpu_shader_3D_polyline_uniform_color)
 GPU_SHADER_CREATE_END()
 
-GPU_SHADER_CREATE_INFO(gpu_shader_3D_polyline_uniform_color_clipped_no_geom)
-METAL_BACKEND_ONLY()
-DO_STATIC_COMPILATION()
-/* TODO(fclem): Put in an UBO to fit the 128byte requirement. */
-PUSH_CONSTANT(MAT4, ModelMatrix)
-PUSH_CONSTANT(VEC4, ClipPlane)
-DEFINE("CLIP")
-ADDITIONAL_INFO(gpu_shader_3D_polyline_uniform_color_no_geom)
-GPU_SHADER_CREATE_END()
-
 GPU_SHADER_CREATE_INFO(gpu_shader_3D_polyline_flat_color)
 DO_STATIC_COMPILATION()
 DEFINE("FLAT")
-VERTEX_IN(1, VEC4, color)
+STORAGE_BUF_FREQ(GPU_SSBO_POLYLINE_COL_BUF_SLOT, READ, float, color[], GEOMETRY)
+PUSH_CONSTANT(IVEC2, gpu_attr_1)
+PUSH_CONSTANT(INT, gpu_attr_1_len)
 ADDITIONAL_INFO(gpu_shader_3D_polyline)
 GPU_SHADER_CREATE_END()
 
-GPU_SHADER_CREATE_INFO(gpu_shader_3D_polyline_flat_color_no_geom)
-METAL_BACKEND_ONLY()
-DO_STATIC_COMPILATION()
-DEFINE("FLAT")
-VERTEX_IN(1, VEC4, color)
-ADDITIONAL_INFO(gpu_shader_3D_polyline_no_geom)
-GPU_SHADER_CREATE_END()
-
 GPU_SHADER_CREATE_INFO(gpu_shader_3D_polyline_smooth_color)
 DO_STATIC_COMPILATION()
 DEFINE("SMOOTH")
-VERTEX_IN(1, VEC4, color)
+STORAGE_BUF_FREQ(GPU_SSBO_POLYLINE_COL_BUF_SLOT, READ, float, color[], GEOMETRY)
+PUSH_CONSTANT(IVEC2, gpu_attr_1)
+PUSH_CONSTANT(INT, gpu_attr_1_len)
 ADDITIONAL_INFO(gpu_shader_3D_polyline)
 GPU_SHADER_CREATE_END()
-
-GPU_SHADER_CREATE_INFO(gpu_shader_3D_polyline_smooth_color_no_geom)
-METAL_BACKEND_ONLY()
-DO_STATIC_COMPILATION()
-DEFINE("SMOOTH")
-VERTEX_IN(1, VEC4, color)
-ADDITIONAL_INFO(gpu_shader_3D_polyline_no_geom)
-GPU_SHADER_CREATE_END()
diff --git a/source/blender/gpu/vulkan/vk_descriptor_set.cc b/source/blender/gpu/vulkan/vk_descriptor_set.cc
index 51834788937..ff0ddc83b1d 100644
--- a/source/blender/gpu/vulkan/vk_descriptor_set.cc
+++ b/source/blender/gpu/vulkan/vk_descriptor_set.cc
@@ -19,10 +19,11 @@ namespace blender::gpu {
 
 void VKDescriptorSetTracker::bind_buffer(VkDescriptorType vk_descriptor_type,
                                          VkBuffer vk_buffer,
+                                         VkDeviceSize buffer_offset,
                                          VkDeviceSize size_in_bytes,
                                          VKDescriptorSet::Location location)
 {
-  vk_descriptor_buffer_infos_.append({vk_buffer, 0, size_in_bytes});
+  vk_descriptor_buffer_infos_.append({vk_buffer, buffer_offset, size_in_bytes});
   vk_write_descriptor_sets_.append({VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
                                     nullptr,
                                     vk_descriptor_set,
@@ -218,13 +219,22 @@ void VKDescriptorSetTracker::bind_storage_buffer_resource(
       vk_device_size = storage_buffer->size_in_bytes();
       break;
     }
+    case BindSpaceStorageBuffers::Type::Buffer: {
+      VKBuffer *buffer = static_cast<VKBuffer *>(elem.resource);
+      vk_buffer = buffer->vk_handle();
+      vk_device_size = buffer->size_in_bytes();
+      break;
+    }
     case BindSpaceStorageBuffers::Type::Unused: {
       BLI_assert_unreachable();
     }
   }
 
-  bind_buffer(
-      VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, vk_buffer, vk_device_size, resource_binding.location);
+  bind_buffer(VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,
+              vk_buffer,
+              elem.offset,
+              vk_device_size - elem.offset,
+              resource_binding.location);
   access_info.buffers.append({vk_buffer, resource_binding.access_mask});
 }
 
@@ -237,6 +247,7 @@ void VKDescriptorSetTracker::bind_uniform_buffer_resource(
   uniform_buffer.ensure_updated();
   bind_buffer(VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER,
               uniform_buffer.vk_handle(),
+              0,
               uniform_buffer.size_in_bytes(),
               resource_binding.location);
   access_info.buffers.append({uniform_buffer.vk_handle(), resource_binding.access_mask});
@@ -254,6 +265,7 @@ void VKDescriptorSetTracker::bind_push_constants(VKPushConstants &push_constants
   const VKUniformBuffer &uniform_buffer = *push_constants.uniform_buffer_get().get();
   bind_buffer(VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER,
               uniform_buffer.vk_handle(),
+              0,
               uniform_buffer.size_in_bytes(),
               push_constants.layout_get().descriptor_set_location_get());
   access_info.buffers.append({uniform_buffer.vk_handle(), VK_ACCESS_UNIFORM_READ_BIT});
diff --git a/source/blender/gpu/vulkan/vk_descriptor_set.hh b/source/blender/gpu/vulkan/vk_descriptor_set.hh
index cef5b0ae0cb..fadb1f5985c 100644
--- a/source/blender/gpu/vulkan/vk_descriptor_set.hh
+++ b/source/blender/gpu/vulkan/vk_descriptor_set.hh
@@ -130,6 +130,7 @@ class VKDescriptorSetTracker {
   void bind_texel_buffer(VkBufferView vk_buffer_view, VKDescriptorSet::Location location);
   void bind_buffer(VkDescriptorType vk_descriptor_type,
                    VkBuffer vk_buffer,
+                   VkDeviceSize buffer_offset,
                    VkDeviceSize size_in_bytes,
                    VKDescriptorSet::Location location);
   void bind_image(VkDescriptorType vk_descriptor_type,
diff --git a/source/blender/gpu/vulkan/vk_immediate.cc b/source/blender/gpu/vulkan/vk_immediate.cc
index 99c2d69b895..7b78ce97e7e 100644
--- a/source/blender/gpu/vulkan/vk_immediate.cc
+++ b/source/blender/gpu/vulkan/vk_immediate.cc
@@ -75,19 +75,38 @@ void VKImmediate::end()
 
   VKContext &context = *VKContext::get();
   BLI_assert(context.shader == unwrap(shader));
-  render_graph::VKResourceAccessInfo &resource_access_info = context.reset_and_get_access_info();
-  vertex_attributes_.update_bindings(*this);
-  context.active_framebuffer_get()->rendering_ensure(context);
+  Shader &shader = *unwrap(this->shader);
+  if (shader.is_polyline) {
+    VKBuffer *buffer = active_buffers_.last().get();
+    VKStateManager &state_manager = context.state_manager_get();
+    state_manager.storage_buffer_bind(BindSpaceStorageBuffers::Type::Buffer,
+                                      buffer,
+                                      GPU_SSBO_POLYLINE_POS_BUF_SLOT,
+                                      buffer_offset_);
+    state_manager.storage_buffer_bind(BindSpaceStorageBuffers::Type::Buffer,
+                                      buffer,
+                                      GPU_SSBO_POLYLINE_COL_BUF_SLOT,
+                                      buffer_offset_);
+    /* Not used. Satisfy the binding. */
+    state_manager.storage_buffer_bind(
+        BindSpaceStorageBuffers::Type::Buffer, buffer, GPU_SSBO_INDEX_BUF_SLOT, buffer_offset_);
+    this->polyline_draw_workaround(0);
+  }
+  else {
+    render_graph::VKResourceAccessInfo &resource_access_info = context.reset_and_get_access_info();
+    vertex_attributes_.update_bindings(*this);
+    context.active_framebuffer_get()->rendering_ensure(context);
 
-  render_graph::VKDrawNode::CreateInfo draw(resource_access_info);
-  draw.node_data.vertex_count = vertex_idx;
-  draw.node_data.instance_count = 1;
-  draw.node_data.first_vertex = 0;
-  draw.node_data.first_instance = 0;
-  vertex_attributes_.bind(draw.node_data.vertex_buffers);
-  context.update_pipeline_data(prim_type, vertex_attributes_, draw.node_data.pipeline_data);
+    render_graph::VKDrawNode::CreateInfo draw(resource_access_info);
+    draw.node_data.vertex_count = vertex_idx;
+    draw.node_data.instance_count = 1;
+    draw.node_data.first_vertex = 0;
+    draw.node_data.first_instance = 0;
+    vertex_attributes_.bind(draw.node_data.vertex_buffers);
+    context.update_pipeline_data(prim_type, vertex_attributes_, draw.node_data.pipeline_data);
 
-  context.render_graph.add_node(draw);
+    context.render_graph.add_node(draw);
+  }
 
   buffer_offset_ += current_subbuffer_len_;
   current_subbuffer_len_ = 0;
@@ -133,7 +152,8 @@ VKBuffer &VKImmediate::ensure_space(size_t bytes_needed)
   VKBuffer &result = *active_buffers_.last();
   result.create(new_buffer_size(bytes_needed),
                 GPU_USAGE_DYNAMIC,
-                VK_BUFFER_USAGE_VERTEX_BUFFER_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT);
+                VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | VK_BUFFER_USAGE_VERTEX_BUFFER_BIT |
+                    VK_BUFFER_USAGE_TRANSFER_DST_BIT);
   debug::object_label(result.vk_handle(), "Immediate");
 
   return result;
diff --git a/source/blender/gpu/vulkan/vk_state_manager.cc b/source/blender/gpu/vulkan/vk_state_manager.cc
index e3fe514d2bf..ef37eef0dce 100644
--- a/source/blender/gpu/vulkan/vk_state_manager.cc
+++ b/source/blender/gpu/vulkan/vk_state_manager.cc
@@ -135,9 +135,10 @@ void VKStateManager::texel_buffer_unbind(VKVertexBuffer &vertex_buffer)
 
 void VKStateManager::storage_buffer_bind(BindSpaceStorageBuffers::Type resource_type,
                                          void *resource,
-                                         int binding)
+                                         int binding,
+                                         VkDeviceSize offset)
 {
-  storage_buffers_.bind(resource_type, resource, binding);
+  storage_buffers_.bind(resource_type, resource, binding, offset);
   is_dirty = true;
 }
 
diff --git a/source/blender/gpu/vulkan/vk_state_manager.hh b/source/blender/gpu/vulkan/vk_state_manager.hh
index 46ab38e4476..8b64cbbc243 100644
--- a/source/blender/gpu/vulkan/vk_state_manager.hh
+++ b/source/blender/gpu/vulkan/vk_state_manager.hh
@@ -115,20 +115,23 @@ class BindSpaceStorageBuffers {
     VertexBuffer,
     IndexBuffer,
     StorageBuffer,
+    Buffer,
   };
   struct Elem {
     Type resource_type;
     void *resource;
+    VkDeviceSize offset;
   };
   Vector<Elem> bound_resources;
 
-  void bind(Type resource_type, void *resource, int binding)
+  void bind(Type resource_type, void *resource, int binding, VkDeviceSize offset)
   {
     if (bound_resources.size() <= binding) {
       bound_resources.resize(binding + 1);
     }
     bound_resources[binding].resource_type = resource_type;
     bound_resources[binding].resource = resource;
+    bound_resources[binding].offset = offset;
   }
 
   const Elem &get(int binding) const
@@ -142,6 +145,7 @@ class BindSpaceStorageBuffers {
       if (bound_resources[index].resource == resource) {
         bound_resources[index].resource = nullptr;
         bound_resources[index].resource_type = Type::Unused;
+        bound_resources[index].offset = 0u;
       }
     }
   }
@@ -234,7 +238,14 @@ class VKStateManager : public StateManager {
 
   void storage_buffer_bind(BindSpaceStorageBuffers::Type resource_type,
                            void *resource,
-                           int binding);
+                           int binding)
+  {
+    storage_buffer_bind(resource_type, resource, binding, 0u);
+  }
+  void storage_buffer_bind(BindSpaceStorageBuffers::Type resource_type,
+                           void *resource,
+                           int binding,
+                           VkDeviceSize offset);
   void storage_buffer_unbind(void *resource);
   void storage_buffer_unbind_all();