From e30389825862fcedfad2f47debaf24f1360bfc3c Mon Sep 17 00:00:00 2001
From: Mike Blumenkrantz <michael.blumenkrantz@gmail.com>
Date: Wed, 11 May 2022 16:04:14 -0400
Subject: [PATCH] zink: manually validate shaders in debug builds

VVL is great, but there's actually cases where it doesn't catch critical
spirv errors, so add in our own validation pass to make sure things are
okay

this is especially useful for running on nvidia, as their compiler will
either crash on or silently drop illegal instructions

Reviewed-by: Dave Airlie <airlied@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/16462>
---
 src/gallium/drivers/zink/zink_compiler.c | 75 ++++++++++++++++++++++++
 1 file changed, 75 insertions(+)

diff --git a/src/gallium/drivers/zink/zink_compiler.c b/src/gallium/drivers/zink/zink_compiler.c
index 599566ca35e..9bb7225bb38 100644
--- a/src/gallium/drivers/zink/zink_compiler.c
+++ b/src/gallium/drivers/zink/zink_compiler.c
@@ -38,6 +38,9 @@
 
 #include "util/u_memory.h"
 
+#include "compiler/spirv/nir_spirv.h"
+#include "vulkan/util/vk_util.h"
+
 bool
 zink_lower_cubemap_to_array(nir_shader *s, uint32_t nonseamless_cube_mask);
 
@@ -1291,6 +1294,78 @@ zink_shader_spirv_compile(struct zink_screen *screen, struct zink_shader *zs, st
    smci.codeSize = spirv->num_words * sizeof(uint32_t);
    smci.pCode = spirv->words;
 
+#ifndef NDEBUG
+   static const struct spirv_to_nir_options spirv_options = {
+      .environment = NIR_SPIRV_VULKAN,
+      .caps = {
+         .float64 = true,
+         .int16 = true,
+         .int64 = true,
+         .tessellation = true,
+         .float_controls = true,
+         .image_ms_array = true,
+         .image_read_without_format = true,
+         .image_write_without_format = true,
+         .storage_image_ms = true,
+         .geometry_streams = true,
+         .storage_8bit = true,
+         .storage_16bit = true,
+         .variable_pointers = true,
+         .stencil_export = true,
+         .post_depth_coverage = true,
+         .transform_feedback = true,
+         .device_group = true,
+         .draw_parameters = true,
+         .shader_viewport_index_layer = true,
+         .multiview = true,
+         .physical_storage_buffer_address = true,
+         .int64_atomics = true,
+         .subgroup_arithmetic = true,
+         .subgroup_basic = true,
+         .subgroup_ballot = true,
+         .subgroup_quad = true,
+         .subgroup_shuffle = true,
+         .subgroup_vote = true,
+         .vk_memory_model = true,
+         .vk_memory_model_device_scope = true,
+         .int8 = true,
+         .float16 = true,
+         .demote_to_helper_invocation = true,
+         .sparse_residency = true,
+         .min_lod = true,
+      },
+      .ubo_addr_format = nir_address_format_32bit_index_offset,
+      .ssbo_addr_format = nir_address_format_32bit_index_offset,
+      .phys_ssbo_addr_format = nir_address_format_64bit_global,
+      .push_const_addr_format = nir_address_format_logical,
+      .shared_addr_format = nir_address_format_32bit_offset,
+   };
+   uint32_t num_spec_entries = 0;
+   struct nir_spirv_specialization *spec_entries = NULL;
+   VkSpecializationInfo sinfo = {0};
+   VkSpecializationMapEntry me[3];
+   uint32_t size[3] = {1,1,1};
+   if (!zs->nir->info.workgroup_size[0]) {
+      sinfo.mapEntryCount = 3;
+      sinfo.pMapEntries = &me[0];
+      sinfo.dataSize = sizeof(uint32_t) * 3;
+      sinfo.pData = size;
+      uint32_t ids[] = {ZINK_WORKGROUP_SIZE_X, ZINK_WORKGROUP_SIZE_Y, ZINK_WORKGROUP_SIZE_Z};
+      for (int i = 0; i < 3; i++) {
+         me[i].size = sizeof(uint32_t);
+         me[i].constantID = ids[i];
+         me[i].offset = i * sizeof(uint32_t);
+      }
+      spec_entries = vk_spec_info_to_nir_spirv(&sinfo, &num_spec_entries);
+   }
+   nir_shader *nir = spirv_to_nir(spirv->words, spirv->num_words,
+                      spec_entries, num_spec_entries,
+                      zs->nir->info.stage, "main", &spirv_options, &screen->nir_options);
+   assert(nir);
+   ralloc_free(nir);
+   free(spec_entries);
+#endif
+
    VkResult ret = VKSCR(CreateShaderModule)(screen->dev, &smci, NULL, &mod);
    bool success = zink_screen_handle_vkresult(screen, ret);
    assert(success);