Copy capture rules for [=, *this]


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


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.

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]

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.

will do thanks

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, it would help with figuring out what’s going on in your case.

1 Like