Book: say whether sync_threads() orders raw-pointer and ptx_asm! shared-memory accesses #1463

What problem does this solve?

The shared-memory chapter of the book explains the hardware rule: sync before reading what another thread wrote. It doesn't say what the compiler may move across thread::sync_threads(). For SharedArray indexing, the answer seems implied. It's less clear for:

1. raw-pointer accesses through SharedArray::as_mut_ptr() or DynamicSharedArray::get()
2. ld.shared / st.shared inside ptx_asm!: with the default (side-effecting) treatment, with clobber("memory"), and with options(register_only)
3. cp.async copies followed by cp.async.wait_group and a barrier

We couldn't tell from the docs whether LLVM or libNVVM may hoist or sink a shared access across the barrier in cases 1 and 2. So in speakrs, every shared load and store is a side-effecting ptx_asm!. That's safe, but it hides the accesses from the optimizer, and it's probably more conservative than needed.

Proposed behavior

A short section in "Shared memory and synchronization" (or the ptx_asm! entry in "Supported features") that states, for each access form above, whether sync_threads() orders it at compile time, and when clobber("memory") is required. A small table would do.

Alternatives considered

Making every shared access a side-effecting ptx_asm!, which is what we do now.

Additional context

Kernels: avencera/speakrs#36 (crates/speakrs-cuda-kernels). We used cuda-oxide 918bbde. I checked that the book on main (a51ef44d) still doesn't cover this.