Merge remote-tracking branch 'origin/sycl' into host_prof

Author: againull

.github/workflows/sycl-linux-precommit.yml

@@ -111,7 +111,7 @@ jobs:
           - name: E2E tests on Intel Ponte Vecchio GPU
             runner: '["Linux", "pvc"]'
             image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
-            target_devices: level_zero:gpu;opencl:gpu
+            target_devices: level_zero:gpu;opencl:gpu;level_zero_v2:gpu
           - name: Dev IGC on Intel Ponte Vecchio GPU
             runner: '["Linux", "pvc"]'
             image: ghcr.io/intel/llvm/ubuntu2404_intel_drivers:devigc
@@ -163,6 +163,7 @@ jobs:
         'false' }}
       # Run only if the PR does not have the 'ci-no-devigc' label.
       skip_run: ${{matrix.use_igc_dev && contains(github.event.pull_request.labels.*.name, 'ci-no-devigc') || 'false'}}
+      env: ${{ contains(needs.detect_changes.outputs.filters, 'esimd') && '{}' || '{"LIT_FILTER_OUT":"ESIMD/"}' }}
 
   test-perf:
     needs: [build, detect_changes]

clang/lib/Driver/ToolChains/Cuda.cpp

@@ -968,6 +968,11 @@ void CudaToolChain::addClangTargetOptions(
                       "--nvptx-prec-sqrtf32=0"});
 
     CC1Args.append({"-mllvm", "-enable-memcpyopt-without-libcalls"});
