| Issue |
180638
|
| Summary |
Deferred diagnostics issued repeatedly with misleading notes attached
|
| Labels |
cuda,
clang:diagnostics,
SYCL
|
| Assignees |
|
| Reporter |
tahonermann
|
The following example is ill-formed because CUDA does not allow code emitted during device compilation to call a host-only function. `hf()` is a host-only function. The lambda in the definition of `l` is a host-device function that may be called by code emitted for either host or device compilation; the lambda erroneously calls a host-only function, but that is only an error if the function is actually called in a device context. `df1()` and `df2()` are both device-only functions.
https://godbolt.org/z/6qGnjrjqY
```
$ cat t.cpp
// Macros locally defined to avoid needing a local CUDA toolkit installation.
#define __global__ __attribute__((global))
#define __host__ __attribute__((host))
#define __device__ __attribute__((device))
__host__ void hf();
__device__ auto l =
[] {
hf(); // #1
hf(); // #2
};
__device__ void df1() {
l();
}
__device__ void df2() {
l();
}
```
Clang produces the following diagnostics. The `-nocudainc` and `-nocudalib` options are used to avoid a dependency on a local CUDA toolkit installation.
```
$ clang -fsyntax-only -nocudainc -nocudalib t.cu
t.cu:8:5: error: reference to __host__ function 'hf' in __host__ __device__ function
8 | hf(); // #1
| ^
t.cu:12:4: note: called by 'df1'
12 | l();
| ^
t.cu:5:15: note: 'hf' declared here
5 | __host__ void hf();
| ^
t.cu:9:5: error: reference to __host__ function 'hf' in __host__ __device__ function
9 | hf(); // #2
| ^
t.cu:5:15: note: 'hf' declared here
5 | __host__ void hf();
| ^
t.cu:8:5: error: reference to __host__ function 'hf' in __host__ __device__ function
8 | hf(); // #1
| ^
t.cu:15:4: note: called by 'df2'
15 | l();
| ^
t.cu:5:15: note: 'hf' declared here
5 | __host__ void hf();
| ^
t.cu:9:5: error: reference to __host__ function 'hf' in __host__ __device__ function
9 | hf(); // #2
| ^
t.cu:5:15: note: 'hf' declared here
5 | __host__ void hf();
| ^
4 errors generated when compiling for sm_52.
```
The diagnostics have the following issues:
1. Diagnostics related to the calls to `hf()` at the locations annotated as `#1` and `#2` are issued multiple times.
2. The "called by" notes are misleading. Consider the "called by 'df1'" case. The error concerns an attempted call to a function named `hf()`. The note suggests the call was made by `df()` and then presents a source line showing a call to `l()`. The note is intended to indicate that `l()` (the lambda) was used as a device function due to it's use in `df1()`. This is useful information, but the presentation is confusing.
3. The "called by" notes are only issued for the first deferred error/warning diagnostic that is issued for a function. Subsequent diagnostics are not noted as being related to use of the function in a device context.
The user experience could be improved by doing one or more of the following:
1. Instead of attaching notes to deferred diagnostics to note the device context, issue distinct diagnostics for each use of a function with deferred diagnostics.
2. Attaching notes that more closely match those used in code synthesis contexts:
- "in instantiation of ... requested here"
- "in implicit initialization of ..."
- "in implicit call to ..."
- "in evaluation of ... needed here"
- "in call to ..."
3. Only issuing deferred diagnostics in the context of the first call from a device context.
4. Only issuing the first deferred diagnostic similar to the diagnostic behavior for errors in template instantiation contexts; see https://godbolt.org/z/cGYEMzTWj.
For example, the follow demonstrates what implementation of options 1 and 3 might look like. The term "device-promoted" is intended to indicate a function that is not declaratively restricted to host or device context but is used in a device context (for CUDA, I believe this only applies to lambdas).
```
t.cu:12:4: error: use of device-promoted function 'lambda at ...' with errors
12 | l();
| ^
t.cu:8:5: error: reference to __host__ function 'hf' in __host__ __device__ function
8 | hf(); // #1
| ^
t.cu:5:15: note: 'hf' declared here
5 | __host__ void hf();
| ^
t.cu:9:5: error: reference to __host__ function 'hf' in __host__ __device__ function
9 | hf(); // #2
| ^
t.cu:5:15: note: 'hf' declared here
5 | __host__ void hf();
| ^
t.cu:15:4: error: use of device-promoted function 'lambda at ...' with errors
15 | l();
| ^
```
These issues were observed while attempting to use deferred diagnostics to implement some SYCL related diagnostics. These issues are not specific to CUDA; they are issues with the deferred diagnostic support.
@yxsamliu and @rjmccall for awareness; the current behavior was implemented in commit 2c31aa2de13a23a00ced87123b92e905f2929c7b for https://reviews.llvm.org/D77028.
![]()
_______________________________________________
llvm-bugs mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
