Merge commit 'a485cbf0d953' into llvmspirv_pulldown

Author: jsji

devops/dependencies-igc-dev.json

@@ -1,10 +1,10 @@
 {
   "linux": {
     "igc_dev": {
-      "github_tag": "igc-dev-81b4de2",
-      "version": "81b4de2",
-      "updated_at": "2025-03-29T17:38:03Z",
-      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/2844658814/zip",
+      "github_tag": "igc-dev-b74b7ab",
+      "version": "b74b7ab",
+      "updated_at": "2025-04-02T18:41:33Z",
+      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/2869865189/zip",
       "root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
     }
   }

devops/scripts/benchmarks/utils/compute_runtime.py

@@ -97,7 +97,7 @@ def build_igc(self, repo, commit):
             options.workdir,
             "vc-intrinsics",
             "https://github.com/intel/vc-intrinsics",
-            "facb2076a2ce6cd6527c1e16570ba0fbaa2f1dba",
+            "b980474c99859f7e4eb157828c5e80202b062177",
         )
         self.llvm_project = git_clone(
             options.workdir,

libdevice/nativecpu_utils.cpp

@@ -318,9 +318,10 @@ DefShuffleINTEL_All(_Float16, f16, _Float16);
 
 Define2ArgForward(uint64_t, __spirv_ocl_u_min, std::min);
 
+#define GET_PROPS __attribute__((pure))
 #define GEN_u32(bname, muxname)                                                \
-  DEVICE_EXTERN_C uint32_t muxname();                                          \
-  DEVICE_EXTERNAL uint32_t bname() { return muxname(); }                       \
+  DEVICE_EXTERN_C GET_PROPS uint32_t muxname();                                \
+  DEVICE_EXTERNAL GET_PROPS uint32_t bname() { return muxname(); }             \
   static_assert(true)
 // subgroup
 GEN_u32(__spirv_SubgroupLocalInvocationId, __mux_get_sub_group_local_id);
@@ -331,8 +332,8 @@ GEN_u32(__spirv_SubgroupSize, __mux_get_sub_group_size);
 
 // I64_I32
 #define GEN_p(bname, muxname, arg)                                             \
-  DEVICE_EXTERN_C uint64_t muxname(uint32_t);                                  \
-  DEVICE_EXTERNAL uint64_t bname() { return muxname(arg); }                    \
+  DEVICE_EXTERN_C GET_PROPS uint64_t muxname(uint32_t);                        \
+  DEVICE_EXTERNAL GET_PROPS uint64_t bname() { return muxname(arg); }          \
   static_assert(true)
 
 #define GEN_xyz(bname, ncpu_name)                                              \
@@ -365,8 +366,8 @@ DefStateSetWithType(set_sub_group_id, SubGroup_id, uint32_t);
 DefStateSetWithType(set_max_sub_group_size, SubGroup_size, uint32_t);
 
 #define DefineStateGetWithType(name, field, type)                              \
-  DEVICE_EXTERNAL_C type __dpcpp_nativecpu_##name(                             \
-      MakeGlobalType<__nativecpu_state> *s) {                                  \
+  DEVICE_EXTERNAL_C GET_PROPS type __dpcpp_nativecpu_##name(                   \
+      MakeGlobalType<const __nativecpu_state> *s) {                            \
     return s->field;                                                           \
   }                                                                            \
   static_assert(true)
@@ -381,8 +382,8 @@ DefineStateGet_U32(get_max_sub_group_size, SubGroup_size);
 DefineStateGet_U32(get_num_sub_groups, NumSubGroups);
 
 #define DefineStateGetWithType2(name, field, rtype, ptype)                     \
-  DEVICE_EXTERNAL_C rtype __dpcpp_nativecpu_##name(                            \
-      ptype dim, MakeGlobalType<__nativecpu_state> *s) {                       \
+  DEVICE_EXTERNAL_C GET_PROPS rtype __dpcpp_nativecpu_##name(                  \
+      ptype dim, MakeGlobalType<const __nativecpu_state> *s) {                 \
     return s->field[dim];                                                      \
   }                                                                            \
   static_assert(true)

llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp

