hal/gl: Allow push constants trough emulation

cts_runner/examples/hello-compute.js

@@ -4,11 +4,12 @@ const numbers = [1, 4, 3, 295];
 
 const device = await adapter.requestDevice();
 
-const shaderCode = `[[block]]
+const shaderCode = `@block
 struct PrimeIndices {
- data: [[stride(4)]] array<u32>;
+ data: @stride(4) array<u32>;
 }; // this is used as both input and output for convenience
-[[group(0), binding(0)]]
+@group(0)
+@binding(0)
 var<storage, read_write> v_indices: PrimeIndices;
 // The Collatz Conjecture states that for any integer n:
 // If n is even, n = n/2
@@ -37,8 +38,9 @@ fn collatz_iterations(n_base: u32) -> u32{
  }
  return i;
 }
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(global_invocation_id)]] global_id: vec3<u32>) {
+@stage(compute)
+@workgroup_size(1)
+fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
  v_indices.data[global_id.x] = collatz_iterations(v_indices.data[global_id.x]);
 }`;
 

player/tests/data/empty.wgsl

@@ -1,3 +1,4 @@
-[[stage(compute), workgroup_size(1)]]
+@stage(compute)
+@workgroup_size(1)
 fn main() {
 }

player/tests/data/quad.wgsl

@@ -1,5 +1,5 @@
-[[stage(vertex)]]
-fn vs_main([[builtin(vertex_index)]] vertex_index: u32) -> [[builtin(position)]] vec4<f32> {
+@stage(vertex)
+fn vs_main(@builtin(vertex_index) vertex_index: u32) -> @builtin(position) vec4<f32> {
  // hacky way to draw a large triangle
  let tmp1 = i32(vertex_index) / 2;
  let tmp2 = i32(vertex_index) & 1;
@@ -10,7 +10,7 @@ fn vs_main([[builtin(vertex_index)]] vertex_index: u32) -> [[builtin(position)]]
  return vec4<f32>(pos, 0.0, 1.0);
 }
 
-[[stage(fragment)]]
-fn fs_main() -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main() -> @location(0) vec4<f32> {
  return vec4<f32>(1.0, 1.0, 1.0, 1.0);
 }

player/tests/data/zero-init-buffer-for-binding.wgsl

@@ -1,11 +1,13 @@
 struct InOutBuffer {
- data: [[stride(4)]] array<u32>;
+ data: @stride(4) array<u32>;
 };
 
-[[group(0), binding(0)]]
+@group(0)
+@binding(0)
 var<storage, read_write> buffer: InOutBuffer;
 
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(global_invocation_id)]] global_id: vec3<u32>) {
+@stage(compute)
+@workgroup_size(1)
+fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
  buffer.data[global_id.x] = buffer.data[global_id.x] + global_id.x;
 }

player/tests/data/zero-init-texture-binding.wgsl

@@ -1,6 +1,7 @@
-[[group(0), binding(0)]] var tex: texture_2d<f32>;
-[[group(0), binding(1)]] var tex_storage: texture_storage_2d<rgba8uint, write>;
+@group(0) @binding(0) var tex: texture_2d<f32>;
+@group(0) @binding(1) var tex_storage: texture_storage_2d<rgba8uint, write>;
 
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(global_invocation_id)]] global_id: vec3<u32>) {
+@stage(compute)
+@workgroup_size(1)
+fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
 }

wgpu-core/Cargo.toml

@@ -38,7 +38,7 @@ thiserror = "1"
 
 [dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "a1840be"
+rev = "81dc674"
 #version = "0.8"
 features = ["span", "validate", "wgsl-in"]
 

wgpu-hal/Cargo.toml

@@ -82,14 +82,14 @@ js-sys = { version = "0.3" }
 
 [dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "a1840be"
+rev = "81dc674"
 #version = "0.8"
 
 # DEV dependencies
 
 [dev-dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "a1840be"
+rev = "81dc674"
 #version = "0.8"
 features = ["wgsl-in"]
 

wgpu-hal/examples/halmark/shader.wgsl

@@ -9,33 +9,37 @@ struct Locals {
  color: u32;
 };
 
-[[group(0), binding(0)]]
+@group(0)
+@binding(0)
 var<uniform> globals: Globals;
 
-[[group(1), binding(0)]]
+@group(1)
+@binding(0)
 var<uniform> locals: Locals;
 
 struct VertexOutput {
- [[builtin(position)]] position: vec4<f32>;
- [[location(0)]] tex_coords: vec2<f32>;
- [[location(1)]] color: vec4<f32>;
+ @builtin(position) position: vec4<f32>;
+ @location(0) tex_coords: vec2<f32>;
+ @location(1) color: vec4<f32>;
 };
 
-[[stage(vertex)]]
-fn vs_main([[builtin(vertex_index)]] vi: u32) -> VertexOutput {
+@stage(vertex)
+fn vs_main(@builtin(vertex_index) vi: u32) -> VertexOutput {
  let tc = vec2<f32>(f32(vi & 1u), 0.5 * f32(vi & 2u));
  let offset = vec2<f32>(tc.x * globals.size.x, tc.y * globals.size.y);
  let pos = globals.mvp * vec4<f32>(locals.position + offset, 0.0, 1.0);
  let color = vec4<f32>((vec4<u32>(locals.color) >> vec4<u32>(0u, 8u, 16u, 24u)) & vec4<u32>(255u)) / 255.0;
  return VertexOutput(pos, tc, color);
 }
 
-[[group(0), binding(1)]]
+@group(0)
+@binding(1)
 var texture: texture_2d<f32>;
-[[group(0), binding(2)]]
+@group(0)
+@binding(2)
 var sam: sampler;
 
-[[stage(fragment)]]
-fn fs_main(in: VertexOutput) -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
  return in.color * textureSampleLevel(texture, sam, in.tex_coords, 0.0);
 }

wgpu-hal/src/gles/adapter.rs

@@ -286,7 +286,8 @@ impl super::Adapter {
 
  let mut features = wgt::Features::empty()
  | wgt::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES
- | wgt::Features::CLEAR_TEXTURE;
+ | wgt::Features::CLEAR_TEXTURE
+ | wgt::Features::PUSH_CONSTANTS;
  features.set(
  wgt::Features::ADDRESS_MODE_CLAMP_TO_BORDER | wgt::Features::ADDRESS_MODE_CLAMP_TO_ZERO,
  extensions.contains("GL_EXT_texture_border_clamp"),
@@ -399,7 +400,7 @@ impl super::Adapter {
  } else {
  !0
  },
- max_push_constant_size: 0,
+ max_push_constant_size: super::MAX_PUSH_CONSTANTS as u32 * 4,
  min_uniform_buffer_offset_alignment,
  min_storage_buffer_offset_alignment,
  max_inter_stage_shader_components: gl.get_parameter_i32(glow::MAX_VARYING_COMPONENTS)

wgpu-hal/src/gles/command.rs

@@ -28,6 +28,7 @@ pub(super) struct State {
  has_pass_label: bool,
  instance_vbuf_mask: usize,
  dirty_vbuf_mask: usize,
+ push_offset_to_uniform: ArrayVec<super::UniformDesc, { super::MAX_PUSH_CONSTANTS }>,
 }
 
 impl super::CommandBuffer {
@@ -43,6 +44,21 @@ impl super::CommandBuffer {
  self.data_bytes.extend(marker.as_bytes());
  start..self.data_bytes.len() as u32
  }
+
+ fn add_push_constant_data(&mut self, data: &[u32]) -> Range<u32> {
+ let data_raw = unsafe {
+ std::slice::from_raw_parts(
+ data.as_ptr() as *const _,
+ data.len() * mem::size_of::<u32>(),
+ )
+ };
+ let start = self.data_bytes.len();
+ assert!(start < u32::MAX as usize);
+ self.data_bytes.extend_from_slice(data_raw);
+ let end = self.data_bytes.len();
+ assert!(end < u32::MAX as usize);
+ (start as u32)..(end as u32)
+ }
 }
 
 impl super::CommandEncoder {
@@ -148,8 +164,10 @@ impl super::CommandEncoder {
  fn set_pipeline_inner(&mut self, inner: &super::PipelineInner) {
  self.cmd_buffer.commands.push(C::SetProgram(inner.program));
 
- //TODO: push constants
- let _ = &inner.uniforms;
+ self.state.push_offset_to_uniform.clear();
+ self.state
+ .push_offset_to_uniform
+ .extend(inner.uniforms.iter().cloned());
 
  // rebind textures, if needed
  let mut dirty_textures = 0u32;
@@ -603,10 +621,25 @@ impl crate::CommandEncoder<super::Api> for super::CommandEncoder {
  &mut self,
  _layout: &super::PipelineLayout,
  _stages: wgt::ShaderStages,
- _offset: u32,
- _data: &[u32],
+ start_offset: u32,
+ data: &[u32],
  ) {
- unimplemented!()
+ let range = self.cmd_buffer.add_push_constant_data(data);
+
+ let end = start_offset + data.len() as u32 * 4;
+ let mut offset = start_offset;
+ while offset < end {
+ let uniform = self.state.push_offset_to_uniform[offset as usize / 4].clone();
+ let size = uniform.size;
+ if uniform.location.is_none() {
+ panic!("No uniform for push constant");
+ }
+ self.cmd_buffer.commands.push(C::SetPushConstants {
+ uniform,
+ offset: range.start + offset,
+ });
+ offset += size;
+ }
  }
 
  unsafe fn insert_debug_marker(&mut self, label: &str) {

wgpu-hal/src/gles/conv.rs

@@ -375,3 +375,59 @@ pub(super) fn map_storage_access(access: wgt::StorageTextureAccess) -> u32 {
  wgt::StorageTextureAccess::ReadWrite => glow::READ_WRITE,
  }
 }
+
+pub(super) fn is_sampler(glsl_uniform_type: u32) -> bool {
+ match glsl_uniform_type {
+ glow::INT_SAMPLER_1D
+ | glow::INT_SAMPLER_1D_ARRAY
+ | glow::INT_SAMPLER_2D
+ | glow::INT_SAMPLER_2D_ARRAY
+ | glow::INT_SAMPLER_2D_MULTISAMPLE
+ | glow::INT_SAMPLER_2D_MULTISAMPLE_ARRAY
+ | glow::INT_SAMPLER_2D_RECT
+ | glow::INT_SAMPLER_3D
+ | glow::INT_SAMPLER_CUBE
+ | glow::INT_SAMPLER_CUBE_MAP_ARRAY
+ | glow::UNSIGNED_INT_SAMPLER_1D
+ | glow::UNSIGNED_INT_SAMPLER_1D_ARRAY
+ | glow::UNSIGNED_INT_SAMPLER_2D
+ | glow::UNSIGNED_INT_SAMPLER_2D_ARRAY
+ | glow::UNSIGNED_INT_SAMPLER_2D_MULTISAMPLE
+ | glow::UNSIGNED_INT_SAMPLER_2D_MULTISAMPLE_ARRAY
+ | glow::UNSIGNED_INT_SAMPLER_2D_RECT
+ | glow::UNSIGNED_INT_SAMPLER_3D
+ | glow::UNSIGNED_INT_SAMPLER_CUBE
+ | glow::UNSIGNED_INT_SAMPLER_CUBE_MAP_ARRAY
+ | glow::SAMPLER_1D
+ | glow::SAMPLER_1D_SHADOW
+ | glow::SAMPLER_1D_ARRAY
+ | glow::SAMPLER_1D_ARRAY_SHADOW
+ | glow::SAMPLER_2D
+ | glow::SAMPLER_2D_SHADOW
+ | glow::SAMPLER_2D_ARRAY
+ | glow::SAMPLER_2D_ARRAY_SHADOW
+ | glow::SAMPLER_2D_MULTISAMPLE
+ | glow::SAMPLER_2D_MULTISAMPLE_ARRAY
+ | glow::SAMPLER_2D_RECT
+ | glow::SAMPLER_2D_RECT_SHADOW
+ | glow::SAMPLER_3D
+ | glow::SAMPLER_CUBE
+ | glow::SAMPLER_CUBE_MAP_ARRAY
+ | glow::SAMPLER_CUBE_MAP_ARRAY_SHADOW
+ | glow::SAMPLER_CUBE_SHADOW => true,
+ _ => false,
+ }
+}
+
+pub(super) fn uniform_byte_size(glsl_uniform_type: u32) -> u32 {
+ match glsl_uniform_type {
+ glow::FLOAT | glow::INT => 4,
+ glow::FLOAT_VEC2 | glow::INT_VEC2 => 8,
+ glow::FLOAT_VEC3 | glow::INT_VEC3 => 12,
+ glow::FLOAT_VEC4 | glow::INT_VEC4 => 16,
+ glow::FLOAT_MAT2 => 16,
+ glow::FLOAT_MAT3 => 36,
+ glow::FLOAT_MAT4 => 64,
+ _ => panic!("Unsupported uniform datatype!"),
+ }
+}

wgpu-hal/src/gles/device.rs

@@ -272,30 +272,35 @@ impl super::Device {
  }
  }
 
- let uniforms = {
- let count = gl.get_active_uniforms(program);
- let mut offset = 0;
- let mut uniforms = Vec::new();
-
- for uniform in 0..count {
- let glow::ActiveUniform { size, utype, name } =
- gl.get_active_uniform(program, uniform).unwrap();
-
- if let Some(location) = gl.get_uniform_location(program, &name) {
- // Sampler2D won't show up in UniformLocation and the only other uniforms
- // should be push constants
- uniforms.push(super::UniformDesc {
- location,
- offset,
- utype,
- });
+ let mut uniforms: [super::UniformDesc; super::MAX_PUSH_CONSTANTS] = Default::default();
+ let count = gl.get_active_uniforms(program);
+ let mut offset = 0;
 
- offset += size as u32;
- }
+ for uniform in 0..count {
+ let glow::ActiveUniform { utype, name, .. } =
+ gl.get_active_uniform(program, uniform).unwrap();
+
+ if conv::is_sampler(utype) {
+ continue;
  }
 
- uniforms.into_boxed_slice()
- };
+ if let Some(location) = gl.get_uniform_location(program, &name) {
+ if uniforms[offset / 4].location.is_some() {
+ panic!("Offset already occupied")
+ }
+
+ // `size` will always be 1 so we need to guess the real size from the type
+ let uniform_size = conv::uniform_byte_size(utype);
+
+ uniforms[offset / 4] = super::UniformDesc {
+ location: Some(location),
+ size: uniform_size,
+ utype,
+ };
+
+ offset += uniform_size as usize;
+ }
+ }
 
  Ok(super::PipelineInner {
  program,
@@ -858,7 +863,6 @@ impl crate::Device<super::Api> for super::Device {
  version: self.shared.shading_language_version,
  writer_flags,
  binding_map,
- push_constant_binding: 0, //TODO?
  },
  })
  }