WGSL: Functions cannot be called at module scope

[aleino-nv]

Failing tests:

• tests/language-feature/constants/static-const-in-generic-interface.slang

Example output:

WGPU error: Error while parsing WGSL: :8:24 error: functions cannot be called at module-scope
const result_0 : i32 = EnsureCompileTimeEval_getValue_0();
                       ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^


 - While validating [ShaderModuleDescriptor]
 - While calling [Device].CreateShaderModule([ShaderModuleDescriptor]).

[aleino-nv]

@csyonghe I haven't looked into this yet, but from the name it sounds like Slang is emitting code where EnsureCompileTimeEval_getValue_0() is supposed to be evaluated by WGSL at compile time.

Is this something Slang does? WGSL has some support for something similar to "constexpr", but I think they have to be explicitly marked as such.

[csyonghe]

I think this is supposed to be inlined and evaluated by slang so it never actually appears in the generated wgsl code.

[aleino-nv]

For some reason I cannot reproduce by the problematic code by just doing a stand-alone compilation:

slangc -o test.wgsl -target wgsl -stage compute -entry computeMain static-const-in-generic-interface.slang

This produces the following code:

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

fn EnsureCompileTimeEval_getValue_0() -> i32
{
    return i32(4);
}

@compute
@workgroup_size(1, 1, 1)
fn computeMain(@builtin(global_invocation_id) dispatchThreadID_0 : vec3<u32>)
{
    outputBuffer_0[i32(0)] = EnsureCompileTimeEval_getValue_0();
    return;
}

I'll debug a bit to see why slang-test somehow comes up with different code, and what the complete generated code looks like.

[aleino-nv]

This is the code that actually gets used to create the WebGPU shader module.
It has the problem that the error message is complaining about

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

fn EnsureCompileTimeEval_getValue_0() -> i32
{
    return i32(4);
}

const result_0 : i32 = EnsureCompileTimeEval_getValue_0();
@compute
@workgroup_size(1, 1, 1)
fn computeMain(@builtin(global_invocation_id) dispatchThreadID_0 : vec3<u32>)
{
    var dispatchThreadID_1 : vec3<i32> = vec3<i32>(dispatchThreadID_0);
    outputBuffer_0[i32(0)] = result_0;
    return;
}

[aleino-nv]

I can reproduce with slangc if I add the -minimum-slang-optimization option! This will make it easier to investigate.

[aleino-nv]

After some debugging based on the above, I have discovered that running simplifyIR at various places in linkAndOptimizeIR resolves the issue.
That feels like overkill though...

The basic issue is that static const int result = EnsureCompileTimeEval<int>.getValue<Impl>(); gets hoisted out to global scope by the time we get to WGSL. (I believe constants must be at global scope in WGSL.)

It seems to me that whatever code is hoisting this variable should trigger some kind of more targeted simplification to avoid illegal function calls at global scope level.

I noticed that in HLSL we do end up with a global function call, though it's allowed:
static const int result_0 = EnsureCompileTimeEval_getValue_0();

For GLSL, it seems the required simplification is running even with -minimum-slang-optimization.
Probably WGSL should be handled in the same way as GLSL, so now I'm investigating how GLSL avoids this issue.

[aleino-nv]

This call here may be difference between WGSL and GLSL in this case:
https://github.com/shader-slang/slang/blob/master/source/slang/slang-emit.cpp#L1466

At least, if I comment this out then GLSL generates code with similar problem (call expression at global scope).

[aleino-nv]

This call here may be difference between WGSL and GLSL in this case: https://github.com/shader-slang/slang/blob/master/source/slang/slang-emit.cpp#L1466

At least, if I comment this out then GLSL generates code with similar problem (call expression at global scope).

Running the same code for WGSL gets rid of the function call at global scope if I just have -minimum-slang-optimization on the command line. However slang-test still ends up generating the function call at global scope.

I believe slang-test is compiling in debug mode. Indeed if I add -g1 then that reintroduces the function call at global scope issue.

