diff --git a/apps/typegpu-docs/tests/individual-example-tests/3d-fish.test.ts b/apps/typegpu-docs/tests/individual-example-tests/3d-fish.test.ts
index dca6c22f4e..6430284a2e 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/3d-fish.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/3d-fish.test.ts
@@ -305,15 +305,6 @@ describe('3d fish example', () => {
         return vertexShader_Output(worldPosition.xyz, worldNormal, canvasPosition, (*currentModelData).variant, _arg_textureUV, (*currentModelData).applySeaFog, (*currentModelData).applySeaDesaturation);
       }
 
-      struct fragmentShader_Input {
-        @location(0) worldPosition: vec3f,
-        @location(1) worldNormal: vec3f,
-        @location(2) variant: f32,
-        @location(3) textureUV: vec2f,
-        @location(4) @interpolate(flat) applySeaFog: u32,
-        @location(5) @interpolate(flat) applySeaDesaturation: u32,
-      }
-
       @group(0) @binding(1) var modelTexture: texture_2d<f32>;
 
       @group(0) @binding(3) var sampler_1: sampler;
@@ -418,6 +409,15 @@ describe('3d fish example', () => {
         return vec3f(r, g, b);
       }
 
+      struct fragmentShader_Input {
+        @location(0) worldPosition: vec3f,
+        @location(1) worldNormal: vec3f,
+        @location(2) variant: f32,
+        @location(3) textureUV: vec2f,
+        @location(4) @interpolate(flat) applySeaFog: u32,
+        @location(5) @interpolate(flat) applySeaDesaturation: u32,
+      }
+
       @fragment fn fragmentShader(_arg_0: fragmentShader_Input) -> @location(0) vec4f {
         var textureColorWithAlpha = textureSample(modelTexture, sampler_1, _arg_0.textureUV);
         var textureColor = textureColorWithAlpha.rgb;
diff --git a/apps/typegpu-docs/tests/individual-example-tests/bitonic-sort.test.ts b/apps/typegpu-docs/tests/individual-example-tests/bitonic-sort.test.ts
index 06b1c586cb..6368fd715f 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/bitonic-sort.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/bitonic-sort.test.ts
@@ -118,12 +118,12 @@ describe('bitonic sort example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
+      @group(0) @binding(0) var<storage, read> data_1: array<u32>;
+
       struct fragmentFn_Input {
         @location(0) uv: vec2f,
       }
 
-      @group(0) @binding(0) var<storage, read> data_1: array<u32>;
-
       @fragment fn fragmentFn(_arg_0: fragmentFn_Input) -> @location(0) vec4f {
         let data = (&data_1);
         let arrayLength_1 = arrayLength(&(*data));
diff --git a/apps/typegpu-docs/tests/individual-example-tests/blur.test.ts b/apps/typegpu-docs/tests/individual-example-tests/blur.test.ts
index 00df4923ed..7f81d96515 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/blur.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/blur.test.ts
@@ -470,14 +470,14 @@ describe('blur example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
-      struct renderFragment_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(0) @binding(0) var renderView: texture_2d<f32>;
 
       @group(0) @binding(1) var sampler_1: sampler;
 
+      struct renderFragment_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn renderFragment(_arg_0: renderFragment_Input) -> @location(0) vec4f {
         return textureSample(renderView, sampler_1, _arg_0.uv);
       }"
diff --git a/apps/typegpu-docs/tests/individual-example-tests/camera-thresholding.test.ts b/apps/typegpu-docs/tests/individual-example-tests/camera-thresholding.test.ts
index 43e10c97b2..783bb59558 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/camera-thresholding.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/camera-thresholding.test.ts
@@ -32,10 +32,6 @@ describe('camera thresholding example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
-      struct mainFrag_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(0) @binding(0) var<uniform> uvTransformUniform: mat2x2f;
 
       @group(1) @binding(0) var inputTexture: texture_external;
@@ -48,6 +44,10 @@ describe('camera thresholding example', () => {
 
       @group(0) @binding(3) var<uniform> thresholdBuffer: f32;
 
+      struct mainFrag_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn mainFrag(_arg_0: mainFrag_Input) -> @location(0) vec4f {
         var uv2 = ((uvTransformUniform * (_arg_0.uv - 0.5f)) + 0.5f);
         var col = textureSampleBaseClampToEdge(inputTexture, sampler_1, uv2);
diff --git a/apps/typegpu-docs/tests/individual-example-tests/caustics.test.ts b/apps/typegpu-docs/tests/individual-example-tests/caustics.test.ts
index a1938d5c34..9d011f7a74 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/caustics.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/caustics.test.ts
@@ -31,10 +31,6 @@ describe('caustics example', () => {
         return mainVertex_Output(vec4f(pos[vertexIndex], 0f, 1f), uv[vertexIndex]);
       }
 
-      struct mainFragment_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(0) @binding(0) var<uniform> tileDensity: f32;
 
       fn tilePattern(uv: vec2f) -> f32 {
@@ -120,6 +116,10 @@ describe('caustics example', () => {
         return mat2x2f(vec2f(cos(angle), sin(angle)), vec2f(-(sin(angle)), cos(angle)));
       }
 
+      struct mainFragment_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn mainFragment(_arg_0: mainFragment_Input) -> @location(0) vec4f {
         var skewMat = mat2x2f(vec2f(0.9800665974617004, 0.19866932928562164), vec2f((-1.9866933079506122f + (_arg_0.uv.x * 3f)), 4.900332889206208f));
         var skewedUv = (skewMat * _arg_0.uv);
diff --git a/apps/typegpu-docs/tests/individual-example-tests/chroma-keying.test.ts b/apps/typegpu-docs/tests/individual-example-tests/chroma-keying.test.ts
index 9b25129771..60175ac939 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/chroma-keying.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/chroma-keying.test.ts
@@ -32,10 +32,6 @@ describe('chroma keying example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
-      struct fragment_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(0) @binding(0) var<uniform> uvTransform: mat2x2f;
 
       @group(1) @binding(0) var inputTexture: texture_external;
@@ -48,6 +44,10 @@ describe('chroma keying example', () => {
 
       @group(0) @binding(3) var<uniform> threshold: f32;
 
+      struct fragment_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn fragment(_arg_0: fragment_Input) -> @location(0) vec4f {
         var uv2 = ((uvTransform * (_arg_0.uv - 0.5f)) + 0.5f);
         var col = textureSampleBaseClampToEdge(inputTexture, sampler_1, uv2);
diff --git a/apps/typegpu-docs/tests/individual-example-tests/confetti.test.ts b/apps/typegpu-docs/tests/individual-example-tests/confetti.test.ts
index 3e39636f7a..921721055b 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/confetti.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/confetti.test.ts
@@ -59,10 +59,10 @@ describe('confetti example', () => {
 
       struct VertexIn {
         @location(0) tilt: f32,
-        @location(1) angle: f32,
-        @location(2) color: vec4f,
-        @location(3) center: vec2f,
         @builtin(vertex_index) vertexIndex: u32,
+        @location(1) angle: f32,
+        @location(2) center: vec2f,
+        @location(3) color: vec4f,
       }
 
       @vertex fn vertex(_arg_0: VertexIn) -> VertexOut {
diff --git a/apps/typegpu-docs/tests/individual-example-tests/cubemap-reflection.test.ts b/apps/typegpu-docs/tests/individual-example-tests/cubemap-reflection.test.ts
index c77faffba2..fe28f8827d 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/cubemap-reflection.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/cubemap-reflection.test.ts
@@ -257,14 +257,14 @@ describe('cubemap reflection example', () => {
         return cubeVertexFn_Output((camera.projection * vec4f(viewPos, 1f)), _arg_position.xyz);
       }
 
-      struct cubeFragmentFn_Input {
-        @location(0) texCoord: vec3f,
-      }
-
       @group(1) @binding(0) var cubemap: texture_cube<f32>;
 
       @group(1) @binding(1) var texSampler: sampler;
 
+      struct cubeFragmentFn_Input {
+        @location(0) texCoord: vec3f,
+      }
+
       @fragment fn cubeFragmentFn(_arg_0: cubeFragmentFn_Input) -> @location(0) vec4f {
         return textureSample(cubemap, texSampler, normalize(_arg_0.texCoord));
       }
@@ -287,11 +287,6 @@ describe('cubemap reflection example', () => {
         return vertexFn_Output((camera.projection * (camera.view * _arg_position)), _arg_normal, _arg_position);
       }
 
-      struct fragmentFn_Input {
-        @location(0) normal: vec4f,
-        @location(1) worldPos: vec4f,
-      }
-
       struct DirectionalLight {
         direction: vec3f,
         color: vec3f,
@@ -314,6 +309,11 @@ describe('cubemap reflection example', () => {
 
       @group(1) @binding(1) var texSampler: sampler;
 
+      struct fragmentFn_Input {
+        @location(0) normal: vec4f,
+        @location(1) worldPos: vec4f,
+      }
+
       @fragment fn fragmentFn(_arg_0: fragmentFn_Input) -> @location(0) vec4f {
         var normalizedNormal = normalize(_arg_0.normal.xyz);
         var normalizedLightDir = normalize(light.direction);
diff --git a/apps/typegpu-docs/tests/individual-example-tests/disco.test.ts b/apps/typegpu-docs/tests/individual-example-tests/disco.test.ts
index 8ea4e8d7f7..ba6d090535 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/disco.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/disco.test.ts
@@ -32,10 +32,6 @@ describe('disco example', () => {
         return mainVertex_Output(vec4f(pos[vertexIndex], 0f, 1f), uv[vertexIndex]);
       }
 
-      struct mainFragment2_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(0) @binding(0) var<uniform> resolutionUniform: vec2f;
 
       fn aspectCorrected(uv: vec2f) -> vec2f {
@@ -65,6 +61,10 @@ describe('disco example', () => {
         return (acc + (col * weight));
       }
 
+      struct mainFragment2_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn mainFragment2(_arg_0: mainFragment2_Input) -> @location(0) vec4f {
         var originalUv = aspectCorrected(_arg_0.uv);
         var aspectUv = originalUv;
@@ -243,10 +243,6 @@ describe('disco example', () => {
         return mainVertex_Output(vec4f(pos[vertexIndex], 0f, 1f), uv[vertexIndex]);
       }
 
-      struct mainFragment1_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(0) @binding(0) var<uniform> resolutionUniform: vec2f;
 
       fn aspectCorrected(uv: vec2f) -> vec2f {
@@ -276,6 +272,10 @@ describe('disco example', () => {
         return (acc + (col * weight));
       }
 
+      struct mainFragment1_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn mainFragment1(_arg_0: mainFragment1_Input) -> @location(0) vec4f {
         var originalUv = aspectCorrected(_arg_0.uv);
         var aspectUv = originalUv;
diff --git a/apps/typegpu-docs/tests/individual-example-tests/fluid-double-buffering.test.ts b/apps/typegpu-docs/tests/individual-example-tests/fluid-double-buffering.test.ts
index ef73e74189..df9620d6ed 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/fluid-double-buffering.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/fluid-double-buffering.test.ts
@@ -558,10 +558,6 @@ describe('fluid double buffering example', () => {
         return vertexMain_Output(vec4f(pos[_arg_idx].x, pos[_arg_idx].y, 0f, 1f), uv[_arg_idx]);
       }
 
-      struct fragmentMain_Input {
-        @location(0) uv: vec2f,
-      }
-
       fn coordsToIndex(x: i32, y: i32) -> i32 {
         return (x + (y * 256i));
       }
@@ -595,6 +591,10 @@ describe('fluid double buffering example', () => {
         return false;
       }
 
+      struct fragmentMain_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn fragmentMain(_arg_0: fragmentMain_Input) -> @location(0) vec4f {
         let x = i32((_arg_0.uv.x * 256f));
         let y = i32((_arg_0.uv.y * 256f));
diff --git a/apps/typegpu-docs/tests/individual-example-tests/game-of-life.test.ts b/apps/typegpu-docs/tests/individual-example-tests/game-of-life.test.ts
index 3028c19bf5..516692dc41 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/game-of-life.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/game-of-life.test.ts
@@ -214,10 +214,6 @@ describe('game of life example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
-      struct displayFragment_Input {
-        @location(0) uv: vec2f,
-      }
-
       struct ZoomParams {
         enabled: u32,
         level: f32,
@@ -242,6 +238,10 @@ describe('game of life example', () => {
 
       @group(0) @binding(2) var<uniform> viewModeUniform: u32;
 
+      struct displayFragment_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn displayFragment(_arg_0: displayFragment_Input) -> @location(0) vec4f {
         let zoom = (&zoomUniform);
         let gs = f32(gameSizeUniform);
diff --git a/apps/typegpu-docs/tests/individual-example-tests/gradient-tiles.test.ts b/apps/typegpu-docs/tests/individual-example-tests/gradient-tiles.test.ts
index f0d3016198..d3c209eb18 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/gradient-tiles.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/gradient-tiles.test.ts
@@ -32,12 +32,12 @@ describe('gradient tiles example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
+      @group(0) @binding(0) var<uniform> spanUniform: vec2f;
+
       struct fragment_Input {
         @location(0) uv: vec2f,
       }
 
-      @group(0) @binding(0) var<uniform> spanUniform: vec2f;
-
       @fragment fn fragment(_arg_0: fragment_Input) -> @location(0) vec4f {
         let red = (floor((_arg_0.uv.x * spanUniform.x)) / spanUniform.x);
         let green = (floor((_arg_0.uv.y * spanUniform.y)) / spanUniform.y);
diff --git a/apps/typegpu-docs/tests/individual-example-tests/gravity.test.ts b/apps/typegpu-docs/tests/individual-example-tests/gravity.test.ts
index 55ed0d2fb8..899b6cb76e 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/gravity.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/gravity.test.ts
@@ -192,14 +192,14 @@ describe('gravity example', () => {
         return skyBoxVertex_Output((camera.projection * vec4f(viewPos, 1f)), _arg_position.xyz);
       }
 
-      struct skyBoxFragment_Input {
-        @location(0) texCoord: vec3f,
-      }
-
       @group(0) @binding(1) var skyBox: texture_cube<f32>;
 
       @group(0) @binding(2) var sampler_1: sampler;
 
+      struct skyBoxFragment_Input {
+        @location(0) texCoord: vec3f,
+      }
+
       @fragment fn skyBoxFragment(_arg_0: skyBoxFragment_Input) -> @location(0) vec4f {
         return textureSample(skyBox, sampler_1, normalize(_arg_0.texCoord));
       }
@@ -250,6 +250,12 @@ describe('gravity example', () => {
         return mainVertex_Output(positionOnCanvas, _arg_uv, _arg_normal, worldPosition, (*currentBody).textureIndex, (*currentBody).destroyed, (*currentBody).ambientLightFactor);
       }
 
+      @group(1) @binding(0) var celestialBodyTextures: texture_2d_array<f32>;
+
+      @group(0) @binding(1) var sampler_1: sampler;
+
+      @group(0) @binding(2) var<uniform> lightSource: vec3f;
+
       struct mainFragment_Input {
         @location(0) uv: vec2f,
         @location(1) normals: vec3f,
@@ -259,12 +265,6 @@ describe('gravity example', () => {
         @location(5) ambientLightFactor: f32,
       }
 
-      @group(1) @binding(0) var celestialBodyTextures: texture_2d_array<f32>;
-
-      @group(0) @binding(1) var sampler_1: sampler;
-
-      @group(0) @binding(2) var<uniform> lightSource: vec3f;
-
       @fragment fn mainFragment(_arg_0: mainFragment_Input) -> @location(0) vec4f {
         if ((_arg_0.destroyed == 1u)) {
           discard;;
diff --git a/apps/typegpu-docs/tests/individual-example-tests/image-tuning.test.ts b/apps/typegpu-docs/tests/individual-example-tests/image-tuning.test.ts
index a9fa1865c8..059828005e 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/image-tuning.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/image-tuning.test.ts
@@ -37,10 +37,6 @@ describe('image tuning example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
-      struct fragment_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(0) @binding(0) var imageView: texture_2d<f32>;
 
       @group(0) @binding(1) var imageSampler: sampler;
@@ -68,6 +64,10 @@ describe('image tuning example', () => {
 
       @group(0) @binding(4) var<uniform> adjustments: Adjustments;
 
+      struct fragment_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn fragment(_arg_0: fragment_Input) -> @location(0) vec4f {
         var color = textureSample(imageView, imageSampler, _arg_0.uv).rgb;
         let inputLuminance = dot(color, vec3f(0.29899999499320984, 0.5870000123977661, 0.11400000005960464));
diff --git a/apps/typegpu-docs/tests/individual-example-tests/jelly-slider.test.ts b/apps/typegpu-docs/tests/individual-example-tests/jelly-slider.test.ts
index 941ee300b4..976bdbaaf2 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/jelly-slider.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/jelly-slider.test.ts
@@ -66,10 +66,6 @@ describe('jelly-slider example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
-      struct raymarchFn_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(0) @binding(0) var<uniform> randomUniform: vec2f;
 
       var<private> seed: vec2f;
@@ -583,6 +579,10 @@ describe('jelly-slider example', () => {
         return background;
       }
 
+      struct raymarchFn_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn raymarchFn(_arg_0: raymarchFn_Input) -> @location(0) vec4f {
         randSeed2((randomUniform * _arg_0.uv));
         var ndc = vec2f(((_arg_0.uv.x * 2f) - 1f), -(((_arg_0.uv.y * 2f) - 1f)));
@@ -713,14 +713,14 @@ describe('jelly-slider example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
-      struct fragmentMain_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(1) @binding(0) var currentTexture: texture_2d<f32>;
 
       @group(0) @binding(0) var filteringSampler: sampler;
 
+      struct fragmentMain_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn fragmentMain(_arg_0: fragmentMain_Input) -> @location(0) vec4f {
         return textureSample(currentTexture, filteringSampler, _arg_0.uv);
       }
diff --git a/apps/typegpu-docs/tests/individual-example-tests/jelly-switch.test.ts b/apps/typegpu-docs/tests/individual-example-tests/jelly-switch.test.ts
index 7a192b5f54..6e507a2602 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/jelly-switch.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/jelly-switch.test.ts
@@ -37,10 +37,6 @@ describe('jelly switch example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
-      struct raymarchFn_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(0) @binding(0) var<uniform> randomUniform: vec2f;
 
       var<private> seed: vec2f;
@@ -385,6 +381,10 @@ describe('jelly switch example', () => {
         return background;
       }
 
+      struct raymarchFn_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn raymarchFn(_arg_0: raymarchFn_Input) -> @location(0) vec4f {
         randSeed2((randomUniform * _arg_0.uv));
         var ndc = vec2f(((_arg_0.uv.x * 2f) - 1f), -(((_arg_0.uv.y * 2f) - 1f)));
@@ -505,14 +505,14 @@ describe('jelly switch example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
-      struct fragmentMain_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(1) @binding(0) var currentTexture: texture_2d<f32>;
 
       @group(0) @binding(0) var filteringSampler: sampler;
 
+      struct fragmentMain_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn fragmentMain(_arg_0: fragmentMain_Input) -> @location(0) vec4f {
         return textureSample(currentTexture, filteringSampler, _arg_0.uv);
       }"
diff --git a/apps/typegpu-docs/tests/individual-example-tests/jump-flood-distance.test.ts b/apps/typegpu-docs/tests/individual-example-tests/jump-flood-distance.test.ts
index 3fb0e7d685..dcf73b0927 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/jump-flood-distance.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/jump-flood-distance.test.ts
@@ -303,10 +303,6 @@ describe('jump flood (distance) example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
-      struct distanceFrag_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(1) @binding(0) var distTexture: texture_2d<f32>;
 
       @group(1) @binding(1) var sampler_1: sampler;
@@ -322,6 +318,10 @@ describe('jump flood (distance) example', () => {
 
       const insideGradient: array<vec3f, 5> = array<vec3f, 5>(vec3f(0.05000000074505806, 0.05000000074505806, 0.15000000596046448), vec3f(0.10000000149011612, 0.20000000298023224, 0.30000001192092896), vec3f(0.20000000298023224, 0.44999998807907104, 0.550000011920929), vec3f(0.4000000059604645, 0.75, 0.699999988079071), vec3f(0.8999999761581421, 1, 0.949999988079071));
 
+      struct distanceFrag_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn distanceFrag(_arg_0: distanceFrag_Input) -> @location(0) vec4f {
         var size = textureDimensions(distTexture);
         var dist = textureSample(distTexture, sampler_1, _arg_0.uv).x;
diff --git a/apps/typegpu-docs/tests/individual-example-tests/liquid-glass.test.ts b/apps/typegpu-docs/tests/individual-example-tests/liquid-glass.test.ts
index e247d26a22..56a5b4ecf5 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/liquid-glass.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/liquid-glass.test.ts
@@ -64,10 +64,6 @@ describe('liquid-glass example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
-      struct fragmentShader_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(0) @binding(0) var<uniform> mousePosUniform: vec2f;
 
       struct Params {
@@ -137,6 +133,10 @@ describe('liquid-glass example', () => {
         return mix(vec4f(color, 1f), vec4f(tint.color, 1f), tint.strength);
       }
 
+      struct fragmentShader_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn fragmentShader(_arg_0: fragmentShader_Input) -> @location(0) vec4f {
         var posInBoxSpace = (_arg_0.uv - mousePosUniform);
         let sdfDist = sdRoundedBox2d(posInBoxSpace, paramsUniform.rectDims, paramsUniform.radius);
diff --git a/apps/typegpu-docs/tests/individual-example-tests/oklab.test.ts b/apps/typegpu-docs/tests/individual-example-tests/oklab.test.ts
index b8e64d3654..23abf35258 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/oklab.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/oklab.test.ts
@@ -32,10 +32,6 @@ describe('oklab example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
-      struct mainFragment_Input {
-        @location(0) uv: vec2f,
-      }
-
       struct item {
         hue: f32,
         alpha: f32,
@@ -226,6 +222,10 @@ describe('oklab example', () => {
         return 1f;
       }
 
+      struct mainFragment_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn mainFragment(_arg_0: mainFragment_Input) -> @location(0) vec4f {
         var uv = ((_arg_0.uv - 0.5f) * vec2f(2, -2));
         let hue = uniforms.hue;
diff --git a/apps/typegpu-docs/tests/individual-example-tests/phong-reflection.test.ts b/apps/typegpu-docs/tests/individual-example-tests/phong-reflection.test.ts
index 1cfbf878ea..0240d9ae44 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/phong-reflection.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/phong-reflection.test.ts
@@ -54,11 +54,6 @@ describe('phong reflection example', () => {
         return vertexShader_Output(_arg_modelPosition, _arg_modelNormal, canvasPosition);
       }
 
-      struct fragmentShader_Input {
-        @location(0) worldPosition: vec3f,
-        @location(1) worldNormal: vec3f,
-      }
-
       struct ExampleControls {
         lightColor: vec3f,
         lightDirection: vec3f,
@@ -69,6 +64,11 @@ describe('phong reflection example', () => {
 
       @group(0) @binding(1) var<uniform> exampleControlsUniform: ExampleControls;
 
+      struct fragmentShader_Input {
+        @location(0) worldPosition: vec3f,
+        @location(1) worldNormal: vec3f,
+      }
+
       @fragment fn fragmentShader(_arg_0: fragmentShader_Input) -> @location(0) vec4f {
         var lightColor = normalize(exampleControlsUniform.lightColor);
         var lightDirection = normalize(exampleControlsUniform.lightDirection);
diff --git a/apps/typegpu-docs/tests/individual-example-tests/point-light-shadow.test.ts b/apps/typegpu-docs/tests/individual-example-tests/point-light-shadow.test.ts
index 283d8d652b..2ce441c6b2 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/point-light-shadow.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/point-light-shadow.test.ts
@@ -41,12 +41,12 @@ describe('point light shadow example', () => {
         return vertexDepth_Output(pos, worldPos);
       }
 
+      @group(0) @binding(1) var<uniform> lightPosition: vec3f;
+
       struct fragmentDepth_Input {
         @location(0) worldPos: vec3f,
       }
 
-      @group(0) @binding(1) var<uniform> lightPosition: vec3f;
-
       @fragment fn fragmentDepth(_arg_0: fragmentDepth_Input) -> @builtin(frag_depth) f32 {
         let dist = length((_arg_0.worldPos - lightPosition));
         return (dist / 100f);
@@ -74,12 +74,6 @@ describe('point light shadow example', () => {
         return vertexMain_Output(pos, worldPos, uv, worldNormal);
       }
 
-      struct fragmentMain_Input {
-        @location(0) worldPos: vec3f,
-        @location(1) uv: vec2f,
-        @location(2) normal: vec3f,
-      }
-
       @group(1) @binding(3) var<uniform> lightPosition: vec3f;
 
       struct item {
@@ -97,6 +91,12 @@ describe('point light shadow example', () => {
 
       @group(1) @binding(2) var shadowSampler: sampler_comparison;
 
+      struct fragmentMain_Input {
+        @location(0) worldPos: vec3f,
+        @location(1) uv: vec2f,
+        @location(2) normal: vec3f,
+      }
+
       @fragment fn fragmentMain(_arg_0: fragmentMain_Input) -> @location(0) vec4f {
         let lightPos = (&lightPosition);
         var toLight = ((*lightPos) - _arg_0.worldPos);
diff --git a/apps/typegpu-docs/tests/individual-example-tests/ray-marching.test.ts b/apps/typegpu-docs/tests/individual-example-tests/ray-marching.test.ts
index 1d56a2d307..bdfa96686f 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/ray-marching.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/ray-marching.test.ts
@@ -31,10 +31,6 @@ describe('ray-marching example', () => {
         return vertexMain_Output(vec4f(pos[idx], 0f, 1f), uv[idx]);
       }
 
-      struct fragmentMain_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(0) @binding(0) var<uniform> resolution: vec2f;
 
       struct Shape {
@@ -148,6 +144,10 @@ describe('ray-marching example', () => {
         return res;
       }
 
+      struct fragmentMain_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn fragmentMain(_arg_0: fragmentMain_Input) -> @location(0) vec4f {
         var uv = ((_arg_0.uv * 2f) - 1f);
         uv.x *= (resolution.x / resolution.y);
diff --git a/apps/typegpu-docs/tests/individual-example-tests/ripple-cube.test.ts b/apps/typegpu-docs/tests/individual-example-tests/ripple-cube.test.ts
index e73568d898..24dcc7ee1c 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/ripple-cube.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/ripple-cube.test.ts
@@ -765,10 +765,6 @@ describe('ripple-cube example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
-      struct fragmentMain_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(1) @binding(0) var colorTexture: texture_2d<f32>;
 
       @group(1) @binding(2) var sampler_1: sampler;
@@ -782,6 +778,10 @@ describe('ripple-cube example', () => {
 
       @group(0) @binding(0) var<uniform> bloomUniform: BloomParams;
 
+      struct fragmentMain_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn fragmentMain(_arg_0: fragmentMain_Input) -> @location(0) vec4f {
         var color = textureSample(colorTexture, sampler_1, _arg_0.uv);
         var bloomColor = textureSample(bloomTexture, sampler_1, _arg_0.uv);
diff --git a/apps/typegpu-docs/tests/individual-example-tests/simple-shadow.test.ts b/apps/typegpu-docs/tests/individual-example-tests/simple-shadow.test.ts
index f1f6775cc0..e6d2de8222 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/simple-shadow.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/simple-shadow.test.ts
@@ -89,11 +89,6 @@ describe('simple shadow example', () => {
         return mainVert_Output(clipPos, transformedNormal, worldPos.xyz);
       }
 
-      struct mainFrag_Input {
-        @location(0) normal: vec4f,
-        @location(1) worldPos: vec3f,
-      }
-
       struct DirectionalLight {
         direction: vec3f,
         color: vec3f,
@@ -118,6 +113,11 @@ describe('simple shadow example', () => {
 
       @group(0) @binding(3) var<uniform> paramsUniform: VisParams;
 
+      struct mainFrag_Input {
+        @location(0) normal: vec4f,
+        @location(1) worldPos: vec3f,
+      }
+
       @fragment fn mainFrag(_arg_0: mainFrag_Input) -> @location(0) vec4f {
         let instanceInfo_1 = (&instanceInfo);
         var N = normalize(_arg_0.normal.xyz);
diff --git a/apps/typegpu-docs/tests/individual-example-tests/slime-mold-3d.test.ts b/apps/typegpu-docs/tests/individual-example-tests/slime-mold-3d.test.ts
index 16ec9b7e12..daab5433e6 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/slime-mold-3d.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/slime-mold-3d.test.ts
@@ -382,10 +382,6 @@ describe('slime mold 3d example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
-      struct fragmentShader_Input {
-        @location(0) uv: vec2f,
-      }
-
       var<private> seed: vec2f;
 
       fn seed2(value: vec2f) {
@@ -438,6 +434,10 @@ describe('slime mold 3d example', () => {
 
       @group(0) @binding(1) var sampler_1: sampler;
 
+      struct fragmentShader_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn fragmentShader(_arg_0: fragmentShader_Input) -> @location(0) vec4f {
         randSeed2(_arg_0.uv);
         var ndc = vec2f(((_arg_0.uv.x * 2f) - 1f), (1f - (_arg_0.uv.y * 2f)));
diff --git a/apps/typegpu-docs/tests/individual-example-tests/slime-mold.test.ts b/apps/typegpu-docs/tests/individual-example-tests/slime-mold.test.ts
index 44d050241e..4b77dc41f4 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/slime-mold.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/slime-mold.test.ts
@@ -308,14 +308,14 @@ describe('slime mold example', () => {
         return fullScreenTriangle_Output(vec4f(pos[_arg_vertexIndex], 0f, 1f), uv[_arg_vertexIndex]);
       }
 
-      struct fragmentShader_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(1) @binding(0) var state: texture_2d<f32>;
 
       @group(0) @binding(0) var filteringSampler: sampler;
 
+      struct fragmentShader_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn fragmentShader(_arg_0: fragmentShader_Input) -> @location(0) vec4f {
         return textureSample(state, filteringSampler, _arg_0.uv);
       }"
diff --git a/apps/typegpu-docs/tests/individual-example-tests/stable-fluid.test.ts b/apps/typegpu-docs/tests/individual-example-tests/stable-fluid.test.ts
index 0e97b0d9e9..b94b01c01f 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/stable-fluid.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/stable-fluid.test.ts
@@ -259,16 +259,16 @@ describe('stable-fluid example', () => {
         return renderFn_Output(vec4f(vertices[_arg_idx], 0f, 1f), texCoords[_arg_idx]);
       }
 
-      struct fragmentImageFn_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(0) @binding(0) var result: texture_2d<f32>;
 
       @group(0) @binding(2) var linSampler: sampler;
 
       @group(0) @binding(1) var background: texture_2d<f32>;
 
+      struct fragmentImageFn_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn fragmentImageFn(_arg_0: fragmentImageFn_Input) -> @location(0) vec4f {
         const pixelStep = 0.001953125f;
         let leftSample = textureSample(result, linSampler, vec2f((_arg_0.uv.x - pixelStep), _arg_0.uv.y)).x;
diff --git a/apps/typegpu-docs/tests/individual-example-tests/uniformity.test.ts b/apps/typegpu-docs/tests/individual-example-tests/uniformity.test.ts
index 21ccd861f6..11a61fe114 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/uniformity.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/uniformity.test.ts
@@ -35,10 +35,6 @@ describe('uniformity test example', () => {
         return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]);
       }
 
-      struct fragmentShader_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(0) @binding(0) var<uniform> canvasRatioUniform: f32;
 
       @group(0) @binding(1) var<uniform> gridSizeUniform: f32;
@@ -65,6 +61,10 @@ describe('uniformity test example', () => {
         return sample();
       }
 
+      struct fragmentShader_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn fragmentShader(_arg_0: fragmentShader_Input) -> @location(0) vec4f {
         var uv = (((_arg_0.uv + 1f) / 2f) * vec2f(canvasRatioUniform, 1f));
         var gridedUV = floor((uv * gridSizeUniform));
@@ -72,10 +72,6 @@ describe('uniformity test example', () => {
         return vec4f(vec3f(randFloat01()), 1f);
       }
 
-      struct fragmentShader_Input_1 {
-        @location(0) uv: vec2f,
-      }
-
       var<private> seed_1: u32;
 
       fn seed2_1(value: vec2f) {
@@ -102,6 +98,10 @@ describe('uniformity test example', () => {
         return sample_1();
       }
 
+      struct fragmentShader_Input_1 {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn fragmentShader_1(_arg_0: fragmentShader_Input_1) -> @location(0) vec4f {
         var uv = (((_arg_0.uv + 1f) / 2f) * vec2f(canvasRatioUniform, 1f));
         var gridedUV = floor((uv * gridSizeUniform));
diff --git a/apps/typegpu-docs/tests/individual-example-tests/vaporrave.test.ts b/apps/typegpu-docs/tests/individual-example-tests/vaporrave.test.ts
index d7834afa66..dafd303f91 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/vaporrave.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/vaporrave.test.ts
@@ -79,10 +79,6 @@ describe('vaporrave example', () => {
         return vertexMain_Output(vec4f(pos[idx], 0f, 1f), uv[idx]);
       }
 
-      struct fragmentMain_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(0) @binding(0) var<uniform> resolutionUniform: vec2f;
 
       struct Ray {
@@ -219,6 +215,10 @@ describe('vaporrave example', () => {
 
       @group(0) @binding(5) var<uniform> glowIntensityUniform: f32;
 
+      struct fragmentMain_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn fragmentMain(_arg_0: fragmentMain_Input) -> @location(0) vec4f {
         var uv = ((_arg_0.uv * 2f) - 1f);
         uv.x *= (resolutionUniform.x / resolutionUniform.y);
diff --git a/apps/typegpu-docs/tests/individual-example-tests/xor-dev-centrifuge-2.test.ts b/apps/typegpu-docs/tests/individual-example-tests/xor-dev-centrifuge-2.test.ts
index 9abffaf6b8..894d6e4d67 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/xor-dev-centrifuge-2.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/xor-dev-centrifuge-2.test.ts
@@ -30,10 +30,6 @@ describe('xor dev centrifuge example', () => {
         return vertexMain_Output(vec4f(pos[_arg_vertexIndex], 0f, 1f), pos[_arg_vertexIndex]);
       }
 
-      struct fragmentMain_Input {
-        @location(0) uv: vec2f,
-      }
-
       struct Params {
         time: f32,
         aspectRatio: f32,
@@ -51,6 +47,10 @@ describe('xor dev centrifuge example', () => {
         return select(tanh(v), sign(v), (abs(v) > vec3f(10)));
       }
 
+      struct fragmentMain_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn fragmentMain(_arg_0: fragmentMain_Input) -> @location(0) vec4f {
         let params = (&paramsUniform);
         var ratio = vec2f((*params).aspectRatio, 1f);
diff --git a/apps/typegpu-docs/tests/individual-example-tests/xor-dev-runner.test.ts b/apps/typegpu-docs/tests/individual-example-tests/xor-dev-runner.test.ts
index 419c95807c..ec9b1c27a0 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/xor-dev-runner.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/xor-dev-runner.test.ts
@@ -34,10 +34,6 @@ describe('xor dev runner example', () => {
         return vertexMain_Output(vec4f(pos[_arg_vertexIndex], 0f, 1f), pos[_arg_vertexIndex]);
       }
 
-      struct fragmentMain_Input {
-        @location(0) uv: vec2f,
-      }
-
       @group(0) @binding(0) var<uniform> colorUniform: vec3f;
 
       struct Camera {
@@ -82,6 +78,10 @@ describe('xor dev runner example', () => {
         return select(tanh(v), sign(v), (abs(v) > 10f));
       }
 
+      struct fragmentMain_Input {
+        @location(0) uv: vec2f,
+      }
+
       @fragment fn fragmentMain(_arg_0: fragmentMain_Input) -> @location(0) vec4f {
         var icolor = (colorUniform * 4f);
         var ray = getRayForUV(_arg_0.uv);
diff --git a/packages/typegpu-gl/src/glOptions.ts b/packages/typegpu-gl/src/glOptions.ts
new file mode 100644
index 0000000000..1e65a593db
--- /dev/null
+++ b/packages/typegpu-gl/src/glOptions.ts
@@ -0,0 +1,7 @@
+import glslGenerator from './glslGenerator.ts';
+
+export function glOptions() {
+  return {
+    unstable_shaderGenerator: glslGenerator,
+  };
+}
diff --git a/packages/typegpu-gl/src/glslGenerator.ts b/packages/typegpu-gl/src/glslGenerator.ts
index eee8206edf..4c549bdac7 100644
--- a/packages/typegpu-gl/src/glslGenerator.ts
+++ b/packages/typegpu-gl/src/glslGenerator.ts
@@ -1,10 +1,17 @@
-import { d, ShaderGenerator, WgslGenerator } from 'typegpu';
+import { NodeTypeCatalog as NODE } from 'tinyest';
+import type { Return } from 'tinyest';
+import tgpu, { d, ShaderGenerator, WgslGenerator } from 'typegpu';
+
+type ResolutionCtx = ShaderGenerator.ResolutionCtx;
+
+const UnknownData: typeof ShaderGenerator.UnknownData = ShaderGenerator.UnknownData;
 
 // ----------
 // WGSL → GLSL type name mapping
 // ----------
 
 const WGSL_TO_GLSL_TYPE: Record<string, string> = {
+  void: 'void',
   f32: 'float',
   u32: 'uint',
   i32: 'int',
@@ -37,13 +44,43 @@ export function translateWgslTypeToGlsl(wgslType: string): string {
   return WGSL_TO_GLSL_TYPE[wgslType] ?? wgslType;
 }
 
+/**
+ * Resolves a struct and adds its declaration to the resolution context.
+ * @param ctx - The resolution context.
+ * @param struct - The struct to resolve.
+ *
+ * @returns The resolved struct name.
+ */
+function resolveStruct(ctx: ResolutionCtx, struct: d.WgslStruct) {
+  const id = ctx.makeUniqueIdentifier(ShaderGenerator.getName(struct), 'global');
+
+  ctx.addDeclaration(`\
+struct ${id} {
+${Object.entries(struct.propTypes)
+  .map(([prop, type]) => `  ${ctx.resolve(type).value} ${prop};\n`)
+  .join('')}\
+};`);
+
+  return id;
+}
+
+const gl_PositionSnippet = tgpu['~unstable'].rawCodeSnippet('gl_Position', d.vec4f, 'private');
+
+interface EntryFnState {
+  structPropToVarMap: Record<string, string>;
+  outVars: { varName: string; propName: string }[];
+}
+
 /**
  * A GLSL ES 3.0 shader generator that extends WgslGenerator.
  * Overrides `dataType` to emit GLSL type names instead of WGSL ones,
  * and overrides variable declaration emission to use `type name = rhs` syntax.
  */
 export class GlslGenerator extends WgslGenerator {
-  public override typeAnnotation(data: d.BaseData): string {
+  #functionType: ShaderGenerator.TgpuShaderStage | 'normal' | undefined;
+  #entryFnState: EntryFnState | undefined;
+
+  override typeAnnotation(data: d.BaseData): string {
     // For WGSL identity types (scalars, vectors, common matrices), map to GLSL directly.
     if (!d.isLooseData(data)) {
       const glslName = WGSL_TO_GLSL_TYPE[data.type];
@@ -52,20 +89,153 @@ export class GlslGenerator extends WgslGenerator {
       }
     }
 
+    if (d.isWgslStruct(data)) {
+      return resolveStruct(this.ctx, data);
+    }
+
     // For all other types (structs, arrays, etc.) delegate to WGSL resolution.
     return super.typeAnnotation(data);
   }
 
-  protected override emitVarDecl(
-    pre: string,
+  override _emitVarDecl(
     _keyword: 'var' | 'let' | 'const',
     name: string,
     dataType: d.BaseData | ShaderGenerator.UnknownData,
     rhsStr: string,
   ): string {
-    const glslTypeName =
-      dataType !== ShaderGenerator.UnknownData ? this.typeAnnotation(dataType) : 'auto';
-    return `${pre}${glslTypeName} ${name} = ${rhsStr};`;
+    const glslTypeName = dataType !== UnknownData ? this.ctx.resolve(dataType).value : 'auto';
+    return `${this.ctx.pre}${glslTypeName} ${name} = ${rhsStr};`;
+  }
+
+  override _return(statement: Return): string {
+    const exprNode = statement[1];
+
+    if (exprNode === undefined) {
+      // Default behavior
+      return super._return(statement);
+    }
+
+    if (this.#functionType !== 'normal') {
+      // oxlint-disable-next-line no-non-null-assertion
+      const entryFnState = this.#entryFnState!;
+      const expectedReturnType = this.ctx.topFunctionReturnType;
+
+      if (typeof exprNode === 'object' && exprNode[0] === NODE.objectExpr) {
+        const transformed = Object.entries(exprNode[1]).map(([prop, rhsNode]) => {
+          let name: string | undefined = entryFnState.structPropToVarMap[prop];
+          if (name === undefined) {
+            if (
+              prop === '$position' ||
+              (expectedReturnType &&
+                d.isWgslStruct(expectedReturnType) &&
+                expectedReturnType.propTypes[prop] === d.builtin.position)
+            ) {
+              name = 'gl_Position';
+            } else {
+              name = this.ctx.makeUniqueIdentifier(prop, 'global');
+              entryFnState.outVars.push({ varName: name, propName: prop });
+            }
+            entryFnState.structPropToVarMap[prop] = name;
+          }
+          const rhsExpr = this._expression(rhsNode);
+          const type = rhsExpr.dataType as d.BaseData;
+
+          const snippet = tgpu['~unstable'].rawCodeSnippet(name, type as d.AnyData, 'private');
+
+          return {
+            name,
+            snippet,
+            assignment: [NODE.assignmentExpr, name, '=', rhsNode],
+          } as const;
+        });
+
+        const block = super._block(
+          [NODE.block, [...transformed.map((t) => t.assignment), [NODE.return]]],
+          Object.fromEntries(
+            transformed.map(({ name, snippet }) => {
+              return [name, snippet.$] as const;
+            }),
+          ),
+        );
+
+        return `${this.ctx.pre}${block}`;
+      } else {
+        // Resolving the expression to inspect it's type
+        // We will resolve it again as part of the modifed statement
+        const expr = expectedReturnType
+          ? this._typedExpression(exprNode, expectedReturnType)
+          : this._expression(exprNode);
+
+        if (expr.dataType === UnknownData) {
+          // Unknown data type, don't know what to do
+          return super._return(statement);
+        }
+
+        if (expr.dataType.type.startsWith('vec')) {
+          const block = super._block(
+            [NODE.block, [[NODE.assignmentExpr, 'gl_Position', '=', exprNode], [NODE.return]]],
+            { gl_Position: gl_PositionSnippet.$ },
+          );
+
+          return `${this.ctx.pre}${block}`;
+        }
+      }
+    }
+
+    return super._return(statement);
+  }
+
+  override functionDefinition(options: ShaderGenerator.FunctionDefinitionOptions): string {
+    if (options.functionType !== 'normal') {
+      this.ctx.reserveIdentifier('gl_Position', 'global');
+    }
+
+    // Function body
+    let lastFunctionType = this.#functionType;
+    this.#functionType = options.functionType;
+    if (options.functionType !== 'normal') {
+      if (this.#entryFnState) {
+        throw new Error('Cannot nest entry functions');
+      }
+      this.#entryFnState = { structPropToVarMap: {}, outVars: [] };
+    }
+
+    try {
+      const body = this._block(options.body);
+
+      // Only after generating the body can we determine the return type
+      const returnType = options.determineReturnType();
+
+      if (options.functionType !== 'normal') {
+        // oxlint-disable-next-line no-non-null-assertion
+        const entryFnState = this.#entryFnState!;
+        if (d.isWgslStruct(returnType)) {
+          for (const { varName, propName } of entryFnState.outVars) {
+            const dataType = returnType.propTypes[propName];
+            if (dataType && d.isDecorated(dataType)) {
+              const location = (dataType.attribs as d.AnyAttribute[]).find(
+                (a) => a.type === '@location',
+              )?.params[0];
+              this.ctx.addDeclaration(`layout(location = ${location}) out ${varName};`);
+            }
+          }
+        }
+        return `void main() ${body}`;
+      }
+
+      const argList = options.args
+        // Stripping out unused arguments in entry functions
+        .filter((arg) => arg.used || options.functionType === 'normal')
+        .map((arg) => {
+          return `${this.ctx.resolve(arg.decoratedType).value} ${arg.name}`;
+        })
+        .join(', ');
+
+      return `${this.ctx.resolve(returnType).value} ${options.name}(${argList}) ${body}`;
+    } finally {
+      this.#functionType = lastFunctionType;
+      this.#entryFnState = undefined;
+    }
   }
 }
 
diff --git a/packages/typegpu-gl/src/index.ts b/packages/typegpu-gl/src/index.ts
index ad06f0ab81..5a3dbb29e2 100644
--- a/packages/typegpu-gl/src/index.ts
+++ b/packages/typegpu-gl/src/index.ts
@@ -1,2 +1,3 @@
 export { initWithGL } from './initWithGL.ts';
 export { initWithGLFallback } from './initWithGLFallback.ts';
+export { glOptions } from './glOptions.ts';
diff --git a/packages/typegpu-gl/tests/glslGenerator.test.ts b/packages/typegpu-gl/tests/glslGenerator.test.ts
index 1d8625882c..a396d3df94 100644
--- a/packages/typegpu-gl/tests/glslGenerator.test.ts
+++ b/packages/typegpu-gl/tests/glslGenerator.test.ts
@@ -1,6 +1,7 @@
 import { describe, expect, it } from 'vitest';
 import tgpu, { d } from 'typegpu';
-import glslGenerator, { translateWgslTypeToGlsl } from '../src/glslGenerator.ts';
+import { glOptions } from '@typegpu/gl';
+import { translateWgslTypeToGlsl } from '../src/glslGenerator.ts';
 
 describe('translateWgslTypeToGlsl', () => {
   it('translates scalar types', () => {
@@ -48,18 +49,14 @@ describe('translateWgslTypeToGlsl', () => {
 
 describe('GlslGenerator - variable declarations', () => {
   it('generates GLSL-style variable declarations for JS function', () => {
-    // 'use gpu' function - uses the TGSL code generation path, which calls functionDefinition
-    const fragFn = tgpu.fragmentFn({
-      in: { uv: d.vec2f },
-      out: d.vec4f,
-    })((input) => {
+    const main = () => {
       'use gpu';
       // A variable that uses a vector type
-      const color = d.vec4f(input.uv[0], input.uv[1], 0, 1);
+      const color = d.vec4f(1, 0, 0, 1);
       return color;
-    });
+    };
 
-    const result = tgpu.resolveWithContext([fragFn], { unstable_shaderGenerator: glslGenerator });
+    const result = tgpu.resolveWithContext([main], glOptions());
     // Should contain the resolved function code
     expect(result.code).toBeDefined();
     expect(result.code.length).toBeGreaterThan(0);
@@ -78,28 +75,36 @@ describe('GlslGenerator - variable declarations', () => {
       return d.vec4f(x, 0, 0, 1);
     });
 
-    const result = tgpu.resolveWithContext([fragFn], { unstable_shaderGenerator: glslGenerator });
+    const result = tgpu.resolveWithContext([fragFn], glOptions());
     expect(result.code).toBeDefined();
     // Variable declaration for f32 should be `float`
     expect(result.code).toContain('float ');
   });
 });
 
-describe('GlslGenerator - functionDefinition post-processing', () => {
-  it('translates WGSL type constructor calls in JS function body to GLSL', () => {
-    const fragFn = tgpu.fragmentFn({
-      in: { uv: d.vec2f },
-      out: d.vec4f,
-    })((input) => {
+describe('GlslGenerator - function definitions', () => {
+  it('generates proper function signatures', () => {
+    function add(a: number, b: number) {
       'use gpu';
-      return d.vec4f(input.uv[0], 0.0, 0.0, 1.0);
-    });
+      return a + b;
+    }
 
-    const result = tgpu.resolveWithContext([fragFn], { unstable_shaderGenerator: glslGenerator });
-    // vec4f(...) in the body should become vec4(...)
-    expect(result.code).toContain('vec4(');
-    // Should not contain the WGSL-style constructor in the body
-    expect(result.code).not.toMatch(/\bvec4f\s*\(/);
+    function main() {
+      'use gpu';
+      return add(1.5, 1.2);
+    }
+
+    const result = tgpu.resolveWithContext([main], glOptions());
+
+    expect(result.code).toMatchInlineSnapshot(`
+      "float add(float a, float b) {
+        return (a + b);
+      }
+
+      float main() {
+        return add(1.5f, 1.2f);
+      }"
+    `);
   });
 
   it('translates vec3f to vec3 in function body', () => {
@@ -111,11 +116,44 @@ describe('GlslGenerator - functionDefinition post-processing', () => {
       return d.vec4f(color[0], color[1], color[2], 1.0);
     });
 
-    const result = tgpu.resolveWithContext([fragFn], { unstable_shaderGenerator: glslGenerator });
+    const result = tgpu.resolveWithContext([fragFn], glOptions());
     expect(result.code).toContain('vec3(');
     expect(result.code).not.toMatch(/\bvec3f\s*\(/);
     expect(result.code).toContain('vec4(');
   });
+
+  it('generates proper struct definition', () => {
+    const Boid = d.struct({
+      pos: d.vec3f,
+      vel: d.vec3f,
+    });
+
+    function createBoid() {
+      'use gpu';
+      return Boid({ pos: d.vec3f(), vel: d.vec3f(0, 1, 0) });
+    }
+
+    function main() {
+      'use gpu';
+      const boid = createBoid();
+    }
+
+    const result = tgpu.resolve([main], glOptions());
+    expect(result).toMatchInlineSnapshot(`
+      "struct Boid {
+        vec3 pos;
+        vec3 vel;
+      };
+
+      Boid createBoid() {
+        return Boid(vec3(), vec3(0, 1, 0));
+      }
+
+      void main() {
+        Boid boid = createBoid();
+      }"
+    `);
+  });
 });
 
 describe('GlslGenerator - entry point generation with JS functions', () => {
@@ -127,7 +165,7 @@ describe('GlslGenerator - entry point generation with JS functions', () => {
       return Out({ pos: d.vec4f(0.0, 0.0, 0.0, 1.0) });
     });
 
-    const result = tgpu.resolveWithContext([vertFn], { unstable_shaderGenerator: glslGenerator });
+    const result = tgpu.resolveWithContext([vertFn], glOptions());
     expect(result.code).toBeDefined();
     expect(result.code.length).toBeGreaterThan(0);
     // The body should have translated type names
@@ -136,32 +174,74 @@ describe('GlslGenerator - entry point generation with JS functions', () => {
 
     expect(result.code).toMatchInlineSnapshot(`
       "struct vertFn_Output {
-        @builtin(position) pos: vec4,
-      }
+        vec4 pos;
+      };
 
-      @vertex fn vertFn() -> vertFn_Output {
+      void main() {
         return vertFn_Output(vec4(0, 0, 0, 1));
       }"
     `);
   });
 
+  it('resolves a vertex function returning a builtin and varying', () => {
+    const vertFn = tgpu.vertexFn({
+      out: {
+        position: d.builtin.position,
+        uv: d.vec2f,
+      },
+    })(() => {
+      'use gpu';
+      const position = d.vec4f();
+      const uv = d.vec2f();
+
+      // NOTE: Don't wrap when assigning variables
+      // is valid and allowed at most once
+      return {
+        position: d.vec4f(position),
+        uv: d.vec2f(uv),
+      };
+    });
+
+    const result = tgpu.resolve([vertFn], glOptions());
+
+    expect(result).toMatchInlineSnapshot(`
+      "layout(location = 0) out uv_1;
+
+      void main() {
+        vec4 position = vec4();
+        vec2 uv = vec2();
+        {
+          gl_Position = position;
+          uv_1 = uv;
+          return;
+        }
+      }"
+    `);
+  });
+
   it('resolves a fragment function returning a color using GLSL generator', () => {
     const fragFn = tgpu.fragmentFn({
       out: d.vec4f,
     })(() => {
       'use gpu';
+      // This variable should get renamed to not conflict with
+      // the global.
+      const gl_Position = 1;
       return d.vec4f(1.0, 0.0, 0.0, 1.0);
     });
 
-    const result = tgpu.resolveWithContext([fragFn], { unstable_shaderGenerator: glslGenerator });
+    const result = tgpu.resolveWithContext([fragFn], glOptions());
     expect(result.code).toBeDefined();
-    // The body should have translated vec4f → vec4
     expect(result.code).toContain('vec4(');
     expect(result.code).not.toMatch(/\bvec4f\s*\(/);
 
     expect(result.code).toMatchInlineSnapshot(`
-      "@fragment fn fragFn() -> @location(0) vec4 {
-        return vec4(1, 0, 0, 1);
+      "void main() {
+        int gl_Position_1 = 1;
+        {
+          gl_Position = vec4(1, 0, 0, 1);
+          return;
+        }
       }"
     `);
   });
diff --git a/packages/typegpu/src/core/buffer/bufferUsage.ts b/packages/typegpu/src/core/buffer/bufferUsage.ts
index 24f28dcbca..d8a1028e78 100644
--- a/packages/typegpu/src/core/buffer/bufferUsage.ts
+++ b/packages/typegpu/src/core/buffer/bufferUsage.ts
@@ -118,7 +118,7 @@ class TgpuFixedBufferImpl<TData extends BaseData, TUsage extends BindableBufferU
 
   [$resolve](ctx: ResolutionCtx): ResolvedSnippet {
     const dataType = this.buffer.dataType;
-    const id = ctx.getUniqueName(this);
+    const id = ctx.makeUniqueIdentifier(getName(this), 'global');
     const { group, binding } = ctx.allocateFixedEntry(
       this.usage === 'uniform' ? { uniform: dataType } : { storage: dataType, access: this.usage },
       this.buffer,
@@ -241,7 +241,7 @@ export class TgpuLaidOutBufferImpl<TData extends BaseData, TUsage extends Bindab
   }
 
   [$resolve](ctx: ResolutionCtx): ResolvedSnippet {
-    const id = ctx.getUniqueName(this);
+    const id = ctx.makeUniqueIdentifier(getName(this), 'global');
     const group = ctx.allocateLayoutEntry(this.#membership.layout);
     const usage = usageToVarTemplateMap[this.usage];
 
diff --git a/packages/typegpu/src/core/constant/tgpuConstant.ts b/packages/typegpu/src/core/constant/tgpuConstant.ts
index 4efbf7bef7..3d11ca79bc 100644
--- a/packages/typegpu/src/core/constant/tgpuConstant.ts
+++ b/packages/typegpu/src/core/constant/tgpuConstant.ts
@@ -113,7 +113,7 @@ class TgpuConstImpl<TDataType extends BaseData> implements TgpuConst<TDataType>,
   }
 
   [$resolve](ctx: ResolutionCtx): ResolvedSnippet {
-    const id = ctx.getUniqueName(this);
+    const id = ctx.makeUniqueIdentifier(getName(this), 'global');
     const resolvedDataType = ctx.resolve(this.dataType).value;
     const resolvedValue = ctx.resolve(this.#value, this.dataType).value;
 
diff --git a/packages/typegpu/src/core/function/autoIO.ts b/packages/typegpu/src/core/function/autoIO.ts
index 6feece51d5..b35a092b1a 100644
--- a/packages/typegpu/src/core/function/autoIO.ts
+++ b/packages/typegpu/src/core/function/autoIO.ts
@@ -85,7 +85,7 @@ export class AutoFragmentFn implements SelfResolvable {
     if (!getName(impl)) {
       setName(impl, 'fragmentFn');
     }
-    this.#core = createFnCore(impl, '@fragment ');
+    this.#core = createFnCore(impl, 'fragment');
     this.autoIn = new AutoStruct({ ...builtinFragmentIn, ...varyings }, undefined, locations);
     setName(this.autoIn, 'FragmentIn');
     this.autoOut = new AutoStruct(builtinFragmentOut, vec4f);
@@ -131,7 +131,7 @@ export class AutoVertexFn implements SelfResolvable {
     if (!getName(impl)) {
       setName(impl, 'vertexFn');
     }
-    this.#core = createFnCore(impl, '@vertex ');
+    this.#core = createFnCore(impl, 'vertex');
     this.autoIn = new AutoStruct({ ...builtinVertexIn, ...attribs }, undefined, locations);
     setName(this.autoIn, 'VertexIn');
     this.autoOut = new AutoStruct(builtinVertexOut, undefined);
diff --git a/packages/typegpu/src/core/function/entryInputRouter.ts b/packages/typegpu/src/core/function/entryInputRouter.ts
index 8ba6669372..b756366c4b 100644
--- a/packages/typegpu/src/core/function/entryInputRouter.ts
+++ b/packages/typegpu/src/core/function/entryInputRouter.ts
@@ -1,12 +1,7 @@
-import { undecorate } from '../../data/dataTypes.ts';
-import { snip, type Snippet } from '../../data/snippet.ts';
+import { type Snippet } from '../../data/snippet.ts';
 import { $internal, $repr } from '../../shared/symbols.ts';
 import { type BaseData, isWgslStruct } from '../../data/wgslTypes.ts';
-
-interface PositionalArgEntry {
-  argName: string;
-  type: BaseData;
-}
+import type { FunctionArgumentAccess } from '../../types.ts';
 
 /**
  * Routes `(input) => { input.x }` style property access to the correct WGSL
@@ -18,38 +13,32 @@ export class EntryInputRouter implements BaseData {
   readonly type = 'entry-input-router' as const;
   // Type-token only, not present at runtime:
   declare readonly [$repr]: never;
-  readonly structArgName: string;
-  readonly dataSchema: BaseData | undefined;
+
+  readonly structArg: FunctionArgumentAccess | undefined;
   /** Maps schemaKey → { WGSL arg name, type } */
-  readonly positionalArgsMap: Map<string, PositionalArgEntry>;
+  readonly positionalArgsMap: Map<string, FunctionArgumentAccess>;
 
   constructor(
-    structArgName: string,
-    dataSchema: BaseData | undefined,
-    positionalArgs: { schemaKey: string; argName: string; type: BaseData }[],
+    structArg: FunctionArgumentAccess | undefined,
+    positionalArgs: { schemaKey: string; arg: FunctionArgumentAccess }[],
   ) {
-    this.structArgName = structArgName;
-    this.dataSchema = dataSchema;
-    this.positionalArgsMap = new Map(
-      positionalArgs.map((a) => [a.schemaKey, { argName: a.argName, type: a.type }]),
-    );
+    this.structArg = structArg;
+    this.positionalArgsMap = new Map(positionalArgs.map((a) => [a.schemaKey, a.arg]));
   }
 
   toString(): string {
     return 'entry-input-router';
   }
 
-  accessProp(propName: string): Snippet | undefined {
+  accessProp(propName: string): Snippet | { target: Snippet; prop: string } | undefined {
     const positionalEntry = this.positionalArgsMap.get(propName);
     if (positionalEntry) {
-      return snip(positionalEntry.argName, positionalEntry.type, 'argument');
+      return positionalEntry();
     }
 
-    if (this.dataSchema && isWgslStruct(this.dataSchema)) {
-      const propType = this.dataSchema.propTypes[propName];
-      if (propType) {
-        return snip(`${this.structArgName}.${propName}`, undecorate(propType), 'argument');
-      }
+    const structSnippet = this.structArg?.();
+    if (structSnippet && isWgslStruct(structSnippet.dataType)) {
+      return { target: structSnippet, prop: propName };
     }
 
     return undefined;
diff --git a/packages/typegpu/src/core/function/fnCore.ts b/packages/typegpu/src/core/function/fnCore.ts
index 66429067e5..4d32ccc518 100644
--- a/packages/typegpu/src/core/function/fnCore.ts
+++ b/packages/typegpu/src/core/function/fnCore.ts
@@ -3,9 +3,10 @@ import { undecorate } from '../../data/dataTypes.ts';
 import { type ResolvedSnippet, snip } from '../../data/snippet.ts';
 import { type BaseData, isWgslData, isWgslStruct, Void } from '../../data/wgslTypes.ts';
 import { MissingLinksError } from '../../errors.ts';
+import { isValidIdentifier } from '../../nameUtils.ts';
 import { getMetaData, getName } from '../../shared/meta.ts';
 import { $getNameForward } from '../../shared/symbols.ts';
-import type { ResolutionCtx } from '../../types.ts';
+import type { ResolutionCtx, TgpuShaderStage } from '../../types.ts';
 import { applyExternals, type ExternalMap, replaceExternalsInWgsl } from '../resolve/externals.ts';
 import { extractArgs } from './extractArgs.ts';
 import type { Implementation, SeparatedEntryArgs } from './fnTypes.ts';
@@ -32,7 +33,11 @@ export interface FnCore {
   ): ResolvedSnippet;
 }
 
-export function createFnCore(implementation: Implementation, fnAttribute = ''): FnCore {
+export function createFnCore(
+  implementation: Implementation,
+  functionType: 'normal' | TgpuShaderStage,
+  workgroupSize?: number[],
+): FnCore {
   /**
    * External application has to be deferred until resolution because
    * some externals can reference the owner function which has not been
@@ -57,25 +62,38 @@ export function createFnCore(implementation: Implementation, fnAttribute = ''):
     ): ResolvedSnippet {
       const externalMap: ExternalMap = {};
 
+      let attributes = '';
+      if (functionType === 'compute') {
+        attributes = `@compute @workgroup_size(${workgroupSize?.join(', ')}) `;
+      } else if (functionType === 'vertex') {
+        attributes = `@vertex `;
+      } else if (functionType === 'fragment') {
+        attributes = `@fragment `;
+      }
+
       for (const externals of externalsToApply) {
         applyExternals(externalMap, externals);
       }
 
-      const id = ctx.getUniqueName(this);
+      const id = ctx.makeUniqueIdentifier(getName(this), 'global');
 
       if (typeof implementation === 'string') {
         if (!returnType) {
           throw new Error('Explicit return type is required for string implementation');
         }
 
-        const validArgNames = entryInput
-          ? Object.fromEntries(
-              entryInput.positionalArgs.map((a) => [a.schemaKey, ctx.makeNameValid(a.schemaKey)]),
-            )
-          : undefined;
+        if (entryInput) {
+          for (const arg of entryInput.positionalArgs) {
+            if (!isValidIdentifier(arg.schemaKey)) {
+              throw new Error(`Invalid argument name: ${arg.schemaKey}`);
+            }
+          }
 
-        if (validArgNames && Object.keys(validArgNames).length > 0) {
-          applyExternals(externalMap, { in: validArgNames });
+          applyExternals(externalMap, {
+            in: Object.fromEntries(
+              entryInput.positionalArgs.map((a) => [a.schemaKey, a.schemaKey]),
+            ),
+          });
         }
 
         const replacedImpl = replaceExternalsInWgsl(ctx, externalMap, implementation);
@@ -83,15 +101,15 @@ export function createFnCore(implementation: Implementation, fnAttribute = ''):
         let header = '';
         let body = '';
 
-        if (fnAttribute !== '' && entryInput && validArgNames) {
+        if (functionType !== 'normal' && entryInput) {
           const { dataSchema, positionalArgs } = entryInput;
           const parts: string[] = [];
           if (dataSchema && isArgUsedInBody('in', replacedImpl)) {
             parts.push(`in: ${ctx.resolve(dataSchema).value}`);
           }
           for (const a of positionalArgs) {
-            const argName = validArgNames[a.schemaKey] ?? '';
-            if (argName !== '' && isArgUsedInBody(argName, replacedImpl)) {
+            const argName = a.schemaKey;
+            if (isArgUsedInBody(argName, replacedImpl)) {
               parts.push(`${getAttributesString(a.type)}${argName}: ${ctx.resolve(a.type).value}`);
             }
           }
@@ -140,7 +158,8 @@ export function createFnCore(implementation: Implementation, fnAttribute = ''):
           body = replacedImpl.slice(providedArgs.range.end);
         }
 
-        ctx.addDeclaration(`${fnAttribute}fn ${id}${header}${body}`);
+        ctx.addDeclaration(`${attributes}fn ${id}${header}${body}`);
+
         return snip(id, returnType, /* origin */ 'runtime');
       }
 
@@ -178,7 +197,7 @@ export function createFnCore(implementation: Implementation, fnAttribute = ''):
       // If an entrypoint implementation has a second argument, it represents the output schema.
       // We look at the identifier chosen by the user and add it to externals.
       const maybeSecondArg = ast.params[1];
-      if (maybeSecondArg && maybeSecondArg.type === 'i' && fnAttribute !== '') {
+      if (maybeSecondArg && maybeSecondArg.type === 'i' && functionType !== 'normal') {
         applyExternals(externalMap, {
           // oxlint-disable-next-line typescript/no-non-null-assertion -- entry functions cannot be shellless
           [maybeSecondArg.name]: undecorate(returnType!),
@@ -187,18 +206,10 @@ export function createFnCore(implementation: Implementation, fnAttribute = ''):
 
       // generate wgsl string
 
-      const {
-        head,
-        body,
-        returnType: actualReturnType,
-      } = ctx.fnToWgsl({
-        functionType: fnAttribute.includes('@compute')
-          ? 'compute'
-          : fnAttribute.includes('@vertex')
-            ? 'vertex'
-            : fnAttribute.includes('@fragment')
-              ? 'fragment'
-              : 'normal',
+      const { code, returnType: actualReturnType } = ctx.resolveFunction({
+        functionType,
+        name: id,
+        workgroupSize,
         argTypes,
         entryInput,
         params: ast.params,
@@ -207,9 +218,7 @@ export function createFnCore(implementation: Implementation, fnAttribute = ''):
         externalMap,
       });
 
-      ctx.addDeclaration(
-        `${fnAttribute}fn ${id}${ctx.resolve(head).value}${ctx.resolve(body).value}`,
-      );
+      ctx.addDeclaration(code);
 
       return snip(id, actualReturnType, /* origin */ 'runtime');
     },
diff --git a/packages/typegpu/src/core/function/shelllessImpl.ts b/packages/typegpu/src/core/function/shelllessImpl.ts
index 23b60bdde1..9263303a31 100644
--- a/packages/typegpu/src/core/function/shelllessImpl.ts
+++ b/packages/typegpu/src/core/function/shelllessImpl.ts
@@ -31,7 +31,7 @@ export function createShelllessImpl(
   argTypes: BaseData[],
   implementation: (...args: never[]) => unknown,
 ): ShelllessImpl {
-  const core = createFnCore(implementation, '');
+  const core = createFnCore(implementation, 'normal');
 
   return {
     [$internal]: true,
diff --git a/packages/typegpu/src/core/function/tgpuComputeFn.ts b/packages/typegpu/src/core/function/tgpuComputeFn.ts
index 6f48cbe643..9cef2d49cc 100644
--- a/packages/typegpu/src/core/function/tgpuComputeFn.ts
+++ b/packages/typegpu/src/core/function/tgpuComputeFn.ts
@@ -114,10 +114,7 @@ function createComputeFn<ComputeIn extends IORecord<AnyComputeBuiltin>>(
       [$getNameForward]: FnCore;
     };
 
-  const core = createFnCore(
-    implementation,
-    `@compute @workgroup_size(${workgroupSize.join(', ')}) `,
-  );
+  const core = createFnCore(implementation, 'compute', workgroupSize);
 
   const result: This = {
     shell,
diff --git a/packages/typegpu/src/core/function/tgpuFn.ts b/packages/typegpu/src/core/function/tgpuFn.ts
index 65ea52de65..40f6548922 100644
--- a/packages/typegpu/src/core/function/tgpuFn.ts
+++ b/packages/typegpu/src/core/function/tgpuFn.ts
@@ -173,7 +173,7 @@ function createFn<ImplSchema extends AnyFn>(
     implementation = _implementation;
   }
 
-  const core = createFnCore(implementation as Implementation, '');
+  const core = createFnCore(implementation as Implementation, 'normal');
 
   const fnBase = {
     shell,
diff --git a/packages/typegpu/src/core/function/tgpuFragmentFn.ts b/packages/typegpu/src/core/function/tgpuFragmentFn.ts
index 631617bcdc..cb3b4c990f 100644
--- a/packages/typegpu/src/core/function/tgpuFragmentFn.ts
+++ b/packages/typegpu/src/core/function/tgpuFragmentFn.ts
@@ -183,7 +183,7 @@ function createFragmentFn(
       [$getNameForward]: FnCore;
     };
 
-  const core = createFnCore(implementation, '@fragment ');
+  const core = createFnCore(implementation, 'fragment');
   const outputType = shell.returnType;
   if (typeof implementation === 'string') {
     addReturnTypeToExternals(implementation, outputType, (externals) =>
diff --git a/packages/typegpu/src/core/function/tgpuVertexFn.ts b/packages/typegpu/src/core/function/tgpuVertexFn.ts
index a261634b3d..8c5b230891 100644
--- a/packages/typegpu/src/core/function/tgpuVertexFn.ts
+++ b/packages/typegpu/src/core/function/tgpuVertexFn.ts
@@ -158,7 +158,7 @@ function createVertexFn(
       [$getNameForward]: FnCore;
     };
 
-  const core = createFnCore(implementation, '@vertex ');
+  const core = createFnCore(implementation, 'vertex');
   const entryInput: SeparatedEntryArgs = separateAllAsPositional(shell.in ?? {});
 
   const result: This = {
diff --git a/packages/typegpu/src/core/resolve/namespace.ts b/packages/typegpu/src/core/resolve/namespace.ts
index f112cd4fb0..90fcba08f7 100644
--- a/packages/typegpu/src/core/resolve/namespace.ts
+++ b/packages/typegpu/src/core/resolve/namespace.ts
@@ -1,6 +1,5 @@
 import type { ResolvedSnippet } from '../../data/snippet.ts';
-import { type NameRegistry, RandomNameRegistry, StrictNameRegistry } from '../../nameRegistry.ts';
-import { getName } from '../../shared/meta.ts';
+import { bannedTokens, builtins } from '../../nameUtils.ts';
 import { $internal } from '../../shared/symbols.ts';
 import { ShelllessRepository } from '../../tgsl/shellless.ts';
 import type { TgpuLazy, TgpuSlot } from '../slot/slotTypes.ts';
@@ -8,8 +7,9 @@ import type { TgpuLazy, TgpuSlot } from '../slot/slotTypes.ts';
 type SlotToValueMap = Map<TgpuSlot<unknown>, unknown>;
 
 export interface NamespaceInternal {
-  readonly nameRegistry: NameRegistry;
+  readonly takenGlobalIdentifiers: Set<string>;
   readonly shelllessRepo: ShelllessRepository;
+  readonly strategy: 'random' | 'strict';
 
   memoizedResolves: WeakMap<
     // WeakMap because if the item does not exist anymore,
@@ -24,73 +24,32 @@ export interface NamespaceInternal {
     TgpuLazy<unknown>,
     { slotToValueMap: SlotToValueMap; result: unknown }[]
   >;
-
-  listeners: {
-    [K in keyof NamespaceEventMap]: Set<(event: NamespaceEventMap[K]) => void>;
-  };
 }
 
-type NamespaceEventMap = {
-  name: { target: object; name: string };
-};
-
-type DetachListener = () => void;
-
 export interface Namespace {
   readonly [$internal]: NamespaceInternal;
-
-  on<TEvent extends keyof NamespaceEventMap>(
-    event: TEvent,
-    listener: (event: NamespaceEventMap[TEvent]) => void,
-  ): DetachListener;
 }
 
 class NamespaceImpl implements Namespace {
   readonly [$internal]: NamespaceInternal;
 
-  constructor(nameRegistry: NameRegistry) {
+  constructor(strategy: 'random' | 'strict') {
     this[$internal] = {
-      nameRegistry,
+      strategy,
+      takenGlobalIdentifiers: new Set([...bannedTokens, ...builtins]),
       shelllessRepo: new ShelllessRepository(),
       memoizedResolves: new WeakMap(),
       memoizedLazy: new WeakMap(),
-      listeners: {
-        name: new Set(),
-      },
     };
   }
-
-  on<TEvent extends keyof NamespaceEventMap>(
-    event: TEvent,
-    listener: (event: NamespaceEventMap[TEvent]) => void,
-  ): DetachListener {
-    if (event === 'name') {
-      const listeners = this[$internal].listeners.name;
-      listeners.add(listener);
-
-      return () => listeners.delete(listener);
-    }
-
-    throw new Error(`Unsupported event: ${event}`);
-  }
 }
 
 export interface NamespaceOptions {
   names?: 'random' | 'strict' | undefined;
 }
 
-export function getUniqueName(namespace: NamespaceInternal, resource: object): string {
-  const name = namespace.nameRegistry.makeUnique(getName(resource), true);
-  for (const listener of namespace.listeners.name) {
-    listener({ target: resource, name });
-  }
-  return name;
-}
-
 export function namespace(options?: NamespaceOptions): Namespace {
   const { names = 'strict' } = options ?? {};
 
-  return new NamespaceImpl(
-    names === 'strict' ? new StrictNameRegistry() : new RandomNameRegistry(),
-  );
+  return new NamespaceImpl(names);
 }
diff --git a/packages/typegpu/src/core/resolve/resolveData.ts b/packages/typegpu/src/core/resolve/resolveData.ts
index 19e5da347d..bc5320c80d 100644
--- a/packages/typegpu/src/core/resolve/resolveData.ts
+++ b/packages/typegpu/src/core/resolve/resolveData.ts
@@ -38,6 +38,7 @@ import type {
   WgslArray,
   WgslStruct,
 } from '../../data/wgslTypes.ts';
+import { getName } from '../../shared/meta.ts';
 import { $internal } from '../../shared/symbols.ts';
 import { assertExhaustive } from '../../shared/utilityTypes.ts';
 import type { ResolutionCtx } from '../../types.ts';
@@ -127,7 +128,7 @@ function resolveStruct(ctx: ResolutionCtx, struct: WgslStruct) {
   if (struct[$internal].isAbstruct) {
     throw new Error('Cannot resolve abstract struct types to WGSL.');
   }
-  const id = ctx.getUniqueName(struct);
+  const id = ctx.makeUniqueIdentifier(getName(struct), 'global');
 
   ctx.addDeclaration(`\
 struct ${id} {
@@ -155,7 +156,7 @@ ${Object.entries(struct.propTypes)
  * ```
  */
 function resolveUnstruct(ctx: ResolutionCtx, unstruct: Unstruct) {
-  const id = ctx.getUniqueName(unstruct);
+  const id = ctx.makeUniqueIdentifier(getName(unstruct), 'global');
 
   ctx.addDeclaration(`\
 struct ${id} {
diff --git a/packages/typegpu/src/core/sampler/sampler.ts b/packages/typegpu/src/core/sampler/sampler.ts
index 4daaf1c917..adf10ee8ce 100644
--- a/packages/typegpu/src/core/sampler/sampler.ts
+++ b/packages/typegpu/src/core/sampler/sampler.ts
@@ -99,7 +99,7 @@ export class TgpuLaidOutSamplerImpl<
   }
 
   [$resolve](ctx: ResolutionCtx): ResolvedSnippet {
-    const id = ctx.getUniqueName(this);
+    const id = ctx.makeUniqueIdentifier(getName(this), 'global');
     const group = ctx.allocateLayoutEntry(this.#membership.layout);
 
     ctx.addDeclaration(
@@ -186,7 +186,7 @@ class TgpuFixedSamplerImpl<T extends WgslSampler | WgslComparisonSampler>
   }
 
   [$resolve](ctx: ResolutionCtx): ResolvedSnippet {
-    const id = ctx.getUniqueName(this);
+    const id = ctx.makeUniqueIdentifier(getName(this), 'global');
 
     const { group, binding } = ctx.allocateFixedEntry(
       this.schema.type === 'sampler_comparison'
diff --git a/packages/typegpu/src/core/texture/externalTexture.ts b/packages/typegpu/src/core/texture/externalTexture.ts
index c9d457c48b..78ae4cd582 100644
--- a/packages/typegpu/src/core/texture/externalTexture.ts
+++ b/packages/typegpu/src/core/texture/externalTexture.ts
@@ -37,7 +37,7 @@ export class TgpuExternalTextureImpl implements TgpuExternalTexture, SelfResolva
   }
 
   [$resolve](ctx: ResolutionCtx): ResolvedSnippet {
-    const id = ctx.getUniqueName(this);
+    const id = ctx.makeUniqueIdentifier(getName(this), 'global');
     const group = ctx.allocateLayoutEntry(this.#membership.layout);
 
     ctx.addDeclaration(
diff --git a/packages/typegpu/src/core/texture/texture.ts b/packages/typegpu/src/core/texture/texture.ts
index 27923995d6..fbe7295d83 100644
--- a/packages/typegpu/src/core/texture/texture.ts
+++ b/packages/typegpu/src/core/texture/texture.ts
@@ -600,7 +600,7 @@ class TgpuFixedTextureViewImpl<T extends WgslTexture | WgslStorageTexture>
   }
 
   [$resolve](ctx: ResolutionCtx): ResolvedSnippet {
-    const id = ctx.getUniqueName(this);
+    const id = ctx.makeUniqueIdentifier(getName(this), 'global');
     const { group, binding } = ctx.allocateFixedEntry(
       isWgslStorageTexture(this.schema)
         ? {
@@ -642,7 +642,7 @@ export class TgpuLaidOutTextureViewImpl<T extends WgslTexture | WgslStorageTextu
   }
 
   [$resolve](ctx: ResolutionCtx): ResolvedSnippet {
-    const id = ctx.getUniqueName(this);
+    const id = ctx.makeUniqueIdentifier(getName(this), 'global');
     const group = ctx.allocateLayoutEntry(this.#membership.layout);
 
     ctx.addDeclaration(
diff --git a/packages/typegpu/src/core/variable/tgpuVariable.ts b/packages/typegpu/src/core/variable/tgpuVariable.ts
index 2f34a523bd..b730bfe690 100644
--- a/packages/typegpu/src/core/variable/tgpuVariable.ts
+++ b/packages/typegpu/src/core/variable/tgpuVariable.ts
@@ -87,7 +87,7 @@ class TgpuVarImpl<TScope extends VariableScope, TDataType extends BaseData>
   }
 
   [$resolve](ctx: ResolutionCtx): ResolvedSnippet {
-    const id = ctx.getUniqueName(this);
+    const id = ctx.makeUniqueIdentifier(getName(this), 'global');
     const pre = `var<${this.#scope}> ${id}: ${ctx.resolve(this.#dataType).value}`;
 
     if (this.#initialValue) {
diff --git a/packages/typegpu/src/data/autoStruct.ts b/packages/typegpu/src/data/autoStruct.ts
index 647d13b79c..c227b565fe 100644
--- a/packages/typegpu/src/data/autoStruct.ts
+++ b/packages/typegpu/src/data/autoStruct.ts
@@ -1,5 +1,5 @@
 import { createIoSchema } from '../core/function/ioSchema.ts';
-import { isValidProp } from '../nameRegistry.ts';
+import { isValidProp } from '../nameUtils.ts';
 import { getName, setName } from '../shared/meta.ts';
 import { $internal, $repr, $resolve } from '../shared/symbols.ts';
 import type { ResolutionCtx, SelfResolvable } from '../types.ts';
diff --git a/packages/typegpu/src/data/snippet.ts b/packages/typegpu/src/data/snippet.ts
index fe658ef99d..9c2e954f0d 100644
--- a/packages/typegpu/src/data/snippet.ts
+++ b/packages/typegpu/src/data/snippet.ts
@@ -74,14 +74,9 @@ export interface Snippet {
   readonly origin: Origin;
 }
 
-export interface ResolvedSnippet {
+export interface ResolvedSnippet extends Snippet {
   readonly value: string;
-  /**
-   * The type that `value` is assignable to (not necessary exactly inferred as).
-   * E.g. `1.1` is assignable to `f32`, but `1.1` itself is an abstract float
-   */
   readonly dataType: BaseData;
-  readonly origin: Origin;
 }
 
 export type MapValueToSnippet<T> = { [K in keyof T]: Snippet };
diff --git a/packages/typegpu/src/data/struct.ts b/packages/typegpu/src/data/struct.ts
index c044da5f8f..c2d024acc3 100644
--- a/packages/typegpu/src/data/struct.ts
+++ b/packages/typegpu/src/data/struct.ts
@@ -1,4 +1,4 @@
-import { isValidProp } from '../nameRegistry.ts';
+import { isValidProp } from '../nameUtils.ts';
 import { getName, setName } from '../shared/meta.ts';
 import { $internal } from '../shared/symbols.ts';
 import { schemaCallWrapper } from './schemaCallWrapper.ts';
diff --git a/packages/typegpu/src/nameRegistry.ts b/packages/typegpu/src/nameUtils.ts
similarity index 58%
rename from packages/typegpu/src/nameRegistry.ts
rename to packages/typegpu/src/nameUtils.ts
index aa2cd1ada0..f66bfbb79a 100644
--- a/packages/typegpu/src/nameRegistry.ts
+++ b/packages/typegpu/src/nameUtils.ts
@@ -1,6 +1,4 @@
-import { invariant } from './errors.ts';
-
-const bannedTokens = new Set([
+export const bannedTokens = new Set([
   // keywords
   'alias',
   'break',
@@ -181,7 +179,7 @@ const bannedTokens = new Set([
   'storage',
 ]);
 
-const builtins = new Set([
+export const builtins = new Set([
   // constructors
   'array',
   'bool',
@@ -361,37 +359,8 @@ const builtins = new Set([
   'quadSwapY',
 ]);
 
-export interface NameRegistry {
-  /**
-   * Creates a valid WGSL identifier, each guaranteed to be unique
-   * in the lifetime of a single resolution process
-   * (excluding non-global identifiers from popped scopes).
-   * Should append "_" to primer, followed by some id.
-   * @param primer Used in the generation process, makes the identifier more recognizable.
-   * @param global Whether the name should be registered in the global scope (true), or in the current function scope (false)
-   */
-  makeUnique(primer: string | undefined, global: boolean): string;
-
-  /**
-   * Creates a valid WGSL identifier.
-   * Renames identifiers that are WGSL reserved words.
-   * @param primer Used in the generation process.
-   *
-   * @example
-   * makeValid("notAKeyword"); // "notAKeyword"
-   * makeValid("struct"); // makeUnique("struct")
-   * makeValid("struct_1"); // makeUnique("struct_1") (to avoid potential name collisions)
-   * makeValid("_"); // ERROR (too difficult to make valid to care)
-   */
-  makeValid(primer: string): string;
-
-  pushFunctionScope(): void;
-  popFunctionScope(): void;
-  pushBlockScope(): void;
-  popBlockScope(): void;
-}
-
-function sanitizePrimer(primer: string | undefined) {
+/*#__NO_SIDE_EFFECTS__*/
+export function sanitizePrimer(primer: string | undefined) {
   if (primer) {
     // sanitizing
     return primer
@@ -411,7 +380,8 @@ function sanitizePrimer(primer: string | undefined) {
  * isValidIdentifier("_"); // ERROR
  * isValidIdentifier("my variable"); // ERROR
  */
-function isValidIdentifier(ident: string): boolean {
+/*#__NO_SIDE_EFFECTS__*/
+export function isValidIdentifier(ident: string): boolean {
   if (ident === '_' || ident.startsWith('__') || /\s/.test(ident)) {
     throw new Error(
       `Invalid identifier '${ident}'. Choose an identifier without whitespaces or leading underscores.`,
@@ -424,6 +394,7 @@ function isValidIdentifier(ident: string): boolean {
 /**
  * Same as `isValidIdentifier`, except does not check for builtin clashes.
  */
+/*#__NO_SIDE_EFFECTS__*/
 export function isValidProp(ident: string): boolean {
   if (ident === '_' || ident.startsWith('__') || /\s/.test(ident)) {
     throw new Error(
@@ -433,126 +404,3 @@ export function isValidProp(ident: string): boolean {
   const prefix = ident.split('_')[0] as string;
   return !bannedTokens.has(prefix);
 }
-type FunctionScopeLayer = {
-  type: 'functionScope';
-};
-
-type BlockScopeLayer = {
-  type: 'blockScope';
-  usedBlockScopeNames: Set<string>;
-};
-
-type ScopeLayer = FunctionScopeLayer | BlockScopeLayer;
-
-abstract class NameRegistryImpl implements NameRegistry {
-  abstract getUniqueVariant(base: string): string;
-
-  readonly #usedNames: Set<string>;
-  readonly #scopeStack: ScopeLayer[];
-
-  constructor() {
-    this.#usedNames = new Set<string>([...bannedTokens, ...builtins]);
-    this.#scopeStack = [];
-  }
-
-  get #usedBlockScopeNames(): Set<string> | undefined {
-    return (this.#scopeStack[this.#scopeStack.length - 1] as BlockScopeLayer | undefined)
-      ?.usedBlockScopeNames;
-  }
-
-  makeUnique(primer: string | undefined, global: boolean): string {
-    const sanitizedPrimer = sanitizePrimer(primer);
-    const name = this.getUniqueVariant(sanitizedPrimer);
-
-    if (global) {
-      this.#usedNames.add(name);
-    } else {
-      this.#usedBlockScopeNames?.add(name);
-    }
-
-    return name;
-  }
-
-  #isUsedInBlocksBefore(name: string): boolean {
-    const functionScopeIndex = this.#scopeStack.findLastIndex(
-      (scope) => scope.type === 'functionScope',
-    );
-    return this.#scopeStack
-      .slice(functionScopeIndex + 1)
-      .some((scope) => (scope as BlockScopeLayer).usedBlockScopeNames.has(name));
-  }
-
-  makeValid(primer: string): string {
-    if (
-      isValidIdentifier(primer) &&
-      !this.#usedNames.has(primer) &&
-      !this.#isUsedInBlocksBefore(primer)
-    ) {
-      this.#usedBlockScopeNames?.add(primer);
-      return primer;
-    }
-    return this.makeUnique(primer, false);
-  }
-
-  isUsed(name: string): boolean {
-    return this.#usedNames.has(name) || this.#isUsedInBlocksBefore(name);
-  }
-
-  pushFunctionScope(): void {
-    this.#scopeStack.push({ type: 'functionScope' });
-    this.#scopeStack.push({
-      type: 'blockScope',
-      usedBlockScopeNames: new Set(),
-    });
-  }
-
-  popFunctionScope(): void {
-    const functionScopeIndex = this.#scopeStack.findLastIndex(
-      (scope) => scope.type === 'functionScope',
-    );
-
-    if (functionScopeIndex === -1) {
-      throw new Error('Tried to pop function scope when no scope was present.');
-    }
-
-    this.#scopeStack.splice(functionScopeIndex);
-  }
-
-  pushBlockScope(): void {
-    this.#scopeStack.push({
-      type: 'blockScope',
-      usedBlockScopeNames: new Set(),
-    });
-  }
-  popBlockScope(): void {
-    invariant(
-      this.#scopeStack[this.#scopeStack.length - 1]?.type === 'blockScope',
-      'Tried to pop block scope, but it is not present',
-    );
-    this.#scopeStack.pop();
-  }
-}
-
-export class RandomNameRegistry extends NameRegistryImpl {
-  #lastUniqueId = 0;
-
-  getUniqueVariant(base: string): string {
-    let name = `${base}_${this.#lastUniqueId++}`;
-    while (this.isUsed(name)) {
-      name = `${base}_${this.#lastUniqueId++}`;
-    }
-    return name;
-  }
-}
-
-export class StrictNameRegistry extends NameRegistryImpl {
-  getUniqueVariant(base: string): string {
-    let index = 0;
-    let name = base;
-    while (this.isUsed(name)) {
-      index++;
-      name = `${base}_${index}`;
-    }
-    return name;
-  }
-}
diff --git a/packages/typegpu/src/resolutionCtx.ts b/packages/typegpu/src/resolutionCtx.ts
index 5b3daa0067..ced6abe106 100644
--- a/packages/typegpu/src/resolutionCtx.ts
+++ b/packages/typegpu/src/resolutionCtx.ts
@@ -1,5 +1,5 @@
 import { isTgpuFn } from './core/function/tgpuFn.ts';
-import { getUniqueName, type Namespace, type NamespaceInternal } from './core/resolve/namespace.ts';
+import type { Namespace, NamespaceInternal } from './core/resolve/namespace.ts';
 import { stitch } from './core/resolve/stitch.ts';
 import { ConfigurableImpl } from './core/root/configurableImpl.ts';
 import type { Configurable, ExperimentalTgpuRoot } from './core/root/rootTypes.ts';
@@ -12,10 +12,9 @@ import {
   type TgpuLazy,
   type TgpuSlot,
 } from './core/slot/slotTypes.ts';
-import { getAttributesString } from './data/attributes.ts';
-import { isData, undecorate, UnknownData } from './data/dataTypes.ts';
+import { isData, UnknownData } from './data/dataTypes.ts';
 import { bool } from './data/numeric.ts';
-import { type ResolvedSnippet, snip, type Snippet } from './data/snippet.ts';
+import { type Origin, type ResolvedSnippet, snip, type Snippet } from './data/snippet.ts';
 import { type BaseData, isPtr, isWgslArray, isWgslStruct, Void } from './data/wgslTypes.ts';
 import { invariant, MissingSlotValueError, ResolutionError, WgslTypeError } from './errors.ts';
 import { provideCtx, topLevelState } from './execMode.ts';
@@ -38,9 +37,11 @@ import { coerceToSnippet, concretize, numericLiteralToSnippet } from './tgsl/gen
 import type { ShaderGenerator } from './tgsl/shaderGenerator.ts';
 import wgslGenerator from './tgsl/wgslGenerator.ts';
 import type {
+  BlockScopeLayer,
   ExecMode,
   ExecState,
-  FnToWgslOptions,
+  ResolveFunctionOptions,
+  FunctionArgumentAccess,
   FunctionScopeLayer,
   ItemLayer,
   ItemStateStack,
@@ -58,6 +59,8 @@ import { createIoSchema } from './core/function/ioSchema.ts';
 import type { IOData } from './core/function/fnTypes.ts';
 import { AutoStruct } from './data/autoStruct.ts';
 import { EntryInputRouter } from './core/function/entryInputRouter.ts';
+import type { FunctionArgument } from './tgsl/shaderGenerator_members.ts';
+import { isValidIdentifier, sanitizePrimer } from './nameUtils.ts';
 
 /**
  * Inserted into bind group entry definitions that belong
@@ -98,6 +101,10 @@ class ItemStateStackImpl implements ItemStateStack {
     return this._stack.findLast((e) => e.type === 'functionScope');
   }
 
+  get topBlockScope(): BlockScopeLayer | undefined {
+    return this._stack.findLast((e) => e.type === 'blockScope');
+  }
+
   pushItem() {
     this._itemDepth++;
     this._stack.push({
@@ -115,16 +122,14 @@ class ItemStateStackImpl implements ItemStateStack {
 
   pushFunctionScope(
     functionType: 'normal' | TgpuShaderStage,
-    args: Snippet[],
-    argAliases: Record<string, Snippet>,
+    argAccess: Record<string, FunctionArgumentAccess>,
     returnType: BaseData | undefined,
     externalMap: Record<string, unknown>,
   ): FunctionScopeLayer {
     const scope: FunctionScopeLayer = {
       type: 'functionScope',
       functionType,
-      args,
-      argAliases,
+      argAccess,
       returnType,
       externalMap,
       reportedReturnTypes: new Set(),
@@ -137,6 +142,7 @@ class ItemStateStackImpl implements ItemStateStack {
   pushBlockScope() {
     this._stack.push({
       type: 'blockScope',
+      takenLocalIdentifiers: new Set(),
       declarations: new Map(),
       externals: new Map(),
     });
@@ -184,13 +190,9 @@ class ItemStateStackImpl implements ItemStateStack {
       const layer = this._stack[i];
 
       if (layer?.type === 'functionScope') {
-        const arg = layer.args.find((a) => a.value === id);
-        if (arg !== undefined) {
-          return arg;
-        }
-
-        if (layer.argAliases[id]) {
-          return layer.argAliases[id];
+        const access = layer.argAccess[id];
+        if (access) {
+          return access();
         }
 
         const external = layer.externalMap[id];
@@ -218,6 +220,26 @@ class ItemStateStackImpl implements ItemStateStack {
     return undefined;
   }
 
+  isIdentifierTakenLocally(id: string): boolean {
+    for (let i = this._stack.length - 1; i >= 0; --i) {
+      const layer = this._stack[i];
+
+      if (layer?.type === 'functionScope') {
+        // Since functions cannot access resources from the calling scope, we
+        // return early here.
+        return false;
+      }
+
+      if (layer?.type === 'blockScope') {
+        if (layer.takenLocalIdentifiers.has(id)) {
+          return true;
+        }
+      }
+    }
+
+    return false;
+  }
+
   defineBlockVariable(id: string, snippet: Snippet): void {
     if (snippet.dataType === UnknownData) {
       throw Error(`Tried to define variable '${id}' of unknown type`);
@@ -311,6 +333,39 @@ interface FixedBindingConfig {
   resource: object;
 }
 
+function createArgument(
+  name: string,
+  type: BaseData,
+  origin: Origin = 'argument',
+): FunctionArgument {
+  let used = false;
+
+  return {
+    name,
+    access: () => {
+      used = true;
+      return snip(name, type, origin);
+    },
+    decoratedType: type,
+    get used() {
+      return used;
+    },
+  };
+}
+
+function createArgumentPropAccess(
+  argAccess: FunctionArgumentAccess,
+  prop: string,
+): FunctionArgumentAccess {
+  return () => {
+    const argSnippet = argAccess();
+    if (!argSnippet) {
+      return undefined;
+    }
+    return accessProp(argSnippet, prop);
+  };
+}
+
 export class ResolutionCtxImpl implements ResolutionCtx {
   readonly #namespaceInternal: NamespaceInternal;
 
@@ -347,6 +402,11 @@ export class ResolutionCtxImpl implements ResolutionCtx {
   public readonly enableExtensions: WgslExtension[] | undefined;
   public expectedType: BaseData | undefined;
 
+  /**
+   * A counter used to generate unique identifiers for globally-scoped definitions in the 'random' strategy.
+   */
+  #lastUniqueId = 0;
+
   constructor(opts: ResolutionCtxImplOptions) {
     this.enableExtensions = opts.enableExtensions;
     this.gen = opts.shaderGenerator ?? wgslGenerator;
@@ -354,12 +414,42 @@ export class ResolutionCtxImpl implements ResolutionCtx {
     this.#namespaceInternal = opts.namespace[$internal];
   }
 
-  getUniqueName(resource: object): string {
-    return getUniqueName(this.#namespaceInternal, resource);
+  isIdentifierTaken(name: string): boolean {
+    return (
+      this.#namespaceInternal.takenGlobalIdentifiers.has(name) ||
+      this._itemStateStack.isIdentifierTakenLocally(name)
+    );
+  }
+
+  makeUniqueIdentifier(primer: string = 'item', scope: 'global' | 'block'): string {
+    if (scope === 'block' && isValidIdentifier(primer) && !this.isIdentifierTaken(primer)) {
+      // Preserving local definitions as they are, provided they are valid and not already taken.
+      this.reserveIdentifier(primer, 'block');
+      return primer;
+    }
+
+    const base = sanitizePrimer(primer);
+    let index = 0;
+    const random = this.#namespaceInternal.strategy === 'random';
+    let name = random ? `${base}_${this.#lastUniqueId++}` : base;
+    while (this.isIdentifierTaken(name)) {
+      name = random ? `${base}_${this.#lastUniqueId++}` : `${base}_${++index}`;
+    }
+
+    this.reserveIdentifier(name, scope);
+    return name;
   }
 
-  makeNameValid(name: string): string {
-    return this.#namespaceInternal.nameRegistry.makeValid(name);
+  reserveIdentifier(name: string, scope: 'global' | 'block'): void {
+    if (scope === 'block') {
+      const blockScope = this._itemStateStack.topBlockScope;
+      if (blockScope) {
+        blockScope.takenLocalIdentifiers.add(name);
+        return;
+      }
+      // Fall through if no block scope is present, treating as global.
+    }
+    this.#namespaceInternal.takenGlobalIdentifiers.add(name);
   }
 
   get pre(): string {
@@ -413,12 +503,10 @@ export class ResolutionCtxImpl implements ResolutionCtx {
   }
 
   pushBlockScope() {
-    this.#namespaceInternal.nameRegistry.pushBlockScope();
     this._itemStateStack.pushBlockScope();
   }
 
   popBlockScope() {
-    this.#namespaceInternal.nameRegistry.popBlockScope();
     this._itemStateStack.pop('blockScope');
   }
 
@@ -438,28 +526,29 @@ export class ResolutionCtxImpl implements ResolutionCtx {
     return this.#logGenerator.logResources;
   }
 
-  fnToWgsl(options: FnToWgslOptions): { head: Wgsl; body: Wgsl; returnType: BaseData } {
-    let fnScopePushed = false;
-
+  resolveFunction(options: ResolveFunctionOptions): { code: string; returnType: BaseData } {
     try {
-      this.#namespaceInternal.nameRegistry.pushFunctionScope();
-      const args: Snippet[] = [];
-      const argAliases: [string, Snippet][] = [];
-      // For entry functions: collect pending header entries to be filtered after body generation.
-      const pendingHeaderEntries: { argName: string; header: string }[] = [];
+      const scope = this._itemStateStack.pushFunctionScope(
+        options.functionType,
+        {},
+        options.returnType,
+        options.externalMap,
+      );
+      // Pushing a block scope as well, so that any identifiers declared at this point will be scoped to the function body.
+      this._itemStateStack.pushBlockScope();
+
+      const args: FunctionArgument[] = [];
 
       if (options.entryInput) {
         const { dataSchema, positionalArgs } = options.entryInput;
         const firstParam = options.params[0];
 
-        const structArgName = this.makeNameValid('_arg_0');
-        const structArg = dataSchema ? snip(structArgName, dataSchema, 'argument') : undefined;
+        const structArg = dataSchema
+          ? createArgument(this.makeUniqueIdentifier('_arg_0', 'block'), dataSchema)
+          : undefined;
+
         if (structArg) {
           args.push(structArg);
-          pendingHeaderEntries.push({
-            argName: structArgName,
-            header: `${structArgName}: ${this.resolve(dataSchema).value}`,
-          });
         }
 
         if (firstParam?.type === FuncParameterType.destructuredObject) {
@@ -467,45 +556,31 @@ export class ResolutionCtxImpl implements ResolutionCtx {
           for (const { name, alias } of firstParam.props) {
             const argInfo = positionalArgs.find((a) => a.schemaKey === name);
             if (argInfo) {
-              const argName = this.makeNameValid(alias);
-              const argSnippet = snip(argName, argInfo.type, 'argument');
-              args.push(argSnippet);
-              argAliases.push([alias, argSnippet]);
-              pendingHeaderEntries.push({
-                argName,
-                header: `${getAttributesString(argInfo.type)}${argName}: ${this.resolve(undecorate(argInfo.type)).value}`,
-              });
+              const arg = createArgument(this.makeUniqueIdentifier(alias, 'block'), argInfo.type);
+              args.push(arg);
+              scope.argAccess[alias] = arg.access;
             } else if (structArg) {
-              const propSnippet = accessProp(structArg, name);
-              if (propSnippet) {
-                argAliases.push([alias, propSnippet]);
-              }
+              scope.argAccess[alias] = createArgumentPropAccess(structArg.access, name);
             }
           }
         } else if (firstParam?.type === FuncParameterType.identifier) {
           // Create named arg snippets, then a proxy for property access routing.
-          const proxyEntries: Array<{ schemaKey: string; argName: string; type: BaseData }> = [];
+          const proxyEntries: Array<{ schemaKey: string; arg: FunctionArgumentAccess }> = [];
           for (const a of positionalArgs) {
-            const argName = this.makeNameValid(`_arg_${a.schemaKey}`);
-            const s = snip(argName, a.type, 'argument');
-            args.push(s);
-            proxyEntries.push({ schemaKey: a.schemaKey, argName, type: a.type });
-            pendingHeaderEntries.push({
-              argName,
-              header: `${getAttributesString(a.type)}${argName}: ${this.resolve(undecorate(a.type)).value}`,
-            });
+            const argName = this.makeUniqueIdentifier(`_arg_${a.schemaKey}`, 'block');
+            const arg = createArgument(argName, a.type);
+            args.push(arg);
+            proxyEntries.push({ schemaKey: a.schemaKey, arg: arg.access });
           }
-          const router = new EntryInputRouter(structArgName, dataSchema, proxyEntries);
-          argAliases.push([firstParam.name, snip(firstParam.name, router, 'argument')]);
+          const router = new EntryInputRouter(structArg?.access, proxyEntries);
+          scope.argAccess[firstParam.name] = () => snip('N/A', router, 'argument');
         } else {
           // No first param: push positional args with schema key names.
           for (const a of positionalArgs) {
-            const argName = this.makeNameValid(`_arg_${a.schemaKey}`);
-            args.push(snip(argName, a.type, 'argument'));
-            pendingHeaderEntries.push({
-              argName,
-              header: `${getAttributesString(a.type)}${argName}: ${this.resolve(undecorate(a.type)).value}`,
-            });
+            const argName = this.makeUniqueIdentifier(`_arg_${a.schemaKey}`, 'block');
+            const arg = createArgument(argName, a.type);
+            args.push(arg);
+            scope.argAccess[argName] = arg.access;
           }
         }
       } else {
@@ -528,22 +603,25 @@ export class ResolutionCtxImpl implements ResolutionCtx {
 
           switch (astParam?.type) {
             case FuncParameterType.identifier: {
-              const rawName = astParam.name;
-              const snippet = snip(this.makeNameValid(rawName), argType, origin);
-              args.push(snippet);
-              if (snippet.value !== rawName) {
-                argAliases.push([rawName, snippet]);
-              }
+              const arg = createArgument(
+                this.makeUniqueIdentifier(astParam.name, 'block'),
+                argType,
+                origin,
+              );
+              args.push(arg);
+              scope.argAccess[astParam.name] = arg.access;
               break;
             }
             case FuncParameterType.destructuredObject: {
-              const objSnippet = snip(`_arg_${i}`, argType, origin);
-              args.push(objSnippet);
-              argAliases.push(
-                ...astParam.props.map(
-                  ({ name, alias }) => [alias, accessProp(objSnippet, name)] as [string, Snippet],
-                ),
+              const objArg = createArgument(
+                this.makeUniqueIdentifier(`_arg_${i}`, 'block'),
+                argType,
+                origin,
               );
+              args.push(objArg);
+              for (const { name, alias } of astParam.props) {
+                scope.argAccess[alias] = createArgumentPropAccess(objArg.access, name);
+              }
               break;
             }
             case undefined: {
@@ -551,84 +629,88 @@ export class ResolutionCtxImpl implements ResolutionCtx {
               // If we're not using an auto-struct, it's not going to
               // have any properties anyway.
               if (!(argType instanceof AutoStruct)) {
-                args.push(snip(`_arg_${i}`, argType, origin));
+                args.push({
+                  name: this.makeUniqueIdentifier(`_arg_${i}`, 'block'),
+                  access: () => {
+                    throw new Error(
+                      `Unreachable: Accessing an argument that wasn't named in the function signature`,
+                    );
+                  },
+                  decoratedType: argType,
+                  used: false,
+                });
               }
             }
           }
         }
       }
 
-      const scope = this._itemStateStack.pushFunctionScope(
-        options.functionType,
-        args,
-        Object.fromEntries(argAliases),
-        options.returnType,
-        options.externalMap,
-      );
-      fnScopePushed = true;
+      let returnType: BaseData | undefined;
 
-      const body = this.gen.functionDefinition(options.body);
-
-      let returnType = options.returnType;
-      if (returnType instanceof AutoStruct) {
-        // We're expecting an "auto" return type, so if there were structs returned,
-        // we accept the struct, otherwise we let the rest of the code unify on a
-        // primitive type.
-        if (isWgslStruct(scope.reportedReturnTypes.values().next().value)) {
-          returnType = returnType.completeStruct;
-        } else {
-          returnType = undefined;
-        }
-      }
+      const code = this.gen.functionDefinition({
+        functionType: options.functionType,
+        name: options.name,
+        workgroupSize: options.workgroupSize,
+        args,
+        body: options.body,
+        determineReturnType: () => {
+          if (returnType) {
+            // Already determined
+            return returnType;
+          }
 
-      if (!returnType) {
-        const returnTypes = [...scope.reportedReturnTypes];
-        if (returnTypes.length === 0) {
-          returnType = Void;
-        } else {
-          const conversion = getBestConversion(returnTypes);
-          if (conversion && !conversion.hasImplicitConversions) {
-            returnType = conversion.targetType;
+          returnType = options.returnType;
+          if (returnType instanceof AutoStruct) {
+            // We're expecting an "auto" return type, so if there were structs returned,
+            // we accept the struct, otherwise we let the rest of the code unify on a
+            // primitive type.
+            if (isWgslStruct(scope.reportedReturnTypes.values().next().value)) {
+              returnType = returnType.completeStruct;
+            } else {
+              returnType = undefined;
+            }
           }
-        }
 
-        if (!returnType) {
-          throw new Error(
-            `Expected function to have a single return type, got [${returnTypes.join(
-              ', ',
-            )}]. Cast explicitly to the desired type.`,
-          );
-        }
+          if (!returnType) {
+            const returnTypes = [...scope.reportedReturnTypes];
+            if (returnTypes.length === 0) {
+              returnType = Void;
+            } else {
+              const conversion = getBestConversion(returnTypes);
+              if (conversion && !conversion.hasImplicitConversions) {
+                returnType = conversion.targetType;
+              }
+            }
 
-        returnType = concretize(returnType);
+            if (!returnType) {
+              throw new Error(
+                `Expected function to have a single return type, got [${returnTypes.join(
+                  ', ',
+                )}]. Cast explicitly to the desired type.`,
+              );
+            }
 
-        if (options.functionType === 'vertex' || options.functionType === 'fragment') {
-          returnType = createIoSchema(returnType as IOData);
-        }
-      }
+            returnType = concretize(returnType);
 
-      if (options.entryInput) {
-        const headerParts = pendingHeaderEntries
-          .filter(({ argName }) => isArgUsedInBody(argName, body))
-          .map(({ header }) => header);
-        const argList = headerParts.join(', ');
-        const returnStr =
-          returnType.type !== 'void'
-            ? `-> ${getAttributesString(returnType)}${this.resolve(returnType).value} `
-            : '';
-        return { head: `(${argList}) ${returnStr}`, body, returnType };
+            if (options.functionType === 'vertex' || options.functionType === 'fragment') {
+              returnType = createIoSchema(returnType as IOData);
+            }
+          }
+          return returnType;
+        },
+      });
+
+      if (!returnType) {
+        throw new Error(`Failed to determine return type`);
       }
 
       return {
-        head: resolveFunctionHeader(this, args, returnType),
-        body,
+        code,
         returnType,
       };
     } finally {
-      if (fnScopePushed) {
-        this._itemStateStack.pop('functionScope');
-      }
-      this.#namespaceInternal.nameRegistry.popFunctionScope();
+      this._itemStateStack.pop('blockScope');
+      this._itemStateStack.pop('functionScope');
     }
   }
 
@@ -1079,17 +1161,3 @@ export function resolve(item: Wgsl, options: ResolutionCtxImplOptions): Resoluti
     logResources: ctx.logResources,
   };
 }
-
-function isArgUsedInBody(argName: string, body: string): boolean {
-  return new RegExp(`\\b${argName}\\b`).test(body);
-}
-
-function resolveFunctionHeader(ctx: ResolutionCtx, args: Snippet[], returnType: BaseData) {
-  const argList = args
-    .map((arg) => `${arg.value}: ${ctx.resolve(arg.dataType as BaseData).value}`)
-    .join(', ');
-
-  return returnType.type !== 'void'
-    ? `(${argList}) -> ${getAttributesString(returnType)}${ctx.resolve(returnType).value} `
-    : `(${argList}) `;
-}
diff --git a/packages/typegpu/src/tgsl/accessProp.ts b/packages/typegpu/src/tgsl/accessProp.ts
index b264a154ba..3410d3e383 100644
--- a/packages/typegpu/src/tgsl/accessProp.ts
+++ b/packages/typegpu/src/tgsl/accessProp.ts
@@ -10,7 +10,7 @@ import {
 } from '../data/dataTypes.ts';
 import { abstractInt, bool, f16, f32, i32, u32 } from '../data/numeric.ts';
 import { derefSnippet } from '../data/ref.ts';
-import { isEphemeralSnippet, snip, type Snippet } from '../data/snippet.ts';
+import { isEphemeralSnippet, isSnippet, snip, type Snippet } from '../data/snippet.ts';
 import {
   vec2b,
   vec2f,
@@ -160,7 +160,14 @@ export function accessProp(target: Snippet, propName: string): Snippet | undefin
   }
 
   if (target.dataType instanceof EntryInputRouter) {
-    return target.dataType.accessProp(propName);
+    const result = target.dataType.accessProp(propName);
+    if (isSnippet(result)) {
+      return result;
+    }
+    if (result) {
+      return accessProp(result.target, result.prop);
+    }
+    return undefined;
   }
 
   if (isPtr(target.dataType)) {
diff --git a/packages/typegpu/src/tgsl/conversion.ts b/packages/typegpu/src/tgsl/conversion.ts
index f410e53e6d..b9525ce87e 100644
--- a/packages/typegpu/src/tgsl/conversion.ts
+++ b/packages/typegpu/src/tgsl/conversion.ts
@@ -308,7 +308,7 @@ export function convertToCommonType<T extends Snippet[]>(
   if ((TEST || DEV) && verbose && conversion.hasImplicitConversions) {
     console.warn(
       `Implicit conversions from [\n${values
-        .map((v) => `  ${v.value}: ${safeStringify(v.dataType)}`)
+        .map((v) => `  ${ctx.resolveSnippet(v).value}: ${safeStringify(v.dataType)}`)
         .join(',\n')}\n] to ${conversion.targetType.type} are supported, but not recommended.
 Consider using explicit conversions instead.`,
     );
diff --git a/packages/typegpu/src/tgsl/makeDereferencable.ts b/packages/typegpu/src/tgsl/makeDereferencable.ts
new file mode 100644
index 0000000000..4c2e56b325
--- /dev/null
+++ b/packages/typegpu/src/tgsl/makeDereferencable.ts
@@ -0,0 +1,49 @@
+import { snip, type Origin } from '../data/snippet.ts';
+import type { BaseData } from '../data/wgslTypes.ts';
+import { $gpuValueOf, $internal, $ownSnippet, $resolve } from '../shared/symbols.ts';
+import { valueProxyHandler } from '../core/valueProxyUtils.ts';
+import type { SelfResolvable } from '../types.ts';
+import { inCodegenMode } from '../execMode.ts';
+
+export function makeDereferencable<T extends SelfResolvable>(
+  value: T,
+  options: makeDereferencable.Options<T>,
+): T {
+  Object.defineProperty(value, $gpuValueOf, {
+    get() {
+      const [dataType, origin] = options.getDataTypeAndOrigin.apply(this);
+
+      return new Proxy(
+        {
+          [$internal]: true,
+          get [$ownSnippet]() {
+            return snip(this, dataType, origin);
+          },
+          [$resolve]: (ctx) => ctx.resolve(this),
+          toString: () => `${this.toString()}.$`,
+        },
+        valueProxyHandler,
+      );
+    },
+  });
+
+  Object.defineProperty(value, '$', {
+    get() {
+      if (inCodegenMode()) {
+        return this[$gpuValueOf];
+      }
+      // TODO: Add proper error message
+      throw new Error(
+        'Cannot read WebGL uniform outside of shader code. Use `.write()` to update it.',
+      );
+    },
+  });
+
+  return value;
+}
+
+export namespace makeDereferencable {
+  export interface Options<T> {
+    getDataTypeAndOrigin(this: T): [dataType: BaseData, origin: Origin];
+  }
+}
diff --git a/packages/typegpu/src/tgsl/makeResolvable.ts b/packages/typegpu/src/tgsl/makeResolvable.ts
new file mode 100644
index 0000000000..5fcacc8c46
--- /dev/null
+++ b/packages/typegpu/src/tgsl/makeResolvable.ts
@@ -0,0 +1,52 @@
+import { snip, type Origin, type ResolvedSnippet } from '../data/snippet.ts';
+import { type BaseData } from '../data/wgslTypes.ts';
+import { $internal, $resolve, isMarkedInternal } from '../shared/symbols.ts';
+import type { ResolutionCtx, SelfResolvable } from '../types.ts';
+
+/**
+ *
+ * @param value
+ * @param options
+ * @returns
+ */
+export function makeResolvable<T>(
+  value: T,
+  options: makeResolvable.Options<T>,
+): T & SelfResolvable {
+  if (!isMarkedInternal(value)) {
+    Object.defineProperty(value, $internal, {
+      value: true,
+    });
+  }
+
+  Object.defineProperty(value, 'toString', {
+    value() {
+      return options.asString.apply(this);
+    },
+  });
+
+  Object.defineProperty(value, $resolve, {
+    value(ctx: ResolutionCtx): ResolvedSnippet {
+      const protoSnippet = options.resolve.apply(this, [ctx]);
+
+      return snip(protoSnippet.value, protoSnippet.dataType, protoSnippet.origin);
+    },
+  });
+
+  return value as T & SelfResolvable;
+}
+
+interface ProtoSnippet {
+  value: string;
+  dataType: BaseData;
+  origin: Origin;
+}
+
+export namespace makeResolvable {
+  export interface Options<T> {
+    resolve(this: T, ctx: ResolutionCtx): ProtoSnippet;
+    asString(this: T): string;
+  }
+
+  export type Resolvable = SelfResolvable;
+}
diff --git a/packages/typegpu/src/tgsl/shaderGenerator.ts b/packages/typegpu/src/tgsl/shaderGenerator.ts
index e482b03cc8..adde48bf00 100644
--- a/packages/typegpu/src/tgsl/shaderGenerator.ts
+++ b/packages/typegpu/src/tgsl/shaderGenerator.ts
@@ -1,7 +1,7 @@
-import type { Block } from 'tinyest';
 import type { BaseData } from '../data/wgslTypes.ts';
 import type { GenerationCtx } from './generationHelpers.ts';
 import type { ResolvedSnippet, Snippet } from '../data/snippet.ts';
+import type { FunctionDefinitionOptions } from './shaderGenerator_members.ts';
 
 /**
  * **NOTE: This is an unstable API and may change in the future.**
@@ -12,7 +12,7 @@ import type { ResolvedSnippet, Snippet } from '../data/snippet.ts';
 export interface ShaderGenerator {
   initGenerator(ctx: GenerationCtx): void;
 
-  functionDefinition(body: Block): string;
+  functionDefinition(options: FunctionDefinitionOptions): string;
   typeInstantiation(schema: BaseData, args: readonly Snippet[]): ResolvedSnippet;
   typeAnnotation(schema: BaseData): string;
 }
diff --git a/packages/typegpu/src/tgsl/shaderGenerator_members.ts b/packages/typegpu/src/tgsl/shaderGenerator_members.ts
index 1eb11ae0c9..5fa67ce0c6 100644
--- a/packages/typegpu/src/tgsl/shaderGenerator_members.ts
+++ b/packages/typegpu/src/tgsl/shaderGenerator_members.ts
@@ -1,6 +1,23 @@
+import type { Block } from 'tinyest';
+import type { BaseData } from '../data/wgslTypes.ts';
+import type { FunctionArgument, TgpuShaderStage } from '../types.ts';
+
 export { UnknownData } from '../data/dataTypes.ts';
+export { getName } from '../shared/meta.ts';
+export { makeDereferencable } from './makeDereferencable.ts';
+export { makeResolvable } from './makeResolvable.ts';
 
 // types
-export type { ResolutionCtx } from '../types.ts';
+export type { ResolutionCtx, FunctionArgument, TgpuShaderStage } from '../types.ts';
 export type { Snippet } from '../data/snippet.ts';
 export type { Origin } from '../data/snippet.ts';
+
+export interface FunctionDefinitionOptions {
+  readonly functionType: 'normal' | TgpuShaderStage;
+  readonly name: string;
+  readonly workgroupSize?: readonly number[] | undefined;
+  readonly args: readonly FunctionArgument[];
+  readonly body: Block;
+
+  determineReturnType(): BaseData;
+}
diff --git a/packages/typegpu/src/tgsl/wgslGenerator.ts b/packages/typegpu/src/tgsl/wgslGenerator.ts
index 38d2a2ee44..fbd00724db 100644
--- a/packages/typegpu/src/tgsl/wgslGenerator.ts
+++ b/packages/typegpu/src/tgsl/wgslGenerator.ts
@@ -51,6 +51,8 @@ import { mathToStd } from './math.ts';
 import type { ExternalMap } from '../core/resolve/externals.ts';
 import * as forOfUtils from './forOfUtils.ts';
 import { isTgpuRange } from '../std/range.ts';
+import type { FunctionDefinitionOptions } from './shaderGenerator_members.ts';
+import { getAttributesString } from '../data/attributes.ts';
 
 const { NodeTypeCatalog: NODE } = tinyest;
 
@@ -214,7 +216,7 @@ export class WgslGenerator implements ShaderGenerator {
     return this.#ctx;
   }
 
-  public _block([_, statements]: tinyest.Block, externalMap?: ExternalMap): string {
+  protected _block([_, statements]: tinyest.Block, externalMap?: ExternalMap): string {
     this.ctx.pushBlockScope();
 
     if (externalMap) {
@@ -239,8 +241,12 @@ ${this.ctx.pre}}`;
     }
   }
 
+  protected _blockStatement(block: tinyest.Block, externalMap?: ExternalMap): string {
+    return `${this.ctx.pre}${this._block(block, externalMap)}`;
+  }
+
   public refVariable(id: string, dataType: wgsl.StorableData): string {
-    const varName = this.ctx.makeNameValid(id);
+    const varName = this.ctx.makeUniqueIdentifier(id, 'block');
     const ptrType = ptrFn(dataType);
     const snippet = snip(
       new RefOperator(snip(varName, dataType, 'function'), ptrType),
@@ -278,22 +284,25 @@ ${this.ctx.pre}}`;
       varOrigin = 'runtime';
     }
 
-    const snippet = snip(this.ctx.makeNameValid(id), dataType, /* origin */ varOrigin);
+    const snippet = snip(
+      this.ctx.makeUniqueIdentifier(id, 'block'),
+      dataType,
+      /* origin */ varOrigin,
+    );
     this.ctx.defineVariable(id, snippet);
     return snippet;
   }
 
-  protected emitVarDecl(
-    pre: string,
+  protected _emitVarDecl(
     keyword: 'var' | 'let' | 'const',
     name: string,
     _dataType: wgsl.BaseData | UnknownData,
     rhsStr: string,
   ): string {
-    return `${pre}${keyword} ${name} = ${rhsStr};`;
+    return `${this.ctx.pre}${keyword} ${name} = ${rhsStr};`;
   }
 
-  public _identifier(id: string): Snippet {
+  protected _identifier(id: string): Snippet {
     if (!id) {
       throw new Error('Cannot resolve an empty identifier');
     }
@@ -314,7 +323,7 @@ ${this.ctx.pre}}`;
    * A wrapper for `generateExpression` that updates `ctx.expectedType`
    * and tries to convert the result when it does not match the expected type.
    */
-  public _typedExpression(
+  protected _typedExpression(
     expression: tinyest.Expression,
     expectedType: wgsl.BaseData | wgsl.BaseData[],
   ) {
@@ -335,6 +344,7 @@ ${this.ctx.pre}}`;
     }
   }
 
+  // TODO: Make protected once we don't test it directly
   public _expression(expression: tinyest.Expression): Snippet {
     if (typeof expression === 'string') {
       return this._identifier(expression);
@@ -887,8 +897,39 @@ ${this.ctx.pre}}`;
     assertExhaustive(expression);
   }
 
-  public functionDefinition(body: tinyest.Block): string {
-    return this._block(body);
+  public functionDefinition(options: FunctionDefinitionOptions): string {
+    // Function body
+    const body = this._block(options.body);
+
+    // Only after generating the body can we determine the return type
+    const returnType = options.determineReturnType();
+
+    const argList = options.args
+      // Stripping out unused arguments in entry functions
+      .filter((arg) => arg.used || options.functionType === 'normal')
+      .map((arg) => {
+        return `${getAttributesString(arg.decoratedType)}${arg.name}: ${this.ctx.resolve(arg.decoratedType).value}`;
+      })
+      .join(', ');
+
+    const head =
+      returnType.type !== 'void'
+        ? `(${argList}) -> ${getAttributesString(returnType)}${this.ctx.resolve(returnType).value} `
+        : `(${argList}) `;
+
+    let attributes = '';
+    if (options.functionType === 'compute') {
+      if (!options.workgroupSize) {
+        throw new Error('Compute shaders must have a workgroup size');
+      }
+      attributes = `@compute @workgroup_size(${options.workgroupSize.join(', ')}) `;
+    } else if (options.functionType === 'vertex') {
+      attributes = `@vertex `;
+    } else if (options.functionType === 'fragment') {
+      attributes = `@fragment `;
+    }
+
+    return `${attributes}fn ${options.name}${head}${body}`;
   }
 
   /**
@@ -910,7 +951,7 @@ ${this.ctx.pre}}`;
     return snip(stitch`${this.ctx.resolve(schema).value}(${args})`, schema, 'runtime');
   }
 
-  public _return(statement: tinyest.Return): string {
+  protected _return(statement: tinyest.Return): string {
     const returnNode = statement[1];
 
     if (returnNode !== undefined) {
@@ -979,7 +1020,7 @@ Try 'return ${typeStr}(${str});' instead.
     return `${this.ctx.pre}return;`;
   }
 
-  public _statement(statement: tinyest.Statement): string {
+  protected _statement(statement: tinyest.Statement): string {
     if (typeof statement === 'string') {
       const id = this._identifier(statement);
       const resolved = id.value && this.ctx.resolve(id.value).value;
@@ -1017,7 +1058,7 @@ Try 'return ${typeStr}(${str});' instead.
           return this._statement(node);
         }
         // simplify 'if (true) {A} else {B}' to '{A}'
-        return `${this.ctx.pre}${this._block(blockifySingleStatement(node))}`;
+        return this._blockStatement(blockifySingleStatement(node));
       }
 
       const consequent = this._block(blockifySingleStatement(consNode));
@@ -1140,17 +1181,11 @@ ${this.ctx.pre}else ${alternate}`;
       const snippet = this.blockVariable(varType, rawId, concretize(dataType), eq.origin);
       const rhsSnippet = tryConvertSnippet(this.ctx, eq, dataType, false);
       const rhsStr = this.ctx.resolve(rhsSnippet.value, rhsSnippet.dataType).value;
-      return this.emitVarDecl(
-        this.ctx.pre,
-        varType,
-        snippet.value as string,
-        concretize(dataType),
-        rhsStr,
-      );
+      return this._emitVarDecl(varType, snippet.value as string, concretize(dataType), rhsStr);
     }
 
     if (statement[0] === NODE.block) {
-      return `${this.ctx.pre}${this._block(statement)}`;
+      return this._blockStatement(statement);
     }
 
     if (statement[0] === NODE.for) {
@@ -1246,12 +1281,9 @@ ${this.ctx.pre}else ${alternate}`;
 
           const blocks = elements.map(
             (e, i) =>
-              `${this.ctx.pre}// unrolled iteration #${i}\n${this.ctx.pre}${this._block(
-                blockified,
-                {
-                  [originalLoopVarName]: e,
-                },
-              )}`,
+              `${this.ctx.pre}// unrolled iteration #${i}\n${this._blockStatement(blockified, {
+                [originalLoopVarName]: e,
+              })}`,
           );
 
           return blocks.join('\n');
@@ -1259,7 +1291,7 @@ ${this.ctx.pre}else ${alternate}`;
 
         this.#unrolling = false;
 
-        const index = this.ctx.makeNameValid('i');
+        const index = this.ctx.makeUniqueIdentifier('i', 'block');
 
         const forHeaderStr = stitch`${this.ctx.pre}for (var ${index} = ${range.start}; ${index} ${range.comparison} ${range.end}; ${index} += ${range.step})`;
 
@@ -1272,7 +1304,7 @@ ${this.ctx.pre}else ${alternate}`;
         } else {
           this.ctx.indent();
           ctxIndent = true;
-          const loopVarName = this.ctx.makeNameValid(originalLoopVarName);
+          const loopVarName = this.ctx.makeUniqueIdentifier(originalLoopVarName, 'block');
           const elementSnippet = forOfUtils.getElementSnippet(
             iterableSnippet,
             snip(index, u32, 'runtime'),
@@ -1286,7 +1318,7 @@ ${this.ctx.pre}else ${alternate}`;
             false,
           )};`;
 
-          bodyStr = `{\n${loopVarDeclStr}\n${this.ctx.pre}${this._block(blockified, {
+          bodyStr = `{\n${loopVarDeclStr}\n${this._blockStatement(blockified, {
             [originalLoopVarName]: snip(loopVarName, elementType, elementSnippet.origin),
           })}\n`;
           this.ctx.dedent();
diff --git a/packages/typegpu/src/types.ts b/packages/typegpu/src/types.ts
index 1c92f8a43f..9b39c20138 100644
--- a/packages/typegpu/src/types.ts
+++ b/packages/typegpu/src/types.ts
@@ -76,8 +76,10 @@ export type Wgsl = Eventual<string | number | boolean | ResolvableObject>;
 
 export type TgpuShaderStage = 'compute' | 'vertex' | 'fragment';
 
-export interface FnToWgslOptions {
+export interface ResolveFunctionOptions {
   functionType: 'normal' | TgpuShaderStage;
+  workgroupSize?: readonly number[] | undefined;
+  name: string;
   argTypes: BaseData[];
   /**
    * The return type of the function. If undefined, the type should be inferred
@@ -99,11 +101,19 @@ export type ItemLayer = {
   usedSlots: Set<TgpuSlot<unknown>>;
 };
 
+export type FunctionArgumentAccess = () => Snippet | undefined;
+
+export interface FunctionArgument {
+  name: string;
+  access: FunctionArgumentAccess;
+  decoratedType: BaseData;
+  used: boolean;
+}
+
 export type FunctionScopeLayer = {
   type: 'functionScope';
   functionType: 'normal' | 'compute' | 'vertex' | 'fragment';
-  args: Snippet[];
-  argAliases: Record<string, Snippet>;
+  argAccess: Record<string, FunctionArgumentAccess>;
   externalMap: Record<string, unknown>;
   /**
    * The return type of the function. If undefined, the type should be inferred
@@ -123,6 +133,7 @@ export type SlotBindingLayer = {
 
 export type BlockScopeLayer = {
   type: 'blockScope';
+  takenLocalIdentifiers: Set<string>;
   declarations: Map<string, Snippet>;
   externals: Map<string, Snippet>;
 };
@@ -132,14 +143,14 @@ export type StackLayer = ItemLayer | SlotBindingLayer | FunctionScopeLayer | Blo
 export interface ItemStateStack {
   readonly itemDepth: number;
   readonly topItem: ItemLayer;
+  readonly topBlockScope: BlockScopeLayer | undefined;
   readonly topFunctionScope: FunctionScopeLayer | undefined;
 
   pushItem(): void;
   pushSlotBindings(pairs: SlotValuePair[]): void;
   pushFunctionScope(
     functionType: 'normal' | TgpuShaderStage,
-    args: Snippet[],
-    argAliases: Record<string, Snippet>,
+    argAccess: Record<string, FunctionArgumentAccess>,
     /**
      * The return type of the function. If undefined, the type should be inferred
      * from the implementation (relevant for shellless functions).
@@ -304,9 +315,8 @@ export interface ResolutionCtx {
    */
   resolveSnippet(snippet: Snippet): ResolvedSnippet;
 
-  fnToWgsl(options: FnToWgslOptions): {
-    head: Wgsl;
-    body: Wgsl;
+  resolveFunction(options: ResolveFunctionOptions): {
+    code: string;
     returnType: BaseData;
   };
 
@@ -324,8 +334,27 @@ export interface ResolutionCtx {
    */
   withRenamed<T>(item: object, name: string | undefined, callback: () => T): T;
 
-  getUniqueName(resource: object): string;
-  makeNameValid(name: string): string;
+  /**
+   * @param primer The basis for the unique identifier. Depending on the strategy, or
+   *               the names already taken, this may be modified to ensure uniqueness.
+   * @param scope The scope in which to generate the identifier. 'global' means
+   *              the identifier is meant to be unique across the entire program, while
+   *              'block' means it cannot shadow any existing identifiers visible from
+   *              within the current block. After the block is popped, any identifiers
+   *              defined within it are no longer visible.
+   * @returns an identifier that is unique within the given scope
+   */
+  makeUniqueIdentifier(primer: string | undefined, scope: 'global' | 'block'): string;
+
+  isIdentifierTaken(name: string): boolean;
+
+  /**
+   * Makes sure the given identifier cannot be generated by {@link makeUniqueIdentifier}
+   * within the given scope.
+   * @param name The name to reserve
+   * @param scope See {@link makeUniqueIdentifier} for a description of the scope parameter.
+   */
+  reserveIdentifier(name: string, scope: 'global' | 'block'): void;
 }
 
 /**
diff --git a/packages/typegpu/tests/namespace.test.ts b/packages/typegpu/tests/namespace.test.ts
index ee6fe33099..7ac6407b43 100644
--- a/packages/typegpu/tests/namespace.test.ts
+++ b/packages/typegpu/tests/namespace.test.ts
@@ -1,4 +1,4 @@
-import { describe, expect, vi } from 'vitest';
+import { describe, expect } from 'vitest';
 import tgpu, { d } from '../src/index.js';
 import { it } from 'typegpu-testing-utility';
 
@@ -71,34 +71,6 @@ describe('tgpu.namespace', () => {
     `);
   });
 
-  it('fires "name" event', () => {
-    const Boid = d.struct({
-      pos: d.vec3f,
-    });
-
-    const names = tgpu['~unstable'].namespace();
-
-    const listener = vi.fn((event) => {});
-    names.on('name', listener);
-
-    const code = tgpu.resolve([Boid], { names });
-
-    expect(listener).toHaveBeenCalledTimes(1);
-    expect(listener).toHaveBeenCalledWith({ name: 'Boid', target: Boid });
-
-    expect(code).toMatchInlineSnapshot(`
-      "struct Boid {
-        pos: vec3f,
-      }"
-    `);
-
-    const code2 = tgpu.resolve([Boid], { names });
-
-    // No more events
-    expect(listener).toHaveBeenCalledTimes(1);
-    expect(code2).toMatchInlineSnapshot(`""`);
-  });
-
   it('handles name collision', () => {
     let code1: string, code2: string;
     const names = tgpu['~unstable'].namespace();
diff --git a/packages/typegpu/tests/renderPipeline.test.ts b/packages/typegpu/tests/renderPipeline.test.ts
index 316f2d0e02..be0c2888e4 100644
--- a/packages/typegpu/tests/renderPipeline.test.ts
+++ b/packages/typegpu/tests/renderPipeline.test.ts
@@ -277,13 +277,6 @@ describe('root.withVertex(...).withFragment(...)', () => {
           return vertexMain_Output(vec3f(), vec3f(), vec3f(), 0f, 0u, vec4f());
         }
 
-        struct fragmentMain_Input {
-          @location(3) baz3: u32,
-          @location(1) bar: vec3f,
-          @location(2) foo: vec3f,
-          @location(5) baz2: f32,
-        }
-
         @fragment fn fragmentMain() -> @location(0) vec4f {
           return vec4f();
         }"
@@ -1568,8 +1561,8 @@ describe('root.createRenderPipeline', () => {
       }
 
       struct VertexIn {
-        @builtin(vertex_index) vertexIndex: u32,
         @location(0) localPos: vec3f,
+        @builtin(vertex_index) vertexIndex: u32,
       }
 
       @vertex fn vertex(_arg_0: VertexIn) -> VertexOut {
@@ -1712,18 +1705,13 @@ describe('root.createRenderPipeline', () => {
                 {
                   "arrayStride": 16,
                   "attributes": [
-                    {
-                      "format": "float32",
-                      "offset": 12,
-                      "shaderLocation": 0,
-                    },
                     {
                       "format": "float32x3",
                       "offset": 0,
-                      "shaderLocation": 1,
+                      "shaderLocation": 0,
                     },
                   ],
-                  "stepMode": "instance",
+                  "stepMode": "vertex",
                 },
                 {
                   "arrayStride": 16,
@@ -1731,10 +1719,15 @@ describe('root.createRenderPipeline', () => {
                     {
                       "format": "float32x3",
                       "offset": 0,
+                      "shaderLocation": 1,
+                    },
+                    {
+                      "format": "float32",
+                      "offset": 12,
                       "shaderLocation": 2,
                     },
                   ],
-                  "stepMode": "vertex",
+                  "stepMode": "instance",
                 },
               ],
               "module": "mockShaderModule",
@@ -1751,10 +1744,10 @@ describe('root.createRenderPipeline', () => {
           {
             "destroy": [MockFunction],
             "getMappedRange": [MockFunction],
-            "label": "instanceBuffer",
+            "label": "vertexBuffer",
             "mapAsync": [MockFunction],
             "mapState": "unmapped",
-            "size": 16,
+            "size": 48,
             "unmap": [MockFunction],
             "usage": 44,
           },
@@ -1766,10 +1759,10 @@ describe('root.createRenderPipeline', () => {
           {
             "destroy": [MockFunction],
             "getMappedRange": [MockFunction],
-            "label": "vertexBuffer",
+            "label": "instanceBuffer",
             "mapAsync": [MockFunction],
             "mapState": "unmapped",
-            "size": 48,
+            "size": 16,
             "unmap": [MockFunction],
             "usage": 44,
           },
diff --git a/packages/typegpu/tests/resolve.test.ts b/packages/typegpu/tests/resolve.test.ts
index 254730def2..83e8e1e787 100644
--- a/packages/typegpu/tests/resolve.test.ts
+++ b/packages/typegpu/tests/resolve.test.ts
@@ -1,6 +1,6 @@
 import { describe, expect, vi } from 'vitest';
 import tgpu, { d } from '../src/index.js';
-import { setName } from '../src/shared/meta.ts';
+import { getName, setName } from '../src/shared/meta.ts';
 import { $gpuValueOf, $internal, $ownSnippet, $resolve } from '../src/shared/symbols.ts';
 import type { ResolutionCtx } from '../src/types.ts';
 import { it } from 'typegpu-testing-utility';
@@ -55,7 +55,7 @@ describe('tgpu resolve', () => {
       } as unknown as number,
 
       [$resolve](ctx: ResolutionCtx) {
-        const name = ctx.getUniqueName(this);
+        const name = ctx.makeUniqueIdentifier(getName(this), 'global');
         ctx.addDeclaration(`@group(0) @binding(0) var<uniform> ${name}: f32;`);
         return snip(name, d.f32, /* origin */ 'runtime');
       },
diff --git a/packages/typegpu/tests/tgsl/extensionEnabled.test.ts b/packages/typegpu/tests/tgsl/extensionEnabled.test.ts
index 727e9fed70..58150fbee5 100644
--- a/packages/typegpu/tests/tgsl/extensionEnabled.test.ts
+++ b/packages/typegpu/tests/tgsl/extensionEnabled.test.ts
@@ -17,14 +17,14 @@ describe('extension based pruning', () => {
     });
 
     expect(tgpu.resolve([someFn], { enableExtensions: ['f16'] })).toMatchInlineSnapshot(`
-        "enable f16;
+      "enable f16;
 
-        fn someFn() -> f32 {
-          {
-            return 6.599609375f;
-          }
-        }"
-      `);
+      fn someFn() -> f32 {
+        {
+          return 6.599609375f;
+        }
+      }"
+    `);
 
     expect(tgpu.resolve([someFn])).toMatchInlineSnapshot(`
       "fn someFn() -> f32 {
diff --git a/packages/typegpu/tests/tgsl/multiplication.test.ts b/packages/typegpu/tests/tgsl/multiplication.test.ts
index 30b3b589c8..1a1e335463 100644
--- a/packages/typegpu/tests/tgsl/multiplication.test.ts
+++ b/packages/typegpu/tests/tgsl/multiplication.test.ts
@@ -31,7 +31,7 @@ test('multiplying i32 with a float literal should implicitly convert to an f32',
     [
       [
         "Implicit conversions from [
-      1: i32
+      1i: i32
     ] to f32 are supported, but not recommended.
     Consider using explicit conversions instead.",
       ],
@@ -43,7 +43,7 @@ test('multiplying i32 with a float literal should implicitly convert to an f32',
       ],
       [
         "Implicit conversions from [
-      1: i32
+      1i: i32
     ] to f32 are supported, but not recommended.
     Consider using explicit conversions instead.",
       ],
@@ -78,7 +78,7 @@ test('multiplying u32 with a float literal should implicitly convert to an f32',
     [
       [
         "Implicit conversions from [
-      10: u32
+      10u: u32
     ] to f32 are supported, but not recommended.
     Consider using explicit conversions instead.",
       ],
@@ -90,7 +90,7 @@ test('multiplying u32 with a float literal should implicitly convert to an f32',
       ],
       [
         "Implicit conversions from [
-      1: u32
+      1u: u32
     ] to f32 are supported, but not recommended.
     Consider using explicit conversions instead.",
       ],
diff --git a/packages/typegpu/tests/tgsl/nameClashes.test.ts b/packages/typegpu/tests/tgsl/nameClashes.test.ts
index 86167245ed..586e8e3bf1 100644
--- a/packages/typegpu/tests/tgsl/nameClashes.test.ts
+++ b/packages/typegpu/tests/tgsl/nameClashes.test.ts
@@ -223,14 +223,14 @@ test('should allow duplicate name after block end', () => {
   };
 
   expect(tgpu.resolve([main])).toMatchInlineSnapshot(`
-      "fn main() -> u32 {
-        for (var i = 0; (i < 3i); i++) {
-          let foo = (i + 1i);
-        }
-        const foo = 7u;
-        return foo;
-      }"
-    `);
+    "fn main() -> u32 {
+      for (var i = 0; (i < 3i); i++) {
+        let foo = (i + 1i);
+      }
+      const foo = 7u;
+      return foo;
+    }"
+  `);
 });
 
 test('should give declarations new names when they are shadowed', () => {
diff --git a/packages/typegpu/tests/tgsl/ternaryOperator.test.ts b/packages/typegpu/tests/tgsl/ternaryOperator.test.ts
index fd37903d12..c41cc6b27c 100644
--- a/packages/typegpu/tests/tgsl/ternaryOperator.test.ts
+++ b/packages/typegpu/tests/tgsl/ternaryOperator.test.ts
@@ -18,14 +18,14 @@ describe('ternary operator', () => {
         myFn.with(mySlot, false).$name('falseFn'),
       ]),
     ).toMatchInlineSnapshot(`
-        "fn trueFn() -> u32 {
-          return 10u;
-        }
-
-        fn falseFn() -> u32 {
-          return 20u;
-        }"
-      `);
+      "fn trueFn() -> u32 {
+        return 10u;
+      }
+
+      fn falseFn() -> u32 {
+        return 20u;
+      }"
+    `);
   });
 
   it('should work for different comptime known expressions', () => {
@@ -72,22 +72,22 @@ describe('ternary operator', () => {
         myFn.with(mySlot, 3).$name('threeFn'),
       ]),
     ).toMatchInlineSnapshot(`
-        "fn myFn() -> u32 {
-          return -1u;
-        }
-
-        fn oneFn() -> u32 {
-          return 10u;
-        }
-
-        fn twoFn() -> u32 {
-          return 20u;
-        }
-
-        fn threeFn() -> u32 {
-          return 30u;
-        }"
-      `);
+      "fn myFn() -> u32 {
+        return -1u;
+      }
+
+      fn oneFn() -> u32 {
+        return 10u;
+      }
+
+      fn twoFn() -> u32 {
+        return 20u;
+      }
+
+      fn threeFn() -> u32 {
+        return 30u;
+      }"
+    `);
   });
 
   it('should not include unused dependencies', ({ root }) => {
@@ -103,20 +103,20 @@ describe('ternary operator', () => {
     });
 
     expect(tgpu.resolve([myFn.with(mySlot, true).$name('trueFn')])).toMatchInlineSnapshot(`
-        "@group(0) @binding(0) var<uniform> myUniform: u32;
+      "@group(0) @binding(0) var<uniform> myUniform: u32;
 
-        fn trueFn() -> u32 {
-          return myUniform;
-        }"
-      `);
+      fn trueFn() -> u32 {
+        return myUniform;
+      }"
+    `);
 
     expect(tgpu.resolve([myFn.with(mySlot, false).$name('falseFn')])).toMatchInlineSnapshot(`
-        "@group(0) @binding(0) var<storage, read> myReadonly: u32;
+      "@group(0) @binding(0) var<storage, read> myReadonly: u32;
 
-        fn falseFn() -> u32 {
-          return myReadonly;
-        }"
-      `);
+      fn falseFn() -> u32 {
+        return myReadonly;
+      }"
+    `);
   });
 
   it('should handle undefined', ({ root }) => {
diff --git a/packages/typegpu/tests/tgsl/wgslGenerator.test.ts b/packages/typegpu/tests/tgsl/wgslGenerator.test.ts
index e1129649ac..43f7b6ce77 100644
--- a/packages/typegpu/tests/tgsl/wgslGenerator.test.ts
+++ b/packages/typegpu/tests/tgsl/wgslGenerator.test.ts
@@ -3,19 +3,17 @@ import { beforeEach, describe, expect, vi } from 'vitest';
 import { namespace } from '../../src/core/resolve/namespace.ts';
 import * as d from '../../src/data/index.ts';
 import { abstractFloat, abstractInt } from '../../src/data/numeric.ts';
-import { snip } from '../../src/data/snippet.ts';
-import { Void, type WgslArray } from '../../src/data/wgslTypes.ts';
+import { type WgslArray } from '../../src/data/wgslTypes.ts';
 import { provideCtx } from '../../src/execMode.ts';
 import tgpu from '../../src/index.js';
 import { ResolutionCtxImpl } from '../../src/resolutionCtx.ts';
 import { getMetaData } from '../../src/shared/meta.ts';
-import { $internal } from '../../src/shared/symbols.ts';
 import * as std from '../../src/std/index.ts';
 import wgslGenerator from '../../src/tgsl/wgslGenerator.ts';
 import { CodegenState } from '../../src/types.ts';
 import { it } from 'typegpu-testing-utility';
 import { ArrayExpression } from '../../src/tgsl/generationHelpers.ts';
-import { extractSnippetFromFn } from '../utils/parseResolved.ts';
+import { expectDataTypeOf, extractSnippetFromFn } from '../utils/parseResolved.ts';
 
 const { NodeTypeCatalog: NODE } = tinyest;
 
@@ -40,19 +38,11 @@ describe('wgslGenerator', () => {
       return true;
     };
 
-    const parsedBody = getMetaData(main)?.ast?.body as tinyest.Block;
-
-    expect(JSON.stringify(parsedBody)).toMatchInlineSnapshot(`"[0,[[10,true]]]"`);
-
-    provideCtx(ctx, () => {
-      ctx[$internal].itemStateStack.pushFunctionScope('normal', [], {}, d.bool, {});
-      const gen = wgslGenerator.functionDefinition(parsedBody);
-      expect(gen).toMatchInlineSnapshot(`
-        "{
-          return true;
-        }"
-      `);
-    });
+    expect(tgpu.resolve([main])).toMatchInlineSnapshot(`
+      "fn main() -> bool {
+        return true;
+      }"
+    `);
   });
 
   it('creates a function body', () => {
@@ -63,23 +53,13 @@ describe('wgslGenerator', () => {
       return a;
     };
 
-    const parsedBody = getMetaData(main)?.ast?.body as tinyest.Block;
-
-    expect(JSON.stringify(parsedBody)).toMatchInlineSnapshot(
-      `"[0,[[12,"a",[5,"12"]],[2,"a","+=",[5,"21"]],[10,"a"]]]"`,
-    );
-
-    provideCtx(ctx, () => {
-      ctx[$internal].itemStateStack.pushFunctionScope('normal', [], {}, d.i32, {});
-      const gen = wgslGenerator.functionDefinition(parsedBody);
-      expect(gen).toMatchInlineSnapshot(`
-        "{
-          var a = 12;
-          a += 21i;
-          return a;
-        }"
-      `);
-    });
+    expect(tgpu.resolve([main])).toMatchInlineSnapshot(`
+      "fn main() -> i32 {
+        var a = 12;
+        a += 21i;
+        return a;
+      }"
+    `);
   });
 
   it('creates correct resources for numeric literals', () => {
@@ -134,57 +114,22 @@ describe('wgslGenerator', () => {
     });
 
     const testBuffer = root.createBuffer(TestStruct).$usage('storage');
-
     const testUsage = testBuffer.as('mutable');
 
-    const testFn = tgpu.fn(
-      [],
-      d.u32,
-    )(() => {
-      return testUsage.$.a + testUsage.$.b.x;
-    });
-
-    const astInfo = getMetaData(
-      testFn[$internal].implementation as (...args: unknown[]) => unknown,
-    );
-    if (!astInfo) {
-      throw new Error('Expected prebuilt AST to be present');
-    }
+    expectDataTypeOf(() => {
+      'use gpu';
+      return testUsage.$.a;
+    }).toStrictEqual(d.u32);
 
-    expect(JSON.stringify(astInfo.ast?.body)).toMatchInlineSnapshot(
-      `"[0,[[10,[1,[7,[7,"testUsage","$"],"a"],"+",[7,[7,[7,"testUsage","$"],"b"],"x"]]]]]"`,
-    );
-    ctx[$internal].itemStateStack.pushFunctionScope(
-      'normal',
-      [],
-      {},
-      d.u32,
-      (astInfo.externals as () => Record<string, unknown>)() ?? {},
-    );
+    expectDataTypeOf(() => {
+      'use gpu';
+      return testUsage.$.b.x;
+    }).toStrictEqual(d.u32);
 
-    provideCtx(ctx, () => {
-      // Check for: return testUsage.$.a + testUsage.$.b.x;
-      //                   ^ this should be a u32
-      const res1 = wgslGenerator._expression(
-        ((astInfo.ast?.body[1][0] as tinyest.Return)[1] as tinyest.BinaryExpression)[1],
-      );
-
-      expect(res1.dataType).toStrictEqual(d.u32);
-
-      // Check for: return testUsage.$.a + testUsage.$.b.x;
-      //                                       ^ this should be a u32
-      const res2 = wgslGenerator._expression(
-        ((astInfo.ast?.body[1][0] as tinyest.Return)[1] as tinyest.BinaryExpression)[3],
-      );
-      expect(res2.dataType).toStrictEqual(d.u32);
-
-      // Check for: return testUsage.$.a + testUsage.$.b.x;
-      //            ^ this should be a u32
-      const sum = wgslGenerator._expression(
-        (astInfo.ast?.body[1][0] as tinyest.Return)[1] as tinyest.Expression,
-      );
-      expect(sum.dataType).toStrictEqual(d.u32);
-    });
+    expectDataTypeOf(() => {
+      'use gpu';
+      return testUsage.$.a + testUsage.$.b.x;
+    }).toStrictEqual(d.u32);
   });
 
   it('generates correct resources for external resource array index access', ({ root }) => {
@@ -192,42 +137,10 @@ describe('wgslGenerator', () => {
 
     const testUsage = testBuffer.as('uniform');
 
-    const testFn = tgpu.fn(
-      [],
-      d.u32,
-    )(() => {
-      return testUsage.$[3] as number;
-    });
-
-    const astInfo = getMetaData(
-      testFn[$internal].implementation as (...args: unknown[]) => unknown,
-    );
-
-    if (!astInfo) {
-      throw new Error('Expected prebuilt AST to be present');
-    }
-
-    expect(JSON.stringify(astInfo.ast?.body)).toMatchInlineSnapshot(
-      `"[0,[[10,[8,[7,"testUsage","$"],[5,"3"]]]]]"`,
-    );
-
-    provideCtx(ctx, () => {
-      ctx[$internal].itemStateStack.pushFunctionScope(
-        'normal',
-        [],
-        {},
-        d.u32,
-        (astInfo.externals as () => Record<string, unknown>)() ?? {},
-      );
-
-      // Check for: return testUsage.$[3];
-      //                   ^ this should be a u32
-      const res = wgslGenerator._expression(
-        (astInfo.ast?.body[1][0] as tinyest.Return)[1] as tinyest.Expression,
-      );
-
-      expect(res.dataType).toStrictEqual(d.u32);
-    });
+    expectDataTypeOf(() => {
+      'use gpu';
+      return testUsage.$[3];
+    }).toStrictEqual(d.u32);
   });
 
   it('generates correct resources for nested struct with atomics in a complex expression', ({
@@ -253,156 +166,29 @@ describe('wgslGenerator', () => {
 
     const testUsage = testBuffer.as('mutable');
 
-    const testFn = tgpu.fn(
-      [d.u32],
-      d.vec4f,
-    )((idx) => {
-      const value = std.atomicLoad(testUsage.$.b.aa[idx]!.y);
-      const vec = std.mix(d.vec4f(), testUsage.$.a, value);
-      std.atomicStore(testUsage.$.b.aa[idx]!.x, vec.y);
-      return vec;
-    });
-
-    const astInfo = getMetaData(
-      testFn[$internal].implementation as (...args: unknown[]) => unknown,
-    );
-
-    if (!astInfo?.ast) {
-      throw new Error('Expected prebuilt AST to be present');
-    }
-
-    expect(JSON.stringify(astInfo.ast.body)).toMatchInlineSnapshot(
-      `"[0,[[13,"value",[6,[7,"std","atomicLoad"],[[7,[8,[7,[7,[7,"testUsage","$"],"b"],"aa"],"idx"],"y"]]]],[13,"vec",[6,[7,"std","mix"],[[6,[7,"d","vec4f"],[]],[7,[7,"testUsage","$"],"a"],"value"]]],[6,[7,"std","atomicStore"],[[7,[8,[7,[7,[7,"testUsage","$"],"b"],"aa"],"idx"],"x"],[7,"vec","y"]]],[10,"vec"]]]"`,
-    );
-
-    if (astInfo.ast.params.filter((arg) => arg.type !== 'i').length > 0) {
-      throw new Error('Expected arguments as identifier names in ast');
-    }
-
-    const args = astInfo.ast.params.map((arg) =>
-      snip((arg as { type: 'i'; name: string }).name, d.u32, /* origin */ 'runtime'),
-    );
-
-    provideCtx(ctx, () => {
-      ctx[$internal].itemStateStack.pushFunctionScope(
-        'normal',
-        args,
-        {},
-        d.vec4f,
-        (astInfo.externals as () => Record<string, unknown>)() ?? {},
-      );
-
-      // Check for: const value = std.atomicLoad(testUsage.$.b.aa[idx]!.y);
-      //                           ^ this part should be a i32
-      const res = wgslGenerator._expression(
-        (astInfo.ast?.body[1][0] as tinyest.Const)[2] as tinyest.Expression,
-      );
-
-      expect(res.dataType).toStrictEqual(d.i32);
-
-      // Check for: const vec = std.mix(d.vec4f(), testUsage.$.a, value);
-      //                        ^ this part should be a vec4f
-      ctx[$internal].itemStateStack.pushBlockScope();
-      wgslGenerator.blockVariable('var', 'value', d.i32, 'runtime');
-      const res2 = wgslGenerator._expression(
-        (astInfo.ast?.body[1][1] as tinyest.Const)[2] as tinyest.Expression,
-      );
-      ctx[$internal].itemStateStack.pop('blockScope');
-
-      expect(res2.dataType).toStrictEqual(d.vec4f);
-
-      // Check for: std.atomicStore(testUsage.$.b.aa[idx]!.x, vec.y);
-      //                            ^ this part should be an atomic u32
-      //            ^ this part should be void
-      ctx[$internal].itemStateStack.pushBlockScope();
-      wgslGenerator.blockVariable('var', 'vec', d.vec4f, 'function');
-      const res3 = wgslGenerator._expression(
-        (astInfo.ast?.body[1][2] as tinyest.Call)[2][0] as tinyest.Expression,
-      );
-      const res4 = wgslGenerator._expression(astInfo.ast?.body[1][2] as tinyest.Expression);
-      ctx[$internal].itemStateStack.pop('blockScope');
-
-      expect(res3.dataType).toStrictEqual(d.atomic(d.u32));
-      expect(res4.dataType).toStrictEqual(Void);
-    });
-  });
-
-  it('creates correct code for for statements', () => {
-    const main = () => {
+    // Check for: const value = std.atomicLoad(testUsage.$.b.aa[idx]!.y);
+    //                           ^ this part should be a i32
+    expectDataTypeOf(() => {
       'use gpu';
-      for (let i = 0; i < 10; i += 1) {
-        continue;
-      }
-    };
-
-    const parsed = getMetaData(main)?.ast?.body as tinyest.Block;
-
-    expect(JSON.stringify(parsed)).toMatchInlineSnapshot(
-      `"[0,[[14,[12,"i",[5,"0"]],[1,"i","<",[5,"10"]],[2,"i","+=",[5,"1"]],[0,[[16]]]]]]"`,
-    );
-
-    const gen = provideCtx(ctx, () => wgslGenerator.functionDefinition(parsed));
+      const idx = d.u32(0);
+      return std.atomicLoad(testUsage.$.b.aa[idx]!.y);
+    }).toStrictEqual(d.i32);
 
-    expect(gen).toMatchInlineSnapshot(`
-      "{
-        for (var i = 0; (i < 10i); i += 1i) {
-          continue;
-        }
-      }"
-    `);
-  });
-
-  it('creates correct code for for statements with outside init', () => {
-    const main = () => {
+    // Check for: const vec = std.mix(d.vec4f(), testUsage.$.a, value);
+    //                        ^ this part should be a vec4f
+    expectDataTypeOf(() => {
       'use gpu';
-      let i = 0;
-      for (; i < 10; i += 1) {
-        continue;
-      }
-    };
-
-    const parsed = getMetaData(main)?.ast?.body as tinyest.Block;
+      const value = std.atomicLoad(testUsage.$.b.aa[0]!.y);
+      return std.mix(d.vec4f(), testUsage.$.a, value);
+    }).toStrictEqual(d.vec4f);
 
-    expect(JSON.stringify(parsed)).toMatchInlineSnapshot(
-      `"[0,[[12,"i",[5,"0"]],[14,null,[1,"i","<",[5,"10"]],[2,"i","+=",[5,"1"]],[0,[[16]]]]]]"`,
-    );
-
-    const gen = provideCtx(ctx, () => wgslGenerator.functionDefinition(parsed));
-
-    expect(gen).toMatchInlineSnapshot(`
-      "{
-        var i = 0;
-        for (; (i < 10i); i += 1i) {
-          continue;
-        }
-      }"
-    `);
-  });
-
-  it('creates correct code for while statements', () => {
-    const main = () => {
+    // Check for: std.atomicStore(testUsage.$.b.aa[idx]!.x, vec.y);
+    //                            ^ this part should be an atomic u32
+    expectDataTypeOf(() => {
       'use gpu';
-      let i = 0;
-      while (i < 10) {
-        i += 1;
-      }
-    };
-
-    const parsed = getMetaData(main)?.ast?.body as tinyest.Block;
-    expect(JSON.stringify(parsed)).toMatchInlineSnapshot(
-      `"[0,[[12,"i",[5,"0"]],[15,[1,"i","<",[5,"10"]],[0,[[2,"i","+=",[5,"1"]]]]]]]"`,
-    );
-
-    const gen = provideCtx(ctx, () => wgslGenerator.functionDefinition(parsed));
-
-    expect(gen).toMatchInlineSnapshot(`
-      "{
-        var i = 0;
-        while ((i < 10i)) {
-          i += 1i;
-        }
-      }"
-    `);
+      const idx = d.u32(0);
+      return testUsage.$.b.aa[idx]!.x;
+    }).toStrictEqual(d.atomic(d.u32));
   });
 
   it('parses correctly "for ... of ..." statements', () => {
@@ -1003,6 +789,11 @@ describe('wgslGenerator', () => {
   });
 
   it('creates correct resources for lazy values and slots', () => {
+    expectDataTypeOf(() => {
+      'use gpu';
+      return lazyV4u.$;
+    }).toStrictEqual(d.vec4u);
+
     const testFn = tgpu.fn([], d.vec4u)(() => lazyV4u.$);
 
     expect(tgpu.resolve([testFn])).toMatchInlineSnapshot(`
@@ -1010,76 +801,14 @@ describe('wgslGenerator', () => {
         return vec4u(44, 88, 132, 176);
       }"
     `);
-
-    const astInfo = getMetaData(
-      testFn[$internal].implementation as (...args: unknown[]) => unknown,
-    );
-
-    if (!astInfo) {
-      throw new Error('Expected prebuilt AST to be present');
-    }
-
-    expect(JSON.stringify(astInfo.ast?.body)).toMatchInlineSnapshot(
-      `"[0,[[10,[7,"lazyV4u","$"]]]]"`,
-    );
-
-    provideCtx(ctx, () => {
-      ctx[$internal].itemStateStack.pushFunctionScope(
-        'normal',
-        [],
-        {},
-        d.vec4u,
-        (astInfo.externals as () => Record<string, unknown>)() ?? {},
-      );
-
-      wgslGenerator.initGenerator(ctx);
-      // Check for: return lazyV4u.$;
-      //                      ^ this should be a vec4u
-      const res = wgslGenerator._expression(
-        (astInfo.ast?.body[1][0] as tinyest.Return)[1] as tinyest.Expression,
-      );
-
-      expect(res.dataType).toStrictEqual(d.vec4u);
-    });
   });
 
   it('creates correct resources for indexing into a lazy value', () => {
-    const testFn = tgpu.fn(
-      [d.u32],
-      d.f32,
-    )((idx) => {
-      return lazyV2f.$[idx] as number;
-    });
-
-    const astInfo = getMetaData(
-      testFn[$internal].implementation as (...args: unknown[]) => unknown,
-    );
-
-    if (!astInfo) {
-      throw new Error('Expected prebuilt AST to be present');
-    }
-
-    expect(JSON.stringify(astInfo.ast?.body)).toMatchInlineSnapshot(
-      `"[0,[[10,[8,[7,"lazyV2f","$"],"idx"]]]]"`,
-    );
-
-    provideCtx(ctx, () => {
-      ctx[$internal].itemStateStack.pushFunctionScope(
-        'normal',
-        [snip('idx', d.u32, /* origin */ 'runtime')],
-        {},
-        d.f32,
-        (astInfo.externals as () => Record<string, unknown>)() ?? {},
-      );
-
-      // Check for: return lazyV2f.$[idx];
-      //                      ^ this should be a f32
-      const res = wgslGenerator._expression(
-        (astInfo.ast?.body[1][0] as tinyest.Return)[1] as tinyest.Expression,
-      );
-
-      expect(res.dataType).toStrictEqual(d.f32);
-    });
+    expectDataTypeOf(() => {
+      'use gpu';
+      const idx = d.u32(0);
+      return lazyV2f.$[idx];
+    }).toStrictEqual(d.f32);
   });
 
   it('creates intermediate representation for array expression', () => {
@@ -1105,44 +834,11 @@ describe('wgslGenerator', () => {
     });
 
     expect(tgpu.resolve([testFn])).toMatchInlineSnapshot(`
-        "fn testFn() -> u32 {
-          var arr = array<u32, 3>(1u, 2u, 3u);
-          return arr[1i];
-        }"
-      `);
-
-    const astInfo = getMetaData(
-      testFn[$internal].implementation as (...args: unknown[]) => unknown,
-    );
-
-    if (!astInfo) {
-      throw new Error('Expected prebuilt AST to be present');
-    }
-
-    expect(JSON.stringify(astInfo.ast?.body)).toMatchInlineSnapshot(
-      `"[0,[[13,"arr",[100,[[6,[7,"d","u32"],[[5,"1"]]],[5,"2"],[5,"3"]]]],[10,[8,"arr",[5,"1"]]]]]"`,
-    );
-
-    provideCtx(ctx, () => {
-      ctx[$internal].itemStateStack.pushFunctionScope(
-        'normal',
-        [],
-        {},
-        d.u32,
-        (astInfo.externals as () => Record<string, unknown>)() ?? {},
-      );
-
-      // Check for: const arr = [1, 2, 3]
-      //                        ^ this should be an array<u32, 3>
-      wgslGenerator.initGenerator(ctx);
-      const res = wgslGenerator._expression(
-        (astInfo.ast?.body[1][0] as tinyest.Const)[2] as unknown as tinyest.Expression,
-      );
-
-      expect(d.isWgslArray(res.dataType)).toBe(true);
-      expect((res.dataType as unknown as WgslArray).elementCount).toBe(3);
-      expect((res.dataType as unknown as WgslArray).elementType).toBe(d.u32);
-    });
+      "fn testFn() -> u32 {
+        var arr = array<u32, 3>(1u, 2u, 3u);
+        return arr[1i];
+      }"
+    `);
   });
 
   it('generates correct code for complex array expressions', () => {
@@ -1164,39 +860,6 @@ describe('wgslGenerator', () => {
         return arr[1i].x;
       }"
     `);
-
-    const astInfo = getMetaData(
-      testFn[$internal].implementation as (...args: unknown[]) => unknown,
-    );
-
-    if (!astInfo) {
-      throw new Error('Expected prebuilt AST to be present');
-    }
-
-    expect(JSON.stringify(astInfo.ast?.body)).toMatchInlineSnapshot(
-      `"[0,[[13,"arr",[100,[[6,[7,"d","vec2u"],[[5,"1"],[5,"2"]]],[6,[7,"d","vec2u"],[[5,"3"],[5,"4"]]],[6,[7,"std","min"],[[6,[7,"d","vec2u"],[[5,"5"],[5,"8"]]],[6,[7,"d","vec2u"],[[5,"7"],[5,"6"]]]]]]]],[10,[7,[8,"arr",[5,"1"]],"x"]]]]"`,
-    );
-
-    provideCtx(ctx, () => {
-      ctx[$internal].itemStateStack.pushFunctionScope(
-        'normal',
-        [],
-        {},
-        d.u32,
-        (astInfo.externals as () => Record<string, unknown>)() ?? {},
-      );
-
-      // Check for: const arr = [1, 2, 3]
-      //                        ^ this should be an array<u32, 3>
-      wgslGenerator.initGenerator(ctx);
-      const res = wgslGenerator._expression(
-        (astInfo.ast?.body[1][0] as tinyest.Const)[2] as unknown as tinyest.Expression,
-      );
-
-      expect(d.isWgslArray(res.dataType)).toBe(true);
-      expect((res.dataType as unknown as WgslArray).elementCount).toBe(3);
-      expect((res.dataType as unknown as WgslArray).elementType).toBe(d.vec2u);
-    });
   });
 
   it('does not autocast lhs of an assignment', () => {
@@ -1247,38 +910,15 @@ describe('wgslGenerator', () => {
       }"
     `);
 
-    const astInfo = getMetaData(
-      testFn[$internal].implementation as (...args: unknown[]) => unknown,
-    );
-
-    if (!astInfo) {
-      throw new Error('Expected prebuilt AST to be present');
-    }
-
-    expect(JSON.stringify(astInfo.ast?.body)).toMatchInlineSnapshot(
-      `"[0,[[13,"arr",[100,[[6,"TestStruct",[[104,{"x":[5,"1"],"y":[5,"2"]}]]],[6,"TestStruct",[[104,{"x":[5,"3"],"y":[5,"4"]}]]]]]],[10,[7,[8,"arr",[5,"1"]],"y"]]]]"`,
-    );
-
-    const res = provideCtx(ctx, () => {
-      ctx[$internal].itemStateStack.pushFunctionScope(
-        'normal',
-        [],
-        {},
-        d.f32,
-        (astInfo.externals as () => Record<string, unknown>)() ?? {},
-      );
-
-      // Check for: const arr = [TestStruct({ x: 1, y: 2 }), TestStruct({ x: 3, y: 4 })];
-      //                        ^ this should be an array<TestStruct, 2>
-      wgslGenerator.initGenerator(ctx);
-      return wgslGenerator._expression(
-        (astInfo.ast?.body[1][0] as tinyest.Const)[2] as tinyest.Expression,
-      );
+    const arraySnippet = extractSnippetFromFn(() => {
+      'use gpu';
+      const arr = [TestStruct({ x: 1, y: 2 }), TestStruct({ x: 3, y: 4 })];
+      return arr;
     });
 
-    expect(d.isWgslArray(res.dataType)).toBe(true);
-    expect((res.dataType as unknown as WgslArray).elementCount).toBe(2);
-    expect((res.dataType as unknown as WgslArray).elementType).toBe(TestStruct);
+    expect(d.isWgslArray(arraySnippet.dataType)).toBe(true);
+    expect((arraySnippet.dataType as unknown as WgslArray).elementCount).toBe(2);
+    expect((arraySnippet.dataType as unknown as WgslArray).elementType).toBe(TestStruct);
   });
 
   it('generates correct code for array expressions with lazy elements', () => {
@@ -1296,37 +936,6 @@ describe('wgslGenerator', () => {
         return arr[1i].y;
       }"
     `);
-
-    const astInfo = getMetaData(
-      testFn[$internal].implementation as (...args: unknown[]) => unknown,
-    );
-
-    if (!astInfo) {
-      throw new Error('Expected prebuilt AST to be present');
-    }
-
-    expect(JSON.stringify(astInfo.ast?.body)).toMatchInlineSnapshot(
-      `"[0,[[13,"arr",[100,[[7,"lazyV2f","$"],[6,[7,"std","mul"],[[7,"lazyV2f","$"],[6,[7,"d","vec2f"],[[5,"2"],[5,"2"]]]]]]]],[10,[7,[8,"arr",[5,"1"]],"y"]]]]"`,
-    );
-
-    const res = provideCtx(ctx, () => {
-      ctx[$internal].itemStateStack.pushFunctionScope(
-        'normal',
-        [],
-        {},
-        d.f32,
-        (astInfo.externals as () => Record<string, unknown>)() ?? {},
-      );
-
-      wgslGenerator.initGenerator(ctx);
-      return wgslGenerator._expression(
-        (astInfo.ast?.body[1][0] as tinyest.Const)[2] as tinyest.Expression,
-      );
-    });
-
-    expect(d.isWgslArray(res.dataType)).toBe(true);
-    expect((res.dataType as unknown as WgslArray).elementCount).toBe(2);
-    expect((res.dataType as unknown as WgslArray).elementType).toBe(d.vec2f);
   });
 
   it('allows for member access on values returned from function calls', () => {
@@ -1364,34 +973,10 @@ describe('wgslGenerator', () => {
       }"
     `);
 
-    const astInfo = getMetaData(fnTwo[$internal].implementation as (...args: unknown[]) => unknown);
-
-    if (!astInfo) {
-      throw new Error('Expected prebuilt AST to be present');
-    }
-
-    expect(JSON.stringify(astInfo.ast?.body)).toMatchInlineSnapshot(
-      `"[0,[[10,[7,[7,[6,"fnOne",[]],"y"],"x"]]]]"`,
-    );
-
-    provideCtx(ctx, () => {
-      ctx[$internal].itemStateStack.pushFunctionScope(
-        'normal',
-        [],
-        {},
-        d.f32,
-        (astInfo.externals as () => Record<string, unknown>)() ?? {},
-      );
-
-      wgslGenerator.initGenerator(ctx);
-      // Check for: return fnOne().y.x;
-      //                   ^ this should be a f32
-      const res = wgslGenerator._expression(
-        (astInfo.ast?.body[1][0] as tinyest.Return)[1] as tinyest.Expression,
-      );
-
-      expect(res.dataType).toStrictEqual(d.f32);
-    });
+    expectDataTypeOf(() => {
+      'use gpu';
+      return fnOne().y.x;
+    }).toStrictEqual(d.f32);
   });
 
   it('generates correct code for conditional with single statement', () => {
@@ -1464,27 +1049,6 @@ describe('wgslGenerator', () => {
     `);
   });
 
-  it('generates correct code for for loops with single statements', () => {
-    const main = () => {
-      'use gpu';
-      for (let i = 0; i < 10; i += 1) {
-        continue;
-      }
-    };
-
-    const gen = provideCtx(ctx, () =>
-      wgslGenerator.functionDefinition(getMetaData(main)?.ast?.body as tinyest.Block),
-    );
-
-    expect(gen).toMatchInlineSnapshot(`
-      "{
-        for (var i = 0; (i < 10i); i += 1i) {
-          continue;
-        }
-      }"
-    `);
-  });
-
   it('generates correct code for while loops with single statements', () => {
     const main = tgpu.fn([])(() => {
       let i = 0;
@@ -1820,59 +1384,51 @@ describe('wgslGenerator', () => {
   it('block externals do not override identifiers', () => {
     const f = () => {
       'use gpu';
-      const y = 100;
-      const x = y;
-      return x;
+      const list = [1];
+      for (const x of tgpu.unroll(list)) {
+        const y = 100;
+        const x = y;
+        return x;
+      }
     };
 
-    const parsed = getMetaData(f)?.ast?.body as tinyest.Block;
-
-    provideCtx(ctx, () => {
-      ctx[$internal].itemStateStack.pushFunctionScope('normal', [], {}, d.u32, {});
-
-      const res = wgslGenerator._block(parsed, { x: 42 });
-
-      expect(res).toMatchInlineSnapshot(`
-          "{
-            const y = 100;
-            const x = y;
-            return u32(x);
-          }"
-        `);
-    });
+    expect(tgpu.resolve([f])).toMatchInlineSnapshot(`
+      "fn f() -> i32 {
+        var list = array<i32, 1>(1);
+        // unrolled iteration #0
+        {
+          const y = 100;
+          const x = y;
+          return x;
+        }
+      }"
+    `);
   });
 
   it('block externals are injected correctly', () => {
     const f = () => {
       'use gpu';
-      for (const x of []) {
+      for (const x of tgpu.unroll([1])) {
         const y = x;
       }
     };
 
-    const parsed = getMetaData(f)?.ast?.body as tinyest.Block;
-
-    provideCtx(ctx, () => {
-      ctx[$internal].itemStateStack.pushFunctionScope('normal', [], {}, d.Void, {});
-
-      const res = wgslGenerator._block((parsed[1][0] as tinyest.ForOf)[3] as tinyest.Block, {
-        x: 67,
-      });
-
-      expect(res).toMatchInlineSnapshot(`
-          "{
-            const y = 67;
-          }"
-        `);
-    });
+    expect(tgpu.resolve([f])).toMatchInlineSnapshot(`
+      "fn f() {
+        // unrolled iteration #0
+        {
+          const y = 1;
+        }
+      }"
+    `);
   });
 
   it('block externals are respected in nested blocks', () => {
     const f = () => {
       'use gpu';
       let result = d.i32(0);
-      const list = d.arrayOf(d.i32, 3)([1, 2, 3]);
-      for (const elem of list) {
+      const list = [1];
+      for (const elem of tgpu.unroll(list)) {
         {
           // We use the `elem` in a nested block
           result += elem;
@@ -1880,24 +1436,18 @@ describe('wgslGenerator', () => {
       }
     };
 
-    const parsed = getMetaData(f)?.ast?.body as tinyest.Block;
-
-    provideCtx(ctx, () => {
-      ctx[$internal].itemStateStack.pushFunctionScope('normal', [], {}, d.Void, {});
-
-      const res = wgslGenerator._block((parsed[1][2] as tinyest.ForOf)[3] as tinyest.Block, {
-        result: snip('result', d.i32, 'function'),
-        elem: 7,
-      });
-
-      expect(res).toMatchInlineSnapshot(`
-        "{
+    expect(tgpu.resolve([f])).toMatchInlineSnapshot(`
+      "fn f() {
+        var result = 0i;
+        var list = array<i32, 1>(1);
+        // unrolled iteration #0
+        {
           {
-            result += 7i;
+            result += list[0u];
           }
-        }"
-      `);
-    });
+        }
+      }"
+    `);
   });
 
   it('prunes comptime if/else', () => {
diff --git a/packages/typegpu/tests/tgslFn.test.ts b/packages/typegpu/tests/tgslFn.test.ts
index 529e9f0cfa..9278a82a00 100644
--- a/packages/typegpu/tests/tgslFn.test.ts
+++ b/packages/typegpu/tests/tgslFn.test.ts
@@ -363,11 +363,7 @@ describe('TGSL tgpu.fn function', () => {
     });
 
     expect(tgpu.resolve([fragmentFn])).toMatchInlineSnapshot(`
-      "struct fragmentFn_Input {
-        @location(0) uv: vec2f,
-      }
-
-      struct fragmentFn_Output {
+      "struct fragmentFn_Output {
         @builtin(sample_mask) sampleMask: u32,
         @builtin(frag_depth) fragDepth: f32,
         @location(0) out: vec4f,
@@ -410,11 +406,7 @@ describe('TGSL tgpu.fn function', () => {
     });
 
     expect(tgpu.resolve([fragmentFn])).toMatchInlineSnapshot(`
-      "struct fragmentFn_Input {
-        @location(0) uv: vec2f,
-      }
-
-      struct fragmentFn_Output {
+      "struct fragmentFn_Output {
         @builtin(sample_mask) sampleMask: u32,
         @builtin(frag_depth) fragDepth: f32,
         @location(0) out: vec4f,
@@ -444,11 +436,7 @@ describe('TGSL tgpu.fn function', () => {
     });
 
     expect(tgpu.resolve([fragmentFn])).toMatchInlineSnapshot(`
-      "struct fragmentFn_Input {
-        @location(0) uv: vec2f,
-      }
-
-      @fragment fn fragmentFn() -> @location(0) vec4f {
+      "@fragment fn fragmentFn() -> @location(0) vec4f {
         var hmm = vec4f(1.25);
         return hmm;
       }"
@@ -481,11 +469,7 @@ describe('TGSL tgpu.fn function', () => {
     });
 
     expect(tgpu.resolve([fragmentFn])).toMatchInlineSnapshot(`
-      "struct fragmentFn_Input {
-        @location(0) uv: vec2f,
-      }
-
-      struct fragmentFn_Output {
+      "struct fragmentFn_Output {
         @builtin(sample_mask) sampleMask: u32,
         @builtin(frag_depth) fragDepth: f32,
         @location(0) out: vec4f,
diff --git a/packages/typegpu/tests/utils/parseResolved.ts b/packages/typegpu/tests/utils/parseResolved.ts
index 7576fedc31..f6dffc8bdd 100644
--- a/packages/typegpu/tests/utils/parseResolved.ts
+++ b/packages/typegpu/tests/utils/parseResolved.ts
@@ -17,10 +17,10 @@ class ExtractingGenerator extends WgslGenerator {
     this.#fnDepth = 0;
   }
 
-  public functionDefinition(body: tinyest.Block): string {
+  public functionDefinition(options: ShaderGenerator.FunctionDefinitionOptions): string {
     this.#fnDepth++;
     try {
-      return super.functionDefinition(body);
+      return super.functionDefinition(options);
     } finally {
       this.#fnDepth--;
     }