[clang] [clang][SYCL] Diagnose reference kernel arguments (PR #192957)

Tom Honermann via cfe-commits cfe-commits at lists.llvm.org
Mon May 4 10:51:45 PDT 2026


tahonermann wrote:

I've been thinking more about the use cases for this visitor and wondering if visiting the `CXXBaseSpecifier` and `FieldDecl` elements of the type suffices for the first use case I listed:
- Identifying subobject types that require special handling for kernel marshaling purposes.

That use case will require producing expressions that access the subobjects of an object. For example, given the following (weird) case involving three subobjects of some SYCL special type:
```
struct [[clang::sycl_special_kernel_parameter]] sycl_special_type {
};
struct base {
  sycl_special_type indirect_sst;
};
struct kernel_type: base, sycl_special_type {
  sycl_special_type direct_sst;
};
void f(sycl::handler &h, kernel_type k) {
  h.single_task<struct KN>(k);
}
```

The eventual call to a `[[clang::sycl_kernel_entry_point]]` attributed function will require synthesizing a call to `sycl_kernel_launch()` that passes the `kernel_type` object, and a call to the result object that passes references to the three `sycl_special_type` subobjects. The call sequence would look something like the following where the three `sycl_special_type` references are passed in depth first order.
```
  sycl_kernel_launch(k)(k.base::indirect_sst, static_cast<sycl_special_type&>(k), k.direct_sst);
```

There may, of course, be multiple base classes to transit to reach a subobject of interest. The question is, given an expression for an lvalue of the kernel object `k`, how straight forward is it to produce the base class or data member accesses using `CXXBaseSpecifier` and `FieldDecl`? I suspect this is relatively straight forward, but it might be worth prototyping or maybe even including a derived visitor/transform class that visits such (created) expressions based on a predicate.

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


More information about the cfe-commits mailing list