fix comments

Author: shiltian

clang/include/clang/Basic/AttrDocs.td

@@ -7549,18 +7549,32 @@ def CUDAClusterDimsAttrDoc : Documentation {
   let Category = DocCatDecl;
   let Content = [{
 In CUDA/HIP programming, the ``__cluster_dims__`` attribute can be applied to a kernel function
-to set the dimensions of a thread block cluster. This allows to group multiple thread blocks into
-a larger unit called a "cluster". `__cluster_dims__` defines the cluster size as ``(X, Y, Z)``,
-where each value is the number of thread blocks in that dimension.
+to set the dimensions of a thread block cluster. ``__cluster_dims__`` defines the cluster size
+as ``(X, Y, Z)``, where each value is the number of thread blocks in that dimension.
+The ``__cluster_dims__`` and `__no_cluster__`` attributes are mutually exclusive.
+
+.. code::
+
+  __global__ __cluster_dims__(2, 1, 1) void kernel(...) {
+    ...
+  }
+
   }];
 }
 
 def CUDANoClusterAttrDoc : Documentation {
   let Category = DocCatDecl;
   let Content = [{
-In CUDA/HIP programming, the ``__no_cluster__`` attribute can be applied to a kernel function to
-indicate that the thread block cluster feature will not be enabled at both compile time and kernel
-launch time. Note: this is a LLVM/Clang only attribute.
+In CUDA/HIP programming, the LLVM/Clang-exclusive ``__no_cluster__`` attribute can be applied to
+a kernel function to indicate that the thread block cluster feature will not be enabled at both
+compile time and kernel launch time. The ``__cluster_dims__`` and `__no_cluster__`` attributes
+are mutually exclusive.
+
+.. code::
+
+  __global__ __no_cluster__ void kernel(...) {
+    ...
+  }
   }];
 }
 

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -13071,7 +13071,7 @@ def warn_cuda_maxclusterrank_sm_90 : Warning<
   "%1 attribute">, InGroup<IgnoredAttributes>;
 
 def err_cuda_cluster_attr_not_supported : Error<
-  "%select{__cluster_dims__|__no_cluster__}0 is not supported for this GPU architecture"
+  "%0 is not supported for this GPU architecture"
 >;
 
 def err_cuda_cluster_dims_too_large : Error<

clang/lib/Sema/SemaDeclAttr.cpp

@@ -5682,11 +5682,11 @@ makeClusterDimsArgExpr(Sema &S, Expr *E, const CUDAClusterDimsAttr &AL,
 
   // Accept template arguments for now as they depend on something else.
   // We'll get to check them when they eventually get instantiated.
-  if (E->isValueDependent())
+  if (E->isInstantiationDependent())
     return {E, 1};
 
-  std::optional<llvm::APSInt> I = llvm::APSInt(64);
-  if (!(I = E->getIntegerConstantExpr(S.Context))) {
+  std::optional<llvm::APSInt> I = E->getIntegerConstantExpr(S.Context);
+  if (!I) {
     S.Diag(E->getExprLoc(), diag::err_attribute_argument_n_type)
         << &AL << Idx << AANT_ArgumentIntegerConstant << E->getSourceRange();
     return {nullptr, 0};
@@ -5754,7 +5754,7 @@ CUDAClusterDimsAttr *Sema::createClusterDimsAttr(const AttributeCommonInfo &CI,
     return nullptr;
   }
 
-  return ::new (Context) CUDAClusterDimsAttr(Context, CI, X, Y, Z);
+  return CUDAClusterDimsAttr::Create(Context, X, Y, Z, CI);
 }
 
 void Sema::addClusterDimsAttr(Decl *D, const AttributeCommonInfo &CI, Expr *X,
@@ -5764,16 +5764,16 @@ void Sema::addClusterDimsAttr(Decl *D, const AttributeCommonInfo &CI, Expr *X,
 }
 
 void Sema::addNoClusterAttr(Decl *D, const AttributeCommonInfo &CI) {
-  if (CUDANoClusterAttr *Attr = ::new (Context) CUDANoClusterAttr(Context, CI))
-    D->addAttr(Attr);
+  D->addAttr(CUDANoClusterAttr::Create(Context, CI));
 }
 
 static void handleClusterDimsAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   auto &TTI = S.Context.getTargetInfo();
   auto Arch = StringToOffloadArch(TTI.getTargetOpts().CPU);
   if ((TTI.getTriple().isNVPTX() && Arch < clang::OffloadArch::SM_90) ||
       (TTI.getTriple().isAMDGPU() && Arch < clang::OffloadArch::GFX1250)) {
-    S.Diag(AL.getLoc(), diag::err_cuda_cluster_attr_not_supported) << 0;
+    S.Diag(AL.getLoc(), diag::err_cuda_cluster_attr_not_supported)
+        << "__cluster_dims__";
     return;
   }
 
@@ -5791,7 +5791,8 @@ static void handleNoClusterAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   auto Arch = StringToOffloadArch(TTI.getTargetOpts().CPU);
   if ((TTI.getTriple().isNVPTX() && Arch < clang::OffloadArch::SM_90) ||
       (TTI.getTriple().isAMDGPU() && Arch < clang::OffloadArch::GFX1250)) {
-    S.Diag(AL.getLoc(), diag::err_cuda_cluster_attr_not_supported) << 1;
+    S.Diag(AL.getLoc(), diag::err_cuda_cluster_attr_not_supported)
+        << "__no_cluster__";
     return;
   }