`::cuda::device::is_object_from(..., ::cuda::device::address_space::cluster_shared)`may be wrong on < SM90

[bernhardmgruber]

Any address in the shared state space is also part of the cluster shared state space. That is, the cluster shared state space is a superset of the shared state space.

Consider the following program:

#include <cuda/memory>

__global__ void k() {
    __shared__ int i;
    _CCCL_VERIFY(::cuda::device::is_object_from(i, ::cuda::device::address_space::cluster_shared), "");
}

int main() {
    k<<<1, 1>>>();
}

This runs successfully on SM90+, but will fail below SM90, because ::cuda::device::is_object_from(..., ::cuda::device::address_space::cluster_shared) just returns false below SM90. I believe this limits the usability of this function and we should just downgrade the check from cluster_shared to shared below SM90.

[fbusato]

I disagree. cuda::device::address_space::cluster_shared should match the result of __isClusterShared. We need a clear distinction between cluster (remote) shared memory and local shared memory.

[bernhardmgruber]

cuda::device::address_space::cluster_shared should match the result of __isClusterShared. We need a clear distinction between cluster (remote) shared memory and local shared memory.

The public PTX ISA documents pretty clearly that the CTA shared memory is part of the cluster shared memory (emphasis mine): https://docs.nvidia.com/cuda/parallel-thread-execution/#shared-state-space

The shared (.shared) state space is a memory that is owned by an executing CTA and is accessible to the threads of all the CTAs within a cluster. An address in shared memory can be read and written by any thread in a CTA cluster.

Additional sub-qualifiers ::cta or ::cluster can be specified on instructions with .shared state space to indicate whether the address belongs to the shared memory window of the executing CTA or of any CTA in the cluster respectively. The addresses in the .shared::cta window also fall within the .shared::cluster window.

From reading through NVBug 4275174 and the referred to diffs, __isClusterShared:

Returns 1 if ptr contains the generic address of an object in shared memory space within the same thread block cluster, otherwise returns 0.

This is also what anyone can deduce from its use in our barrier_block_scope.h, see e.g. this PR: https://github.com/NVIDIA/cccl/pull/4866/files#diff-63c3b9cfc510e8842725a740cf49e29110513e01fd2d35f3f0c9d8de420e2d09. The logic in all the assertions (which we are shipping for years) is that __isClusterShared includes every address that __isShared. Fortunately, __isClusterShared only appears on code paths guarded to be run on at least SM90, so we are side stepping the problem of this issue in current production code. Or at least, that's what we think.

We need a clear distinction between cluster (remote) shared memory and local shared memory.

We have that today. A cluster remote shared memory address (aka. distributed SMEM) is one for which __isClusterShared(p) && !__isShared(p).

Now to the topic of this issue: I would argue that

__shared__ int i;
__isClusterShared(&i);

should give the same answer before and after SM90. It does not do this. Since it's an internal, undocumented compiler built-in, this is fine. But we can do better on our public documented API and make:

__shared__ int i;
::cuda::device::is_object_from(i, ::cuda::device::address_space::cluster_shared);

consistent across SM versions.

Be aware that this is not a hill I am willing to die on. If you argue that this is nonsensical, because there is no cluster shared <SM90, I will also accept that and we can close the issue.

[fbusato]

NVBug 4275174 clarifies a lot of issues. It is OK in this case. We probably need a way to clearly distinguish between local and remove shared memory in the future. @gonzalobg for visibility

[gonzalobg]

AFAICT __isClusterShared is not a public API (e.g. we intentionally don't document anywhere what it does).
If you want to build something on top of it, step 0 is probably to ask and get CUDA Runtime/Driver to specify what it does (there is a discussion in an nvbug about that which hasn't concluded yet).

[bernhardmgruber]

AFAICT __isClusterShared is not a public API [...] If you want to build something on top of it, step 0 is probably to ask and get CUDA Runtime/Driver to specify what it does.

The problem is that our cuda::barrier implementation has been using __isClusterShared for some years already, and when ::cuda::device::is_object_from was created, the uses of __isClusterShared were replaced by it. So we have been using it already for a while.

We recently got a bug about it though (nvbug 5810743) which @davebayer is trying to fix in #7259. But you are right, nobody specifies or guarantees what the intrinsic actually does. We only use it in assertions, so we could also just remove all uses of it.

That still leaves the question what the meaning of ::cuda::device::is_object_from(..., ::cuda::device::address_space::cluster_shared) is and how it can be implemented.