Merge remote-tracking branch 'upstream/sycl' into georgi/hip-undef-cuda-arch

Author: Georgi Mirazchiyski

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

@@ -117,11 +117,11 @@ jobs:
       env: ${{ matrix.env || '{}' }}
 
       # Do not install drivers on AMD and CUDA runners.
-      install_igc_driver: |
+      install_igc_driver: >-
         ${{ !contains(matrix.target_devices, 'ext_oneapi_cuda') &&
         !contains(matrix.target_devices, 'ext_oneapi_hip') &&
         contains(needs.detect_changes.outputs.filters, 'drivers') }}
-      install_dev_igc_driver: |
+      install_dev_igc_driver: >-
         ${{ !contains(matrix.target_devices, 'ext_oneapi_cuda') &&
         !contains(matrix.target_devices, 'ext_oneapi_hip') &&
         matrix.use_igc_dev && contains(needs.detect_changes.outputs.filters, 'devigccfg') ||

clang/include/clang/Basic/LangOptions.def

@@ -307,7 +307,6 @@ LANGOPT(SYCLAllowFuncPtr  , 1, 0, "Allow function pointers in SYCL device code")
 LANGOPT(SYCLStdLayoutKernelParams, 1, 0, "Enable standard layout requirement for SYCL kernel parameters")
 LANGOPT(SYCLUnnamedLambda , 1, 0, "Allow unnamed lambda SYCL kernels")
 LANGOPT(SYCLForceInlineKernelLambda , 1, 0, "Force inline SYCL kernel lambdas in entry point")
-LANGOPT(SYCLAllowAllFeaturesInConstexpr, 1, 0, "Allow all C++ features in SYCL device code in manifestly constant-evaluated expressions")
 LANGOPT(SYCLESIMDForceStatelessMem, 1, 0, "Make accessors use USM memory in ESIMD kernels")
 LANGOPT(SYCLESIMDBuildHostCode, 1, 1, "Build the host implementation of ESIMD functions")
 ENUM_LANGOPT(SYCLVersion  , SYCLMajorVersion, 2, SYCL_None, "Version of the SYCL standard used")

clang/include/clang/Driver/Options.td

@@ -8628,11 +8628,6 @@ def fsycl_is_native_cpu : Flag<["-"], "fsycl-is-native-cpu">,
   HelpText<"Perform device compilation for Native CPU.">,
   Visibility<[CC1Option]>,
   MarshallingInfoFlag<LangOpts<"SYCLIsNativeCPU">>;
-defm sycl_allow_all_features_in_constexpr
-    : BoolFOption<
-          "sycl-allow-all-features-in-constexpr", LangOpts<"SYCLAllowAllFeaturesInConstexpr">,
-          DefaultFalse,
-          PosFlag<SetTrue, [], [CC1Option], "Allow all C++ features in SYCL device code in manifestly constant-evaluated expressions">, NegFlag<SetFalse>>;
 
 } // let Visibility = [CC1Option]
 

clang/lib/Sema/Sema.cpp

@@ -1829,7 +1829,7 @@ class DeferredDiagnosticsEmitter
   }
 
   void VisitDeclStmt(DeclStmt *DS) {
-    if (S.getLangOpts().SYCLAllowAllFeaturesInConstexpr) {
+    if (S.getLangOpts().SYCLIsDevice) {
       if (DS->isSingleDecl()) {
         Decl *D = DS->getSingleDecl();
         if (auto *VD = dyn_cast<VarDecl>(D))
@@ -1850,7 +1850,7 @@ class DeferredDiagnosticsEmitter
   }
 
   void VisitConstantExpr(ConstantExpr *E) {
-    if (S.getLangOpts().SYCLAllowAllFeaturesInConstexpr)
+    if (S.getLangOpts().SYCLIsDevice)
       return;
     this->VisitStmt(E);
   }
@@ -2264,7 +2264,7 @@ void Sema::checkTypeSupport(QualType Ty, SourceLocation Loc, ValueDecl *D) {
 
   CheckType(Ty);
   if (const auto *FD = dyn_cast_if_present<FunctionDecl>(D)) {
-    if (LangOpts.SYCLAllowAllFeaturesInConstexpr && FD->isConsteval())
+    if (LangOpts.SYCLIsDevice && FD->isConsteval())
       return;
     if (const auto *FPTy = dyn_cast<FunctionProtoType>(Ty)) {
       for (const auto &ParamTy : FPTy->param_types())

clang/lib/Sema/SemaDeclCXX.cpp

@@ -18418,7 +18418,7 @@ void Sema::ActOnCXXEnterDeclInitializer(Scope *S, Decl *D) {
 
   if (auto *VD = dyn_cast<VarDecl>(D);
       VD && (VD->mightBeUsableInConstantExpressions(Context)))
-    InConstexprVarInit = LangOpts.SYCLAllowAllFeaturesInConstexpr;
+    InConstexprVarInit = true;
   PushExpressionEvaluationContext(
       ExpressionEvaluationContext::PotentiallyEvaluated, D);
 }

clang/lib/Sema/SemaSYCL.cpp

@@ -782,8 +782,7 @@ class DeviceFunctionTracker {
 
 public:
   DeviceFunctionTracker(SemaSYCL &S) : SemaSYCLRef(S) {
-    if (S.getLangOpts().SYCLAllowAllFeaturesInConstexpr)
-      CG.setSkipConstantExpressions(S.getASTContext());
+    CG.setSkipConstantExpressions(S.getASTContext());
     CG.addToCallGraph(S.getASTContext().getTranslationUnitDecl());
     CollectSyclExternalFuncs();
   }
@@ -5594,14 +5593,12 @@ SemaSYCL::DiagIfDeviceCode(SourceLocation Loc, unsigned DiagID,
       return SemaDiagnosticBuilder::K_ImmediateWithCallStack;
     if (!FD)
       return SemaDiagnosticBuilder::K_Nop;
-    if (SemaRef.getLangOpts().SYCLAllowAllFeaturesInConstexpr &&
-        (SemaRef.isConstantEvaluatedContext() ||
-         SemaRef.currentEvaluationContext().isDiscardedStatementContext()))
+    if (SemaRef.isConstantEvaluatedContext() ||
+        SemaRef.currentEvaluationContext().isDiscardedStatementContext())
       return SemaDiagnosticBuilder::K_Nop;
     // Defer until we know that the variable's intializer is actually a
     // manifestly constant-evaluated expression.
-    if (SemaRef.getLangOpts().SYCLAllowAllFeaturesInConstexpr &&
-        SemaRef.InConstexprVarInit)
+    if (SemaRef.InConstexprVarInit)
       return SemaDiagnosticBuilder::K_Deferred;
     if (SemaRef.getEmissionStatus(FD) ==
         Sema::FunctionEmissionStatus::Emitted) {

clang/test/CodeGenSYCL/free_function_int_header.cpp

@@ -76,6 +76,26 @@ __attribute__((sycl_device))
 
 template void ff_6(Agg S1, Derived S2, int);
 
+constexpr int TestArrSize = 3;
+
+template <int ArrSize>
+struct KArgWithPtrArray {
+  int *data[ArrSize];
+  int start[ArrSize];
+  int end[ArrSize];
+  constexpr int getArrSize() { return ArrSize; }
+};
+
+template <int ArrSize>
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+void ff_7(KArgWithPtrArray<ArrSize> KArg) {
+  for (int j = 0; j < ArrSize; j++)
+    for (int i = KArg.start[j]; i <= KArg.end[j]; i++)
+      KArg.data[j][i] = KArg.start[j] + KArg.end[j];
+}
+
+template void ff_7(KArgWithPtrArray<TestArrSize> KArg);
+
 // CHECK:      const char* const kernel_names[] = {
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_2Piii
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_2Piiii
@@ -84,6 +104,7 @@ template void ff_6(Agg S1, Derived S2, int);
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_3IdEvPT_S0_S0_
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_410NoPointers8Pointers3Agg
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i
+// CHECK-NEXT:   {{.*}}__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
 // CHECK-NEXT:   ""
 // CHECK-NEXT: };
 
@@ -124,6 +145,9 @@ template void ff_6(Agg S1, Derived S2, int);
 // CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 40, 32 },
 // CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 72 },
 
+// CHECK:  //--- _Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 48, 0 },
+
 // CHECK:        { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
 // CHECK-NEXT: };
 
@@ -249,6 +273,26 @@ template void ff_6(Agg S1, Derived S2, int);
 // CHECK-NEXT:   static constexpr bool value = true;
 // CHECK-NEXT: };
 // CHECK-NEXT: }
+//
+// CHECK: Definition of _Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE as a free function kernel
+
+// CHECK: Forward declarations of kernel and its argument types:
+// CHECK: template <int ArrSize> struct KArgWithPtrArray;
+//
+// CHECK: template <int ArrSize> void ff_7(KArgWithPtrArray<ArrSize> KArg);
+// CHECK-NEXT: static constexpr auto __sycl_shim8() {
+// CHECK-NEXT:   return (void (*)(struct KArgWithPtrArray<3>))ff_7<3>;
+// CHECK-NEXT: }
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: template <>
+// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim8()> {
+// CHECK-NEXT:   static constexpr bool value = true;
+// CHECK-NEXT: };
+// CHECK-NEXT: template <>
+// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim8()> {
+// CHECK-NEXT:   static constexpr bool value = true;
+// CHECK-NEXT: };
+// CHECK-NEXT: }
 
 // CHECK: #include <sycl/kernel_bundle.hpp>
 
@@ -307,3 +351,11 @@ template void ff_6(Agg S1, Derived S2, int);
 // CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i"});
 // CHECK-NEXT: }
 // CHECK-NEXT: }
+
+// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: template <>
+// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim8()>() {
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE"});
+// CHECK-NEXT: }
+// CHECK-NEXT: }

clang/test/CodeGenSYCL/free_function_kernel_params.cpp

@@ -26,10 +26,33 @@ __attribute__((sycl_device))
 void ff_4(NoPointers S1, Pointers S2, Agg S3) {
 }
 
+constexpr int TestArrSize = 3;
+
+template <int ArrSize>
+struct KArgWithPtrArray {
+  int *data[ArrSize];
+  int start[ArrSize];
+  int end[ArrSize];
+  constexpr int getArrSize() { return ArrSize; }
+};
+
+template <int ArrSize>
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
+void ff_6(KArgWithPtrArray<ArrSize> KArg) {
+  for (int j = 0; j < ArrSize; j++)
+    for (int i = KArg.start[j]; i <= KArg.end[j]; i++)
+      KArg.data[j][i] = KArg.start[j] + KArg.end[j];
+}
+
+template void ff_6(KArgWithPtrArray<TestArrSize> KArg);
+
 // CHECK: %struct.NoPointers = type { i32 }
 // CHECK: %struct.Pointers = type { ptr addrspace(4), ptr addrspace(4) }
 // CHECK: %struct.Agg = type { %struct.NoPointers, i32, ptr addrspace(4), %struct.Pointers }
 // CHECK: %struct.__generated_Pointers = type { ptr addrspace(1), ptr addrspace(1) }
 // CHECK: %struct.__generated_Agg = type { %struct.NoPointers, i32, ptr addrspace(1), %struct.__generated_Pointers.0 }
 // CHECK: %struct.__generated_Pointers.0 = type { ptr addrspace(1), ptr addrspace(1) }
+// CHECK: %struct.__generated_KArgWithPtrArray = type { [3 x ptr addrspace(1)], [3 x i32], [3 x i32] }
+// CHECK: %struct.KArgWithPtrArray = type { [3 x ptr addrspace(4)], [3 x i32], [3 x i32] }
 // CHECK: define dso_local spir_kernel void @{{.*}}__sycl_kernel{{.*}}(ptr noundef byval(%struct.NoPointers) align 4 %__arg_S1, ptr noundef byval(%struct.__generated_Pointers) align 8 %__arg_S2, ptr noundef byval(%struct.__generated_Agg) align 8 %__arg_S3)
+// CHECK: define dso_local spir_kernel void @{{.*}}__sycl_kernel_ff_6{{.*}}(ptr noundef byval(%struct.__generated_KArgWithPtrArray) align 8 %__arg_KArg)

clang/test/Driver/sycl-device-lib-amdgcn.cpp

@@ -44,7 +44,7 @@
 
 // Check that llvm-link uses the "-only-needed" flag.
 // Not using the flag breaks kernel bundles.
-// RUN: %clangxx -### -nogpulib --sysroot=%S/Inputs/SYCL \
+// RUN: %clangxx -### -nogpulib -fno-sycl-libspirv --sysroot=%S/Inputs/SYCL \
 // RUN: -fsycl -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend --offload-arch=gfx906 %s 2>&1 \
 // RUN: | FileCheck -check-prefix=CHK-ONLY-NEEDED %s
 

clang/test/Driver/sycl-device-lib-nvptx.cpp

@@ -44,7 +44,7 @@
 
 // Check that llvm-link uses the "-only-needed" flag.
 // Not using the flag breaks kernel bundles.
-// RUN: %clangxx -### --sysroot=%S/Inputs/SYCL -fsycl -fsycl-targets=nvptx64-nvidia-cuda %s 2>&1 \
+// RUN: %clangxx -### -nocudalib -fno-sycl-libspirv --sysroot=%S/Inputs/SYCL -fsycl -fsycl-targets=nvptx64-nvidia-cuda %s 2>&1 \
 // RUN: | FileCheck -check-prefix=CHK-ONLY-NEEDED %s
 
 // CHK-ONLY-NEEDED: llvm-link"{{.*}}"-only-needed"{{.*}}"{{.*}}devicelib--cuda.bc"{{.*}}