+
+    if (DriverArgs.hasFlag(options::OPT_fsycl_id_queries_fit_in_int,
+                           options::OPT_fno_sycl_id_queries_fit_in_int, false))
+      CC1Args.append(
+          {"-mllvm", "-nvvm-reflect-add=__CUDA_ID_QUERIES_FIT_IN_INT=1"});
   } else {
     CC1Args.append({"-fcuda-is-device", "-mllvm",
                     "-enable-memcpyopt-without-libcalls",

clang/lib/Sema/SemaSYCL.cpp

@@ -1197,6 +1197,8 @@ bool SemaSYCL::isFreeFunction(const FunctionDecl *FD) {
                    NameValuePair.first == "sycl-single-task-kernel";
           });
       IsFreeFunctionAttr = it != NameValuePairs.end();
+      if (IsFreeFunctionAttr)
+        break;
     }
     if (Redecl->isFirstDecl()) {
       if (IsFreeFunctionAttr)
@@ -6615,9 +6617,18 @@ class FreeFunctionPrinter {
         // function
         NSInserted = true;
       }
+      if (FD->isFunctionTemplateSpecialization() &&
+          FD->isThisDeclarationADefinition())
+        O << "template <> ";
       O << TemplateParameters;
       O << FD->getReturnType().getAsString() << " ";
-      O << FD->getNameAsString() << "(" << Args << ");";
+      FD->printName(O, Policy);
+      if (FD->isFunctionTemplateSpecialization() &&
+          FD->isThisDeclarationADefinition())
+        O << getTemplateSpecializationArgString(
+            FD->getTemplateSpecializationArgs());
+
+      O << "(" << Args << ");";
       if (NSInserted) {
         O << "\n";
         PrintNSClosingBraces(O, FD);
@@ -6639,35 +6650,49 @@ class FreeFunctionPrinter {
     if (NSInserted)
       PrintNamespaces(O, FD, /*isPrintNamesOnly=*/true);
     O << FD->getIdentifier()->getName().data();
-    if (FD->getPrimaryTemplate()) {
-      std::string Buffer;
-      llvm::raw_string_ostream StringStream(Buffer);
-      const TemplateArgumentList *TAL = FD->getTemplateSpecializationArgs();
-      ArrayRef<TemplateArgument> A = TAL->asArray();
-      bool FirstParam = true;
-      for (const auto &X : A) {
-        if (FirstParam)
-          FirstParam = false;
-        else if (X.getKind() == TemplateArgument::Pack) {
-          for (const auto &PackArg : X.pack_elements()) {
-            StringStream << ", ";
-            PackArg.print(Policy, StringStream, true);
-          }
-          continue;
-        } else {
+    if (FD->getPrimaryTemplate())
+      O << getTemplateSpecializationArgString(
+          FD->getTemplateSpecializationArgs());
+  }
+
+private:
+  /// Helper method to get string with template types
+  /// \param TAL The template argument list.
+  /// \returns string Example:
+  /// \code
+  ///  template <typename T1, typename T2>
+  ///  void foo(T1 a, T2 b);
+  /// \endcode
+  /// returns string "<T1, T2>"
+  /// If TAL is nullptr, returns empty string.
+  std::string
+  getTemplateSpecializationArgString(const TemplateArgumentList *TAL) {
+    if (!TAL)
+      return "";
+    std::string Buffer;
+    llvm::raw_string_ostream StringStream(Buffer);
+    ArrayRef<TemplateArgument> A = TAL->asArray();
+    bool FirstParam = true;
+    for (const auto &X : A) {
+      if (FirstParam)
+        FirstParam = false;
+      else if (X.getKind() == TemplateArgument::Pack) {
+        for (const auto &PackArg : X.pack_elements()) {
           StringStream << ", ";
+          PackArg.print(Policy, StringStream, /*IncludeType*/ true);
         }
+        continue;
+      } else
+        StringStream << ", ";
 
-        X.print(Policy, StringStream, true);
-      }
-      StringStream.flush();
-      if (Buffer.front() != '<')
-        Buffer = "<" + Buffer + ">";
-      O << Buffer;
+      X.print(Policy, StringStream, /*IncludeType*/ true);
     }
+    StringStream.flush();
+    if (Buffer.front() != '<')
+      Buffer = "<" + Buffer + ">";
+    return Buffer;
   }
 
-private:
   /// Helper method to get arguments of templated function as a string
   /// \param Parameters Array of parameters of the function.
   /// \param Policy Printing policy.
@@ -7081,6 +7106,10 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
     FreeFunctionPrinter FFPrinter(O, Policy);
     if (FTD) {
       FFPrinter.printFreeFunctionDeclaration(FTD, S);
+      if (const auto kind = K.SyclKernel->getTemplateSpecializationKind();
+          K.SyclKernel->isFunctionTemplateSpecialization() &&
+          kind == TSK_ExplicitSpecialization)
+        FFPrinter.printFreeFunctionDeclaration(K.SyclKernel, ParmListWithNames);
     } else {
       FFPrinter.printFreeFunctionDeclaration(K.SyclKernel, ParmListWithNames);
     }

clang/test/CodeGenSYCL/free_function_int_header.cpp

@@ -254,6 +254,16 @@ __attribute__((sycl_device))
 void ff_20(sycl::accessor<int, 1, sycl::access::mode::read_write> acc) {
 }
 
+[[__sycl_detail__::add_ir_attributes_function("work_group_size", 16)]]
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
+void ff_21(AliasType start, AliasType *ptr) {
+}
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
+[[__sycl_detail__::add_ir_attributes_function("work_group_size", 16)]]
+void ff_22(AliasType start, AliasType *ptr) {
+}
+
 // CHECK:      const char* const kernel_names[] = {
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_2Piii
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_2Piiii
@@ -286,6 +296,8 @@ void ff_20(sycl::accessor<int, 1, sycl::access::mode::read_write> acc) {
 // CHECK-NEXT:   {{.*}}__sycl_kernel_free_functions5tests5ff_18ENS_3AggEPS1_
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_19N14free_functions16KArgWithPtrArrayILi50EEE
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_20N4sycl3_V18accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE
+// CHECK-NEXT:   {{.*}}__sycl_kernel_ff_217DerivedPS_
+// CHECK-NEXT:   {{.*}}__sycl_kernel_ff_227DerivedPS_
 
 // CHECK-NEXT:   ""
 // CHECK-NEXT: };
@@ -479,6 +491,7 @@ void ff_20(sycl::accessor<int, 1, sycl::access::mode::read_write> acc) {
 // CHECK: Definition of _Z18__sycl_kernel_ff_3IdEvPT_S0_S0_ as a free function kernel
 // CHECK: Forward declarations of kernel and its argument types:
 // CHECK: template <typename T> void ff_3(T * ptr, T start, T end);
+// CHECK: template <> void ff_3<double>(double * ptr, double start, double end);
 // CHECK-NEXT: static constexpr auto __sycl_shim5() {
 // CHECK-NEXT:   return (void (*)(double *, double, double))ff_3<double>;
 // CHECK-NEXT: }
@@ -980,6 +993,37 @@ void ff_20(sycl::accessor<int, 1, sycl::access::mode::read_write> acc) {
 // CHECK-NEXT: };
 // CHECK-NEXT: }
 
+
+// CHECK: void ff_21(Derived start, Derived * ptr);
+// CHECK-NEXT: static constexpr auto __sycl_shim30() {
+// CHECK-NEXT:  return (void (*)(struct Derived, struct Derived *))ff_21;
+// CHECK-NEXT: }
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: template <>
+// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim30()> {
+// CHECK-NEXT: static constexpr bool value = true;
+// CHECK-NEXT: };
+// CHECK-NEXT: template <>
+// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim30()> {
+// CHECK-NEXT: static constexpr bool value = true;
+// CHECK-NEXT: };
+// CHECK-NEXT: }
+
+// CHECK: void ff_22(Derived start, Derived * ptr);
+// CHECK-NEXT: static constexpr auto __sycl_shim31() {
+// CHECK-NEXT:  return (void (*)(struct Derived, struct Derived *))ff_22;
+// CHECK-NEXT: }
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: template <>
+// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim31()> {
+// CHECK-NEXT:  static constexpr bool value = true;
+// CHECK-NEXT: };
+// CHECK-NEXT: template <>
+// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim31()> {
+// CHECK-NEXT: static constexpr bool value = true;
+// CHECK-NEXT: };
+// CHECK-NEXT: }
+
 // CHECK: #include <sycl/kernel_bundle.hpp>
 
 // CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_2Piii
@@ -1196,3 +1240,17 @@ void ff_20(sycl::accessor<int, 1, sycl::access::mode::read_write> acc) {
 // CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{"_Z19__sycl_kernel_ff_20N4sycl3_V18accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE"});
 // CHECK-NEXT: }
 // CHECK-NEXT: }
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: template <>
+// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim30()>() {
+// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_Z19__sycl_kernel_ff_217DerivedPS_"});
+// CHECK-NEXT: }
+// CHECK-NEXT: }
+
+// CHECK: namespace sycl {
+// CHECK-NEXT: template <>
+// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim31()>() {
+// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_Z19__sycl_kernel_ff_227DerivedPS_"});
+// CHECK-NEXT: }
+// CHECK-NEXT: }

clang/test/Driver/sycl-nvptx-id-queries-fit-in-int.cpp

@@ -0,0 +1,16 @@
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang -### -nocudalib \
+// RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda %s 2>&1 \
+// RUN: | FileCheck --check-prefix=CHECK-DEFAULT %s
+
+// RUN: %clang -### -nocudalib \
+// RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda -fno-sycl-id-queries-fit-in-int %s 2>&1 \
+// RUN: | FileCheck --check-prefix=CHECK-DEFAULT %s
+
+// RUN: %clang -### -nocudalib \
+// RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda -fsycl-id-queries-fit-in-int %s 2>&1 \
+// RUN: | FileCheck --check-prefix=CHECK-INT %s
+
+// CHECK-INT: "-mllvm" "-nvvm-reflect-add=__CUDA_ID_QUERIES_FIT_IN_INT=1"
+// CHECK-DEFAULT-NOT: "-nvvm-reflect-add=__CUDA_ID_QUERIES_FIT_IN_INT=1"

libclc/libspirv/lib/ptx-nvidiacl/workitem/get_global_id.cl

@@ -8,17 +8,31 @@
 
 #include <libspirv/spirv.h>
 
+extern int __nvvm_reflect_ocl(constant char *);
+
 _CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalInvocationId_x() {
+  if (__nvvm_reflect_ocl("__CUDA_ID_QUERIES_FIT_IN_INT")) {
+    return (uint)__spirv_WorkgroupId_x() * (uint)__spirv_WorkgroupSize_x() +
+           (uint)__spirv_LocalInvocationId_x() + (uint)__spirv_GlobalOffset_x();
+  }
   return __spirv_WorkgroupId_x() * __spirv_WorkgroupSize_x() +
          __spirv_LocalInvocationId_x() + __spirv_GlobalOffset_x();
 }
 
 _CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalInvocationId_y() {
+  if (__nvvm_reflect_ocl("__CUDA_ID_QUERIES_FIT_IN_INT")) {
+    return (uint)__spirv_WorkgroupId_y() * (uint)__spirv_WorkgroupSize_y() +
+           (uint)__spirv_LocalInvocationId_y() + (uint)__spirv_GlobalOffset_y();
+  }
   return __spirv_WorkgroupId_y() * __spirv_WorkgroupSize_y() +
          __spirv_LocalInvocationId_y() + __spirv_GlobalOffset_y();
 }
 
 _CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalInvocationId_z() {
+  if (__nvvm_reflect_ocl("__CUDA_ID_QUERIES_FIT_IN_INT")) {
+    return (uint)__spirv_WorkgroupId_z() * (uint)__spirv_WorkgroupSize_z() +
+           (uint)__spirv_LocalInvocationId_z() + (uint)__spirv_GlobalOffset_z();
+  }
   return __spirv_WorkgroupId_z() * __spirv_WorkgroupSize_z() +
          __spirv_LocalInvocationId_z() + __spirv_GlobalOffset_z();
 }

libdevice/crt_wrapper.cpp

@@ -12,10 +12,8 @@
 
 #include <cstdint>
 
-#ifndef __NVPTX__
 #define RAND_NEXT_LEN 1024
 DeviceGlobal<uint64_t[RAND_NEXT_LEN]> RandNext;
-#endif
 
 #if defined(__SPIR__) || defined(__SPIRV__) || defined(__NVPTX__) ||           \
     defined(__AMDGCN__)
@@ -34,8 +32,6 @@ int memcmp(const void *s1, const void *s2, size_t n) {
   return __devicelib_memcmp(s1, s2, n);
 }
 
-#ifndef __NVPTX__
-
 // This simple rand is for ease of use only, the implementation aligns with
 // LLVM libc rand which is based on xorshift64star pseudo random number
 // generator. If work item number <= 1024, each work item has its own internal
@@ -107,8 +103,6 @@ void srand(unsigned int seed) {
   RAND_NEXT_ACC[gid1] = seed;
 }
 
-#endif
-
 #if defined(_WIN32)
 // Truncates a wide (16 or 32 bit) string (wstr) into an ASCII string (str).
 // Any non-ASCII characters are replaced by question mark '?'.

llvm/lib/SYCLLowerIR/LowerWGScope.cpp

@@ -144,7 +144,9 @@ template <typename T> static unsigned asUInt(T val) {
 
 static IntegerType *getSizeTTy(Module &M) {
   LLVMContext &Ctx = M.getContext();
-  auto PtrSize = M.getDataLayout().getPointerTypeSize(PointerType::getUnqual(Ctx));
+  const DataLayout &DL = M.getDataLayout();
+  auto PtrSize = DL.getPointerTypeSize(
+      PointerType::get(Ctx, DL.getDefaultGlobalsAddressSpace()));
   return PtrSize == 8 ? Type::getInt64Ty(Ctx) : Type::getInt32Ty(Ctx);
 }
 

llvm/test/SYCLLowerIR/convergent_ptr_size_in_addrspace.ll

@@ -0,0 +1,24 @@
+; RUN: opt < %s -passes=LowerWGScope -S | FileCheck %s
+
+; This test checks that pointer size in default global address space is used for
+; size_t type, which is value type of GV __spirv_BuiltInLocalInvocationIndex.
+; Note that pointer size in the default address space is 4.
+
+target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-i64:64-v16:16-v32:32-n16:32:64-G1"
+
+%struct.baz = type { i64 }
+
+; CHECK: @__spirv_BuiltInLocalInvocationIndex = external addrspace(1) constant i64, align 8
+
+define internal void @wibble(ptr byval(%struct.baz) %arg1) !work_group_scope !0 {
+; CHECK:   load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 8
+; CHECK:   call void @_Z22__spirv_ControlBarrieriii(i32 2, i32 2, i32 272)
+  ret void
+}
+
+; CHECK: ; Function Attrs: convergent
+; CHECK: declare void @_Z22__spirv_ControlBarrieriii(i32, i32, i32) #[[ATTR_NUM:[0-9]+]]
+
+; CHECK: attributes #[[ATTR_NUM]] = { convergent }
+
+!0 = !{}