Merge from 'sycl' to 'sycl-web'

Author: iclsrc

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -3438,10 +3438,6 @@ def err_attribute_argument_is_zero : Error<
 def warn_attribute_argument_n_negative : Warning<
   "%0 attribute parameter %1 is negative and will be ignored">,
   InGroup<CudaCompat>;
-def warn_reqd_sub_group_attribute_cuda_n_32
-    : Warning<"attribute argument %0 is invalid and will be ignored; CUDA "
-              "requires sub_group size 32">,
-      InGroup<CudaCompat>;
 def err_property_function_in_objc_container : Error<
   "use of Objective-C property in function nested in Objective-C "
   "container not supported, move function outside its container">;
@@ -3540,6 +3536,10 @@ def warn_dllimport_dropped_from_inline_function : Warning<
 def warn_attribute_on_direct_kernel_callee_only : Warning<"%0 attribute allowed"
   " only on a function directly called from a SYCL kernel function; attribute ignored">,
   InGroup<IgnoredAttributes>;
+def warn_reqd_sub_group_attribute_n
+    : Warning<"attribute argument %0 is invalid and will be ignored; %1 "
+              "requires sub_group size %2">,
+  InGroup<IgnoredAttributes>;
 def warn_nothrow_attribute_ignored : Warning<"'nothrow' attribute conflicts with"
   " exception specification; attribute ignored">,
   InGroup<IgnoredAttributes>;

clang/lib/Sema/SemaDeclAttr.cpp

@@ -4020,9 +4020,25 @@ void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI,
           << CI << /*positive*/ 0;
       return;
     }
-    if (Context.getTargetInfo().getTriple().isNVPTX() && ArgVal != 32) {
-      Diag(E->getExprLoc(), diag::warn_reqd_sub_group_attribute_cuda_n_32)
-          << ArgVal.getSExtValue();
+    auto &TI = Context.getTargetInfo();
+    if (TI.getTriple().isNVPTX() && ArgVal != 32)
+      Diag(E->getExprLoc(), diag::warn_reqd_sub_group_attribute_n)
+          << ArgVal.getSExtValue() << TI.getTriple().getArchName() << 32;
+    if (TI.getTriple().isAMDGPU()) {
+      const auto HasWaveFrontSize64 =
+          TI.getTargetOpts().FeatureMap["wavefrontsize64"];
+      const auto HasWaveFrontSize32 =
+          TI.getTargetOpts().FeatureMap["wavefrontsize32"];
+
+      // CDNA supports only 64 wave front size, for those GPUs allow subgroup
+      // size of 64. Some GPUs support both 32 and 64, for those (and the rest)
+      // only allow 32. Warn on incompatible sizes.
+      const auto SupportedWaveFrontSize =
+          HasWaveFrontSize64 && !HasWaveFrontSize32 ? 64 : 32;
+      if (ArgVal != SupportedWaveFrontSize)
+        Diag(E->getExprLoc(), diag::warn_reqd_sub_group_attribute_n)
+            << ArgVal.getSExtValue() << TI.getTriple().getArchName()
+            << SupportedWaveFrontSize;
     }
 
     // Check to see if there's a duplicate attribute with different values

clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp

@@ -0,0 +1,24 @@
+// RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx1010 -internal-isystem %S/Inputs -std=c++2b -verify %s
+
+// Sub-group size is optimized for 32, warn (and ignore the attribute) if the
+// size is not 32.
+#include "sycl.hpp"
+
+int main() {
+
+  sycl::queue Q;
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class invalid_kernel>([=] [[sycl::reqd_sub_group_size(64)]] {}); // expected-warning {{attribute argument 64 is invalid and will be ignored; amdgcn requires sub_group size 32}}
+  });
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class valid_kernel>([=] [[sycl::reqd_sub_group_size(32)]] {});
+  });
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class invalid_kernel_2>([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; amdgcn requires sub_group size 32}}
+  });
+
+  return 0;
+}

clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp

@@ -0,0 +1,24 @@
+// RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx90a -internal-isystem %S/Inputs -std=c++2b -verify %s
+
+// Sub-group size is optimized for 64, warn (and ignore the attribute) if the
+// size is not 64.
+#include "sycl.hpp"
+
+int main() {
+
+  sycl::queue Q;
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class valid_kernel>([=] [[sycl::reqd_sub_group_size(64)]] {});
+  });
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class invalid_kernel>([=] [[sycl::reqd_sub_group_size(32)]] {}); // expected-warning {{attribute argument 32 is invalid and will be ignored; amdgcn requires sub_group size 64}}
+  });
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class invalid_kernel_2>([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; amdgcn requires sub_group size 64}}
+  });
+
+  return 0;
+}

clang/test/SemaSYCL/reqd-sub-group-size-cuda.cpp

