================
@@ -484,6 +484,88 @@ non-constexpr function, which is by default a host
function.
Users can override the inferred host and device attributes of default
destructors by adding explicit host and device attributes to them.
+Deferred Diagnostics
+====================
+
+In HIP (and CUDA), a ``__host__ __device__`` function can be called from both
+host and device code. Certain operations are not allowed on the device (e.g.,
+calling a host-only function, using variable-length arrays, or throwing
+exceptions). However, a ``__host__ __device__`` function containing such
+operations is only ill-formed if it is actually called from device code.
+
+Clang handles this through *deferred diagnostics*: errors and warnings in
+``__host__ __device__`` functions are recorded during parsing but not emitted
+immediately. They are only emitted if the function turns out to be reachable
+from code that must run on the device.
+
+Device-Promoted Functions
+-------------------------
+
+A *device-promoted function* is a function that is not explicitly restricted to
+device context (it is either ``__host__ __device__`` or, in the case of
+lambdas, implicitly ``__host__ __device__``) but is used from device code,
+forcing it to be compiled for the device. Device-promoted functions are the
+primary source of deferred diagnostics.
+
+Common examples of device-promoted functions:
+
+- Lambdas without explicit ``__host__`` or ``__device__`` attributes
+- ``__host__ __device__`` functions that call host-only functions
+- ``inline __host__ __device__`` helper functions used from device code
+
+When a device-promoted function contains operations that are not valid on the
+device, clang emits the deferred diagnostics along with notes showing how the
+function was reached from device code.
+
+Example
+^^^^^^^
+
+.. code-block:: c++
+
+ __host__ void host_only();
+
+ // This lambda is implicitly __host__ __device__. It is device-promoted
+ // when called from a __device__ function.
+ __device__ auto lambda = [] {
+ host_only(); // error: only emitted if lambda is used from device code
+ };
+
+ __device__ void df1() {
+ lambda(); // triggers deferred diagnostic for lambda
+ }
+
+ __device__ void df2() {
+ lambda(); // same lambda, same error — not duplicated
+ }
+
+Clang emits the error once and lists all device callers:
+
+.. code-block:: text
+
+ error: reference to __host__ function 'host_only' in __host__ __device__
function
+ note: 'host_only' declared here
+ note: called by 'df1'
+ note: called by 'df2'
+
+Call Chain Notes
+^^^^^^^^^^^^^^^^
+
+When a device-promoted function is reached through a chain of intermediate
+functions, clang shows the full call chain. The first note in each chain uses
+"called by" and subsequent notes use "then called by":
+
+.. code-block:: text
+
+ error: reference to __host__ function 'host_only' in __host__ __device__
function
+ note: called by 'helper1'
+ note: then called by 'device_func1'
+ note: called by 'helper2'
+ note: then called by 'device_func2'
----------------
Artem-B wrote:
What's the call chain in this example? I assume it's
`device_func1->helper1->host_only()`.
If that's the case, then "then called by device_func1()" looks odd, because
it's the very first call we make, but "then called by" implies that the call
happened *after* helper1(). Replace "then called by" with "which was called by"
? Or use non-verbal cue to indicate the order. E.g.
```
note: called by 'helper1'
note: + called by 'device_func1'
```
https://github.com/llvm/llvm-project/pull/185926
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