@@ -301,29 +301,20 @@ PreservedAnalyses PrepareSYCLNativeCPUPass::run(Module &M,
 
   llvm::Constant *CurrentStatePointerTLS = nullptr;
 
-  // check if any of the kernels is called by some other function.
-  // This can happen e.g. with OCK, where wrapper functions are
-  // created around the original kernel.
-  bool KernelIsCalled = false;
-  for (auto &K : OldKernels) {
-    for (auto &U : K->uses()) {
-      if (isa<CallBase>(U.getUser())) {
-        KernelIsCalled = true;
-      }
-    }
-  }
+  // Contains the used builtins and kernels that need to be processed to
+  // receive a pointer to the state struct.
+  llvm::SmallVector<std::pair<llvm::Function *, StringRef>>
+      UsedBuiltinsAndKernels;
 
   // Then we iterate over all the supported builtins, find the used ones
-  llvm::SmallVector<std::pair<llvm::Function *, StringRef>> UsedBuiltins;
   for (const auto &Entry : BuiltinNamesMap) {
     auto *Glob = M.getFunction(Entry.first);
     if (!Glob)
       continue;
     if (CurrentStatePointerTLS == nullptr) {
       for (const auto &Use : Glob->uses()) {
         auto *I = cast<CallBase>(Use.getUser());
-        if (KernelIsCalled ||
-            IsNonKernelCalledByNativeCPUKernel(I->getFunction())) {
+        if (IsNonKernelCalledByNativeCPUKernel(I->getFunction())) {
           // only use the threadlocal if we have kernels calling builtins
           // indirectly, or if the kernel is called by some other func.
           CurrentStatePointerTLS = M.getOrInsertGlobal(
@@ -344,7 +335,7 @@ PreservedAnalyses PrepareSYCLNativeCPUPass::run(Module &M,
         }
       }
     }
-    UsedBuiltins.push_back({Glob, Entry.second});
+    UsedBuiltinsAndKernels.push_back({Glob, Entry.second});
   }
 
 #ifdef NATIVECPU_USE_OCK
@@ -407,6 +398,11 @@ PreservedAnalyses PrepareSYCLNativeCPUPass::run(Module &M,
     OldF->replaceAllUsesWith(NewF);
     OldF->eraseFromParent();
     NewKernels.push_back(NewF);
+    if (!CurrentStatePointerTLS && NewF->getNumUses() > 0)
+      // If a thread_local is not used we need to keep track of the called
+      // kernel so we can update its call sites with the pointer to the state
+      // struct like we do for the called builtins.
+      UsedBuiltinsAndKernels.push_back({NewF, ""});
     ModuleChanged = true;
   }
 
@@ -419,13 +415,25 @@ PreservedAnalyses PrepareSYCLNativeCPUPass::run(Module &M,
 
   // Then we iterate over all used builtins and
   // replace them with calls to our Native CPU functions.
-  for (const auto &Entry : UsedBuiltins) {
+  // For the used kernels we need to replace calls to them
+  // with calls receiving the state pointer argument.
+  for (const auto &Entry : UsedBuiltinsAndKernels) {
     SmallVector<std::pair<Instruction *, Instruction *>> ToRemove;
     SmallVector<Function *> ToRemove2;
     Function *const Glob = Entry.first;
     Function *ReplaceFunc = nullptr;
     for (const auto &Use : Glob->uses()) {
       auto I = cast<CallBase>(Use.getUser());
+      if (Entry.second == "") {
+        if (const Function *CF = I->getCalledFunction()) {
+          unsigned numParams = CF->getFunctionType()->getNumParams();
+          auto numArgs = I->arg_size();
+          if (numArgs == numParams)
+            continue;
+          assert(numArgs + 1 == numParams);
+        }
+        ReplaceFunc = Entry.first;
+      }
       Function *const C = I->getFunction();
       if (IsUnusedBuiltinOrPrivateDef(*C)) {
         ToRemove2.push_back(C);
@@ -464,6 +472,10 @@ PreservedAnalyses PrepareSYCLNativeCPUPass::run(Module &M,
     for (auto Temp : ToRemove2)
       Temp->eraseFromParent();
 
+    // Don't erase if it's not a builtin
+    if (Entry.second == "")
+      continue;
+
     // Finally, we erase the builtin from the module
     Glob->eraseFromParent();
   }

sycl/source/detail/queue_impl.hpp

@@ -870,10 +870,7 @@ class queue_impl {
   event finalizeHandlerPostProcess(
       HandlerType &Handler,
       const optional<SubmitPostProcessF> &PostProcessorFunc) {
-    auto HandlerImpl = detail::getSyclObjImpl(Handler);
-    const CGType Type = HandlerImpl->MCGType;
-
-    bool IsKernel = Type == CGType::Kernel;
+    bool IsKernel = Handler.getType() == CGType::Kernel;
     bool KernelUsesAssert = false;
 
     if (IsKernel)

sycl/test-e2e/Basic/image/image_max_size.cpp

@@ -1,5 +1,8 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// RUN: %{build} -o %t.out
+
+// %O0 added because of GSD-10960. Without it, IGC will fail with
+// an access violation error.
+// RUN: %{build} %O0 -o %t.out
 // RUN: %{run} %t.out
 
 // UNSUPPORTED: cuda
@@ -13,100 +16,88 @@
 #include <sycl/detail/core.hpp>
 using namespace sycl;
 
-void init(uint32_t *A, uint32_t *B, size_t NumI32Elts) {
-  for (int I = 0; I < NumI32Elts; I++) {
-    A[I] = I;
-    B[I] = 0;
-  }
-}
+template <int Dimensions> class CopyKernel;
 
-int check(uint32_t *B, size_t NumI32Elts) {
-  for (int I = 0; I < NumI32Elts; I++) {
-    if (B[I] != I) {
-      std::cout << "Failed" << std::endl;
-      std::cerr << "Error for the index: " << I << ", computed: " << B[I]
-                << std::endl;
-      return 1;
-    }
-  }
-  std::cout << "Passed" << std::endl;
-  return 0;
-}
+template <int Dimensions>
+bool testND(queue &Q, size_t XSize, size_t YSize, size_t ZSize = 1) {
 
-int test2D(queue &Q, size_t XSize, size_t YSize) {
-  std::cout << "Starting the test with size = {" << XSize << ", " << YSize
-            << "} ... ";
-  size_t NumI32Elts = XSize * YSize * 4;
-  uint32_t *A = (uint32_t *)malloc(NumI32Elts * sizeof(uint32_t));
-  uint32_t *B = (uint32_t *)malloc(NumI32Elts * sizeof(uint32_t));
-  init(A, B, NumI32Elts);
+  static_assert(Dimensions == 2 || Dimensions == 3,
+                "Only 2D and 3D images are supported.");
 
-  try {
-    image<2> ImgA(A, image_channel_order::rgba,
-                  image_channel_type::unsigned_int32, range<2>{XSize, YSize});
-    image<2> ImgB(B, image_channel_order::rgba,
-                  image_channel_type::unsigned_int32, range<2>{XSize, YSize});
+  if constexpr (Dimensions == 2)
+    std::cout << "Starting the test with size = {" << XSize << ", " << YSize
+              << "} ... ";
+  else
+    std::cout << "Starting the test with size = {" << XSize << ", " << YSize
+              << ", " << ZSize << "} ... ";
 
-    Q.submit([&](handler &CGH) {
-       auto AAcc = ImgA.get_access<uint4, access::mode::read>(CGH);
-       auto BAcc = ImgB.get_access<uint4, access::mode::write>(CGH);
-       CGH.parallel_for<class I2D>(range<2>{XSize, YSize}, [=](id<2> Id) {
-         sycl::int2 Coord(Id[0], Id[1]);
-         BAcc.write(Coord, AAcc.read(Coord));
-       });
-     }).wait();
-  } catch (exception const &e) {
-    std::cout << "Failed" << std::endl;
-    std::cerr << "SYCL Exception caught: " << e.what();
-    return 1;
-  }
+  const size_t NumI32Elts = XSize * YSize * ZSize * 4;
+  range<Dimensions> ImgRange;
+  if constexpr (Dimensions == 2)
+    ImgRange = range<Dimensions>{XSize, YSize};
+  else
+    ImgRange = range<Dimensions>{XSize, YSize, ZSize};
 
-  int NumErrors = check(B, NumI32Elts);
-  free(A);
-  free(B);
-  return NumErrors;
-}
+  // Allocate input buffer and initialize it with some values.
+  uint32_t *Input = (uint32_t *)malloc(NumI32Elts * sizeof(uint32_t));
+  for (int i = 0; i < NumI32Elts; i++)
+    Input[i] = i;
 
-int test3D(queue &Q, size_t XSize, size_t YSize, size_t ZSize) {
-  std::cout << "Starting the test with size = {" << XSize << ", " << YSize
-            << ", " << ZSize << "} ... ";
-  size_t NumI32Elts = XSize * YSize * ZSize * 4;
-  uint32_t *A = (uint32_t *)malloc(NumI32Elts * sizeof(uint32_t));
-  uint32_t *B = (uint32_t *)malloc(NumI32Elts * sizeof(uint32_t));
-  init(A, B, NumI32Elts);
+  // calloc to ensure that the output buffer is initialized to zero.
+  uint32_t *Output = (uint32_t *)calloc(NumI32Elts, sizeof(uint32_t));
 
+  // Create the image and submit the copy kernel.
   try {
-    image<3> ImgA(A, image_channel_order::rgba,
-                  image_channel_type::unsigned_int32,
-                  range<3>{XSize, YSize, ZSize});
-    image<3> ImgB(B, image_channel_order::rgba,
-                  image_channel_type::unsigned_int32,
-                  range<3>{XSize, YSize, ZSize});
+    image<Dimensions> ImgA(Input, image_channel_order::rgba,
+                           image_channel_type::unsigned_int32, ImgRange);
+    image<Dimensions> ImgB(Output, image_channel_order::rgba,
+                           image_channel_type::unsigned_int32, ImgRange);
 
     Q.submit([&](handler &CGH) {
-       auto AAcc = ImgA.get_access<uint4, access::mode::read>(CGH);
-       auto BAcc = ImgB.get_access<uint4, access::mode::write>(CGH);
-       CGH.parallel_for<class I3D>(range<3>{XSize, YSize, ZSize},
-                                   [=](id<3> Id) {
-                                     sycl::int4 Coord(Id[0], Id[1], Id[2], 0);
-                                     BAcc.write(Coord, AAcc.read(Coord));
-                                   });
+       auto AAcc = ImgA.template get_access<uint4, access::mode::read>(CGH);
+       auto BAcc = ImgB.template get_access<uint4, access::mode::write>(CGH);
+       CGH.parallel_for<CopyKernel<Dimensions>>(
+           ImgRange, [=](id<Dimensions> Id) {
+             // Use int2 for 2D and int4 for 3D images.
+             if constexpr (Dimensions == 3) {
+               sycl::int4 Coord(Id[0], Id[1], Id[2], 0);
+               BAcc.write(Coord, AAcc.read(Coord));
+             } else {
+               sycl::int2 Coord(Id[0], Id[1]);
+               BAcc.write(Coord, AAcc.read(Coord));
+             }
+           });
      }).wait();
   } catch (exception const &e) {
+
     std::cout << "Failed" << std::endl;
     std::cerr << "SYCL Exception caught: " << e.what();
+    free(Input);
+    free(Output);
     return 1;
   }
 
-  int NumErrors = check(B, NumI32Elts);
-  free(A);
-  free(B);
-  return NumErrors;
+  // Check the output buffer.
+  bool HasError = false;
+  for (int i = 0; i < NumI32Elts; i++) {
+    if (Output[i] != i) {
+      HasError = true;
+      break;
+    }
+  }
+
+  if (!HasError) {
+    std::cout << "Passed" << std::endl;
+  } else {
+    std::cout << "Failed" << std::endl;
+  }
+
+  free(Input);
+  free(Output);
+  return HasError;
 }
 
 int main() {
-  int NumErrors = 0;
-
   queue Q;
   device Dev = Q.get_device();
   std::cout << "Running on " << Dev.get_info<info::device::name>()
@@ -127,17 +118,18 @@ int main() {
 
   // Using max sizes in one image may require too much memory.
   // Check them one by one.
-  NumErrors += test2D(Q, MaxWidth2D, 2);
-  NumErrors += test2D(Q, 2, MaxHeight2D);
+  bool HasError = false;
+  HasError |= testND<2>(Q, MaxWidth2D, 2);
+  HasError |= testND<2>(Q, 2, MaxHeight2D);
 
-  NumErrors += test3D(Q, MaxWidth3D, 2, 3);
-  NumErrors += test3D(Q, 2, MaxHeight3D, 3);
-  NumErrors += test3D(Q, 2, 3, MaxDepth3D);
+  HasError |= testND<3>(Q, MaxWidth3D, 2, 3);
+  HasError |= testND<3>(Q, 2, MaxHeight3D, 3);
+  HasError |= testND<3>(Q, 2, 3, MaxDepth3D);
 
-  if (NumErrors)
-    std::cerr << "Test failed." << std::endl;
+  if (HasError)
+    std::cout << "Test failed." << std::endl;
   else
     std::cout << "Test passed." << std::endl;
 
-  return NumErrors;
+  return HasError ? 1 : 0;
 }

sycl/test-e2e/Graph/Profiling/event_profiling_info.cpp

@@ -6,6 +6,9 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
+// UNSUPPORTED: level_zero_v2_adapter
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/17760
+
 // This test checks the profiling of an event returned
 // from graph submission with event::get_profiling_info().
 // It first tests a graph made exclusively of memory operations,

sycl/test-e2e/Graph/Profiling/event_profiling_info_usm.cpp

@@ -6,6 +6,9 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
+// UNSUPPORTED: level_zero_v2_adapter
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/17760
+
 // This test checks the profiling of an event returned
 // from graph submission with event::get_profiling_info().
 // It first tests a graph made exclusively of memory operations,

sycl/test-e2e/Graph/Update/lit.local.cfg

@@ -1 +1,3 @@
 config.required_features += ['aspect-ext_oneapi_graph']
+# V2 does not have support for MCL yet
+config.unsupported_features += ['level_zero_v2_adapter']