Debugging further, I found that WGSL emit logic works differently in these two cases: the call is deemed to possibly have side effects when I run via slang-test, and so shouldFoldInstIntoUseSites returns false, whereas it returns true when I run slangc (at least without -g1).

I also noticed that -g1 causes GLSL to generate the call at global scope, so I don't expect to find a solution by looking at how GLSL handles this, because it seemingly doesn't!

[aleino-nv]

I assume it's correct that the call may have side effects: in debug mode any call could have the side effect of e.g. breaking into the debugger.
One way around this would be to say that the call can be folded into usage sites on the condition that it's also possible to then completely inline it so that there is no call anymore: hence no unintended side effects.

[aleino-nv]

The test passes if I comment out the mightHaveSideEffects check and just return true...

// If the instruction is at global scope, then it might represent
// a constant (e.g., the value of an enum case).
//
if (as<IRModuleInst>(inst->getParent()))
{
    if (!inst->mightHaveSideEffects())
        return true;
}

The function has kIROp_ReadNoneDecoration in slangc, which is why mightHaveSideEffects returns false, and in slang-test it does not have this decoration.

[aleino-nv]

Yong pointed out yesterday that static const values get lowered to global constants right away when producing the initial IR, and also that the SPIR-V backend doesn't have this problem.

The reason is that SPIR-V IR legalization contains some code that will inline such global constants.
The task is now to make that code reusable and then run it for WGSL as well. (I also noticed that GLSL seems to have the same issue when you compile with -g1, so it should probably also reuse this legalization.)

[aleino-nv]

When I try to reuse the same kind of legalization as SPIR-V does (https://github.com/aleino-nv/slang/tree/aleino/wgpu), the reference to the
global call gets inlined into the function, but unfortunately the global call remains.

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

fn EnsureCompileTimeEval_getValue_0() -> i32
{
    return i32(4);
}

const result_0 : i32 = EnsureCompileTimeEval_getValue_0();
@compute
@workgroup_size(1, 1, 1)
fn computeMain(@builtin(global_invocation_id) dispatchThreadID_0 : vec3<u32>)
{
    var dispatchThreadID_1 : vec3<i32> = vec3<i32>(dispatchThreadID_0);
    var result_1 : i32 = EnsureCompileTimeEval_getValue_0();
    var result_2 : i32 = EnsureCompileTimeEval_getValue_0();
    outputBuffer_0[i32(0)] = result_2;
    return;
}

I tried adding the same kind of "legalization cleanup" that's done for SPIR-V, but the
same code as above is still generated.

@csyonghe I guess I need some tweaks to the way that inlining is done in this case..

[csyonghe]

You will need to delete the global insts after they have been inlined.

[aleino-nv]

You will need to delete the global insts after they have been inlined.

Sure, in one way or another... but this is now the same logic as for the SPIR-V legalization.
I guess for SPIR-V, the global gets cleaned up by some other means?

[csyonghe]

The spirv emit logic will not emit anything not referenced by the EntryPoint. But wgsl backend will simply emit everything in the module so it doesn’t hurt to just remove the unused global for both backends.

[aleino-nv]

Ok, I added some code to remove any globals that are left unused after the inlining, and now the test is passing.

[aleino-nv]

Here is a PR #5768

[aleino-nv]

There is an issue where after removing the inlined IRInst's we end up deleting an SPIRVAsm instruction for tests/spirv/global-compute.slang.

The instruction is representing static const int vmin = min(1,2); as SPIRVAsm.

Then later on, for some reason spvtools ends up crashing (dereferencing the result type, which is nullptr) when optimizing this instruction.

[aleino-nv]

I think I have a proper fix for the issue in the above PR now -- it's now in review.

[aleino-nv]

I think I have a proper fix for the issue in the above PR now -- it's now in review.

Found out in review that another fix is needed, so this required another attempt which I've just submitted, so the patch is in review again.

[aleino-nv]

It turned out there was a problem with the existing inlining code, that only became apparent when I added my change to delete the inlined nodes. The inlining problem should now be solved, as well.