================
@@ -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
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
