diff --git a/src/panfrost/vulkan/panvk_cmd_buffer.h b/src/panfrost/vulkan/panvk_cmd_buffer.h
index 9b622cb97d706..e3774199410da 100644
--- a/src/panfrost/vulkan/panvk_cmd_buffer.h
+++ b/src/panfrost/vulkan/panvk_cmd_buffer.h
@@ -92,7 +92,11 @@ struct panvk_descriptor_state {
    uint32_t dirty;
    const struct panvk_descriptor_set *sets[MAX_SETS];
    struct panvk_push_descriptor_set *push_sets[MAX_SETS];
-   struct panvk_sysvals sysvals;
+   union {
+      struct panvk_graphics_sysvals gfx;
+      struct panvk_compute_sysvals compute;
+   } sysvals;
+
    struct {
       struct mali_uniform_buffer_packed ubos[MAX_DYNAMIC_UNIFORM_BUFFERS];
       struct panvk_ssbo_addr ssbos[MAX_DYNAMIC_STORAGE_BUFFERS];
diff --git a/src/panfrost/vulkan/panvk_shader.h b/src/panfrost/vulkan/panvk_shader.h
index 3f2c6453773ed..f7aa8ff164218 100644
--- a/src/panfrost/vulkan/panvk_shader.h
+++ b/src/panfrost/vulkan/panvk_shader.h
@@ -24,30 +24,31 @@ struct nir_shader;
 struct pan_blend_state;
 struct panvk_device;
 
-union panvk_sysval_vec4 {
-   float f32[4];
-   uint32_t u32[4];
+struct panvk_graphics_sysvals {
+   struct {
+      struct {
+         float x, y, z;
+      } scale, offset;
+   } viewport;
+
+   struct {
+      float constants[4];
+   } blend;
+
+   struct {
+      uint32_t first_vertex;
+      uint32_t base_vertex;
+      uint32_t base_instance;
+   } vs;
 };
 
-struct panvk_sysvals {
-   union {
-      struct {
-         /* Only for graphics */
-         union panvk_sysval_vec4 viewport_scale;
-         union panvk_sysval_vec4 viewport_offset;
-         union panvk_sysval_vec4 blend_constants;
-
-         uint32_t first_vertex;
-         uint32_t base_vertex;
-         uint32_t base_instance;
-      };
-
-      struct {
-         /* Only for compute */
-         union panvk_sysval_vec4 num_work_groups;
-         union panvk_sysval_vec4 local_group_size;
-      };
-   };
+struct panvk_compute_sysvals {
+   struct {
+      uint32_t x, y, z;
+   } num_work_groups;
+   struct {
+      uint32_t x, y, z;
+   } local_group_size;
 };
 
 struct panvk_shader {
diff --git a/src/panfrost/vulkan/panvk_vX_cmd_buffer.c b/src/panfrost/vulkan/panvk_vX_cmd_buffer.c
index f26b601505fb4..9e77a70040e80 100644
--- a/src/panfrost/vulkan/panvk_vX_cmd_buffer.c
+++ b/src/panfrost/vulkan/panvk_vX_cmd_buffer.c
@@ -256,71 +256,57 @@ panvk_per_arch(cmd_alloc_tls_desc)(struct panvk_cmd_buffer *cmdbuf, bool gfx)
    }
 }
 
-/*
- * Upload the viewport scale. Defined as (px/2, py/2, pz) at the start of
- * section 24.5 ("Controlling the Viewport") of the Vulkan spec. At the end of
- * the section, the spec defines:
- *
- * px = width
- * py = height
- * pz = maxDepth - minDepth
- */
-static void
-panvk_sysval_upload_viewport_scale(const VkViewport *viewport,
-                                   union panvk_sysval_vec4 *data)
-{
-   data->f32[0] = 0.5f * viewport->width;
-   data->f32[1] = 0.5f * viewport->height;
-   data->f32[2] = (viewport->maxDepth - viewport->minDepth);
-}
-
-/*
- * Upload the viewport offset. Defined as (ox, oy, oz) at the start of section
- * 24.5 ("Controlling the Viewport") of the Vulkan spec. At the end of the
- * section, the spec defines:
- *
- * ox = x + width/2
- * oy = y + height/2
- * oz = minDepth
- */
-static void
-panvk_sysval_upload_viewport_offset(const VkViewport *viewport,
-                                    union panvk_sysval_vec4 *data)
-{
-   data->f32[0] = (0.5f * viewport->width) + viewport->x;
-   data->f32[1] = (0.5f * viewport->height) + viewport->y;
-   data->f32[2] = viewport->minDepth;
-}
-
 static void
 panvk_cmd_prepare_draw_sysvals(
    struct panvk_cmd_buffer *cmdbuf,
    struct panvk_cmd_bind_point_state *bind_point_state,
    struct panvk_draw_info *draw)
 {
-   struct panvk_sysvals *sysvals = &bind_point_state->desc_state.sysvals;
+   struct panvk_graphics_sysvals *sysvals =
+      &bind_point_state->desc_state.sysvals.gfx;
 
    unsigned base_vertex = draw->index_size ? draw->vertex_offset : 0;
-   if (sysvals->first_vertex != draw->offset_start ||
-       sysvals->base_vertex != base_vertex ||
-       sysvals->base_instance != draw->first_instance) {
-      sysvals->first_vertex = draw->offset_start;
-      sysvals->base_vertex = base_vertex;
-      sysvals->base_instance = draw->first_instance;
+   if (sysvals->vs.first_vertex != draw->offset_start ||
+       sysvals->vs.base_vertex != base_vertex ||
+       sysvals->vs.base_instance != draw->first_instance) {
+      sysvals->vs.first_vertex = draw->offset_start;
+      sysvals->vs.base_vertex = base_vertex;
+      sysvals->vs.base_instance = draw->first_instance;
       bind_point_state->desc_state.push_uniforms = 0;
    }
 
    if (cmdbuf->state.dirty & PANVK_DYNAMIC_BLEND_CONSTANTS) {
-      memcpy(&sysvals->blend_constants, cmdbuf->state.blend.constants,
+      memcpy(&sysvals->blend.constants, cmdbuf->state.blend.constants,
              sizeof(cmdbuf->state.blend.constants));
       bind_point_state->desc_state.push_uniforms = 0;
    }
 
    if (cmdbuf->state.dirty & PANVK_DYNAMIC_VIEWPORT) {
-      panvk_sysval_upload_viewport_scale(&cmdbuf->state.viewport,
-                                         &sysvals->viewport_scale);
-      panvk_sysval_upload_viewport_offset(&cmdbuf->state.viewport,
-                                          &sysvals->viewport_offset);
+      VkViewport *viewport = &cmdbuf->state.viewport;
+
+      /* Upload the viewport scale. Defined as (px/2, py/2, pz) at the start of
+       * section 24.5 ("Controlling the Viewport") of the Vulkan spec. At the
+       * end of the section, the spec defines:
+       *
+       * px = width
+       * py = height
+       * pz = maxDepth - minDepth
+       */
+      sysvals->viewport.scale.x = 0.5f * viewport->width;
+      sysvals->viewport.scale.y = 0.5f * viewport->height;
+      sysvals->viewport.scale.z = (viewport->maxDepth - viewport->minDepth);
+
+      /* Upload the viewport offset. Defined as (ox, oy, oz) at the start of
+       * section 24.5 ("Controlling the Viewport") of the Vulkan spec. At the
+       * end of the section, the spec defines:
+       *
+       * ox = x + width/2
+       * oy = y + height/2
+       * oz = minDepth
+       */
+      sysvals->viewport.offset.x = (0.5f * viewport->width) + viewport->x;
+      sysvals->viewport.offset.y = (0.5f * viewport->height) + viewport->y;
+      sysvals->viewport.offset.z = viewport->minDepth;
       bind_point_state->desc_state.push_uniforms = 0;
    }
 }
@@ -1735,13 +1721,13 @@ panvk_per_arch(CmdDispatch)(VkCommandBuffer commandBuffer, uint32_t x,
    struct panfrost_ptr job =
       pan_pool_alloc_desc(&cmdbuf->desc_pool.base, COMPUTE_JOB);
 
-   struct panvk_sysvals *sysvals = &desc_state->sysvals;
-   sysvals->num_work_groups.u32[0] = x;
-   sysvals->num_work_groups.u32[1] = y;
-   sysvals->num_work_groups.u32[2] = z;
-   sysvals->local_group_size.u32[0] = pipeline->cs.local_size.x;
-   sysvals->local_group_size.u32[1] = pipeline->cs.local_size.y;
-   sysvals->local_group_size.u32[2] = pipeline->cs.local_size.z;
+   struct panvk_compute_sysvals *sysvals = &desc_state->sysvals.compute;
+   sysvals->num_work_groups.x = x;
+   sysvals->num_work_groups.y = y;
+   sysvals->num_work_groups.z = z;
+   sysvals->local_group_size.x = pipeline->cs.local_size.x;
+   sysvals->local_group_size.y = pipeline->cs.local_size.y;
+   sysvals->local_group_size.z = pipeline->cs.local_size.z;
    desc_state->push_uniforms = 0;
 
    panvk_per_arch(cmd_alloc_tls_desc)(cmdbuf, false);
diff --git a/src/panfrost/vulkan/panvk_vX_shader.c b/src/panfrost/vulkan/panvk_vX_shader.c
index 5c103564ebd6c..ebfd714bcf9ef 100644
--- a/src/panfrost/vulkan/panvk_vX_shader.c
+++ b/src/panfrost/vulkan/panvk_vX_shader.c
@@ -79,28 +79,35 @@ panvk_lower_sysvals(nir_builder *b, nir_instr *instr, void *data)
    nir_def *val = NULL;
    b->cursor = nir_before_instr(instr);
 
-#define SYSVAL(name) offsetof(struct panvk_sysvals, name)
+#define SYSVAL(ptype, name) offsetof(struct panvk_ ## ptype ## _sysvals, name)
    switch (intr->intrinsic) {
    case nir_intrinsic_load_num_workgroups:
-      val = load_sysval_from_push_const(b, intr, SYSVAL(num_work_groups));
+      val =
+         load_sysval_from_push_const(b, intr, SYSVAL(compute, num_work_groups));
       break;
    case nir_intrinsic_load_workgroup_size:
-      val = load_sysval_from_push_const(b, intr, SYSVAL(local_group_size));
+      val = load_sysval_from_push_const(b, intr,
+                                        SYSVAL(compute, local_group_size));
       break;
    case nir_intrinsic_load_viewport_scale:
-      val = load_sysval_from_push_const(b, intr, SYSVAL(viewport_scale));
+      val =
+         load_sysval_from_push_const(b, intr, SYSVAL(graphics, viewport.scale));
       break;
    case nir_intrinsic_load_viewport_offset:
-      val = load_sysval_from_push_const(b, intr, SYSVAL(viewport_offset));
+      val = load_sysval_from_push_const(b, intr,
+                                        SYSVAL(graphics, viewport.offset));
       break;
    case nir_intrinsic_load_first_vertex:
-      val = load_sysval_from_push_const(b, intr, SYSVAL(first_vertex));
+      val = load_sysval_from_push_const(b, intr,
+                                        SYSVAL(graphics, vs.first_vertex));
       break;
    case nir_intrinsic_load_base_vertex:
-      val = load_sysval_from_push_const(b, intr, SYSVAL(base_vertex));
+      val =
+         load_sysval_from_push_const(b, intr, SYSVAL(graphics, vs.base_vertex));
       break;
    case nir_intrinsic_load_base_instance:
-      val = load_sysval_from_push_const(b, intr, SYSVAL(base_instance));
+      val = load_sysval_from_push_const(b, intr,
+                                        SYSVAL(graphics, vs.base_instance));
       break;
    case nir_intrinsic_load_blend_const_color_rgba:
       if (opts->static_blend_constants) {
@@ -113,7 +120,8 @@ panvk_lower_sysvals(nir_builder *b, nir_instr *instr, void *data)
 
          val = nir_build_imm(b, 4, 32, constants);
       } else {
-         val = load_sysval_from_push_const(b, intr, SYSVAL(blend_constants));
+         val = load_sysval_from_push_const(b, intr,
+                                           SYSVAL(graphics, blend.constants));
       }
       break;