GPU codegen issue with cooperative_groups and sync

On my branch matmul_paper_schedules (which currently just contains a few small tweaks to fork opts, this bug is unrelated but those tweaks are needed to get matmul scheduled the way I want it). Building the Juno matmul (I'm using a schedule that creates one thread for each element of the output matrix via n/32 * l/32 blocks and 32 * 32 threads per block) I end up getting errors out of nvcc on miranda that:

/tmp/.tmppAVmvH/matmul.cu(89): error: class "cooperative_groups::__v1::thread_block_tile<1024U, void>" has no member "sync"

/tmp/.tmppAVmvH/matmul.cu(108): error: class "cooperative_groups::__v1::thread_block_tile<1024U, void>" has no member "sync"

/tmp/.tmppAVmvH/matmul.cu(127): error: class "cooperative_groups::__v1::thread_block_tile<1024U, void>" has no member "sync"

NVCC version is 11.5.119 on miranda if that's relevant.

assigned to @rarbore2

So on my machine, with NVCC version 12.8, the example compiles fine. On miranda, I see the following error:

/usr/include/cooperative_groups.h(1510): error: static assertion failed with "Tiled partition with Size > 32 is supported only by cooperative_groups::experimental::tiled_partition available with experimental features enabled"
            detected during instantiation of "cooperative_groups::__v1::thread_block_tile<Size, ParentT> cooperative_groups::__v1::tiled_partition<Size,ParentT>(const ParentT &) [with Size=1024U, ParentT=cooperative_groups::__v1::thread_block]"

What are the benefits/drawbacks of the number of threads per block? It seems like this is saying we either need to enable experimental features or use less threads per block

Well it's a matter of capability - having at most 32 threads per block limits the size of tiles in general. We should really either update CUDA or enable experimental features (I don't actually know how to do this).

Probably resolved by !214 (merged)

closed