Merge branch 'main' into main

Author: eZWALT

clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp

@@ -1160,22 +1160,6 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
   case NVPTX::BI__nvvm_fence_sc_cluster:
     return Builder.CreateCall(
         CGM.getIntrinsic(Intrinsic::nvvm_fence_sc_cluster));
-  case NVPTX::BI__nvvm_bar_sync:
-    return Builder.CreateCall(
-        CGM.getIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all),
-        EmitScalarExpr(E->getArg(0)));
-  case NVPTX::BI__syncthreads:
-    return Builder.CreateCall(
-        CGM.getIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all),
-        Builder.getInt32(0));
-  case NVPTX::BI__nvvm_barrier_sync:
-    return Builder.CreateCall(
-        CGM.getIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all),
-        EmitScalarExpr(E->getArg(0)));
-  case NVPTX::BI__nvvm_barrier_sync_cnt:
-    return Builder.CreateCall(
-        CGM.getIntrinsic(Intrinsic::nvvm_barrier_cta_sync),
-        {EmitScalarExpr(E->getArg(0)), EmitScalarExpr(E->getArg(1))});
   default:
     return nullptr;
   }

clang/lib/Sema/SemaChecking.cpp

@@ -11745,10 +11745,6 @@ static void CheckImplicitArgumentConversions(Sema &S, CallExpr *TheCall,
 
 static void DiagnoseNullConversion(Sema &S, Expr *E, QualType T,
                                    SourceLocation CC) {
-  if (S.Diags.isIgnored(diag::warn_impcast_null_pointer_to_integer,
-                        E->getExprLoc()))
-    return;
-
   // Don't warn on functions which have return type nullptr_t.
   if (isa<CallExpr>(E))
     return;
@@ -11765,6 +11761,10 @@ static void DiagnoseNullConversion(Sema &S, Expr *E, QualType T,
       T->isMemberPointerType() || !T->isScalarType() || T->isNullPtrType())
     return;
 
+  if (S.Diags.isIgnored(diag::warn_impcast_null_pointer_to_integer,
+                        E->getExprLoc()))
+    return;
+
   SourceLocation Loc = E->getSourceRange().getBegin();
 
   // Venture through the macro stacks to get to the source of macro arguments.

clang/test/CodeGen/builtins-nvptx-ptx60.cu

@@ -32,10 +32,10 @@ __device__ void nvvm_sync(unsigned mask, int i, float f, int a, int b,
   // CHECK: call void @llvm.nvvm.bar.warp.sync(i32
   // expected-error@+1 {{'__nvvm_bar_warp_sync' needs target feature ptx60}}
   __nvvm_bar_warp_sync(mask);
-  // CHECK: call void @llvm.nvvm.barrier.cta.sync.all(i32
+  // CHECK: call void @llvm.nvvm.barrier.sync(i32
   // expected-error@+1 {{'__nvvm_barrier_sync' needs target feature ptx60}}
   __nvvm_barrier_sync(mask);
-  // CHECK: call void @llvm.nvvm.barrier.cta.sync(i32
+  // CHECK: call void @llvm.nvvm.barrier.sync.cnt(i32
   // expected-error@+1 {{'__nvvm_barrier_sync_cnt' needs target feature ptx60}}
   __nvvm_barrier_sync_cnt(mask, i);
 

clang/test/CodeGen/builtins-nvptx.c

@@ -198,7 +198,7 @@ __device__ int read_pms() {
 
 __device__ void sync() {
 
-// CHECK: call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
+// CHECK: call void @llvm.nvvm.bar.sync(i32 0)
 
   __nvvm_bar_sync(0);
 
@@ -259,7 +259,7 @@ __device__ void nvvm_math(float f1, float f2, double d1, double d2) {
   __nvvm_membar_gl();
 // CHECK: call void @llvm.nvvm.membar.sys()
   __nvvm_membar_sys();
-// CHECK: call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
+// CHECK: call void @llvm.nvvm.barrier0()
   __syncthreads();
 }
 

clang/test/Headers/gpuintrin.c

@@ -887,7 +887,7 @@ __gpu_kernel void foo() {
 // NVPTX-LABEL: define internal void @__gpu_sync_threads(
 // NVPTX-SAME: ) #[[ATTR0]] {
 // NVPTX-NEXT:  [[ENTRY:.*:]]
-// NVPTX-NEXT:    call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
+// NVPTX-NEXT:    call void @llvm.nvvm.barrier0()
 // NVPTX-NEXT:    ret void
 //
 //

lldb/source/Commands/CommandObjectThreadUtil.cpp

@@ -37,6 +37,8 @@ void CommandObjectIterateOverThreads::DoExecute(Args &command,
   result.SetStatus(m_success_return);
 
   bool all_threads = false;
+  m_unique_stacks = false;
+
   if (command.GetArgumentCount() == 0) {
     Thread *thread = m_exe_ctx.GetThreadPtr();
     if (thread)

lldb/test/API/functionalities/thread/num_threads/TestNumThreads.py

@@ -132,10 +132,32 @@ def is_thread3(thread):
         # Construct our expected back trace string
         expect_string = "10 thread(s)%s" % (expect_threads)
 
+        # There was a bug where if you used 'thread backtrace unique'
+        # we would switch all future backtraces to use the
+        # "frame-format-unique" not the "frame-format".  Make
+        # sure we don't do that...
+        setting_data = self.dbg.GetSetting("frame-format-unique")
+        setting_str = setting_data.GetStringValue(1000)
+        setting_str = "UNIQUE: " + setting_str
+        lldb.SBDebugger.SetInternalVariable(
+            "frame-format-unique", setting_str, self.dbg.GetInstanceName()
+        )
         # Now that we are stopped, we should have 10 threads waiting in the
         # thread3 function. All of these threads should show as one stack.
         self.expect(
             "thread backtrace unique",
             "Backtrace with unique stack shown correctly",
-            substrs=[expect_string, "main.cpp:%d" % self.thread3_before_lock_line],
+            substrs=[
+                expect_string,
+                "UNIQUE:",
+                "main.cpp:%d" % self.thread3_before_lock_line,
+            ],
+        )
+        # Make sure setting the unique flag in the command isn't
+        # persistent:
+        self.expect(
+            "thread backtrace",
+            "Backtrace unique is not sticky",
+            substrs=["UNIQUE:"],
+            matching=False,
         )

llvm/docs/NVPTXUsage.rst

@@ -199,58 +199,21 @@ map in the following way to CUDA builtins:
 Barriers
 --------
 
-'``llvm.nvvm.barrier.cta.*``'
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+'``llvm.nvvm.barrier0``'
+^^^^^^^^^^^^^^^^^^^^^^^^^^^
 
 Syntax:
 """""""
 
 .. code-block:: llvm
 
-  declare void @llvm.nvvm.barrier.cta.sync(i32 %id, i32 %n)
-  declare void @llvm.nvvm.barrier.cta.sync.all(i32 %id)
-  declare void @llvm.nvvm.barrier.cta.arrive(i32 %id, i32 %n)
-
-  declare void @llvm.nvvm.barrier.cta.sync.aligned(i32 %id, i32 %n)
-  declare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 %id)
-  declare void @llvm.nvvm.barrier.cta.arrive.aligned(i32 %id, i32 %n)
+  declare void @llvm.nvvm.barrier0()
 
 Overview:
 """""""""
 
-The '``@llvm.nvvm.barrier.cta.*``' family of intrinsics perform barrier
-synchronization and communication within a CTA. They can be used by the threads
-within the CTA for synchronization and communication.
-
-Semantics:
-""""""""""
-
-Operand %id specifies a logical barrier resource and must fall within the range
-0 through 15. When present, operand %n specifies the number of threads
-participating in the barrier. When specifying a thread count, the value must be
-a multiple of the warp size. With the '``@llvm.nvvm.barrier.cta.sync.*``'
-variants, the '``.all``' suffix indicates that all threads in the CTA should
-participate in the barrier and the %n operand is not present.
-
-All forms of the '``@llvm.nvvm.barrier.cta.*``' intrinsic cause the executing
-thread to wait for all non-exited threads from its warp and then marks the
-warp's arrival at the barrier. In addition to signaling its arrival at the 
-barrier, the '``@llvm.nvvm.barrier.cta.sync.*``' intrinsics cause the executing
-thread to wait for non-exited threads of all other warps participating in the
-barrier to arrive. On the other hand, the '``@llvm.nvvm.barrier.cta.arrive.*``'
-intrinsic does not cause the executing thread to wait for threads of other
-participating warps.
-
-When a barrier completes, the waiting threads are restarted without delay,
-and the barrier is reinitialized so that it can be immediately reused.
-
-The '``@llvm.nvvm.barrier.cta.*``' intrinsic has an optional '``.aligned``'
-modifier to indicate textual alignment of the barrier. When specified, it
-indicates that all threads in the CTA will execute the same
-'``@llvm.nvvm.barrier.cta.*``' instruction. In conditionally executed code, an
-aligned '``@llvm.nvvm.barrier.cta.*``' instruction should only be used if it is
-known that all threads in the CTA evaluate the condition identically, otherwise
-behavior is undefined.
+The '``@llvm.nvvm.barrier0()``' intrinsic emits a PTX ``bar.sync 0``
+instruction, equivalent to the ``__syncthreads()`` call in CUDA.
 
 Electing a thread
 -----------------

llvm/include/llvm/IR/IntrinsicsNVVM.td

@@ -128,12 +128,6 @@
 //   * llvm.nvvm.swap.lo.hi.b64      --> llvm.fshl(x, x, 32)
 //   * llvm.nvvm.atomic.load.inc.32  --> atomicrmw uinc_wrap
 //   * llvm.nvvm.atomic.load.dec.32  --> atomicrmw udec_wrap
-// * llvm.nvvm.barrier0              --> llvm.nvvm.barrier.cta.sync.aligned.all(0)
-// * llvm.nvvm.barrier.n             --> llvm.nvvm.barrier.cta.sync.aligned.all(x)
-// * llvm.nvvm.bar.sync              --> llvm.nvvm.barrier.cta.sync.aligned.all(x)
-// * llvm.nvvm.barrier               --> llvm.nvvm.barrier.cta.sync.aligned(x, y)
-// * llvm.nvvm.barrier.sync          --> llvm.nvvm.barrier.cta.sync.all(x)
-// * llvm.nvvm.barrier.sync.cnt      --> llvm.nvvm.barrier.cta.sync(x, y)
 
 def llvm_global_ptr_ty  : LLVMQualPointerType<1>;         // (global)ptr
 def llvm_shared_ptr_ty  : LLVMQualPointerType<3>;         // (shared)ptr
@@ -1269,28 +1263,35 @@ let TargetPrefix = "nvvm" in {
   defm int_nvvm_atomic_cas_gen_i  : PTXAtomicWithScope3<llvm_anyint_ty>;
 
 // Bar.Sync
+
+  // The builtin for "bar.sync 0" is called __syncthreads.  Unlike most of the
+  // intrinsics in this file, this one is a user-facing API.
+  def int_nvvm_barrier0 : ClangBuiltin<"__syncthreads">,
+      Intrinsic<[], [], [IntrConvergent, IntrNoCallback]>;
+  // Synchronize all threads in the CTA at barrier 'n'.
+  def int_nvvm_barrier_n : ClangBuiltin<"__nvvm_bar_n">,
+      Intrinsic<[], [llvm_i32_ty], [IntrConvergent, IntrNoCallback]>;
+  // Synchronize 'm', a multiple of warp size, (arg 2) threads in
+  // the CTA at barrier 'n' (arg 1).
+  def int_nvvm_barrier : ClangBuiltin<"__nvvm_bar">,
+      Intrinsic<[], [llvm_i32_ty, llvm_i32_ty], [IntrConvergent, IntrNoCallback]>;
   def int_nvvm_barrier0_popc : ClangBuiltin<"__nvvm_bar0_popc">,
       Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrConvergent, IntrNoCallback]>;
   def int_nvvm_barrier0_and : ClangBuiltin<"__nvvm_bar0_and">,
       Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrConvergent, IntrNoCallback]>;
   def int_nvvm_barrier0_or : ClangBuiltin<"__nvvm_bar0_or">,
       Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrConvergent, IntrNoCallback]>;
 
+  def int_nvvm_bar_sync : NVVMBuiltin,
+      Intrinsic<[], [llvm_i32_ty], [IntrConvergent, IntrNoCallback]>;
   def int_nvvm_bar_warp_sync : NVVMBuiltin,
       Intrinsic<[], [llvm_i32_ty], [IntrConvergent, IntrNoCallback]>;
 
-  // barrier{.cta}.sync{.aligned}      a{, b};
-  // barrier{.cta}.arrive{.aligned}    a, b;
-  let IntrProperties = [IntrConvergent, IntrNoCallback] in {
-    foreach align = ["", "_aligned"] in {
-      def int_nvvm_barrier_cta_sync # align # _all :
-          Intrinsic<[], [llvm_i32_ty]>;
-      def int_nvvm_barrier_cta_sync # align :
-          Intrinsic<[], [llvm_i32_ty, llvm_i32_ty]>;
-      def int_nvvm_barrier_cta_arrive # align :
-          Intrinsic<[], [llvm_i32_ty, llvm_i32_ty]>;
-    }
-  }
+  // barrier.sync id[, cnt]
+  def int_nvvm_barrier_sync : NVVMBuiltin,
+      Intrinsic<[], [llvm_i32_ty], [IntrConvergent, IntrNoCallback]>;
+  def int_nvvm_barrier_sync_cnt : NVVMBuiltin,
+      Intrinsic<[], [llvm_i32_ty, llvm_i32_ty], [IntrConvergent, IntrNoCallback]>;
 
   // barrier.cluster.[wait, arrive, arrive.relaxed]
   def int_nvvm_barrier_cluster_arrive :

llvm/lib/IR/AutoUpgrade.cpp

@@ -1343,9 +1343,12 @@ static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn,
         // nvvm.abs.{i,ii}
         Expand =
             Name == "i" || Name == "ll" || Name == "bf16" || Name == "bf16x2";
-      else if (Name.consume_front("fabs."))
+      else if (Name == "fabs.f" || Name == "fabs.ftz.f" || Name == "fabs.d")
         // nvvm.fabs.{f,ftz.f,d}
-        Expand = Name == "f" || Name == "ftz.f" || Name == "d";
+        Expand = true;
+      else if (Name == "clz.ll" || Name == "popc.ll" || Name == "h2f" ||
+               Name == "swap.lo.hi.b64")
+        Expand = true;
       else if (Name.consume_front("max.") || Name.consume_front("min."))
         // nvvm.{min,max}.{i,ii,ui,ull}
         Expand = Name == "s" || Name == "i" || Name == "ll" || Name == "us" ||
@@ -1377,18 +1380,7 @@ static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn,
         Expand = (Name.starts_with("i.") || Name.starts_with("f.") ||
                   Name.starts_with("p."));
       else
-        Expand = StringSwitch<bool>(Name)
-                     .Case("barrier0", true)
-                     .Case("barrier.n", true)
-                     .Case("barrier.sync.cnt", true)
-                     .Case("barrier.sync", true)
-                     .Case("barrier", true)
-                     .Case("bar.sync", true)
-                     .Case("clz.ll", true)
-                     .Case("popc.ll", true)
-                     .Case("h2f", true)
-                     .Case("swap.lo.hi.b64", true)
-                     .Default(false);
+        Expand = false;
 
       if (Expand) {
         NewFn = nullptr;
@@ -2486,20 +2478,6 @@ static Value *upgradeNVVMIntrinsicCall(StringRef Name, CallBase *CI,
     MDNode *MD = MDNode::get(Builder.getContext(), {});
     LD->setMetadata(LLVMContext::MD_invariant_load, MD);
     return LD;
-  } else if (Name == "barrier0" || Name == "barrier.n" || Name == "bar.sync") {
-    Value *Arg =
-        Name.ends_with('0') ? Builder.getInt32(0) : CI->getArgOperand(0);
-    Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
-                                  {}, {Arg});
-  } else if (Name == "barrier") {
-    Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned, {},
-                                  {CI->getArgOperand(0), CI->getArgOperand(1)});
-  } else if (Name == "barrier.sync") {
-    Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
-                                  {CI->getArgOperand(0)});
-  } else if (Name == "barrier.sync.cnt") {
-    Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync, {},
-                                  {CI->getArgOperand(0), CI->getArgOperand(1)});
   } else {
     Intrinsic::ID IID = shouldUpgradeNVPTXBF16Intrinsic(Name);
     if (IID != Intrinsic::not_intrinsic &&