[clang] [CUDA][HIP] Fix host/device context in concept (PR #67721)

Yaxun Liu via cfe-commits cfe-commits at lists.llvm.org
Fri Oct 6 07:23:51 PDT 2023


================
@@ -176,3 +176,34 @@ Predefined Macros
    * - ``HIP_API_PER_THREAD_DEFAULT_STREAM``
      - Alias to ``__HIP_API_PER_THREAD_DEFAULT_STREAM__``. Deprecated.
 
+C++20 Concepts with HIP and CUDA
+--------------------------------
+
+In Clang, when working with HIP or CUDA, it's important to note that all constraints in C++20 concepts are assumed to be for the host side only. This behavior is consistent across both programming models, and developers should be aware of this assumption when writing code that utilizes C++20 concepts.
+
+Example:
+.. code-block:: c++
+
+   template <class T>
+   concept MyConcept = requires(T& obj) {
+     my_function(obj);  // Assumed to be a host-side requirement
----------------
yxsamliu wrote:

currently, we do not support that. 

I would suggest adding an extension to the clang that allows `__host__` and `__device__` attributes on call expressions in concept definition to indicate the required callability for the host or device.

For example,

```
template <class T>
   concept MyConcept = requires(T& obj) {
     __device__ my_function(obj);  // requires my_function(obj) callable on device side
}
```

https://github.com/llvm/llvm-project/pull/67721


More information about the cfe-commits mailing list