================
@@ -4055,6 +4055,115 @@ template<typename T> constexpr T *addressof(T &value) {
}
```
+### `__addrspaceof`
+
+`__addrspaceof` returns a Clang address-space identifier for a type, object, or
+lvalue expression.
+
+This is useful for performance-sensitive code that needs a compile-time value
+to choose an address-space-specific operation, overload, or template
+specialization. For example, GPU code may select different loads for global,
+shared/local, or constant memory.
+
+**Syntax**:
+
+```c++
+int __addrspaceof(type-or-expression)
+```
+
+The parentheses are required. The operand is not evaluated. The result is an
+integer constant expression.
+Reference types are adjusted to their referenced type before the address space
+is queried.
+
+For a type operand, the operator returns its top-level explicit address space,
+or the default address space if it has none.
+
+For an unparenthesized id-expression or unparenthesized member access,
+the operator queries the directly named variable or data member. It first uses
+the top-level explicit address space of the entity's type. For a CUDA/HIP
+entity without a top-level explicit address space, it uses the storage address
+space specified by `__device__`, `__shared__`, or `__constant__`. Otherwise, it
+returns the default address space. An explicit `__device__` `const` global or
+static data member that is promoted to constant memory is reported as the CUDA
+or HIP constant address space.
+
+Other expression operands must be lvalues. The operator returns the top-level
+explicit address space of the expression's type, or the default address space
+if it has none. Prvalue operands, including address-of expressions, are
+ill-formed. As with `decltype`, extra parentheses can change a direct entity
+query into an expression query.
+
+```c++
+__device__ int *p;
+
+static_assert(__addrspaceof(p) == __ADDRSPACE_GLOBAL);
+static_assert(__addrspaceof(*p) == __ADDRSPACE_DEFAULT);
+static_assert(__addrspaceof((p)) == __ADDRSPACE_DEFAULT);
+```
+
+The first query reports where the pointer variable is stored. The second
+reports the address space of the pointed-to object type. The third treats `p`
+as an lvalue expression because of the extra parentheses.
+
+Clang predefines macros for the values returned by this operator. Most language
+address spaces are reported using common memory regions:
+`__ADDRSPACE_DEFAULT`, `__ADDRSPACE_GLOBAL`, `__ADDRSPACE_LOCAL`,
+`__ADDRSPACE_CONSTANT`, `__ADDRSPACE_PRIVATE`, and `__ADDRSPACE_GENERIC`.
+OpenCL, CUDA, HIP, SYCL, and HLSL address spaces that describe the same memory
+region produce the same value. For an address space written explicitly with
+`__attribute__((address_space(N)))`, the result is
+`__ADDRSPACE_TARGET_OFFSET + N`. Deprecated OpenCL and SYCL `global_device` and
----------------
AaronBallman wrote:
Would it be more user-friendly for `__ADDRSPACE_TARGET_OFFSET` to instead be
`__ADDRSPACE_TARGET(OFFSET)` as a function-like macro that does the math for
the user?
https://github.com/llvm/llvm-project/pull/210242
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
