[spv-out] Fix invalid spirv being generated from integer dot products

[PyryM]

Short: this fixes integer dot product SPIR-V generation by moving a constant declaration to the right place.

Longer:
The code generation for integer dot products would (only sometimes, I am not sure why the tests do not catch this) produce SPIR-V that is technically invalid because it tries to declare constant null/0 inside of a function:

; Function 21
%21 = OpFunction %void None %22
%17 = OpLabel
// [...]
// vvv generated integer dot product
%38 = OpConstantNull %uint
%39 = OpCompositeExtract %uint %33 0
%40 = OpCompositeExtract %uint %36 0
%41 = OpIMul %uint %39 %40
%42 = OpIAdd %uint %38 %41
// [...]

Running against the SPIR-V validator (or a sufficiently pedantic driver like those for the Intel Arc Alchemist) causes a validation failure along the lines of:

error: line 55: ConstantNull cannot appear in a function declaration
  %38 = OpConstantNull %uint

Mysterious crashes from this had been noticed before on some hardware (e.g., gfx-rs/wgpu#2694).

The fix is to use the already-provided helper function for declaring constant null/0.

[teoxoy]

Thanks for the PR; it looks great!

[teoxoy]

One thing before I merge this:

Running against the SPIR-V validator (or a sufficiently pedantic driver like those for the Intel Arc Alchemist) causes a validation failure along the lines of:

error: line 55: ConstantNull cannot appear in a function declaration
  %38 = OpConstantNull %uint

How did you get spirv-val to produce this error?

I tried the command below (the one CI uses as well) and can't get spirv-val to error.

cat ./tests/out/spv/functions.spvasm | spirv-as --target-env spv1.0 -o - | spirv-val

[PyryM]

I'm not sure why the specific WGSL shaders in the tests don't exhibit the issue; is there some later optimization pass that is perhaps lifting shared constants and obscuring the issue?

However, this compute shader will provoke the behavior:

@group(0) @binding(0) var<storage, read_write> vals_out: array<u32>;
@compute @workgroup_size(4, 4, 4)
fn cs_main(@builtin(global_invocation_id) global_id: vec3<u32>) {
  vals_out[global_id.x] = dot(global_id, global_id);
}

Run through Naga -> spirv-val produces:

error: line 33: ConstantNull cannot appear in a function declaration
  %22 = OpConstantNull %uint

Run through Naga -> spirv-dis produces the following disassembly where the OpConstantNull does indeed appear
inside a function body:

[...] 
         %13 = OpFunction %void None %14
          %9 = OpLabel
         %12 = OpLoad %v3uint %gl_GlobalInvocationID
         %17 = OpAccessChain %_ptr_StorageBuffer__runtimearr_uint %6 %uint_0
               OpBranch %18
         %18 = OpLabel
         %19 = OpCompositeExtract %uint %12 0
=====>   %22 = OpConstantNull %uint
         %23 = OpCompositeExtract %uint %12 0
         %24 = OpCompositeExtract %uint %12 0
         %25 = OpIMul %uint %23 %24
         %26 = OpIAdd %uint %22 %25
         %27 = OpCompositeExtract %uint %12 1
         %28 = OpCompositeExtract %uint %12 1
         %29 = OpIMul %uint %27 %28
         %30 = OpIAdd %uint %26 %29
         %31 = OpCompositeExtract %uint %12 2
         %32 = OpCompositeExtract %uint %12 2
         %33 = OpIMul %uint %31 %32
         %21 = OpIAdd %uint %30 %33
         %34 = OpAccessChain %_ptr_StorageBuffer_uint %17 %19
               OpStore %34 %21
               OpReturn
               OpFunctionEnd

[PyryM]

Actually, running tests/in/functions.wgsl through CLI Naga will also result in the invalid spir-v being generated; there must be some difference between how the CLI (and also wgpu/wgpu_native) processes wgsl->spv and how the snapshot/test outputs are generated.

[teoxoy]

I see, the assembler of rspirv is lifting those constants out of the functions. We might want to run spirv-val on the binaries directly; writing both the text and binary representations.