[hlsl-out] Invalid hlsl for write to scalar/vector storage buffer

[hasali19]

WGSL:

@group(0)
@binding(0)
var<storage, read_write> buf: vec2<i32>;

@compute
@workgroup_size(1)
fn main() {
    buf = vec2<i32>(1);
}

HLSL:

RWByteAddressBuffer buf : register(u0);

[numthreads(1, 1, 1)]
void main()
{
    buf.Store2(, asuint((1).xx));
    return;
}

The Store2 call seems to have a missing argument.

[hasali19]

Looks like this also affects scalars:

@group(0)
@binding(0)
var<storage, read_write> buf: i32;

@compute
@workgroup_size(1)
fn main() {
    buf = 0;
}

HLSL:

RWByteAddressBuffer buf : register(u0);

[numthreads(1, 1, 1)]
void main()
{
    buf.Store(, asuint(0));
    return;
}

[teoxoy]

The WGSL spec and also naga used to only allow structs for global variables in uniform and storage address spaces.
Following the spec we recently relaxed this requirement (in #1682) but it looks like the Writer::write_storage_address method still assumes there is a non empty chain of accesses.

I think to get this working we just need to write a 0 if the chain is empty here

naga/src/back/hlsl/storage.rs

Lines 41 to 62 in b3d5e6d

	pub(super) fn write_storage_address(
	&mut self,
	module: &crate::Module,
	chain: &[SubAccess],
	func_ctx: &FunctionCtx,
	) -> BackendResult {
	for (i, access) in chain.iter().enumerate() {
	if i != 0 {
	write!(self.out, "+")?;
	}
	match *access {
	SubAccess::Offset(offset) => {
	write!(self.out, "{}", offset)?;
	}
	SubAccess::Index { value, stride } => {
	self.write_expr(module, value, func_ctx)?;
	write!(self.out, "*{}", stride)?;
	}
	}
	}
	Ok(())
	}