When translating WGSL to HLSL, the expected translation of pointer-to-workgroup formal parameters to a user defined function is as an HLSL inout formal parameter.

Example WGSL:

var<workgroup> w: array<i32,100>;

fn helper(pw: ptr<workgroup,array<i32,100>>) {
  (*pw)[5] = 1;
}

@compute @workgroup_size(1)
fn main() {
  helper(&w);
}

Tint's emitted HLSL code:

groupshared int w[100];

void helper(inout int pw[100]) {
  pw[5] = 1;
}

struct tint_symbol_1 {
  uint local_invocation_index : SV_GroupIndex;
};

void main_inner(uint local_invocation_index) {
  {
    [loop] for(uint idx = local_invocation_index; (idx < 100u); idx = (idx + 1u)) {
      const uint i = idx;
      w[i] = 0;
    }
  }
  GroupMemoryBarrierWithGroupSync();
  helper(w);
}

[numthreads(1, 1, 1)]
void main(tint_symbol_1 tint_symbol) {
  main_inner(tint_symbol.local_invocation_index);
  return;
}

In last week's meeting, Tex mentioned that in general, translating to a groupshared inout variable will not work as the programmer intended.

HLSL's inout semantics are copy-in-copy-out. And strictly implementing copy-in-copy-out breaks the programmer expectation. Often the backend compiler will be able to prove to itself that there is no aliasing and elide the copies. But that is not always possible, and is dependent on the sophistication of the backend compiler.

So the pointer-to-workgroup feature for user-defined functions may need to be reconsidered for as long as HLSL is a required underlying language.