Issue with cooperative_groups::memcpy_async

My suggestion would be to file a bug. It’s not obvious to me that the kernel1 case/usage should have any limitations compared to the kernel0 case/usage, for the example you have given.

If you file a bug, I would expect either a technical issue to be surfaced, or else expose a need to improve the docs.

I could mention a few other things; you probably know them already:

  1. the underlying hardware accelerated transfer only applies to 4-byte quantities.
  2. the underlying hardware accelerated transfer only applies to cc 8.0, not 7.0
  3. the underlying hardware accelerated transfer only applies to the global->smem path

The net of that is for the case you have given, I think there should be no difference by using the kernel0 case as a WAR.