fix(hlsl-out): use Interlocked<op> intrinsic for atomic integers

[ErichDonGubler]

Fixes #2284.

---

We currently assume that we are using raw RWByteAddressBuffer methods for all atomic operations (<pointer>.Interlocked<op>(<raw_byte_offset>, …)), which is only true when we use var<storage, read_write> globals. For var<workgroup> globals, we need Interlocked<op>(<pointer>, …), using the original expression as the first argument.

Fix this by branching on the pointer's address space in Atomic statements, and implementing the workgroup address space case with intrinsics.

[ErichDonGubler]

@jimblandy and I are planning on pair-reviewing this draft Soon™. 😊

[ErichDonGubler]

@jimblandy: I've pushed some significant updates to this during my afternoon yesterday. Most of the logic is the same, just renamed or shuffled around a bit.

[ErichDonGubler]

@jimblandy: Just pushed up newest changes branching on address space. TIL a bunch, thanks for the help! ❤️

[teoxoy]

The new test file looks much more complete than what we had before!

naga/tests/in/access.wgsl

Lines 154 to 169 in da8e911

	@compute @workgroup_size(1)
	fn atomics() {
	var tmp: i32;
	let value = atomicLoad(&bar.atom);
	tmp = atomicAdd(&bar.atom, 5);
	tmp = atomicSub(&bar.atom, 5);
	tmp = atomicAnd(&bar.atom, 5);
	tmp = atomicOr(&bar.atom, 5);
	tmp = atomicXor(&bar.atom, 5);
	tmp = atomicMin(&bar.atom, 5);
	tmp = atomicMax(&bar.atom, 5);
	tmp = atomicExchange(&bar.atom, 5);
	// https://github.com/gpuweb/gpuweb/issues/2021
	// tmp = atomicCompareExchangeWeak(&bar.atom, 5, 5);
	atomicStore(&bar.atom, value);
	}

I think we could remove this but the atomicExchange function seems to be missing from the new test file.

[ErichDonGubler]

@teoxoy:

The new test file looks much more complete than what we had before!

😊

I think we could remove this [(access.wgsl:154-169)] but the atomicExchange function seems to be missing from the new test file.

Agreed. Added coverage for atomicExchange with 340a698 and b14bd40, and removed the overlapping contents of access.wgsl with f8c7cb1 (and threw in e310121 for fun). Will autosquash momentarily so this branch's contents are still nice commits.

[jimblandy]

Looks good. Some minor cleanups needed.