Harbor

Changelog 9ad271dcb849

pin gpu workflow docs examples

@sky · 1 month ago · parent a7bc9d349271
0 added 4 modified 0 deleted
gpu/examples/offscreen_ir/main.odin +21 -34 modified
13 unchanged lines hidden
14 14 }
15 15
16 16 VERTEX_SHADER :: `
17 - struct In
18 - @location(0) position: vec2
19 - @location(1) uv: vec2
20 - end
21 -
22 - @varying
23 - struct Varying
24 - @builtin(position) position: vec4
25 - @location(0) uv: vec2
26 - end
27 -
28 - @entry(vertex)
29 - function main(input: In) -> Varying
30 - return Varying{position = vec4(input.position, 0.0, 1.0), uv = input.uv}
17 + vertex main
18 + in position: vec2 @location(0)
19 + in uv: vec2 @location(1)
20 + out position: vec4 @builtin(position)
21 + out uv: vec2
22 + do
23 + position = vec4(position, 0.0, 1.0)
24 + out.uv = uv
31 25 end
32 26 `
33 27
34 28 OFFSCREEN_FRAGMENT_SHADER :: `
35 - @varying
36 - struct Varying
37 - @builtin(position) position: vec4
38 - @location(0) uv: vec2
39 - end
40 -
41 - @entry(fragment)
42 - function main(input: Varying) -> vec4
43 - let band = step(0.5, fract((input.uv.x + input.uv.y) * 8.0))
44 - return mix(vec4(0.05, 0.15, 0.65, 1.0), vec4(input.uv.x, input.uv.y, 0.25, 1.0), band)
29 + fragment main
30 + in uv: vec2
31 + out color: vec4
32 + do
33 + let band = step(0.5, fract((uv.x + uv.y) * 8.0))
34 + color = mix(vec4(0.05, 0.15, 0.65, 1.0), vec4(uv.x, uv.y, 0.25, 1.0), band)
45 35 end
46 36 `
47 37
48 38 SAMPLE_FRAGMENT_SHADER :: `
49 - @group(0) @binding(0)
50 - uniform offscreen_color: sampler2D
51 -
52 - @varying
53 - struct Varying
54 - @builtin(position) position: vec4
55 - @location(0) uv: vec2
39 + group frame = 0
40 + @binding(0) texture offscreen_color: sampler2D
56 41 end
57 42
58 - @entry(fragment)
59 - function main(input: Varying) -> vec4
60 - return sample(offscreen_color, input.uv)
43 + fragment main
44 + in uv: vec2
45 + out color: vec4
46 + do
47 + color = sample(offscreen_color, uv)
61 48 end
62 49 `
63 50
182 unchanged lines hidden
gpu/examples/particles_compute/main.odin +11 -12 modified
43 unchanged lines hidden
44 44 @push_constant
45 45 uniform pc: PC
46 46
47 - @group(0) @binding(0)
48 - buffer storage particles: []Particle
47 + group sim = 0
48 + @binding(0) buffer storage particles: []Particle
49 + end
49 50
50 - @workgroup_size(256, 1, 1)
51 - @entry(compute)
52 - function main(@builtin(global_invocation_id) gid: uvec3)
53 - let i = gid.x
54 - if i >= pc.count then
55 - return
51 + compute main @workgroup_size(256, 1, 1)
52 + do
53 + let i = work.global_id.x
54 + if i < pc.count then
55 + particles[i].position = vec2(400.0, 300.0)
56 + particles[i].velocity = vec2(0.0, 0.0)
57 + particles[i].color = vec4(1.0, 0.7, 0.1, 1.0)
58 + particles[i].life = 1.0
56 59 end
57 - particles[i].position = vec2(400.0, 300.0)
58 - particles[i].velocity = vec2(0.0, 0.0)
59 - particles[i].color = vec4(1.0, 0.7, 0.1, 1.0)
60 - particles[i].life = 1.0
61 60 end
62 61 `
63 62
91 unchanged lines hidden
gpu/shader/README.md +8 -6 modified
48 unchanged lines hidden
49 49 end
50 50 ```
51 51
52 - Current canonical syntax:
52 + Compatibility syntax:
53 53
54 54 ```lua
55 55 struct Uniforms
69 unchanged lines hidden
125 125 end
126 126 ```
127 127
128 - V1 workflow output writes must be top-level assignments (`color = expr` or
129 - `out.color = expr`). Writes inside `if`, `for`, or `while` are rejected until
130 - control-flow definite assignment lands. Missing or duplicate output writes are
131 - compile errors.
128 + Workflow output writes are assignments (`color = expr` or `out.color = expr`).
129 + They may appear in `if`/`elseif`/`else` when every non-discard path definitely
130 + assigns each output before stage exit. Overwrites are allowed. Assignments in
131 + loops are allowed as overwrites but do not prove final assignment after the
132 + loop. `return expr` is only fragment single-output top-level sugar and cannot be
133 + mixed with explicit `out` writes.
132 134
133 135 ### Bindings
134 136
4 unchanged lines hidden
139 141
140 142 group material = 1
141 143 uniform albedo: sampler2D -- group block, auto binding
142 - buffer storage particles: ParticleBuffer @binding(4)
144 + @binding(4) buffer storage particles: ParticleBuffer
143 145 end
144 146 ```
145 147
139 unchanged lines hidden
gpu/shader/docs/language.md +51 -22 modified
103 unchanged lines hidden
104 104
105 105 ## Entry Points
106 106
107 - ### Standard (struct-based I/O)
107 + ### Workflow stages
108 108
109 109 ```lua
110 - @entry(vertex)
111 - function main(input: VertexInput) -> VertexOutput
112 - -- ...
110 + vertex main
111 + in position: vec3 @location(0)
112 + in uv: vec2 @location(1)
113 + out clip: vec4 @builtin(position)
114 + out uv: vec2
115 + do
116 + clip = vec4(position, 1.0)
117 + out.uv = uv
113 118 end
114 - ```
115 119
116 - ### Inline Parameters
117 -
118 - Parameters can carry attributes directly:
119 -
120 - ```lua
121 - @entry(fragment)
122 - function main(@location(0) uv: vec2) -> (color: vec4)
120 + fragment shade
121 + in uv: vec2
122 + out color: vec4
123 + do
123 124 return vec4(uv, 0.0, 1.0)
124 125 end
126 +
127 + compute simulate @workgroup_size(64, 1, 1)
128 + do
129 + let i = work.global_id.x
130 + -- write storage buffers/resources
131 + end
125 132 ```
126 133
127 - ### Single Output Shorthand
134 + `return expr` is only fragment single-output top-level sugar. Vertex stages use
135 + explicit `out` assignments. Compute stages do not return and cannot declare
136 + `out` slots.
128 137
138 + Workflow output assignments can appear in control flow if every non-discard path
139 + definitely assigns every output before stage exit. Overwrites are allowed.
140 +
141 + ### Compatibility entry syntax
142 +
143 + The older annotation syntax remains supported:
144 +
129 145 ```lua
146 + @entry(vertex)
147 + function main(input: VertexInput) -> VertexOutput
148 + -- ...
149 + end
150 +
130 151 @entry(fragment)
131 - function main(@location(0) uv: vec2) -> vec4
152 + function main(@location(0) uv: vec2) -> (color: vec4)
132 153 return vec4(uv, 0.0, 1.0)
133 154 end
134 155 ```
31 unchanged lines hidden
166 187 ### Uniform Buffers
167 188
168 189 ```lua
169 - @group(0) @binding(0)
170 - uniform camera: CameraData
190 + group frame = 0
191 + @binding(0) uniform camera: CameraData
192 + end
171 193 ```
172 194
173 195 ### Storage Buffers
174 196
175 197 ```lua
176 - @group(0) @binding(1)
177 - buffer particles: ParticleBuffer
198 + group sim = 0
199 + @binding(1) buffer particles: ParticleBuffer
200 + end
178 201 ```
179 202
180 203 ### Textures
181 204
182 205 ```lua
183 - @group(1) @binding(0)
184 - uniform diffuse: sampler2D
206 + group material = 1
207 + @binding(0) texture diffuse: sampler2D
208 + end
185 209 ```
186 210
187 211 ### Auto-numbering
5 unchanged lines hidden
193 217 uniform b: TypeB -- group=0, binding=1
194 218 @group(1)
195 219 uniform c: sampler2D -- group=1, binding=0 (texture), binding=1 (sampler)
220 +
221 + group material = 1
222 + texture c: sampler2D
223 + end
196 224 ```
197 225
198 226 ## Shared Memory
3 unchanged lines hidden
202 230 ```lua
203 231 shared tile: [64]float
204 232
205 - @entry(compute) @workgroup_size(64, 1, 1)
206 - function main(@builtin(local_invocation_index) idx: uint)
233 + compute main @workgroup_size(64, 1, 1)
234 + do
235 + let idx = work.local_index
207 236 tile[idx] = float(idx)
208 237 barrier()
209 238 -- all threads can now read tile
81 unchanged lines hidden