Fix texture built-ins where u32 was expected (#2245)

- The Typifier was updated to expect `uint`
- Both `glsl` and `spv` frontends where updated to cast the result to `sint`.
- Both `glsl` and `spv` backends where updated to cast the result to `uint`.
- Remove cast in `msl` backend.

src/back/glsl/mod.rs

@@ -2534,6 +2534,16 @@ impl<'a, W: Write> Writer<'a, W> {
  crate::ImageDimension::D3 => 3,
  crate::ImageDimension::Cube => 2,
  };
+
+ if let crate::ImageQuery::Size { .. } = query {
+ match components {
+ 1 => write!(self.out, "uint(")?,
+ _ => write!(self.out, "uvec{components}(")?,
+ }
+ } else {
+ write!(self.out, "uint(")?;
+ }
+
  match query {
  crate::ImageQuery::Size { level } => {
  match class {
@@ -2593,6 +2603,8 @@ impl<'a, W: Write> Writer<'a, W> {
  write!(self.out, ")",)?;
  }
  }
+
+ write!(self.out, ")")?;
  }
  // `Unary` is pretty straightforward
  // "-" - for `Negate`

src/back/msl/writer.rs

@@ -1453,24 +1453,21 @@ impl<W: Write> Writer<W> {
  self.put_image_size_query(
  image,
  level.map(LevelOfDetail::Direct),
- crate::ScalarKind::Sint,
+ crate::ScalarKind::Uint,
  context,
  )?;
  }
  crate::ImageQuery::NumLevels => {
- write!(self.out, "int(")?;
  self.put_expression(image, context, false)?;
- write!(self.out, ".get_num_mip_levels())")?;
+ write!(self.out, ".get_num_mip_levels()")?;
  }
  crate::ImageQuery::NumLayers => {
- write!(self.out, "int(")?;
  self.put_expression(image, context, false)?;
- write!(self.out, ".get_array_size())")?;
+ write!(self.out, ".get_array_size()")?;
  }
  crate::ImageQuery::NumSamples => {
- write!(self.out, "int(")?;
  self.put_expression(image, context, false)?;
- write!(self.out, ".get_num_samples())")?;
+ write!(self.out, ".get_num_samples()")?;
  }
  },
  crate::Expression::Unary { op, expr } => {

src/back/spv/image.rs

@@ -1032,21 +1032,19 @@ impl<'w> BlockContext<'w> {
  Id::D2 | Id::Cube => 2,
  Id::D3 => 3,
  };
- let extended_size_type_id = {
- let array_coords = usize::from(arrayed);
- let vector_size = match dim_coords + array_coords {
- 2 => Some(crate::VectorSize::Bi),
- 3 => Some(crate::VectorSize::Tri),
- 4 => Some(crate::VectorSize::Quad),
- _ => None,
- };
- self.get_type_id(LookupType::Local(LocalType::Value {
- vector_size,
- kind: crate::ScalarKind::Sint,
- width: 4,
- pointer_space: None,
- }))
+ let array_coords = usize::from(arrayed);
+ let vector_size = match dim_coords + array_coords {
+ 2 => Some(crate::VectorSize::Bi),
+ 3 => Some(crate::VectorSize::Tri),
+ 4 => Some(crate::VectorSize::Quad),
+ _ => None,
  };
+ let extended_size_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
+ vector_size,
+ kind: crate::ScalarKind::Sint,
+ width: 4,
+ pointer_space: None,
+ }));
 
  let (query_op, level_id) = match class {
  Ic::Sampled { multi: true, .. }
@@ -1075,7 +1073,24 @@ impl<'w> BlockContext<'w> {
  }
  block.body.push(inst);
 
- if result_type_id != extended_size_type_id {
+ let bitcast_type_id = self.get_type_id(
+ LocalType::Value {
+ vector_size,
+ kind: crate::ScalarKind::Uint,
+ width: 4,
+ pointer_space: None,
+ }
+ .into(),
+ );
+ let bitcast_id = self.gen_id();
+ block.body.push(Instruction::unary(
+ spirv::Op::Bitcast,
+ bitcast_type_id,
+ bitcast_id,
+ id_extended,
+ ));
+
+ if result_type_id != bitcast_type_id {
  let id = self.gen_id();
  let components = match dim {
  // always pick the first component, and duplicate it for all 3 dimensions
@@ -1085,23 +1100,41 @@ impl<'w> BlockContext<'w> {
  block.body.push(Instruction::vector_shuffle(
  result_type_id,
  id,
- id_extended,
- id_extended,
+ bitcast_id,
+ bitcast_id,
  components,
  ));
+
  id
  } else {
- id_extended
+ bitcast_id
  }
  }
  Iq::NumLevels => {
- let id = self.gen_id();
+ let query_id = self.gen_id();
  block.body.push(Instruction::image_query(
  spirv::Op::ImageQueryLevels,
+ self.get_type_id(
+ LocalType::Value {
+ vector_size: None,
+ kind: crate::ScalarKind::Sint,
+ width: 4,
+ pointer_space: None,
+ }
+ .into(),
+ ),
+ query_id,
+ image_id,
+ ));
+
+ let id = self.gen_id();
+ block.body.push(Instruction::unary(
+ spirv::Op::Bitcast,
  result_type_id,
  id,
- image_id,
+ query_id,
  ));
+
  id
  }
  Iq::NumLayers => {
@@ -1125,23 +1158,58 @@ impl<'w> BlockContext<'w> {
  );
  inst.add_operand(self.get_index_constant(0));
  block.body.push(inst);
- let id = self.gen_id();
+
+ let extract_id = self.gen_id();
  block.body.push(Instruction::composite_extract(
- result_type_id,
- id,
+ self.get_type_id(
+ LocalType::Value {
+ vector_size: None,
+ kind: crate::ScalarKind::Sint,
+ width: 4,
+ pointer_space: None,
+ }
+ .into(),
+ ),
+ extract_id,
  id_extended,
  &[vec_size as u32 - 1],
  ));
+
+ let id = self.gen_id();
+ block.body.push(Instruction::unary(
+ spirv::Op::Bitcast,
+ result_type_id,
+ id,
+ extract_id,
+ ));
+
  id
  }
  Iq::NumSamples => {
- let id = self.gen_id();
+ let query_id = self.gen_id();
  block.body.push(Instruction::image_query(
  spirv::Op::ImageQuerySamples,
+ self.get_type_id(
+ LocalType::Value {
+ vector_size: None,
+ kind: crate::ScalarKind::Sint,
+ width: 4,
+ pointer_space: None,
+ }
+ .into(),
+ ),
+ query_id,
+ image_id,
+ ));
+
+ let id = self.gen_id();
+ block.body.push(Instruction::unary(
+ spirv::Op::Bitcast,
  result_type_id,
  id,
- image_id,
+ query_id,
  ));
+
  id
  }
  };

src/front/glsl/builtins.rs

@@ -1873,7 +1873,7 @@ impl MacroCall {
  name: None,
  inner: TypeInner::Vector {
  size,
- kind: crate::ScalarKind::Sint,
+ kind: crate::ScalarKind::Uint,
  width: 4,
  },
  },
@@ -1883,7 +1883,15 @@ impl MacroCall {
  expr = ctx.add_expression(Expression::Compose { components, ty }, meta, body)
  }
 
- expr
+ ctx.add_expression(
+ Expression::As {
+ expr,
+ kind: Sk::Sint,
+ convert: Some(4),
+ },
+ Span::default(),
+ body,
+ )
  }
  MacroCall::ImageLoad { multi } => {
  let comps =

src/front/spv/image.rs

@@ -684,6 +684,11 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
  image: image_lexp.handle,
  query: crate::ImageQuery::Size { level },
  };
+ let expr = crate::Expression::As {
+ expr: ctx.expressions.append(expr, self.span_from_with_op(start)),
+ kind: crate::ScalarKind::Sint,
+ convert: Some(4),
+ };
  self.lookup_expression.insert(
  result_id,
  LookupExpression {
@@ -714,6 +719,11 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
  image: image_lexp.handle,
  query,
  };
+ let expr = crate::Expression::As {
+ expr: expressions.append(expr, self.span_from_with_op(start)),
+ kind: crate::ScalarKind::Sint,
+ convert: Some(4),
+ };
  self.lookup_expression.insert(
  result_id,
  LookupExpression {

src/front/wgsl/tests.rs

@@ -357,10 +357,10 @@ fn parse_texture_query() {
  "
  var t: texture_multisampled_2d_array<f32>;
  fn foo() {
- var dim: vec2<i32> = textureDimensions(t);
+ var dim: vec2<u32> = textureDimensions(t);
  dim = textureDimensions(t, 0);
- let layers: i32 = textureNumLayers(t);
- let samples: i32 = textureNumSamples(t);
+ let layers: u32 = textureNumLayers(t);
+ let samples: u32 = textureNumSamples(t);
  }
  ",
  )

src/proc/typifier.rs

@@ -525,17 +525,17 @@ impl<'a> ResolveContext<'a> {
  crate::ImageQuery::Size { level: _ } => match *past(image)?.inner_with(types) {
  Ti::Image { dim, .. } => match dim {
  crate::ImageDimension::D1 => Ti::Scalar {
- kind: crate::ScalarKind::Sint,
+ kind: crate::ScalarKind::Uint,
  width: 4,
  },
  crate::ImageDimension::D2 | crate::ImageDimension::Cube => Ti::Vector {
  size: crate::VectorSize::Bi,
- kind: crate::ScalarKind::Sint,
+ kind: crate::ScalarKind::Uint,
  width: 4,
  },
  crate::ImageDimension::D3 => Ti::Vector {
  size: crate::VectorSize::Tri,
- kind: crate::ScalarKind::Sint,
+ kind: crate::ScalarKind::Uint,
  width: 4,
  },
  },
@@ -547,7 +547,7 @@ impl<'a> ResolveContext<'a> {
  crate::ImageQuery::NumLevels
  | crate::ImageQuery::NumLayers
  | crate::ImageQuery::NumSamples => Ti::Scalar {
- kind: crate::ScalarKind::Sint,
+ kind: crate::ScalarKind::Uint,
  width: 4,
  },
  }),

tests/in/binding-arrays.wgsl

@@ -30,8 +30,8 @@ fn main(fragment_in: FragmentIn) -> @location(0) vec4<f32> {
  let uniform_index = uni.index;
  let non_uniform_index = fragment_in.index;
 
- var i1 = 0;
- var i2 = vec2<i32>(0);
+ var u1 = 0u;
+ var u2 = vec2<u32>(0u);
  var v1 = 0.0;
  var v4 = vec4<f32>(0.0);
 
@@ -46,9 +46,9 @@ fn main(fragment_in: FragmentIn) -> @location(0) vec4<f32> {
  let uv = vec2<f32>(0.0);
  let pix = vec2<i32>(0);
 
- i2 += textureDimensions(texture_array_unbounded[0]);
- i2 += textureDimensions(texture_array_unbounded[uniform_index]);
- i2 += textureDimensions(texture_array_unbounded[non_uniform_index]);
+ u2 += textureDimensions(texture_array_unbounded[0]);
+ u2 += textureDimensions(texture_array_unbounded[uniform_index]);
+ u2 += textureDimensions(texture_array_unbounded[non_uniform_index]);
 
  v4 += textureGather(0, texture_array_bounded[0], samp[0], uv);
  v4 += textureGather(0, texture_array_bounded[uniform_index], samp[uniform_index], uv);
@@ -62,17 +62,17 @@ fn main(fragment_in: FragmentIn) -> @location(0) vec4<f32> {
  v4 += textureLoad(texture_array_unbounded[uniform_index], pix, 0);
  v4 += textureLoad(texture_array_unbounded[non_uniform_index], pix, 0);
 
- i1 += textureNumLayers(texture_array_2darray[0]);
- i1 += textureNumLayers(texture_array_2darray[uniform_index]);
- i1 += textureNumLayers(texture_array_2darray[non_uniform_index]);
+ u1 += textureNumLayers(texture_array_2darray[0]);
+ u1 += textureNumLayers(texture_array_2darray[uniform_index]);
+ u1 += textureNumLayers(texture_array_2darray[non_uniform_index]);
 
- i1 += textureNumLevels(texture_array_bounded[0]);
- i1 += textureNumLevels(texture_array_bounded[uniform_index]);
- i1 += textureNumLevels(texture_array_bounded[non_uniform_index]);
+ u1 += textureNumLevels(texture_array_bounded[0]);
+ u1 += textureNumLevels(texture_array_bounded[uniform_index]);
+ u1 += textureNumLevels(texture_array_bounded[non_uniform_index]);
 
- i1 += textureNumSamples(texture_array_multisampled[0]);
- i1 += textureNumSamples(texture_array_multisampled[uniform_index]);
- i1 += textureNumSamples(texture_array_multisampled[non_uniform_index]);
+ u1 += textureNumSamples(texture_array_multisampled[0]);
+ u1 += textureNumSamples(texture_array_multisampled[uniform_index]);
+ u1 += textureNumSamples(texture_array_multisampled[non_uniform_index]);
 
  v4 += textureSample(texture_array_bounded[0], samp[0], uv);
  v4 += textureSample(texture_array_bounded[uniform_index], samp[uniform_index], uv);
@@ -102,7 +102,7 @@ fn main(fragment_in: FragmentIn) -> @location(0) vec4<f32> {
  textureStore(texture_array_storage[uniform_index], pix, v4);
  textureStore(texture_array_storage[non_uniform_index], pix, v4);
 
- let v2 = vec2<f32>(i2 + vec2<i32>(i1));
+ let v2 = vec2<f32>(u2 + vec2<u32>(u1));
 
  return v4 + vec4<f32>(v2.x, v2.y, v2.x, v2.y) + v1;
 }

tests/in/image.wgsl

@@ -18,7 +18,7 @@ var image_dst: texture_storage_1d<r32uint,write>;
 @compute @workgroup_size(16)
 fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
  let dim = textureDimensions(image_storage_src);
- let itc = dim * vec2<i32>(local_id.xy) % vec2<i32>(10, 20);
+ let itc = vec2<i32>(dim * local_id.xy) % vec2<i32>(10, 20);
  // loads with ivec2 coords.
  let value1 = textureLoad(image_mipmapped_src, itc, i32(local_id.z));
  let value2 = textureLoad(image_multisampled_src, itc, i32(local_id.z));
@@ -39,8 +39,8 @@ fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
 
 @compute @workgroup_size(16, 1, 1)
 fn depth_load(@builtin(local_invocation_id) local_id: vec3<u32>) {
- let dim: vec2<i32> = textureDimensions(image_storage_src);
- let itc: vec2<i32> = ((dim * vec2<i32>(local_id.xy)) % vec2<i32>(10, 20));
+ let dim: vec2<u32> = textureDimensions(image_storage_src);
+ let itc: vec2<i32> = (vec2<i32>(dim * local_id.xy) % vec2<i32>(10, 20));
  let val: f32 = textureLoad(image_depth_multisampled_src, itc, i32(local_id.z));
  textureStore(image_dst, itc.x, vec4<u32>(u32(val)));
  return;