https://github.com/shiltian updated https://github.com/llvm/llvm-project/pull/156686
>From 5603cdb1f3ff4c7204105d6e2efa31e12c9e89fd Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Wed, 3 Sep 2025 10:47:52 -0400 Subject: [PATCH 01/10] [Clang][HIP][CUDA] Add `__cluster_dims__` and `__no_cluster__` attribute This PR adds basic frontend support for `__cluster_dims__` and `__no_cluster__` attribute. Co-authored-by: Yaxun (Sam) Liu <[email protected]> Co-authored-by: Jay Foad <[email protected]> --- clang/include/clang/Basic/Attr.td | 17 +++ .../clang/Basic/DiagnosticSemaKinds.td | 8 ++ clang/include/clang/Sema/Sema.h | 8 ++ clang/lib/CodeGen/Targets/AMDGPU.cpp | 26 ++++ .../lib/Headers/__clang_hip_runtime_wrapper.h | 2 + clang/lib/Sema/SemaDeclAttr.cpp | 130 ++++++++++++++++++ .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 37 +++++ clang/test/CodeGenCUDA/Inputs/cuda.h | 4 + clang/test/CodeGenCUDA/cluster_dims.cu | 38 +++++ ...a-attribute-supported-attributes-list.test | 2 + clang/test/SemaCUDA/Inputs/cuda.h | 2 + clang/test/SemaCUDA/cluster_dims.cu | 64 +++++++++ 12 files changed, 338 insertions(+) create mode 100644 clang/test/CodeGenCUDA/cluster_dims.cu create mode 100644 clang/test/SemaCUDA/cluster_dims.cu diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 3c697ed8dd882..faae9ed661fe1 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1572,6 +1572,23 @@ def HIPManaged : InheritableAttr { let Documentation = [HIPManagedAttrDocs]; } +def CUDAClusterDims : InheritableAttr { + let Spellings = [GNU<"cluster_dims">, Declspec<"__cluster_dims__">]; + let Args = [ExprArgument<"X">, ExprArgument<"Y", 1>, ExprArgument<"Z", 1>]; + let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">; + let LangOpts = [CUDA]; + let Documentation = [Undocumented]; +} + +def CUDANoCluster : InheritableAttr { + let Spellings = [GNU<"no_cluster">, Declspec<"__no_cluster__">]; + let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">; + let LangOpts = [CUDA]; + let Documentation = [Undocumented]; +} + +def : MutualExclusions<[CUDAClusterDims, CUDANoCluster]>; + def CUDAInvalidTarget : InheritableAttr { let Spellings = []; let Subjects = SubjectList<[Function]>; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index b157cbb0b8069..5599ee51addb3 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -13067,6 +13067,14 @@ def warn_cuda_maxclusterrank_sm_90 : Warning< "maxclusterrank requires sm_90 or higher, CUDA arch provided: %0, ignoring " "%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" +>; + +def err_cuda_cluster_dims_too_large : Error< + "only a maximum of %0 thread blocks in a cluster is supported" +>; + // VTable pointer authentication errors def err_non_polymorphic_vtable_pointer_auth : Error< "cannot set vtable pointer authentication on monomorphic type %0">; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 265462a86e405..5ac286fd68192 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -5010,6 +5010,14 @@ class Sema final : public SemaBase { void AddLaunchBoundsAttr(Decl *D, const AttributeCommonInfo &CI, Expr *MaxThreads, Expr *MinBlocks, Expr *MaxBlocks); + /// Add a cluster_dims attribute to a particular declaration. + CUDAClusterDimsAttr *createClusterDimsAttr(const AttributeCommonInfo &CI, + Expr *X, Expr *Y, Expr *Z); + void addClusterDimsAttr(Decl *D, const AttributeCommonInfo &CI, Expr *X, + Expr *Y, Expr *Z); + /// Add a no_cluster attribute to a particular declaration. + void addNoClusterAttr(Decl *D, const AttributeCommonInfo &CI); + enum class RetainOwnershipKind { NS, CF, OS }; UuidAttr *mergeUuidAttr(Decl *D, const AttributeCommonInfo &CI, diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index 0fcbf7e458a34..48855ce485f91 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -342,6 +342,9 @@ static bool requiresAMDGPUProtectedVisibility(const Decl *D, void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes( const FunctionDecl *FD, llvm::Function *F, CodeGenModule &M) const { + llvm::StringMap<bool> TargetFetureMap; + M.getContext().getFunctionFeatureMap(TargetFetureMap, FD); + const auto *ReqdWGS = M.getLangOpts().OpenCL ? FD->getAttr<ReqdWorkGroupSizeAttr>() : nullptr; const bool IsOpenCLKernel = @@ -402,6 +405,29 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes( F->addFnAttr("amdgpu-max-num-workgroups", AttrVal.str()); } + + if (auto *Attr = FD->getAttr<CUDAClusterDimsAttr>()) { + uint32_t X = + Attr->getX()->EvaluateKnownConstInt(M.getContext()).getExtValue(); + uint32_t Y = + Attr->getY() + ? Attr->getY()->EvaluateKnownConstInt(M.getContext()).getExtValue() + : 1; + uint32_t Z = + Attr->getZ() + ? Attr->getZ()->EvaluateKnownConstInt(M.getContext()).getExtValue() + : 1; + + llvm::SmallString<32> AttrVal; + llvm::raw_svector_ostream OS(AttrVal); + OS << X << ',' << Y << ',' << Z; + F->addFnAttr("amdgpu-cluster-dims", AttrVal.str()); + } + + // OpenCL doesn't support cluster feature. + if ((IsOpenCLKernel && TargetFetureMap.lookup("gfx1250-insts")) || + FD->getAttr<CUDANoClusterAttr>()) + F->addFnAttr("amdgpu-cluster-dims", "0,0,0"); } void AMDGPUTargetCodeGenInfo::setTargetAttributes( diff --git a/clang/lib/Headers/__clang_hip_runtime_wrapper.h b/clang/lib/Headers/__clang_hip_runtime_wrapper.h index da1e39ac7270e..fb0ece96e1418 100644 --- a/clang/lib/Headers/__clang_hip_runtime_wrapper.h +++ b/clang/lib/Headers/__clang_hip_runtime_wrapper.h @@ -25,6 +25,8 @@ #define __constant__ __attribute__((constant)) #define __managed__ __attribute__((managed)) +#define __cluster_dims__(...) __attribute__((cluster_dims(__VA_ARGS__))) + #if !defined(__cplusplus) || __cplusplus < 201103L #define nullptr NULL; #endif diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 328ccf6694073..5a393084467f5 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5683,6 +5683,130 @@ static void handleLaunchBoundsAttr(Sema &S, Decl *D, const ParsedAttr &AL) { AL.getNumArgs() > 2 ? AL.getArgAsExpr(2) : nullptr); } +static std::pair<Expr *, int> +makeClusterDimsArgExpr(Sema &S, Expr *E, const CUDAClusterDimsAttr &AL, + const unsigned Idx) { + if (S.DiagnoseUnexpandedParameterPack(E)) + return {nullptr, 0}; + + // 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()) + return {E, 1}; + + std::optional<llvm::APSInt> I = llvm::APSInt(64); + if (!(I = E->getIntegerConstantExpr(S.Context))) { + S.Diag(E->getExprLoc(), diag::err_attribute_argument_n_type) + << &AL << Idx << AANT_ArgumentIntegerConstant << E->getSourceRange(); + return {nullptr, 0}; + } + // Make sure we can fit it in 4 bits. + if (!I->isIntN(4)) { + S.Diag(E->getExprLoc(), diag::err_ice_too_large) + << toString(*I, 10, false) << 4 << /* Unsigned */ 1; + return {nullptr, 0}; + } + if (*I < 0) + S.Diag(E->getExprLoc(), diag::warn_attribute_argument_n_negative) + << &AL << Idx << E->getSourceRange(); + + // We may need to perform implicit conversion of the argument. + InitializedEntity Entity = InitializedEntity::InitializeParameter( + S.Context, S.Context.getConstType(S.Context.IntTy), /*consume*/ false); + ExprResult ValArg = S.PerformCopyInitialization(Entity, SourceLocation(), E); + assert(!ValArg.isInvalid() && + "Unexpected PerformCopyInitialization() failure."); + + return {ValArg.getAs<Expr>(), I->getZExtValue()}; +} + +CUDAClusterDimsAttr *Sema::createClusterDimsAttr(const AttributeCommonInfo &CI, + Expr *X, Expr *Y, Expr *Z) { + CUDAClusterDimsAttr TmpAttr(Context, CI, X, Y, Z); + + int ValX = 1; + int ValY = 1; + int ValZ = 1; + + std::tie(X, ValX) = makeClusterDimsArgExpr(*this, X, TmpAttr, /*Idx=*/0); + if (!X) + return nullptr; + + if (Y) { + std::tie(Y, ValY) = makeClusterDimsArgExpr(*this, Y, TmpAttr, /*Idx=*/1); + if (!Y) + return nullptr; + } + + if (Z) { + std::tie(Z, ValZ) = makeClusterDimsArgExpr(*this, Z, TmpAttr, /*Idx=*/2); + if (!Z) + return nullptr; + } + + int FlatDim = ValX * ValY * ValZ; + auto TT = (!Context.getLangOpts().CUDAIsDevice && Context.getAuxTargetInfo()) + ? Context.getAuxTargetInfo()->getTriple() + : Context.getTargetInfo().getTriple(); + int MaxDim = 1; + if (TT.isNVPTX()) + MaxDim = 8; + else if (TT.isAMDGPU()) + MaxDim = 16; + else + return nullptr; + + // A maximum of 8 thread blocks in a cluster is supported as a portable + // cluster size in CUDA. The number is 16 for AMDGPU. + if (FlatDim > MaxDim) { + Diag(CI.getLoc(), diag::err_cuda_cluster_dims_too_large) << MaxDim; + return nullptr; + } + + return ::new (Context) CUDAClusterDimsAttr(Context, CI, X, Y, Z); +} + +void Sema::addClusterDimsAttr(Decl *D, const AttributeCommonInfo &CI, Expr *X, + Expr *Y, Expr *Z) { + if (auto *Attr = createClusterDimsAttr(CI, X, Y, Z)) + D->addAttr(Attr); +} + +void Sema::addNoClusterAttr(Decl *D, const AttributeCommonInfo &CI) { + if (CUDANoClusterAttr *Attr = ::new (Context) CUDANoClusterAttr(Context, CI)) + D->addAttr(Attr); +} + +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; + return; + } + + if (!AL.checkAtLeastNumArgs(S, /*Num=*/1) || + !AL.checkAtMostNumArgs(S, /*Num=*/3)) + return; + + S.addClusterDimsAttr(D, AL, AL.getArgAsExpr(0), + AL.getNumArgs() > 1 ? AL.getArgAsExpr(1) : nullptr, + AL.getNumArgs() > 2 ? AL.getArgAsExpr(2) : nullptr); +} + +static void handleNoClusterAttr(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) << 1; + return; + } + + S.addNoClusterAttr(D, AL); +} + static void handleArgumentWithTypeTagAttr(Sema &S, Decl *D, const ParsedAttr &AL) { if (!AL.isArgIdent(0)) { @@ -7145,6 +7269,12 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL, case ParsedAttr::AT_CUDALaunchBounds: handleLaunchBoundsAttr(S, D, AL); break; + case ParsedAttr::AT_CUDAClusterDims: + handleClusterDimsAttr(S, D, AL); + break; + case ParsedAttr::AT_CUDANoCluster: + handleNoClusterAttr(S, D, AL); + break; case ParsedAttr::AT_Restrict: handleRestrictAttr(S, D, AL); break; diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 3819f775811e5..081a720916fef 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -707,6 +707,38 @@ static void instantiateDependentAMDGPUMaxNumWorkGroupsAttr( S.AMDGPU().addAMDGPUMaxNumWorkGroupsAttr(New, Attr, XExpr, YExpr, ZExpr); } +static void instantiateDependentCUDAClusterDimsAttr( + Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, + const CUDAClusterDimsAttr &Attr, Decl *New) { + EnterExpressionEvaluationContext Unevaluated( + S, Sema::ExpressionEvaluationContext::ConstantEvaluated); + + Expr *XExpr = nullptr; + Expr *YExpr = nullptr; + Expr *ZExpr = nullptr; + + if (Attr.getX()) { + ExprResult ResultX = S.SubstExpr(Attr.getX(), TemplateArgs); + if (ResultX.isUsable()) + XExpr = ResultX.getAs<Expr>(); + } + + if (Attr.getY()) { + ExprResult ResultY = S.SubstExpr(Attr.getY(), TemplateArgs); + if (ResultY.isUsable()) + YExpr = ResultY.getAs<Expr>(); + } + + if (Attr.getZ()) { + ExprResult ResultZ = S.SubstExpr(Attr.getZ(), TemplateArgs); + if (ResultZ.isUsable()) + ZExpr = ResultZ.getAs<Expr>(); + } + + if (XExpr) + S.addClusterDimsAttr(New, Attr, XExpr, YExpr, ZExpr); +} + // This doesn't take any template parameters, but we have a custom action that // needs to happen when the kernel itself is instantiated. We need to run the // ItaniumMangler to mark the names required to name this kernel. @@ -921,6 +953,11 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, *this, TemplateArgs, *AMDGPUMaxNumWorkGroups, New); } + if (const auto *CUDAClusterDims = dyn_cast<CUDAClusterDimsAttr>(TmplAttr)) { + instantiateDependentCUDAClusterDimsAttr(*this, TemplateArgs, + *CUDAClusterDims, New); + } + if (const auto *ParamAttr = dyn_cast<HLSLParamModifierAttr>(TmplAttr)) { instantiateDependentHLSLParamModifierAttr(*this, TemplateArgs, ParamAttr, New); diff --git a/clang/test/CodeGenCUDA/Inputs/cuda.h b/clang/test/CodeGenCUDA/Inputs/cuda.h index dc85eae0c5178..e7ad784335027 100644 --- a/clang/test/CodeGenCUDA/Inputs/cuda.h +++ b/clang/test/CodeGenCUDA/Inputs/cuda.h @@ -13,6 +13,8 @@ #endif #define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__))) #define __grid_constant__ __attribute__((grid_constant)) +#define __cluster_dims__(...) __attribute__((cluster_dims(__VA_ARGS__))) +#define __no_cluster__ __attribute__((no_cluster)) #else #define __constant__ #define __device__ @@ -22,6 +24,8 @@ #define __managed__ #define __launch_bounds__(...) #define __grid_constant__ +#define __cluster_dims__(...) +#define __no_cluster__ #endif struct dim3 { diff --git a/clang/test/CodeGenCUDA/cluster_dims.cu b/clang/test/CodeGenCUDA/cluster_dims.cu new file mode 100644 index 0000000000000..00635e3572a7f --- /dev/null +++ b/clang/test/CodeGenCUDA/cluster_dims.cu @@ -0,0 +1,38 @@ +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx1250 -fcuda-is-device -emit-llvm -x hip -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -aux-triple amdgcn-amd-amdhsa -emit-llvm -x hip -o - %s | FileCheck --check-prefix=HOST %s + +#include "Inputs/cuda.h" + +const int constint = 4; + +// HOST-NOT: "amdgpu-cluster-dims" + +// CHECK: "amdgpu-cluster-dims"="2,2,2" +__global__ void __cluster_dims__(2, 2, 2) test_literal_3d() {} + +// CHECK: "amdgpu-cluster-dims"="2,2,1" +__global__ void __cluster_dims__(2, 2) test_literal_2d() {} + +// CHECK: "amdgpu-cluster-dims"="4,1,1" +__global__ void __cluster_dims__(4) test_literal_1d() {} + +// CHECK: "amdgpu-cluster-dims"="4,2,1" +__global__ void __cluster_dims__(constint, constint / 2, 1) test_constant() {} + +// CHECK: "amdgpu-cluster-dims"="0,0,0" +__global__ void __no_cluster__ test_no_cluster() {} + +// CHECK: "amdgpu-cluster-dims"="7,1,1" +template<unsigned a> +__global__ void __cluster_dims__(a) test_template_1d() {} +template __global__ void test_template_1d<7>(); + +// CHECK: "amdgpu-cluster-dims"="2,6,1" +template<unsigned a, unsigned b> +__global__ void __cluster_dims__(a, b) test_template_2d() {} +template __global__ void test_template_2d<2, 6>(); + +// CHECK: "amdgpu-cluster-dims"="1,2,3" +template<unsigned a, unsigned b, unsigned c> +__global__ void __cluster_dims__(a, b, c) test_template_3d() {} +template __global__ void test_template_3d<1, 2, 3>(); diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index 73d4cb1769ed5..693d54159804e 100644 --- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test @@ -35,6 +35,7 @@ // CHECK-NEXT: CFUnknownTransfer (SubjectMatchRule_function) // CHECK-NEXT: CPUDispatch (SubjectMatchRule_function) // CHECK-NEXT: CPUSpecific (SubjectMatchRule_function) +// CHECK-NEXT: CUDAClusterDims (SubjectMatchRule_function) // CHECK-NEXT: CUDAConstant (SubjectMatchRule_variable) // CHECK-NEXT: CUDADevice (SubjectMatchRule_function, SubjectMatchRule_variable) // CHECK-NEXT: CUDADeviceBuiltinSurfaceType (SubjectMatchRule_record) @@ -43,6 +44,7 @@ // CHECK-NEXT: CUDAGridConstant (SubjectMatchRule_variable_is_parameter) // CHECK-NEXT: CUDAHost (SubjectMatchRule_function) // CHECK-NEXT: CUDALaunchBounds (SubjectMatchRule_objc_method, SubjectMatchRule_hasType_functionType) +// CHECK-NEXT: CUDANoCluster (SubjectMatchRule_function) // CHECK-NEXT: CUDAShared (SubjectMatchRule_variable) // CHECK-NEXT: CXX11NoReturn (SubjectMatchRule_function) // CHECK-NEXT: CallableWhen (SubjectMatchRule_function_is_member) diff --git a/clang/test/SemaCUDA/Inputs/cuda.h b/clang/test/SemaCUDA/Inputs/cuda.h index 10db947d8246c..2bf45e03d91c7 100644 --- a/clang/test/SemaCUDA/Inputs/cuda.h +++ b/clang/test/SemaCUDA/Inputs/cuda.h @@ -13,6 +13,8 @@ #define __managed__ __attribute__((managed)) #define __grid_constant__ __attribute__((grid_constant)) #define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__))) +#define __cluster_dims__(...) __attribute__((cluster_dims(__VA_ARGS__))) +#define __no_cluster__ __attribute__((no_cluster)) struct dim3 { unsigned x, y, z; diff --git a/clang/test/SemaCUDA/cluster_dims.cu b/clang/test/SemaCUDA/cluster_dims.cu new file mode 100644 index 0000000000000..3cd0e0197c29b --- /dev/null +++ b/clang/test/SemaCUDA/cluster_dims.cu @@ -0,0 +1,64 @@ +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -ast-print -x hip -verify=NS,all %s +// RUN: %clang_cc1 -triple nvptx-nvidia-cuda -fcuda-is-device -ast-print -x hip -verify=NS,all %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx1250 -fcuda-is-device -ast-print -x hip -verify=amd,common,all %s | FileCheck -check-prefixes=CHECK %s +// RUN: %clang_cc1 -triple nvptx-nvidia-cuda -target-cpu sm_90 -fcuda-is-device -ast-print -x hip -verify=cuda,common,all %s | FileCheck -check-prefixes=CHECK %s +// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -aux-triple amdgcn-amd-amdhsa -ast-print -x hip -verify=amd,common,all %s | FileCheck -check-prefixes=CHECK %s + +#include "Inputs/cuda.h" + +const int constint = 4; + +// CHECK: __attribute__((global)) __attribute__((cluster_dims(2, 2, 2))) void test_literal_3d() +__global__ void __cluster_dims__(2, 2, 2) test_literal_3d() {} //NS-error {{__cluster_dims__ is not supported for this GPU architecture}} + +// CHECK: __attribute__((global)) __attribute__((cluster_dims(2, 2))) void test_literal_2d() +__global__ void __cluster_dims__(2, 2) test_literal_2d() {} //NS-error {{__cluster_dims__ is not supported for this GPU architecture}} + +// CHECK: __attribute__((global)) __attribute__((cluster_dims(4))) void test_literal_1d() +__global__ void __cluster_dims__(4) test_literal_1d() {} //NS-error {{__cluster_dims__ is not supported for this GPU architecture}} + +// CHECK: __attribute__((global)) __attribute__((cluster_dims(constint, constint / 4, 1))) void test_constant() +__global__ void __cluster_dims__(constint, constint / 4, 1) test_constant() {} //NS-error {{__cluster_dims__ is not supported for this GPU architecture}} + +// CHECK: template <int x, int y, int z> void test_template() __attribute__((cluster_dims(x, y, z))) +template <int x, int y, int z> void test_template(void) __cluster_dims__(x, y, z){} //NS-error {{__cluster_dims__ is not supported for this GPU architecture}} + +// CHECK: template <int x, int y, int z> void test_template_expr() __attribute__((cluster_dims(x + constint, y, z))) +template <int x, int y, int z> void test_template_expr(void) __cluster_dims__(x + constint, y, z) {} //NS-error {{__cluster_dims__ is not supported for this GPU architecture}} + +//NS-error@+1 {{__cluster_dims__ is not supported for this GPU architecture}} +__global__ void __cluster_dims__(32, 2, 4) test_too_large_dim_0() {} // common-error {{integer constant expression evaluates to value 32 that cannot be represented in a 4-bit unsigned integer type}} + +// cuda-error@+2 {{only a maximum of 8 thread blocks in a cluster is supported}} +// amd-error@+1 {{only a maximum of 16 thread blocks in a cluster is supported}} +__global__ void __cluster_dims__(4, 4, 4) test_too_large_dim_1() {} // NS-error {{__cluster_dims__ is not supported for this GPU architecture}} + +// cuda-error@+3 {{only a maximum of 8 thread blocks in a cluster is supported}} +// amd-error@+2 {{only a maximum of 16 thread blocks in a cluster is supported}} +template<unsigned a, unsigned b, unsigned c> +__global__ void __cluster_dims__(a, b, c) test_too_large_dim_template() {} // NS-error {{__cluster_dims__ is not supported for this GPU architecture}} +template __global__ void test_too_large_dim_template<4, 4, 4>(); // common-note {{in instantiation of function template specialization 'test_too_large_dim_template<4U, 4U, 4U>' requested here}} + +int none_const_int = 4; + +//NS-error@+1 {{__cluster_dims__ is not supported for this GPU architecture}} +__global__ void __cluster_dims__(none_const_int, 2, 4) test_non_constant_0() {} // common-error {{'cluster_dims' attribute requires parameter 0 to be an integer constant}} + +//NS-error@+1 {{__cluster_dims__ is not supported for this GPU architecture}} +__global__ void __cluster_dims__(8, none_const_int / 2, 4) test_non_constant_1() {} // common-error {{'cluster_dims' attribute requires parameter 1 to be an integer constant}} + +//NS-error@+1 {{__cluster_dims__ is not supported for this GPU architecture}} +__global__ void __cluster_dims__(8, 2, none_const_int / 4) test_non_constant_2() {} // common-error {{'cluster_dims' attribute requires parameter 2 to be an integer constant}} + +//NS-error@+1 {{__no_cluster__ is not supported for this GPU architecture}} +__global__ void __no_cluster__ test_no_cluster() {} + +//NS-error@+2 {{__no_cluster__ is not supported for this GPU architecture}} +//NS-error@+1 {{__cluster_dims__ is not supported for this GPU architecture}} +__global__ void __no_cluster__ __cluster_dims__(2,2,2) test_have_both() {} // common-error {{'cluster_dims' and 'no_cluster' attributes are not compatible}} common-note {{conflicting attribute is here}} + +template <int... args> +__cluster_dims__(args) void test_template_variadic_args(void) {} // all-error {{expression contains unexpanded parameter pack 'args'}} + +template <int... args> +__cluster_dims__(1, args) void test_template_variadic_args_2(void) {} // all-error {{expression contains unexpanded parameter pack 'args'}} >From 69aba1539e05ddb4c70e72cc99e9699acbc56754 Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Wed, 3 Sep 2025 18:44:16 -0400 Subject: [PATCH 02/10] add documentation --- clang/include/clang/Basic/Attr.td | 4 ++-- clang/include/clang/Basic/AttrDocs.td | 19 +++++++++++++++++++ 2 files changed, 21 insertions(+), 2 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index faae9ed661fe1..6ab76e8721434 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1577,14 +1577,14 @@ def CUDAClusterDims : InheritableAttr { let Args = [ExprArgument<"X">, ExprArgument<"Y", 1>, ExprArgument<"Z", 1>]; let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">; let LangOpts = [CUDA]; - let Documentation = [Undocumented]; + let Documentation = [CUDAClusterDimsAttrDoc]; } def CUDANoCluster : InheritableAttr { let Spellings = [GNU<"no_cluster">, Declspec<"__no_cluster__">]; let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">; let LangOpts = [CUDA]; - let Documentation = [Undocumented]; + let Documentation = [CUDANoClusterAttrDoc]; } def : MutualExclusions<[CUDAClusterDims, CUDANoCluster]>; diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 20a52b49a8f10..0ca31cbdc048b 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -7533,6 +7533,25 @@ A managed variable can be accessed in both device and host code. }]; } +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. + }]; +} + +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. + }]; +} + def LifetimeOwnerDocs : Documentation { let Category = DocCatDecl; let Content = [{ >From cdcaa8b31810fca73eecb557d0a56bc1812216b7 Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Wed, 10 Sep 2025 12:03:42 -0400 Subject: [PATCH 03/10] fix comments --- clang/include/clang/Basic/AttrDocs.td | 26 ++++++++++++++----- .../clang/Basic/DiagnosticSemaKinds.td | 2 +- clang/lib/Sema/SemaDeclAttr.cpp | 17 ++++++------ 3 files changed, 30 insertions(+), 15 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 0ca31cbdc048b..fcc4f5cbdefd3 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -7537,18 +7537,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(...) { + ... + } }]; } diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 5599ee51addb3..ce37075cb6184 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -13068,7 +13068,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< diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 5a393084467f5..5f4e7ff2f9b04 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5691,11 +5691,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}; @@ -5763,7 +5763,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, @@ -5773,8 +5773,7 @@ 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) { @@ -5782,7 +5781,8 @@ static void handleClusterDimsAttr(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) << 0; + S.Diag(AL.getLoc(), diag::err_cuda_cluster_attr_not_supported) + << "__cluster_dims__"; return; } @@ -5800,7 +5800,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; } >From 12a904894588c25d628f5835003dbcb4ef26d852 Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Wed, 24 Sep 2025 13:00:38 -0400 Subject: [PATCH 04/10] minor improvements --- clang/lib/Sema/SemaDeclAttr.cpp | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 5f4e7ff2f9b04..5af7e7c42af55 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5687,32 +5687,33 @@ static std::pair<Expr *, int> makeClusterDimsArgExpr(Sema &S, Expr *E, const CUDAClusterDimsAttr &AL, const unsigned Idx) { if (S.DiagnoseUnexpandedParameterPack(E)) - return {nullptr, 0}; + return {}; // Accept template arguments for now as they depend on something else. // We'll get to check them when they eventually get instantiated. if (E->isInstantiationDependent()) - return {E, 1}; + return {}; 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}; + return {}; } // Make sure we can fit it in 4 bits. if (!I->isIntN(4)) { S.Diag(E->getExprLoc(), diag::err_ice_too_large) - << toString(*I, 10, false) << 4 << /* Unsigned */ 1; - return {nullptr, 0}; + << toString(*I, 10, false) << 4 << /*Unsigned=*/1; + return {}; } - if (*I < 0) + if (*I < 0) { S.Diag(E->getExprLoc(), diag::warn_attribute_argument_n_negative) << &AL << Idx << E->getSourceRange(); + } // We may need to perform implicit conversion of the argument. InitializedEntity Entity = InitializedEntity::InitializeParameter( - S.Context, S.Context.getConstType(S.Context.IntTy), /*consume*/ false); + S.Context, S.Context.getConstType(S.Context.IntTy), /*consume=*/false); ExprResult ValArg = S.PerformCopyInitialization(Entity, SourceLocation(), E); assert(!ValArg.isInvalid() && "Unexpected PerformCopyInitialization() failure."); >From 4e715e75c7f62eb2ff9cde1475a1fa7da0d53c25 Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Wed, 24 Sep 2025 13:26:26 -0400 Subject: [PATCH 05/10] fix an error --- clang/lib/Sema/SemaDeclAttr.cpp | 12 +++--------- 1 file changed, 3 insertions(+), 9 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 5af7e7c42af55..e3cd74964a25b 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5692,7 +5692,7 @@ 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->isInstantiationDependent()) - return {}; + return {E, 1}; std::optional<llvm::APSInt> I = E->getIntegerConstantExpr(S.Context); if (!I) { @@ -5711,14 +5711,8 @@ makeClusterDimsArgExpr(Sema &S, Expr *E, const CUDAClusterDimsAttr &AL, << &AL << Idx << E->getSourceRange(); } - // We may need to perform implicit conversion of the argument. - InitializedEntity Entity = InitializedEntity::InitializeParameter( - S.Context, S.Context.getConstType(S.Context.IntTy), /*consume=*/false); - ExprResult ValArg = S.PerformCopyInitialization(Entity, SourceLocation(), E); - assert(!ValArg.isInvalid() && - "Unexpected PerformCopyInitialization() failure."); - - return {ValArg.getAs<Expr>(), I->getZExtValue()}; + return {ConstantExpr::Create(S.getASTContext(), E, APValue(*I)), + I->getZExtValue()}; } CUDAClusterDimsAttr *Sema::createClusterDimsAttr(const AttributeCommonInfo &CI, >From 9b346bf22df0c3548c46f4b9bb1d9b05ccfed387 Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Thu, 25 Sep 2025 19:29:02 -0400 Subject: [PATCH 06/10] fix another comment --- clang/include/clang/Basic/Attr.td | 4 ++-- .../test/Misc/pragma-attribute-supported-attributes-list.test | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 6ab76e8721434..7204bb4a9367a 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1575,14 +1575,14 @@ def HIPManaged : InheritableAttr { def CUDAClusterDims : InheritableAttr { let Spellings = [GNU<"cluster_dims">, Declspec<"__cluster_dims__">]; let Args = [ExprArgument<"X">, ExprArgument<"Y", 1>, ExprArgument<"Z", 1>]; - let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">; + let Subjects = SubjectList<[ObjCMethod, FunctionLike]>; let LangOpts = [CUDA]; let Documentation = [CUDAClusterDimsAttrDoc]; } def CUDANoCluster : InheritableAttr { let Spellings = [GNU<"no_cluster">, Declspec<"__no_cluster__">]; - let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">; + let Subjects = SubjectList<[ObjCMethod, FunctionLike]>; let LangOpts = [CUDA]; let Documentation = [CUDANoClusterAttrDoc]; } diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index 693d54159804e..ab4153a64f028 100644 --- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test @@ -35,7 +35,7 @@ // CHECK-NEXT: CFUnknownTransfer (SubjectMatchRule_function) // CHECK-NEXT: CPUDispatch (SubjectMatchRule_function) // CHECK-NEXT: CPUSpecific (SubjectMatchRule_function) -// CHECK-NEXT: CUDAClusterDims (SubjectMatchRule_function) +// CHECK-NEXT: CUDAClusterDims (SubjectMatchRule_objc_method, SubjectMatchRule_hasType_functionType) // CHECK-NEXT: CUDAConstant (SubjectMatchRule_variable) // CHECK-NEXT: CUDADevice (SubjectMatchRule_function, SubjectMatchRule_variable) // CHECK-NEXT: CUDADeviceBuiltinSurfaceType (SubjectMatchRule_record) @@ -44,7 +44,7 @@ // CHECK-NEXT: CUDAGridConstant (SubjectMatchRule_variable_is_parameter) // CHECK-NEXT: CUDAHost (SubjectMatchRule_function) // CHECK-NEXT: CUDALaunchBounds (SubjectMatchRule_objc_method, SubjectMatchRule_hasType_functionType) -// CHECK-NEXT: CUDANoCluster (SubjectMatchRule_function) +// CHECK-NEXT: CUDANoCluster (SubjectMatchRule_objc_method, SubjectMatchRule_hasType_functionType) // CHECK-NEXT: CUDAShared (SubjectMatchRule_variable) // CHECK-NEXT: CXX11NoReturn (SubjectMatchRule_function) // CHECK-NEXT: CallableWhen (SubjectMatchRule_function_is_member) >From 646ec7f7f81c919bacdddc11716bf80dfdb0324f Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Fri, 26 Sep 2025 00:41:39 -0400 Subject: [PATCH 07/10] more information on no_dims --- clang/include/clang/Basic/AttrDocs.td | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index fcc4f5cbdefd3..2cbdc394d9fb6 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -7553,10 +7553,14 @@ The ``__cluster_dims__`` and `__no_cluster__`` attributes are mutually exclusive def CUDANoClusterAttrDoc : Documentation { let Category = DocCatDecl; let Content = [{ -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. +In CUDA/HIP programming, a kernel function can still be launched with the cluster feature +enabled at runtime, even without the ``__cluster_dims__`` attribute. The LLVM/Clang-exclusive +``__no_cluster__`` attribute can be applied to a kernel function to explicitly indicate that +the cluster feature will not be enabled either at compile time or at kernel launch time. This +allows the compiler to apply certain optimizations without assuming that clustering could be +enabled at runtime. It is undefined behavior to launch a kernel with the ``__no_cluster__`` +attribute if the cluster feature is enabled at runtime. The ``__cluster_dims__`` and +``__no_cluster__`` attributes are mutually exclusive. .. code:: >From 0208c0ccbd36de284fac14ef9057633b04851187 Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Tue, 30 Sep 2025 14:37:50 -0400 Subject: [PATCH 08/10] drop `__` in `Declspec` --- clang/include/clang/Basic/Attr.td | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 7204bb4a9367a..fc9c8a0f2cdbc 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1573,7 +1573,7 @@ def HIPManaged : InheritableAttr { } def CUDAClusterDims : InheritableAttr { - let Spellings = [GNU<"cluster_dims">, Declspec<"__cluster_dims__">]; + let Spellings = [GNU<"cluster_dims">, Declspec<"cluster_dims">]; let Args = [ExprArgument<"X">, ExprArgument<"Y", 1>, ExprArgument<"Z", 1>]; let Subjects = SubjectList<[ObjCMethod, FunctionLike]>; let LangOpts = [CUDA]; @@ -1581,7 +1581,7 @@ def CUDAClusterDims : InheritableAttr { } def CUDANoCluster : InheritableAttr { - let Spellings = [GNU<"no_cluster">, Declspec<"__no_cluster__">]; + let Spellings = [GNU<"no_cluster">, Declspec<"no_cluster">]; let Subjects = SubjectList<[ObjCMethod, FunctionLike]>; let LangOpts = [CUDA]; let Documentation = [CUDANoClusterAttrDoc]; >From e1e11580c43e227965ae6086bca9ec99e0434303 Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Tue, 30 Sep 2025 17:59:55 -0400 Subject: [PATCH 09/10] fix comments --- clang/include/clang/Basic/Attr.td | 2 +- clang/lib/CodeGen/Targets/AMDGPU.cpp | 19 +++++++------------ .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 6 +++--- 3 files changed, 11 insertions(+), 16 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index fc9c8a0f2cdbc..42583421ca9bc 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1574,7 +1574,7 @@ def HIPManaged : InheritableAttr { def CUDAClusterDims : InheritableAttr { let Spellings = [GNU<"cluster_dims">, Declspec<"cluster_dims">]; - let Args = [ExprArgument<"X">, ExprArgument<"Y", 1>, ExprArgument<"Z", 1>]; + let Args = [ExprArgument<"X">, ExprArgument<"Y", /*opt=*/1>, ExprArgument<"Z", /*opt=*/1>]; let Subjects = SubjectList<[ObjCMethod, FunctionLike]>; let LangOpts = [CUDA]; let Documentation = [CUDAClusterDimsAttrDoc]; diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index 48855ce485f91..ee4d2aa660269 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -407,20 +407,15 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes( } if (auto *Attr = FD->getAttr<CUDAClusterDimsAttr>()) { - uint32_t X = - Attr->getX()->EvaluateKnownConstInt(M.getContext()).getExtValue(); - uint32_t Y = - Attr->getY() - ? Attr->getY()->EvaluateKnownConstInt(M.getContext()).getExtValue() - : 1; - uint32_t Z = - Attr->getZ() - ? Attr->getZ()->EvaluateKnownConstInt(M.getContext()).getExtValue() - : 1; - + auto GetExprVal = [&](const auto &E) { + return E ? E->EvaluateKnownConstInt(M.getContext()).getExtValue() : 1; + }; + unsigned X = GetExprVal(Attr->getX()); + unsigned Y = GetExprVal(Attr->getY()); + unsigned Z = GetExprVal(Attr->getZ()); llvm::SmallString<32> AttrVal; llvm::raw_svector_ostream OS(AttrVal); - OS << X << ',' << Y << ',' << Z; + OS << X << ", " << Y << ", " << Z; F->addFnAttr("amdgpu-cluster-dims", AttrVal.str()); } diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 081a720916fef..7505bdd64edb6 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -720,19 +720,19 @@ static void instantiateDependentCUDAClusterDimsAttr( if (Attr.getX()) { ExprResult ResultX = S.SubstExpr(Attr.getX(), TemplateArgs); if (ResultX.isUsable()) - XExpr = ResultX.getAs<Expr>(); + XExpr = ResultX.get(); } if (Attr.getY()) { ExprResult ResultY = S.SubstExpr(Attr.getY(), TemplateArgs); if (ResultY.isUsable()) - YExpr = ResultY.getAs<Expr>(); + YExpr = ResultY.get(); } if (Attr.getZ()) { ExprResult ResultZ = S.SubstExpr(Attr.getZ(), TemplateArgs); if (ResultZ.isUsable()) - ZExpr = ResultZ.getAs<Expr>(); + ZExpr = ResultZ.get(); } if (XExpr) >From 4f0237379d3403fb5c1ba8c6c0dd044f8e6cc509 Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Mon, 6 Oct 2025 13:53:51 -0400 Subject: [PATCH 10/10] refine target feature lookup; fix comments --- clang/include/clang/Basic/Attr.td | 4 ++-- clang/include/clang/Basic/AttrDocs.td | 7 +++--- .../clang/Basic/DiagnosticSemaKinds.td | 2 +- clang/lib/CodeGen/Targets/AMDGPU.cpp | 8 ++++--- clang/lib/Sema/SemaDeclAttr.cpp | 24 +++++++++++-------- .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 3 +-- clang/test/SemaCUDA/cluster_dims.cu | 8 +++---- 7 files changed, 31 insertions(+), 25 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 42583421ca9bc..212c209317748 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1573,7 +1573,7 @@ def HIPManaged : InheritableAttr { } def CUDAClusterDims : InheritableAttr { - let Spellings = [GNU<"cluster_dims">, Declspec<"cluster_dims">]; + let Spellings = [GNU<"cluster_dims">]; let Args = [ExprArgument<"X">, ExprArgument<"Y", /*opt=*/1>, ExprArgument<"Z", /*opt=*/1>]; let Subjects = SubjectList<[ObjCMethod, FunctionLike]>; let LangOpts = [CUDA]; @@ -1581,7 +1581,7 @@ def CUDAClusterDims : InheritableAttr { } def CUDANoCluster : InheritableAttr { - let Spellings = [GNU<"no_cluster">, Declspec<"no_cluster">]; + let Spellings = [GNU<"no_cluster">]; let Subjects = SubjectList<[ObjCMethod, FunctionLike]>; let LangOpts = [CUDA]; let Documentation = [CUDANoClusterAttrDoc]; diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 2cbdc394d9fb6..de9e84dd8fa39 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -7537,9 +7537,10 @@ 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. ``__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. +to set the dimensions of a thread block cluster, which is an optional level of hierarchy and made +up of thread blocks. ``__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:: diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index ce37075cb6184..0491580771911 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -13072,7 +13072,7 @@ def err_cuda_cluster_attr_not_supported : Error< >; def err_cuda_cluster_dims_too_large : Error< - "only a maximum of %0 thread blocks in a cluster is supported" + "cluster does not support more than %0 thread blocks; %1 provided" >; // VTable pointer authentication errors diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index ee4d2aa660269..bb14d4602894b 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -415,13 +415,15 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes( unsigned Z = GetExprVal(Attr->getZ()); llvm::SmallString<32> AttrVal; llvm::raw_svector_ostream OS(AttrVal); - OS << X << ", " << Y << ", " << Z; + OS << X << ',' << Y << ',' << Z; F->addFnAttr("amdgpu-cluster-dims", AttrVal.str()); } // OpenCL doesn't support cluster feature. - if ((IsOpenCLKernel && TargetFetureMap.lookup("gfx1250-insts")) || - FD->getAttr<CUDANoClusterAttr>()) + const TargetInfo &TTI = M.getContext().getTargetInfo(); + if ((IsOpenCLKernel && + TTI.hasFeatureEnabled(TTI.getTargetOpts().FeatureMap, "clusters")) || + FD->hasAttr<CUDANoClusterAttr>()) F->addFnAttr("amdgpu-cluster-dims", "0,0,0"); } diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index e3cd74964a25b..11ae42dd399ca 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5740,9 +5740,10 @@ CUDAClusterDimsAttr *Sema::createClusterDimsAttr(const AttributeCommonInfo &CI, } int FlatDim = ValX * ValY * ValZ; - auto TT = (!Context.getLangOpts().CUDAIsDevice && Context.getAuxTargetInfo()) - ? Context.getAuxTargetInfo()->getTriple() - : Context.getTargetInfo().getTriple(); + const llvm::Triple TT = + (!Context.getLangOpts().CUDAIsDevice && Context.getAuxTargetInfo()) + ? Context.getAuxTargetInfo()->getTriple() + : Context.getTargetInfo().getTriple(); int MaxDim = 1; if (TT.isNVPTX()) MaxDim = 8; @@ -5754,7 +5755,8 @@ CUDAClusterDimsAttr *Sema::createClusterDimsAttr(const AttributeCommonInfo &CI, // A maximum of 8 thread blocks in a cluster is supported as a portable // cluster size in CUDA. The number is 16 for AMDGPU. if (FlatDim > MaxDim) { - Diag(CI.getLoc(), diag::err_cuda_cluster_dims_too_large) << MaxDim; + Diag(CI.getLoc(), diag::err_cuda_cluster_dims_too_large) + << MaxDim << FlatDim; return nullptr; } @@ -5772,10 +5774,11 @@ void Sema::addNoClusterAttr(Decl *D, const AttributeCommonInfo &CI) { } static void handleClusterDimsAttr(Sema &S, Decl *D, const ParsedAttr &AL) { - auto &TTI = S.Context.getTargetInfo(); - auto Arch = StringToOffloadArch(TTI.getTargetOpts().CPU); + const TargetInfo &TTI = S.Context.getTargetInfo(); + OffloadArch Arch = StringToOffloadArch(TTI.getTargetOpts().CPU); if ((TTI.getTriple().isNVPTX() && Arch < clang::OffloadArch::SM_90) || - (TTI.getTriple().isAMDGPU() && Arch < clang::OffloadArch::GFX1250)) { + (TTI.getTriple().isAMDGPU() && + !TTI.hasFeatureEnabled(TTI.getTargetOpts().FeatureMap, "clusters"))) { S.Diag(AL.getLoc(), diag::err_cuda_cluster_attr_not_supported) << "__cluster_dims__"; return; @@ -5791,10 +5794,11 @@ static void handleClusterDimsAttr(Sema &S, Decl *D, const ParsedAttr &AL) { } static void handleNoClusterAttr(Sema &S, Decl *D, const ParsedAttr &AL) { - auto &TTI = S.Context.getTargetInfo(); - auto Arch = StringToOffloadArch(TTI.getTargetOpts().CPU); + const TargetInfo &TTI = S.Context.getTargetInfo(); + OffloadArch Arch = StringToOffloadArch(TTI.getTargetOpts().CPU); if ((TTI.getTriple().isNVPTX() && Arch < clang::OffloadArch::SM_90) || - (TTI.getTriple().isAMDGPU() && Arch < clang::OffloadArch::GFX1250)) { + (TTI.getTriple().isAMDGPU() && + !TTI.hasFeatureEnabled(TTI.getTargetOpts().FeatureMap, "clusters"))) { S.Diag(AL.getLoc(), diag::err_cuda_cluster_attr_not_supported) << "__no_cluster__"; return; diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 7505bdd64edb6..75aae2ca180ca 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -735,8 +735,7 @@ static void instantiateDependentCUDAClusterDimsAttr( ZExpr = ResultZ.get(); } - if (XExpr) - S.addClusterDimsAttr(New, Attr, XExpr, YExpr, ZExpr); + S.addClusterDimsAttr(New, Attr, XExpr, YExpr, ZExpr); } // This doesn't take any template parameters, but we have a custom action that diff --git a/clang/test/SemaCUDA/cluster_dims.cu b/clang/test/SemaCUDA/cluster_dims.cu index 3cd0e0197c29b..a4797ef007eaa 100644 --- a/clang/test/SemaCUDA/cluster_dims.cu +++ b/clang/test/SemaCUDA/cluster_dims.cu @@ -29,12 +29,12 @@ template <int x, int y, int z> void test_template_expr(void) __cluster_dims__(x //NS-error@+1 {{__cluster_dims__ is not supported for this GPU architecture}} __global__ void __cluster_dims__(32, 2, 4) test_too_large_dim_0() {} // common-error {{integer constant expression evaluates to value 32 that cannot be represented in a 4-bit unsigned integer type}} -// cuda-error@+2 {{only a maximum of 8 thread blocks in a cluster is supported}} -// amd-error@+1 {{only a maximum of 16 thread blocks in a cluster is supported}} +// cuda-error@+2 {{cluster does not support more than 8 thread blocks; 64 provided}} +// amd-error@+1 {{cluster does not support more than 16 thread blocks; 64 provided}} __global__ void __cluster_dims__(4, 4, 4) test_too_large_dim_1() {} // NS-error {{__cluster_dims__ is not supported for this GPU architecture}} -// cuda-error@+3 {{only a maximum of 8 thread blocks in a cluster is supported}} -// amd-error@+2 {{only a maximum of 16 thread blocks in a cluster is supported}} +// cuda-error@+3 {{cluster does not support more than 8 thread blocks; 64 provided}} +// amd-error@+2 {{cluster does not support more than 16 thread blocks; 64 provided}} template<unsigned a, unsigned b, unsigned c> __global__ void __cluster_dims__(a, b, c) test_too_large_dim_template() {} // NS-error {{__cluster_dims__ is not supported for this GPU architecture}} template __global__ void test_too_large_dim_template<4, 4, 4>(); // common-note {{in instantiation of function template specialization 'test_too_large_dim_template<4U, 4U, 4U>' requested here}} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
