merge main into amd-staging

Author: ronlieb

bolt/lib/Passes/BinaryPasses.cpp

@@ -60,7 +60,12 @@ static cl::opt<DynoStatsSortOrder> DynoStatsSortOrderOpt(
     "print-sorted-by-order",
     cl::desc("use ascending or descending order when printing functions "
              "ordered by dyno stats"),
-    cl::init(DynoStatsSortOrder::Descending), cl::cat(BoltOptCategory));
+    cl::init(DynoStatsSortOrder::Descending),
+    cl::values(clEnumValN(DynoStatsSortOrder::Ascending, "ascending",
+                          "Ascending order"),
+               clEnumValN(DynoStatsSortOrder::Descending, "descending",
+                          "Descending order")),
+    cl::cat(BoltOptCategory));
 
 cl::list<std::string>
 HotTextMoveSections("hot-text-move-sections",

bolt/test/AArch64/print-sorted-by-order.s

@@ -0,0 +1,31 @@
+# Check that --print-sorted-by-order=<ascending/descending> option works properly in llvm-bolt
+#
+# RUN: llvm-mc -filetype=obj -triple aarch64-unknown-unknown %s -o %t.o
+# RUN: %clang %cflags -fPIC -pie %t.o -o %t.exe -nostdlib -Wl,-q
+# RUN: link_fdata %s %t.o %t.fdata
+# RUN: llvm-bolt %t.exe -o %t.bolt --print-sorted-by=all --print-sorted-by-order=ascending \
+# RUN:   --data %t.fdata | FileCheck %s -check-prefix=CHECK-ASCEND
+# RUN: llvm-bolt %t.exe -o %t.bolt --print-sorted-by=all --print-sorted-by-order=descending \
+# RUN:   --data %t.fdata | FileCheck %s -check-prefix=CHECK-DESCEND
+
+# CHECK-ASCEND: BOLT-INFO: top functions sorted by dyno stats are:
+# CHECK-ASCEND-NEXT: bar
+# CHECK-ASCEND-NEXT: foo
+# CHECK-DESCEND: BOLT-INFO: top functions sorted by dyno stats are:
+# CHECK-DESCEND-NEXT: foo
+# CHECK-DESCEND-NEXT: bar
+
+  .text
+  .align 4
+  .global bar
+  .type bar, %function
+bar:
+  mov w0, wzr
+  ret
+
+  .global foo
+  .type foo, %function
+foo:
+# FDATA: 1 foo 0 1 bar 0 0 1
+  bl bar
+  ret

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

@@ -1429,9 +1429,9 @@ void CGOpenMPRuntimeGPU::emitParallelCall(
   if (!CGF.HaveInsertPoint())
     return;
 
-  auto &&ParallelGen = [this, Loc, OutlinedFn, CapturedVars, IfCond, NumThreads,
-                        NumThreadsModifier, Severity, Message](
-                           CodeGenFunction &CGF, PrePostActionTy &Action) {
+  auto &&ParallelGen = [this, Loc, OutlinedFn, CapturedVars, IfCond,
+                        NumThreads](CodeGenFunction &CGF,
+                                    PrePostActionTy &Action) {
     CGBuilderTy &Bld = CGF.Builder;
     llvm::Value *NumThreadsVal = NumThreads;
     llvm::Function *WFn = WrapperFunctionsMap[OutlinedFn];
@@ -1479,34 +1479,21 @@ void CGOpenMPRuntimeGPU::emitParallelCall(
       NumThreadsVal = Bld.CreateZExtOrTrunc(NumThreadsVal, CGF.Int32Ty);
 
     assert(IfCondVal && "Expected a value");
-    RuntimeFunction FnID = OMPRTL___kmpc_parallel_51;
     llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
-    if (CGM.getLangOpts().OpenMPNoNestedParallelism &&
-        CGM.IsSPMDExecutionMode()) {
-      llvm::Value *Args[] = {
-          RTLoc, NumThreadsVal, FnPtr,
-          Bld.CreateBitOrPointerCast(CapturedVarsAddrs.emitRawPointer(CGF),
-                                     CGF.VoidPtrPtrTy),
-          llvm::ConstantInt::get(CGM.SizeTy, CapturedVars.size())};
-      CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
-                              CGM.getModule(), OMPRTL___kmpc_parallel_spmd),
-                          Args);
-    } else {
-      llvm::SmallVector<llvm::Value *, 10> Args(
-        {RTLoc, getThreadID(CGF, Loc), IfCondVal, NumThreadsVal,
-         llvm::ConstantInt::get(CGF.Int32Ty, -1), FnPtr, ID,
-         Bld.CreateBitOrPointerCast(CapturedVarsAddrs.emitRawPointer(CGF),
-                                    CGF.VoidPtrPtrTy),
-         llvm::ConstantInt::get(CGM.SizeTy, CapturedVars.size())});
-    if (NumThreadsModifier == OMPC_NUMTHREADS_strict) {
-      FnID = OMPRTL___kmpc_parallel_60;
-      Args.append({llvm::ConstantInt::get(CGM.Int32Ty, true),
-                   emitSeverityClause(Severity),
-                   emitMessageClause(CGF, Message)});
-    }
-    CGF.EmitRuntimeCall(
-        OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(), FnID), Args);
-   }
+    llvm::Value *Args[] = {
+        RTLoc,
+        getThreadID(CGF, Loc),
+        IfCondVal,
+        NumThreadsVal,
+        llvm::ConstantInt::get(CGF.Int32Ty, -1),
+        FnPtr,
+        ID,
+        Bld.CreateBitOrPointerCast(CapturedVarsAddrs.emitRawPointer(CGF),
+                                   CGF.VoidPtrPtrTy),
+        llvm::ConstantInt::get(CGM.SizeTy, CapturedVars.size())};
+    CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
+                            CGM.getModule(), OMPRTL___kmpc_parallel_51),
+                        Args);
   };
 
   RegionCodeGenTy RCG(ParallelGen);

