[clang] [Clang][HIP][CUDA] Add `__cluster_dims__` and `__no_cluster__` attribute (PR #156686)
Erich Keane via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 16 18:51:22 PDT 2025
================
@@ -7545,6 +7545,45 @@ 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, conventionally exposed as
+``__cluster_dims__`` macro, can be applied to a kernel function 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::
+
+ __global__ __cluster_dims__(2, 1, 1) void kernel(...) {
+ ...
+ }
+
+ }];
+}
+
+def CUDANoClusterAttrDoc : Documentation {
+ let Category = DocCatDecl;
+ let Content = [{
+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
----------------
erichkeane wrote:
Either `cluster_dims` here, or find some way to not say 'attribute? Perhaps, 'even without marking it `__cluster_dims__`'?
https://github.com/llvm/llvm-project/pull/156686
More information about the cfe-commits
mailing list