[hlsl-out] Arrays in Structs Are Incorrectly Loaded When Loaded as A Whole Struct

[cwfitzgerald]

WGSL

struct Ah {
    inner: array<f32, 2>,
};
@group(0) @binding(0)
var<storage> ah: Ah;

@compute @workgroup_size(1)
fn cs_main() {
    _ = ah;
}

HLSL

struct Ah {
    float inner[2];
};

ByteAddressBuffer ah : register(t0);

Ah ConstructAh(float arg0[2]) {
    Ah ret = (Ah)0;
    ret.inner = arg0;
    return ret;
}

[numthreads(1, 1, 1)]
void cs_main()
{
    // ERROR: Unexpected {
    Ah unnamed = ConstructAh({asfloat(ah.Load(0+0)), asfloat(ah.Load(0+4))});
}

Solution

As far as I can tell, the only real solution is to make the array part of the initialization a temporary. I wish there was a way to initialize an array inline, but I don't think you can. I'm not entirely sure how to do it but this code is the code that needs to be assigned to a temporary.

naga/src/back/hlsl/storage.rs

Lines 145 to 156 in 2754a01

	crate::TypeInner::Array {
	base,
	size: crate::ArraySize::Constant(const_handle),
	..
	} => {
	write!(self.out, "{{")?;
	let count = module.constants[const_handle].to_array_length().unwrap();
	let stride = module.types[base].inner.size(&module.constants);
	let iter = (0..count).map(|i| (TypeResolution::Handle(base), stride * i));
	self.write_storage_load_sequence(module, var_handle, iter, func_ctx)?;
	write!(self.out, "}}")?;
	}

.

[numthreads(1, 1, 1)]
void cs_main()
{
    float temp[2] = {asfloat(ah.Load(0+0)), asfloat(ah.Load(0+4))};
    Ah unnamed = ConstructAh(temp);
}

[teoxoy]

We can inject a function to do inline init of arrays (we have this mechanism already).

[ErichDonGubler]

I'll take a look at this in the next few days weeks. EDIT: duration

[teoxoy]

The mechanism for this is already in place but we only use it for structs right now.

naga/src/back/hlsl/help.rs

Lines 876 to 881 in 8a72b7a

	let constructor = WrappedConstructor { ty };
	if !writer.wrapped.constructors.contains(&constructor) {
	writer
	.write_wrapped_constructor_function(module, constructor)?;
	writer.wrapped.constructors.insert(constructor);
	}

this is where we detect and inject the function that constructs the type

naga/src/back/hlsl/storage.rs

Lines 158 to 161 in 8a72b7a

	let constructor = super::help::WrappedConstructor {
	ty: result_ty.handle().unwrap(),
	};
	self.write_wrapped_constructor_function_name(module, constructor)?;

this is where we inject a call to the function that constructs the type

We need to more or less copy-paste the code above for arrays.