anv: wire up jay

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40835>

src/intel/vulkan/anv_physical_device.c

@@ -6,6 +6,7 @@
 #include "anv_api_version.h"
 #include "anv_measure.h"
 
+#include "dev/intel_debug.h"
 #include "i915/anv_device.h"
 #include "xe/anv_device.h"
 
@@ -133,7 +134,8 @@ static void
 get_device_extensions(const struct anv_physical_device *device,
                       struct vk_device_extension_table *ext)
 {
-   const bool rt_enabled = ANV_SUPPORT_RT && device->info.has_ray_tracing;
+   const bool rt_enabled = ANV_SUPPORT_RT && device->info.has_ray_tracing &&
+                           !intel_use_jay_any_stage(&device->info);
    const bool hw_video_encode_supported = device->info.verx10 < 125;
    const bool video_encode_enabled = hw_video_encode_supported &&
                                      (device->instance->debug & ANV_DEBUG_VIDEO_ENCODE);
@@ -2779,7 +2781,8 @@ anv_physical_device_try_create(struct vk_instance *vk_instance,
 
    device->has_cooperative_matrix =
       (device->info.has_systolic || debug_get_bool_option("INTEL_LOWER_DPAS", false)) &&
-      device->info.cooperative_matrix_configurations[0].scope != INTEL_CMAT_SCOPE_NONE;
+      device->info.cooperative_matrix_configurations[0].scope != INTEL_CMAT_SCOPE_NONE &&
+      !intel_use_jay_any_stage(&device->info);
 
    if (is_virtio) {
       struct util_sync_provider *sync = intel_virtio_sync_provider(fd);

src/intel/vulkan/anv_shader_compile.c

@@ -15,6 +15,7 @@
 #include "compiler/brw/brw_nir.h"
 #include "compiler/brw/brw_nir_rt.h"
 #include "compiler/intel_nir.h"
+#include "compiler/jay/jay.h"
 
 #include "git_sha1.h"
 
@@ -833,6 +834,7 @@ anv_shader_compile_vs(struct anv_device *device,
                       char **error_str)
 {
    const struct brw_compiler *compiler = device->physical->compiler;
+   const struct intel_device_info *devinfo = compiler->devinfo;
    nir_shader *nir = shader_data->info->nir;
 
    shader_data->num_stats = 1;
@@ -850,7 +852,17 @@ anv_shader_compile_vs(struct anv_device *device,
       .prog_data = &shader_data->prog_data.vs,
    };
 
-   shader_data->code = (void *)brw_compile_vs(compiler, &params);
+   if (intel_use_jay(devinfo, nir->info.stage)) {
+      struct jay_shader_bin *bin =
+         jay_compile(devinfo, mem_ctx, nir,
+                     (union brw_any_prog_data *) params.prog_data,
+                     (union brw_any_prog_key *) params.key);
+
+      shader_data->code = (void *) bin->kernel;
+   } else {
+      shader_data->code = (void *) brw_compile_vs(compiler, &params);
+   }
+
    *error_str = params.base.error_str;
 }
 
@@ -1040,6 +1052,7 @@ anv_shader_compile_fs(struct anv_device *device,
                       char **error_str)
 {
    const struct brw_compiler *compiler = device->physical->compiler;
+   const struct intel_device_info *devinfo = compiler->devinfo;
    nir_shader *nir = shader_data->info->nir;
 
    /* When using Primitive Replication for multiview, each view gets its own
@@ -1072,7 +1085,17 @@ anv_shader_compile_fs(struct anv_device *device,
       .max_polygons = UCHAR_MAX,
    };
 
-   shader_data->code = (void *)brw_compile_fs(compiler, &params);
+   if (intel_use_jay(devinfo, nir->info.stage)) {
+      struct jay_shader_bin *bin =
+         jay_compile(devinfo, mem_ctx, nir,
+                     (union brw_any_prog_data *) params.prog_data,
+                     (union brw_any_prog_key *) params.key);
+
+      shader_data->code = (void *) bin->kernel;
+   } else {
+      shader_data->code = (void *) brw_compile_fs(compiler, &params);
+   }
+
    *error_str = params.base.error_str;
 
    shader_data->num_stats = (uint32_t)!!shader_data->prog_data.fs.dispatch_multi +
@@ -1101,6 +1124,7 @@ anv_shader_compile_cs(struct anv_device *device,
                       char **error_str)
 {
    const struct brw_compiler *compiler = device->physical->compiler;
+   const struct intel_device_info *devinfo = compiler->devinfo;
    nir_shader *nir = shader_data->info->nir;
 
    shader_data->num_stats = 1;
@@ -1118,7 +1142,21 @@ anv_shader_compile_cs(struct anv_device *device,
       .prog_data = &shader_data->prog_data.cs,
    };
 
-   shader_data->code = (void *)brw_compile_cs(compiler, &params);
+   if (intel_use_jay(devinfo, nir->info.stage)) {
+      struct jay_shader_bin *bin = jay_compile(devinfo, mem_ctx, nir,
+                             (union brw_any_prog_data*)params.prog_data,
+                             (union brw_any_prog_key*)params.key);
+
+       shader_data->code = (void*)bin->kernel;
+       shader_data->stats[0] = bin->stats;
+
+       params.prog_data->local_size[0] = nir->info.workgroup_size[0];
+       params.prog_data->local_size[1] = nir->info.workgroup_size[1];
+       params.prog_data->local_size[2] = nir->info.workgroup_size[2];
+   } else {
+       shader_data->code = (void*)brw_compile_cs(compiler, &params);
+   }
+
    *error_str = params.base.error_str;
 }
 

src/intel/vulkan/meson.build

@@ -277,7 +277,8 @@ libvulkan_intel = shared_library(
     idep_nir, idep_genxml, idep_vulkan_util, idep_vulkan_wsi,
     idep_vulkan_runtime, idep_mesautil, idep_xmlconfig,
     idep_intel_driver_ds, idep_intel_dev, idep_intel_blorp,
-    idep_intel_compiler_brw, idep_libvulkan_intel_decoder
+    idep_intel_compiler_brw, idep_libvulkan_intel_decoder,
+    idep_intel_compiler_jay,
   ],
   c_args : anv_flags,
   gnu_symbol_visibility : 'hidden',
@@ -315,7 +316,7 @@ if with_tests
       dep_thread, dep_dl, dep_m, anv_deps,
       idep_nir, idep_vulkan_util, idep_vulkan_wsi, idep_vulkan_runtime,
       idep_mesautil, idep_intel_dev, idep_intel_blorp,
-      idep_intel_compiler_brw, idep_libvulkan_intel_decoder,
+      idep_intel_compiler_brw, idep_intel_compiler_jay, idep_libvulkan_intel_decoder,
     ],
     c_args : anv_flags,
     gnu_symbol_visibility : 'hidden',