Copy capture rules for [=, *this]

Clang Frontend
,
1

Hi

I’m wondering if someone can point me at a resource that explains what what is allowable (or should work) for copy capturing a copy of *this on device lambdas

I have some simple things working but more complex things do not

Thanks

2

Which language are you asking about? CUDA, SYCL, OpenMP? Each language defines device copyability requirements for objects copied from host to device. Generally, there is a trivially copyable requirement (e.g., memcpy()) or a protocol.

3

Sorry yes it’s specifically C++17 cuda device lambdas

I have a situation where a simple class with some simple members (int/double) does get copied properly via [*this] but it fails at runtime with something more complex

it’s strange though if I expand the capture to just copy the actual variables contained within the (*this) it works, eg [a=this->a, b=this->b]

4

I wouldn’t be surprised if this is just a compiler bug. I checked existing bug reports, but didn’t find one that looked like a match. I suggest filing a new bug report at Issues · llvm/llvm-project · GitHub.

5

will do thanks

6

There’s nothing particularly special about lambdas in CUDA code on the language level (modulo giving them implicit __host__ __device__ attributes).

However, it also means that we do not do anything special about transferring lambda objects across CPU/GPU boundary. If the captured data contains a pointer, there’s a good chance that it will not work on the GPU.

If you can create a reduced reproducer for the issue on godbolt.org, it would help with figuring out what’s going on in your case.

1 Like