Quick work around for SPIRV runner (#4762)

chengjunlu · web-flow · commit a1d1057973cc · 2025-07-22T14:23:26.000+02:00
The kernel calling convention has been changed for supporting the global
scratch space.

This is a quick work around for SPIRV runner to support the global
scratch since there is no real global scratch space required by Triton
XPU for now.

Signed-off-by: Lu,Chengjun &lt;chengjun.lu@intel.com&gt;
diff --git a/utils/SPIRVRunner/SPIRVRunner.cpp b/utils/SPIRVRunner/SPIRVRunner.cpp
@@ -285,6 +285,10 @@ static void sycl_kernel_launch(sycl::queue &stream, sycl::kernel &kernel_ptr,
             "Type entry is missing in JSON argument_list\n");
       }
     }
+    if (narg != expected_num_params) {
+      // global scratch.
+      cgh.set_arg(narg++, nullptr);
+    }
     if (triton_args.shared_memory) {
       using share_mem_t = sycl::local_accessor<int8_t, 1>;
       share_mem_t local_buffer = share_mem_t(triton_args.shared_memory, cgh);

Original file line number	Diff line number	Diff line change
`@@ -285,6 +285,10 @@ static void sycl_kernel_launch(sycl::queue &stream, sycl::kernel &kernel_ptr,`
`285`	`285`	`"Type entry is missing in JSON argument_list\n");`
`286`	`286`	`}`
`287`	`287`	`}`
	`288`	`+ if (narg != expected_num_params) {`
	`289`	`+ // global scratch.`
	`290`	`+ cgh.set_arg(narg++, nullptr);`
	`291`	`+ }`
`288`	`292`	`if (triton_args.shared_memory) {`
`289`	`293`	`using share_mem_t = sycl::local_accessor<int8_t, 1>;`
`290`	`294`	`share_mem_t local_buffer = share_mem_t(triton_args.shared_memory, cgh);`