clang/lib/CodeGen/CGOpenMPRuntimeGPU.h

@@ -224,11 +224,6 @@ class CGOpenMPRuntimeGPU : public CGOpenMPRuntime {
   /// Emits call to void __kmpc_push_num_threads(ident_t *loc, kmp_int32
   /// global_tid, kmp_int32 num_threads) to generate code for 'num_threads'
   /// clause.
-  /// If the modifier 'strict' is given:
-  /// Emits call to void __kmpc_push_num_threads_strict(ident_t *loc, kmp_int32
-  /// global_tid, kmp_int32 num_threads, int severity, const char *message) to
-  /// generate code for 'num_threads' clause with 'strict' modifier.
-  /// \param NumThreads An integer value of threads.
   void emitNumThreadsClause(
       CodeGenFunction &CGF, llvm::Value *NumThreads, SourceLocation Loc,
       OpenMPNumThreadsClauseModifier Modifier = OMPC_NUMTHREADS_unknown,
@@ -297,11 +292,11 @@ class CGOpenMPRuntimeGPU : public CGOpenMPRuntime {
   /// \param NumThreads The value corresponding to the num_threads clause, if
   /// any, or nullptr.
   /// \param NumThreadsModifier The modifier of the num_threads clause, if
-  /// any, ignored otherwise.
+  /// any, ignored otherwise. Currently unused on the device.
   /// \param Severity The severity corresponding to the num_threads clause, if
-  /// any, ignored otherwise.
+  /// any, ignored otherwise. Currently unused on the device.
   /// \param Message The message string corresponding to the num_threads clause,
-  /// if any, or nullptr.
+  /// if any, or nullptr. Currently unused on the device.
   void emitParallelCall(CodeGenFunction &CGF, SourceLocation Loc,
                         llvm::Function *OutlinedFn,
                         ArrayRef<llvm::Value *> CapturedVars,

clang/test/AST/ByteCode/openmp.cpp

@@ -17,12 +17,12 @@ extern int omp_get_thread_num(void);
 
 int test2() {
   int x = 0;
-  int device_result[N] = {0};
+  int result[N] = {0};
 
-  #pragma omp target parallel loop num_threads(strict: N) severity(warning) message("msg")
+  #pragma omp parallel loop num_threads(strict: N) severity(warning) message("msg")
   for (int i = 0; i < N; i++) {
     x = omp_get_thread_num();
-    device_result[i] = i + x;
+    result[i] = i + x;
   }
 }