Merge from 'sycl' to 'sycl-web' (8 commits)

  CONFLICT (content): Merge conflict in clang/lib/AST/MicrosoftMangle.cpp
  CONFLICT (content): Merge conflict in clang/lib/CodeGen/CodeGenFunction.cpp
  CONFLICT (content): Merge conflict in clang/lib/CodeGen/CodeGenModule.cpp
  CONFLICT (content): Merge conflict in clang/lib/CodeGen/Targets/AMDGPU.cpp

Author: vmaksimo

clang/lib/CodeGen/TargetInfo.cpp

@@ -117,7 +117,7 @@ unsigned TargetCodeGenInfo::getDeviceKernelCallingConv() const {
     // to multiple function arguments etc.
     return llvm::CallingConv::SPIR_KERNEL;
   }
-  if (getABIInfo().getContext().getLangOpts().SYCLIsNativeCPU) {
+  if (getABIInfo().getContext().getTargetInfo().getTriple().isNativeCPU()) {
     return llvm::CallingConv::SPIR_KERNEL;
   }
   llvm_unreachable("Unknown kernel calling convention");

clang/lib/CodeGen/Targets/AMDGPU.cpp

@@ -305,6 +305,8 @@ class AMDGPUTargetCodeGenInfo : public TargetCodeGenInfo {
   void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
                            CodeGen::CodeGenModule &M) const override;
   unsigned getDeviceKernelCallingConv() const override;
+  void
+  setOCLKernelStubCallingConvention(const FunctionType *&FT) const override;
 
   llvm::Constant *getNullPointer(const CodeGen::CodeGenModule &CGM,
       llvm::PointerType *T, QualType QT) const override;
@@ -435,6 +437,14 @@ unsigned AMDGPUTargetCodeGenInfo::getDeviceKernelCallingConv() const {
   return llvm::CallingConv::AMDGPU_KERNEL;
 }
 
+void AMDGPUTargetCodeGenInfo::setOCLKernelStubCallingConvention(
+    const FunctionType *&FT) const {
+  bool IsSYCL = getABIInfo().getContext().getLangOpts().isSYCL();
+  FT = getABIInfo().getContext().adjustFunctionType(
+      FT,
+      FT->getExtInfo().withCallingConv(!IsSYCL ? CC_C : CC_DeviceKernel));
+}
+
 // Currently LLVM assumes null pointers always have value 0,
 // which results in incorrectly transformed IR. Therefore, instead of
 // emitting null pointers in private and local address spaces, a null

clang/test/CodeGenSYCL/spirv-builtins-addr-space.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -fsycl-is-device -fdeclare-spirv-builtins -emit-llvm -o - -O0 | FileCheck %s
+// RUN: %clang_cc1 %s -triple spir64 -fsycl-is-device -fdeclare-spirv-builtins -emit-llvm -o - -O0 | FileCheck %s
 //
 // Check that SPIR-V builtins are declared with SYCL address spaces rather
 // than OpenCL address spaces when using them with SYCL. OpenCL address spaces
@@ -10,19 +10,19 @@
 
 #include "Inputs/sycl.hpp"
 
-// CHECK: __spirv_ocl_modf{{.*}}SYglobal
+// CHECK: __spirv_ocl_modf{{.*}}AS1f
 void modf_global(float a) {
   __attribute__((opencl_global)) float *ptr = nullptr;
   sycl::kernel_single_task<class fake_kernel>([=]() { __spirv_ocl_modf(a, ptr); });
 }
 
-// CHECK: __spirv_ocl_modf{{.*}}SYlocal
+// CHECK: __spirv_ocl_modf{{.*}}AS3f
 void modf_local(float a) {
   __attribute__((opencl_local)) float *ptr = nullptr;
   sycl::kernel_single_task<class fake_kernel>([=]() { __spirv_ocl_modf(a, ptr); });
 }
 
-// CHECK: __spirv_ocl_modf{{.*}}SYprivate
+// CHECK: __spirv_ocl_modf{{.*}}AS0f
 void modf_private(float a) {
   __attribute__((opencl_private)) float *ptr = nullptr;
   sycl::kernel_single_task<class fake_kernel>([=]() { __spirv_ocl_modf(a, ptr); });

clang/test/Frontend/sycl-propagate-aspect-warning.cpp

@@ -1,5 +1,4 @@
-// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -mconstructor-aliases -emit-codegen-only -verify %s
-
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -triple native_cpu -fsycl-is-device -mconstructor-aliases -emit-llvm %s -o /dev/null
 // Tests for warnings when propagated aspects do not match the aspects available
 // in a function, as specified through the 'sycl::device_has' attribute.
 

devops/dependencies-igc-dev.json

@@ -1,10 +1,10 @@
 {
   "linux": {
     "igc_dev": {
-      "github_tag": "igc-dev-cee4ac4",
-      "version": "cee4ac4",
-      "updated_at": "2025-07-17T03:17:43Z",
-      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/3551081814/zip",
+      "github_tag": "igc-dev-420b632",
+      "version": "420b632",
+      "updated_at": "2025-07-20T08:48:29Z",
+      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/3571853551/zip",
       "root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
     }
   }

devops/scripts/benchmarks/benches/base.py

@@ -199,7 +199,8 @@ def benchmarks(self) -> list[Benchmark]:
     def name(self) -> str:
         pass
 
-    def setup(self):
+    @abstractmethod
+    def setup(self) -> None:
         return
 
     def additional_metadata(self) -> dict[str, BenchmarkMetadata]:

devops/scripts/benchmarks/benches/benchdnn.py

@@ -58,7 +58,7 @@ def benchmarks(self) -> list:
                 )
         return benchmarks
 
-    def setup(self):
+    def setup(self) -> None:
         if options.sycl is None:
             return
 

devops/scripts/benchmarks/benches/compute.py

@@ -52,7 +52,7 @@ def git_url(self) -> str:
     def git_hash(self) -> str:
         return "83b9ae3ebb3563552409f3a317cdc1cf3d3ca6bd"
 
-    def setup(self):
+    def setup(self) -> None:
         if options.sycl is None:
             return
 

devops/scripts/benchmarks/benches/gromacs.py

@@ -50,7 +50,7 @@ def benchmarks(self) -> list[Benchmark]:
             # GromacsBenchmark(self, "0192", "rf", "eager"),
         ]
 
-    def setup(self):
+    def setup(self) -> None:
         self.gromacs_src = git_clone(
             self.directory,
             "gromacs-repo",

devops/scripts/benchmarks/benches/llamacpp.py

@@ -28,7 +28,7 @@ def git_url(self) -> str:
     def git_hash(self) -> str:
         return "916c83bfe7f8b08ada609c3b8e583cf5301e594b"
 
-    def setup(self):
+    def setup(self) -> None:
         if options.sycl is None:
             return