@@ -9,7 +9,7 @@ int main() {
   sycl::queue Q;
 
   Q.submit([&](sycl::handler &h) {
-    h.single_task<class invalid_kernel>([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; CUDA requires sub_group size 32}}
+    h.single_task<class invalid_kernel>([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; nvptx requires sub_group size 32}}
   });
 
   Q.submit([&](sycl::handler &h) {

libclc/amdgcn-amdhsa/libspirv/workitem/get_local_size.cl

@@ -23,17 +23,19 @@
 CONST_AS char * __clc_amdgcn_dispatch_ptr(void) __asm("llvm.amdgcn.dispatch.ptr");
 #endif
 
+// Mimic `EmitAMDGPUWorkGroupSize` in `clang/lib/CodeGen/CGBuiltin.cpp`.
+
 _CLC_DEF _CLC_OVERLOAD size_t __spirv_WorkgroupSize_x() {
-    CONST_AS uint * ptr = (CONST_AS uint *) __dispatch_ptr();
-    return ptr[1] & 0xffffu;
+    CONST_AS ushort * ptr = (CONST_AS ushort *) __dispatch_ptr();
+    return ptr[2];
 }
 
 _CLC_DEF _CLC_OVERLOAD size_t __spirv_WorkgroupSize_y() {
-    CONST_AS uint * ptr = (CONST_AS uint *) __dispatch_ptr();
-    return ptr[1] >> 16;
+    CONST_AS ushort * ptr = (CONST_AS ushort *) __dispatch_ptr();
+    return ptr[3];
 }
 
 _CLC_DEF _CLC_OVERLOAD size_t __spirv_WorkgroupSize_z() {
-    CONST_AS uint * ptr = (CONST_AS uint *) __dispatch_ptr();
-    return ptr[2] & 0xffffu;
+    CONST_AS ushort * ptr = (CONST_AS ushort *) __dispatch_ptr();
+    return ptr[4];
 }

sycl-fusion/jit-compiler/lib/KernelFusion.cpp

@@ -100,12 +100,6 @@ FusionResult KernelFusion::fuseKernels(
         "Fusion output target format not supported by this build");
   }
 
-  if (TargetFormat != BinaryFormat::SPIRV &&
-      TargetFormat != BinaryFormat::PTX && IsHeterogeneousList) {
-    return FusionResult{
-        "Heterogeneous ND ranges not supported for this target"};
-  }
-
   bool CachingEnabled = ConfigHelper::get<option::JITEnableCaching>();
   CacheKeyT CacheKey{KernelsToFuse,
                      Identities,

sycl-fusion/passes/kernel-fusion/Builtins.cpp

@@ -36,10 +36,9 @@ static raw_ostream &operator<<(raw_ostream &Os, const NDRange &ND) {
             << ND.getLocalSize();
 }
 
-/// Will generate a unique function name so that it can be reused in further
-/// stages.
-static std::string getFunctionName(BuiltinKind K, const NDRange &SrcNDRange,
-                                   const NDRange &FusedNDRange) {
+std::string Remapper::getFunctionName(BuiltinKind K, const NDRange &SrcNDRange,
+                                      const NDRange &FusedNDRange,
+                                      uint32_t Idx) {
   std::string Res;
   raw_string_ostream S{Res};
   S << "__" <<
@@ -63,6 +62,8 @@ static std::string getFunctionName(BuiltinKind K, const NDRange &SrcNDRange,
         llvm_unreachable("Unhandled kind");
       }()
     << "_remapper_" << SrcNDRange << "_" << FusedNDRange;
+  if (Idx != (uint32_t)-1)
+    S << "_" << static_cast<char>('x' + Idx);
   return S.str();
 }
 
@@ -339,13 +340,8 @@ jit_compiler::Remapper::remapBuiltins(Function *F, const NDRange &SrcNDRange,
         // If the builtin should not be remapped, return the original function.
         return F;
 
-      // Remap given builtin.
-      const auto Name = getFunctionName(K, SrcNDRange, FusedNDRange);
-      auto *M = F->getParent();
-      assert(!M->getFunction(Name) && "Function name should be unique");
-
       return Cached = TargetInfo.createRemapperFunction(
-                 *this, K, F->getName(), Name, M, SrcNDRange, FusedNDRange);
+                 *this, K, F, F->getParent(), SrcNDRange, FusedNDRange);
     }
     if (TargetInfo.isSafeToNotRemapBuiltin(F)) {
       // No need to remap.
@@ -375,20 +371,9 @@ jit_compiler::Remapper::remapBuiltins(Function *F, const NDRange &SrcNDRange,
 
   // Set Cached to support recursive functions.
   Cached = Clone;
-  for (auto &I : instructions(Clone)) {
-    if (auto *Call = dyn_cast<CallBase>(&I)) {
-      // Recursive call
-      auto *OldF = Call->getCalledFunction();
-      auto ErrOrNewF = remapBuiltins(OldF, SrcNDRange, FusedNDRange);
-      if (auto Err = ErrOrNewF.takeError()) {
-        return std::move(Err);
-      }
-      // Override called function.
-      auto *NewF = *ErrOrNewF;
-      Call->setCalledFunction(NewF);
-      Call->setCallingConv(NewF->getCallingConv());
-      Call->setAttributes(NewF->getAttributes());
-    }
+  if (auto Err = TargetInfo.scanForBuiltinsToRemap(Clone, *this, SrcNDRange,
+                                                   FusedNDRange)) {
+    return Err;
   }
   return Clone;
 }

sycl-fusion/passes/kernel-fusion/Builtins.h

@@ -47,6 +47,12 @@ class Remapper {
   explicit Remapper(const llvm::TargetFusionInfo &TargetInfo)
       : TargetInfo(TargetInfo) {}
 
+  ///
+  /// Generate a unique function name for a remapper function.
+  static std::string getFunctionName(BuiltinKind K, const NDRange &SrcNDRange,
+                                     const NDRange &FusedNDRange,
+                                     uint32_t Idx = -1);
+
   ///
   /// Recursively remap index space getters builtins.
   llvm::Expected<llvm::Function *> remapBuiltins(llvm::Function *F,