> and to get back good debug info because I needed to move data from VRAM to normal RAM. I was trying to learn at the time so efficiency wasn't really a factor

> and to get back good debug info because I needed to move data from VRAM to normal RAM. I was trying to learn at the time so efficiency wasn't really a factor I mean you're literally using the calvinball programming language - it's just a metafunctor in the language of macros, what's the big deal? literally my (sanitized) code from 10 years ago ```cuda __global__ void kernel_myKernel(int num_items, myTask_t * task_item, #if VALIDATION == 1 myDebug1_t * output_debug1_arr, myDebug2_t * output_debug2_arr, #endif myOutput_t * output_arr, int current_iteration) ``` those myDebug_t pointers are pointers to host memory, so transferring the data back to host memory is transparently managed by CUDA or Thrust, and you just have an #IF or #IFDEF block in the code that dumps intermediate state to the pointer if validation is enabled in the build. And then you can run whatever test suite, on the actual intermediate data that's happening inside your kernel, so you can test your invariants/assertions. Test-suite lite edition/just the part you need - but you're testing actual kernel invocations, not just theoretical. And you can go to the level of dumping intermediate state of prefix_sums/reductions (or generalized "intermediate work items") every warp/grid iteration if you want - it's debug mode. but granted this relies on the ability to have data dynamically piped between device and host memory spaces... which was a novel bit of syntactic sugar they added to the CUDA toolkit back in like 2012 (it was big around the Kepler era and it was a big feature on the Jetson TK1 too) so AMD might well have not had it. Or they might have said they had it but it just broke horribly if you used it. but this is literally like three lines of code, if you had just used NVIDIA instead of AMD, because the feature was there when you needed it. Technically can be accomplished by just allocating extra VRAM for the buffer and copying it back afterwards, but granted, more work there. What's the price of getting your work done? What's your time worth? That's always been the only moat of relevance that CUDA has had... and it's always been worth the expense.