github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-04-22 03:02:01 -04:00
Code Issues Packages Projects Releases Wiki Activity

[glsl-out] Es: convert 1D images to 2D (#1436)

Browse Source
This commit is contained in:
João Capucho
2021-10-04 04:01:44 +01:00
committed by GitHub
parent b34fa5f2cc
commit 71a75d727d
14 changed files with 540 additions and 282 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
src/back/glsl/features.rs
View File

@@ -22,8 +22,6 @@ bitflags::bitflags! {
/// Adds support for image load and early depth tests
const IMAGE_LOAD_STORE = 1 << 8;
const CONSERVATIVE_DEPTH = 1 << 9;
/// Isn't supported in ES
const TEXTURE_1D = 1 << 10;
/// Interpolation and auxiliary qualifiers. Perspective, Flat, and
/// Centroid are available in all GLSL versions we support.
const NOPERSPECTIVE_QUALIFIER = 1 << 11;
@@ -92,9 +90,6 @@ impl FeaturesManager {
check_feature!(IMAGE_LOAD_STORE, 130, 310);
check_feature!(CONSERVATIVE_DEPTH, 130, 300);
check_feature!(CONSERVATIVE_DEPTH, 130, 300);
// 1D textures are supported by all core versions and aren't supported by an es versions
// so use 0 that way the check will always be false and can be optimized away
check_feature!(TEXTURE_1D, 0);
check_feature!(NOPERSPECTIVE_QUALIFIER, 130);
check_feature!(SAMPLE_QUALIFIER, 400, 320);
// gl_ClipDistance is supported by core versions > 1.3 and aren't supported by an es versions without extensions
@@ -298,8 +293,6 @@ impl<'a, W> Writer<'a, W> {
} => {
if arrayed && dim == ImageDimension::Cube {
self.features.request(Features::CUBE_TEXTURES_ARRAY)
} else if dim == ImageDimension::D1 {
self.features.request(Features::TEXTURE_1D)
}
match class {

52
src/back/glsl/mod.rs
View File

@@ -514,7 +514,7 @@ impl<'a, W: Write> Writer<'a, W> {
// We treat images separately because they might require
// writing the storage format
TypeInner::Image {
dim,
mut dim,
arrayed,
class,
} => {
@@ -526,6 +526,11 @@ impl<'a, W: Write> Writer<'a, W> {
} => Some((format, access)),
_ => None,
};
if dim == crate::ImageDimension::D1 && es {
dim = crate::ImageDimension::D2
}
// Gether the location if needed
let layout_binding = if self.options.version.supports_explicit_locations() {
let br = global.binding.as_ref().unwrap();
@@ -1954,12 +1959,12 @@ impl<'a, W: Write> Writer<'a, W> {
// We need to get the coordinates vector size to later build a vector that's `size + 1`
// if `depth_ref` is some, if it isn't a vector we panic as that's not a valid expression
let size = match *ctx.info[coordinate].ty.inner_with(&self.module.types) {
TypeInner::Vector { size, .. } => size,
let mut coord_dim = match *ctx.info[coordinate].ty.inner_with(&self.module.types) {
TypeInner::Vector { size, .. } => size as u8,
TypeInner::Scalar { .. } => 1,
_ => unreachable!(),
};
let mut coord_dim = size as u8;
if array_index.is_some() {
coord_dim += 1;
}
@@ -1968,9 +1973,16 @@ impl<'a, W: Write> Writer<'a, W> {
coord_dim += 1;
}
let tex_1d_hack = dim == crate::ImageDimension::D1 && self.options.version.is_es();
let is_vec = tex_1d_hack || coord_dim != 1;
// Compose a new texture coordinates vector
write!(self.out, "vec{}(", coord_dim)?;
if is_vec {
write!(self.out, "vec{}(", coord_dim + tex_1d_hack as u8)?;
}
self.write_expr(coordinate, ctx)?;
if tex_1d_hack {
write!(self.out, ", 0.0")?;
}
if let Some(expr) = array_index {
write!(self.out, ", ")?;
self.write_expr(expr, ctx)?;
@@ -1981,7 +1993,9 @@ impl<'a, W: Write> Writer<'a, W> {
self.write_expr(expr, ctx)?;
}
}
write!(self.out, ")")?;
if is_vec {
write!(self.out, ")")?;
}
if cube_array_shadow {
if let Some(expr) = depth_ref {
@@ -2025,7 +2039,13 @@ impl<'a, W: Write> Writer<'a, W> {
if let Some(constant) = offset {
write!(self.out, ", ")?;
if tex_1d_hack {
write!(self.out, "ivec2(")?;
}
self.write_constant(constant)?;
if tex_1d_hack {
write!(self.out, ", 0)")?;
}
}
// End the function
@@ -2103,11 +2123,11 @@ impl<'a, W: Write> Writer<'a, W> {
ImageClass::Sampled { .. } | ImageClass::Depth { .. } => {
write!(self.out, "textureSize(")?;
self.write_expr(image, ctx)?;
write!(self.out, ",")?;
write!(self.out, ", ")?;
if let Some(expr) = level {
self.write_expr(expr, ctx)?;
} else {
write!(self.out, "0",)?;
write!(self.out, "0")?;
}
}
ImageClass::Storage { .. } => {
@@ -2115,7 +2135,10 @@ impl<'a, W: Write> Writer<'a, W> {
self.write_expr(image, ctx)?;
}
}
write!(self.out, ").{}", &"xyz"[..components])?;
write!(self.out, ")")?;
if components != 1 || self.options.version.is_es() {
write!(self.out, ".{}", &"xyz"[..components])?;
}
}
crate::ImageQuery::NumLevels => {
write!(self.out, "textureQueryLevels(",)?;
@@ -2129,7 +2152,9 @@ impl<'a, W: Write> Writer<'a, W> {
};
write!(self.out, "{}(", fun_name)?;
self.write_expr(image, ctx)?;
write!(self.out, ",0).{}", back::COMPONENTS[components])?;
if components != 1 || self.options.version.is_es() {
write!(self.out, ", 0).{}", back::COMPONENTS[components])?;
}
}
crate::ImageQuery::NumSamples => {
// assumes ARB_shader_texture_image_samples
@@ -2473,7 +2498,14 @@ impl<'a, W: Write> Writer<'a, W> {
write!(self.out, ")")?;
}
None => {
let tex_1d_hack = dim == IDim::D1 && self.options.version.is_es();
if tex_1d_hack {
write!(self.out, "ivec2(")?;
}
self.write_expr(coordinate, ctx)?;
if tex_1d_hack {
write!(self.out, ", 0.0)")?;
}
}
}
Ok(())

1
tests/in/image.param.ron
View File

@@ -3,4 +3,5 @@
version: (1, 1),
debug: true,
),
glsl_blacklist: ["depth_load", "levels_queries"]
)

42
tests/in/image.wgsl
View File

@@ -23,10 +23,18 @@ fn main(
let itc = dim * vec2<i32>(local_id.xy) % vec2<i32>(10, 20);
let value1 = textureLoad(image_mipmapped_src, itc, i32(local_id.z));
let value2 = textureLoad(image_multisampled_src, itc, i32(local_id.z));
let value3 = textureLoad(image_depth_multisampled_src, itc, i32(local_id.z));
let value4 = textureLoad(image_storage_src, itc);
let value5 = textureLoad(image_array_src, itc, i32(local_id.z), i32(local_id.z) + 1);
textureStore(image_dst, itc.x, value1 + value2 + u32(value3) + value4 + value5);
textureStore(image_dst, itc.x, value1 + value2 + value4 + value5);
}
[[stage(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 val: f32 = textureLoad(image_depth_multisampled_src, itc, i32(local_id.z));
textureStore(image_dst, itc.x, vec4<u32>(u32(val)));
return;
}
[[group(0), binding(0)]]
@@ -48,27 +56,34 @@ var image_aa: texture_multisampled_2d<f32>;
fn queries() -> [[builtin(position)]] vec4<f32> {
let dim_1d = textureDimensions(image_1d);
let dim_2d = textureDimensions(image_2d);
let num_levels_2d = textureNumLevels(image_2d);
let dim_2d_lod = textureDimensions(image_2d, 1);
let dim_2d_array = textureDimensions(image_2d_array);
let num_levels_2d_array = textureNumLevels(image_2d_array);
let dim_2d_array_lod = textureDimensions(image_2d_array, 1);
let num_layers_2d = textureNumLayers(image_2d_array);
let dim_cube = textureDimensions(image_cube);
let num_levels_cube = textureNumLevels(image_cube);
let dim_cube_lod = textureDimensions(image_cube, 1);
let dim_cube_array = textureDimensions(image_cube_array);
let num_levels_cube_array = textureNumLevels(image_cube_array);
let dim_cube_array_lod = textureDimensions(image_cube_array, 1);
let num_layers_cube = textureNumLayers(image_cube_array);
let dim_3d = textureDimensions(image_3d);
let num_levels_3d = textureNumLevels(image_3d);
let dim_3d_lod = textureDimensions(image_3d, 1);
let num_samples_aa = textureNumSamples(image_aa);
let sum = dim_1d + dim_2d.y + dim_2d_lod.y + dim_2d_array.y + dim_2d_array_lod.y +
num_layers_2d + dim_cube.y + dim_cube_lod.y + dim_cube_array.y + dim_cube_array_lod.y +
num_layers_cube + dim_3d.z + dim_3d_lod.z + num_samples_aa +
dim_cube.y + dim_cube_lod.y + dim_cube_array.y + dim_cube_array_lod.y +
dim_3d.z + dim_3d_lod.z;
return vec4<f32>(f32(sum));
}
[[stage(vertex)]]
fn levels_queries() -> [[builtin(position)]] vec4<f32> {
let num_levels_2d = textureNumLevels(image_2d);
let num_levels_2d_array = textureNumLevels(image_2d_array);
let num_layers_2d = textureNumLayers(image_2d_array);
let num_levels_cube = textureNumLevels(image_cube);
let num_levels_cube_array = textureNumLevels(image_cube_array);
let num_layers_cube = textureNumLayers(image_cube_array);
let num_levels_3d = textureNumLevels(image_3d);
let num_samples_aa = textureNumSamples(image_aa);
let sum = num_layers_2d + num_layers_cube + num_samples_aa +
num_levels_2d + num_levels_2d_array + num_levels_3d + num_levels_cube + num_levels_cube_array;
return vec4<f32>(f32(sum));
}
@@ -80,11 +95,12 @@ var sampler_reg: sampler;
fn sample() -> [[location(0)]] vec4<f32> {
let tc = vec2<f32>(0.5);
let level = 2.3;
let s1d = textureSample(image_1d, sampler_reg, tc.x);
let s2d = textureSample(image_2d, sampler_reg, tc);
let s2d_offset = textureSample(image_2d, sampler_reg, tc, vec2<i32>(3, 1));
let s2d_level = textureSampleLevel(image_2d, sampler_reg, tc, level);
let s2d_level_offset = textureSampleLevel(image_2d, sampler_reg, tc, level, vec2<i32>(3, 1));
return s2d + s2d_offset + s2d_level + s2d_level_offset;
return s1d + s2d + s2d_offset + s2d_level + s2d_level_offset;
}
[[group(1), binding(1)]]

31
tests/out/glsl/image.main.Compute.glsl Normal file
View File

@@ -0,0 +1,31 @@
#version 310 es
#extension GL_EXT_texture_cube_map_array : require
precision highp float;
precision highp int;
layout(local_size_x = 16, local_size_y = 1, local_size_z = 1) in;
uniform highp usampler2D _group_0_binding_0;
uniform highp usampler2DMS _group_0_binding_3;
layout(rgba8ui) readonly uniform highp uimage2D _group_0_binding_1;
uniform highp usampler2DArray _group_0_binding_5;
layout(r32ui) writeonly uniform highp uimage2D _group_0_binding_2;
void main() {
uvec3 local_id = gl_LocalInvocationID;
ivec2 dim = imageSize(_group_0_binding_1).xy;
ivec2 itc = ((dim * ivec2(local_id.xy)) % ivec2(10, 20));
uvec4 value1_ = texelFetch(_group_0_binding_0, itc, int(local_id.z));
uvec4 value2_ = texelFetch(_group_0_binding_3, itc, int(local_id.z));
uvec4 value4_ = imageLoad(_group_0_binding_1, itc);
uvec4 value5_ = texelFetch(_group_0_binding_5, ivec3(itc, int(local_id.z)), (int(local_id.z) + 1));
imageStore(_group_0_binding_2, ivec2(itc.x, 0.0), (((value1_ + value2_) + value4_) + value5_));
return;
}

37
tests/out/glsl/image.queries.Vertex.glsl Normal file
View File

@@ -0,0 +1,37 @@
#version 310 es
#extension GL_EXT_texture_cube_map_array : require
precision highp float;
precision highp int;
uniform highp sampler2D _group_0_binding_0;
uniform highp sampler2D _group_0_binding_1;
uniform highp sampler2DArray _group_0_binding_2;
uniform highp samplerCube _group_0_binding_3;
uniform highp samplerCubeArray _group_0_binding_4;
uniform highp sampler3D _group_0_binding_5;
void main() {
int dim_1d = textureSize(_group_0_binding_0, 0).x;
ivec2 dim_2d = textureSize(_group_0_binding_1, 0).xy;
ivec2 dim_2d_lod = textureSize(_group_0_binding_1, 1).xy;
ivec2 dim_2d_array = textureSize(_group_0_binding_2, 0).xy;
ivec2 dim_2d_array_lod = textureSize(_group_0_binding_2, 1).xy;
ivec2 dim_cube = textureSize(_group_0_binding_3, 0).xy;
ivec2 dim_cube_lod = textureSize(_group_0_binding_3, 1).xy;
ivec2 dim_cube_array = textureSize(_group_0_binding_4, 0).xy;
ivec2 dim_cube_array_lod = textureSize(_group_0_binding_4, 1).xy;
ivec3 dim_3d = textureSize(_group_0_binding_5, 0).xyz;
ivec3 dim_3d_lod = textureSize(_group_0_binding_5, 1).xyz;
int sum = ((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z);
gl_Position = vec4(float(sum));
gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w);
return;
}

23
tests/out/glsl/image.sample.Fragment.glsl Normal file
View File

@@ -0,0 +1,23 @@
#version 310 es
#extension GL_EXT_texture_cube_map_array : require
precision highp float;
precision highp int;
uniform highp sampler2D _group_0_binding_0;
uniform highp sampler2D _group_0_binding_1;
layout(location = 0) out vec4 _fs2p_location0;
void main() {
vec2 tc = vec2(0.5);
vec4 s1d = texture(_group_0_binding_0, vec2(tc.x, 0.0));
vec4 s2d = texture(_group_0_binding_1, vec2(tc));
vec4 s2d_offset = textureOffset(_group_0_binding_1, vec2(tc), ivec2(3, 1));
vec4 s2d_level = textureLod(_group_0_binding_1, vec2(tc), 2.299999952316284);
vec4 s2d_level_offset = textureLodOffset(_group_0_binding_1, vec2(tc), 2.299999952316284, ivec2(3, 1));
_fs2p_location0 = ((((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset);
return;
}

18
tests/out/glsl/image.sample_comparison.Fragment.glsl Normal file
View File

@@ -0,0 +1,18 @@
#version 310 es
#extension GL_EXT_texture_cube_map_array : require
precision highp float;
precision highp int;
uniform highp sampler2DShadow _group_1_binding_2;
layout(location = 0) out float _fs2p_location0;
void main() {
vec2 tc = vec2(0.5);
float s2d_depth = texture(_group_1_binding_2, vec3(tc, 0.5));
float s2d_depth_level = textureLod(_group_1_binding_2, vec3(tc, 0.5), 0.0);
_fs2p_location0 = (s2d_depth + s2d_depth_level);
return;
}

142
tests/out/hlsl/image.hlsl
View File

@@ -31,10 +31,19 @@ void main(uint3 local_id : SV_GroupThreadID)
int2 itc = ((dim * int2(local_id.xy)) % int2(10, 20));
uint4 value1_ = image_mipmapped_src.Load(int3(itc, int(local_id.z)));
uint4 value2_ = image_multisampled_src.Load(itc, int(local_id.z));
float value3_ = image_depth_multisampled_src.Load(itc, int(local_id.z)).x;
uint4 value4_ = image_storage_src.Load(itc);
uint4 value5_ = image_array_src.Load(int4(itc, int(local_id.z), (int(local_id.z) + 1)));
image_dst[itc.x] = ((((value1_ + value2_) + uint4(uint(value3_).xxxx)) + value4_) + value5_);
image_dst[itc.x] = (((value1_ + value2_) + value4_) + value5_);
return;
}
[numthreads(16, 1, 1)]
void depth_load(uint3 local_id1 : SV_GroupThreadID)
{
int2 dim1 = NagaRWDimensions2D(image_storage_src);
int2 itc1 = ((dim1 * int2(local_id1.xy)) % int2(10, 20));
float val = image_depth_multisampled_src.Load(itc1, int(local_id1.z)).x;
image_dst[itc1.x] = uint4(uint(val).xxxx);
return;
}
@@ -52,13 +61,6 @@ int2 NagaDimensions2D(Texture2D<float4> tex)
return ret.xy;
}
int NagaNumLevels2D(Texture2D<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z);
return ret.z;
}
int2 NagaMipDimensions2D(Texture2D<float4> tex, uint mip_level)
{
uint4 ret;
@@ -73,13 +75,6 @@ int2 NagaDimensions2DArray(Texture2DArray<float4> tex)
return ret.xy;
}
int NagaNumLevels2DArray(Texture2DArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int2 NagaMipDimensions2DArray(Texture2DArray<float4> tex, uint mip_level)
{
uint4 ret;
@@ -87,13 +82,6 @@ int2 NagaMipDimensions2DArray(Texture2DArray<float4> tex, uint mip_level)
return ret.xy;
}
int NagaNumLayers2DArray(Texture2DArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int2 NagaDimensionsCube(TextureCube<float4> tex)
{
uint4 ret;
@@ -101,13 +89,6 @@ int2 NagaDimensionsCube(TextureCube<float4> tex)
return ret.xy;
}
int NagaNumLevelsCube(TextureCube<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z);
return ret.z;
}
int2 NagaMipDimensionsCube(TextureCube<float4> tex, uint mip_level)
{
uint4 ret;
@@ -122,13 +103,6 @@ int2 NagaDimensionsCubeArray(TextureCubeArray<float4> tex)
return ret.xy;
}
int NagaNumLevelsCubeArray(TextureCubeArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int2 NagaMipDimensionsCubeArray(TextureCubeArray<float4> tex, uint mip_level)
{
uint4 ret;
@@ -136,13 +110,6 @@ int2 NagaMipDimensionsCubeArray(TextureCubeArray<float4> tex, uint mip_level)
return ret.xy;
}
int NagaNumLayersCubeArray(TextureCubeArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int3 NagaDimensions3D(Texture3D<float4> tex)
{
uint4 ret;
@@ -150,18 +117,77 @@ int3 NagaDimensions3D(Texture3D<float4> tex)
return ret.xyz;
}
int NagaNumLevels3D(Texture3D<float4> tex)
int3 NagaMipDimensions3D(Texture3D<float4> tex, uint mip_level)
{
uint4 ret;
tex.GetDimensions(mip_level, ret.x, ret.y, ret.z, ret.w);
return ret.xyz;
}
float4 queries() : SV_Position
{
int dim_1d = NagaDimensions1D(image_1d);
int2 dim_2d = NagaDimensions2D(image_2d);
int2 dim_2d_lod = NagaMipDimensions2D(image_2d, 1);
int2 dim_2d_array = NagaDimensions2DArray(image_2d_array);
int2 dim_2d_array_lod = NagaMipDimensions2DArray(image_2d_array, 1);
int2 dim_cube = NagaDimensionsCube(image_cube);
int2 dim_cube_lod = NagaMipDimensionsCube(image_cube, 1);
int2 dim_cube_array = NagaDimensionsCubeArray(image_cube_array);
int2 dim_cube_array_lod = NagaMipDimensionsCubeArray(image_cube_array, 1);
int3 dim_3d = NagaDimensions3D(image_3d);
int3 dim_3d_lod = NagaMipDimensions3D(image_3d, 1);
int sum = ((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z);
return float4(float(sum).xxxx);
}
int NagaNumLevels2D(Texture2D<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z);
return ret.z;
}
int NagaNumLevels2DArray(Texture2DArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int3 NagaMipDimensions3D(Texture3D<float4> tex, uint mip_level)
int NagaNumLayers2DArray(Texture2DArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(mip_level, ret.x, ret.y, ret.z, ret.w);
return ret.xyz;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int NagaNumLevelsCube(TextureCube<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z);
return ret.z;
}
int NagaNumLevelsCubeArray(TextureCubeArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int NagaNumLayersCubeArray(TextureCubeArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int NagaNumLevels3D(Texture3D<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int NagaMSNumSamples2D(Texture2DMS<float4> tex)
@@ -171,39 +197,29 @@ int NagaMSNumSamples2D(Texture2DMS<float4> tex)
return ret.z;
}
float4 queries() : SV_Position
float4 levels_queries() : SV_Position
{
int dim_1d = NagaDimensions1D(image_1d);
int2 dim_2d = NagaDimensions2D(image_2d);
int num_levels_2d = NagaNumLevels2D(image_2d);
int2 dim_2d_lod = NagaMipDimensions2D(image_2d, 1);
int2 dim_2d_array = NagaDimensions2DArray(image_2d_array);
int num_levels_2d_array = NagaNumLevels2DArray(image_2d_array);
int2 dim_2d_array_lod = NagaMipDimensions2DArray(image_2d_array, 1);
int num_layers_2d = NagaNumLayers2DArray(image_2d_array);
int2 dim_cube = NagaDimensionsCube(image_cube);
int num_levels_cube = NagaNumLevelsCube(image_cube);
int2 dim_cube_lod = NagaMipDimensionsCube(image_cube, 1);
int2 dim_cube_array = NagaDimensionsCubeArray(image_cube_array);
int num_levels_cube_array = NagaNumLevelsCubeArray(image_cube_array);
int2 dim_cube_array_lod = NagaMipDimensionsCubeArray(image_cube_array, 1);
int num_layers_cube = NagaNumLayersCubeArray(image_cube_array);
int3 dim_3d = NagaDimensions3D(image_3d);
int num_levels_3d = NagaNumLevels3D(image_3d);
int3 dim_3d_lod = NagaMipDimensions3D(image_3d, 1);
int num_samples_aa = NagaMSNumSamples2D(image_aa);
int sum = ((((((((((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + num_layers_2d) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + num_layers_cube) + dim_3d.z) + dim_3d_lod.z) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array);
return float4(float(sum).xxxx);
int sum1 = (((((((num_layers_2d + num_layers_cube) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array);
return float4(float(sum1).xxxx);
}
float4 sample1() : SV_Target0
{
float2 tc = float2(0.5.xx);
float4 s1d = image_1d.Sample(sampler_reg, tc.x);
float4 s2d = image_2d.Sample(sampler_reg, tc);
float4 s2d_offset = image_2d.Sample(sampler_reg, tc, int2(3, 1));
float4 s2d_level = image_2d.SampleLevel(sampler_reg, tc, 2.299999952316284);
float4 s2d_level_offset = image_2d.SampleLevel(sampler_reg, tc, 2.299999952316284, int2(3, 1));
return (((s2d + s2d_offset) + s2d_level) + s2d_level_offset);
return ((((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset);
}
float sample_comparison() : SV_Target0

4
tests/out/hlsl/image.hlsl.config
View File

@@ -1,3 +1,3 @@
vertex=(queries:vs_5_1 )
vertex=(queries:vs_5_1 levels_queries:vs_5_1 )
fragment=(sample1:ps_5_1 sample_comparison:ps_5_1 )
compute=(main:cs_5_1 )
compute=(main:cs_5_1 depth_load:cs_5_1 )

67
tests/out/msl/image.msl
View File

@@ -10,7 +10,6 @@ kernel void main1(
metal::uint3 local_id [[thread_position_in_threadgroup]]
, metal::texture2d<uint, metal::access::sample> image_mipmapped_src [[user(fake0)]]
, metal::texture2d_ms<uint, metal::access::read> image_multisampled_src [[user(fake0)]]
, metal::depth2d_ms<float, metal::access::read> image_depth_multisampled_src [[user(fake0)]]
, metal::texture2d<uint, metal::access::read> image_storage_src [[user(fake0)]]
, metal::texture2d_array<uint, metal::access::sample> image_array_src [[user(fake0)]]
, metal::texture1d<uint, metal::access::write> image_dst [[user(fake0)]]
@@ -19,16 +18,31 @@ kernel void main1(
metal::int2 itc = (dim * static_cast<metal::int2>(local_id.xy)) % metal::int2(10, 20);
metal::uint4 value1_ = image_mipmapped_src.read(metal::uint2(itc), static_cast<int>(local_id.z));
metal::uint4 value2_ = image_multisampled_src.read(metal::uint2(itc), static_cast<int>(local_id.z));
float value3_ = image_depth_multisampled_src.read(metal::uint2(itc), static_cast<int>(local_id.z));
metal::uint4 value4_ = image_storage_src.read(metal::uint2(itc));
metal::uint4 value5_ = image_array_src.read(metal::uint2(itc), static_cast<int>(local_id.z), static_cast<int>(local_id.z) + 1);
image_dst.write((((value1_ + value2_) + metal::uint4(static_cast<uint>(value3_))) + value4_) + value5_, metal::uint(itc.x));
image_dst.write(((value1_ + value2_) + value4_) + value5_, metal::uint(itc.x));
return;
}
struct depth_loadInput {
};
kernel void depth_load(
metal::uint3 local_id1 [[thread_position_in_threadgroup]]
, metal::depth2d_ms<float, metal::access::read> image_depth_multisampled_src [[user(fake0)]]
, metal::texture2d<uint, metal::access::read> image_storage_src [[user(fake0)]]
, metal::texture1d<uint, metal::access::write> image_dst [[user(fake0)]]
) {
metal::int2 dim1 = int2(image_storage_src.get_width(), image_storage_src.get_height());
metal::int2 itc1 = (dim1 * static_cast<metal::int2>(local_id1.xy)) % metal::int2(10, 20);
float val = image_depth_multisampled_src.read(metal::uint2(itc1), static_cast<int>(local_id1.z));
image_dst.write(metal::uint4(static_cast<uint>(val)), metal::uint(itc1.x));
return;
}
struct queriesOutput {
metal::float4 member1 [[position]];
metal::float4 member2 [[position]];
};
vertex queriesOutput queries(
metal::texture1d<float, metal::access::sample> image_1d [[user(fake0)]]
@@ -37,50 +51,67 @@ vertex queriesOutput queries(
, metal::texturecube<float, metal::access::sample> image_cube [[user(fake0)]]
, metal::texturecube_array<float, metal::access::sample> image_cube_array [[user(fake0)]]
, metal::texture3d<float, metal::access::sample> image_3d [[user(fake0)]]
, metal::texture2d_ms<float, metal::access::read> image_aa [[user(fake0)]]
) {
int dim_1d = int(image_1d.get_width());
metal::int2 dim_2d = int2(image_2d.get_width(), image_2d.get_height());
int num_levels_2d = int(image_2d.get_num_mip_levels());
metal::int2 dim_2d_lod = int2(image_2d.get_width(1), image_2d.get_height(1));
metal::int2 dim_2d_array = int2(image_2d_array.get_width(), image_2d_array.get_height());
int num_levels_2d_array = int(image_2d_array.get_num_mip_levels());
metal::int2 dim_2d_array_lod = int2(image_2d_array.get_width(1), image_2d_array.get_height(1));
int num_layers_2d = int(image_2d_array.get_array_size());
metal::int2 dim_cube = int2(image_cube.get_width());
int num_levels_cube = int(image_cube.get_num_mip_levels());
metal::int2 dim_cube_lod = int2(image_cube.get_width(1));
metal::int2 dim_cube_array = int2(image_cube_array.get_width());
int num_levels_cube_array = int(image_cube_array.get_num_mip_levels());
metal::int2 dim_cube_array_lod = int2(image_cube_array.get_width(1));
int num_layers_cube = int(image_cube_array.get_array_size());
metal::int3 dim_3d = int3(image_3d.get_width(), image_3d.get_height(), image_3d.get_depth());
int num_levels_3d = int(image_3d.get_num_mip_levels());
metal::int3 dim_3d_lod = int3(image_3d.get_width(1), image_3d.get_height(1), image_3d.get_depth(1));
int num_samples_aa = int(image_aa.get_num_samples());
int sum = (((((((((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + num_layers_2d) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + num_layers_cube) + dim_3d.z) + dim_3d_lod.z) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array;
int sum = (((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z;
return queriesOutput { metal::float4(static_cast<float>(sum)) };
}
struct levels_queriesOutput {
metal::float4 member3 [[position]];
};
vertex levels_queriesOutput levels_queries(
metal::texture2d<float, metal::access::sample> image_2d [[user(fake0)]]
, metal::texture2d_array<float, metal::access::sample> image_2d_array [[user(fake0)]]
, metal::texturecube<float, metal::access::sample> image_cube [[user(fake0)]]
, metal::texturecube_array<float, metal::access::sample> image_cube_array [[user(fake0)]]
, metal::texture3d<float, metal::access::sample> image_3d [[user(fake0)]]
, metal::texture2d_ms<float, metal::access::read> image_aa [[user(fake0)]]
) {
int num_levels_2d = int(image_2d.get_num_mip_levels());
int num_levels_2d_array = int(image_2d_array.get_num_mip_levels());
int num_layers_2d = int(image_2d_array.get_array_size());
int num_levels_cube = int(image_cube.get_num_mip_levels());
int num_levels_cube_array = int(image_cube_array.get_num_mip_levels());
int num_layers_cube = int(image_cube_array.get_array_size());
int num_levels_3d = int(image_3d.get_num_mip_levels());
int num_samples_aa = int(image_aa.get_num_samples());
int sum1 = ((((((num_layers_2d + num_layers_cube) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array;
return levels_queriesOutput { metal::float4(static_cast<float>(sum1)) };
}
struct sampleOutput {
metal::float4 member2 [[color(0)]];
metal::float4 member4 [[color(0)]];
};
fragment sampleOutput sample(
metal::texture2d<float, metal::access::sample> image_2d [[user(fake0)]]
metal::texture1d<float, metal::access::sample> image_1d [[user(fake0)]]
, metal::texture2d<float, metal::access::sample> image_2d [[user(fake0)]]
, metal::sampler sampler_reg [[user(fake0)]]
) {
metal::float2 tc = metal::float2(0.5);
metal::float4 s1d = image_1d.sample(sampler_reg, tc.x);
metal::float4 s2d = image_2d.sample(sampler_reg, tc);
metal::float4 s2d_offset = image_2d.sample(sampler_reg, tc, const_type8_);
metal::float4 s2d_level = image_2d.sample(sampler_reg, tc, metal::level(2.299999952316284));
metal::float4 s2d_level_offset = image_2d.sample(sampler_reg, tc, metal::level(2.299999952316284), const_type8_);
return sampleOutput { ((s2d + s2d_offset) + s2d_level) + s2d_level_offset };
return sampleOutput { (((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset };
}
struct sample_comparisonOutput {
float member3 [[color(0)]];
float member5 [[color(0)]];
};
fragment sample_comparisonOutput sample_comparison(
metal::sampler sampler_cmp [[user(fake0)]]

357
tests/out/spv/image.spvasm
View File

@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 215
; Bound: 244
OpCapability SampledCubeArray
OpCapability ImageQuery
OpCapability Image1D
@@ -10,12 +10,15 @@ OpCapability Sampled1D
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %69 "main" %66
OpEntryPoint Vertex %115 "queries" %113
OpEntryPoint Fragment %182 "sample" %181
OpEntryPoint Fragment %203 "sample_comparison" %201
OpEntryPoint GLCompute %107 "depth_load" %105
OpEntryPoint Vertex %128 "queries" %126
OpEntryPoint Vertex %176 "levels_queries" %175
OpEntryPoint Fragment %205 "sample" %204
OpEntryPoint Fragment %232 "sample_comparison" %230
OpExecutionMode %69 LocalSize 16 1 1
OpExecutionMode %182 OriginUpperLeft
OpExecutionMode %203 OriginUpperLeft
OpExecutionMode %107 LocalSize 16 1 1
OpExecutionMode %205 OriginUpperLeft
OpExecutionMode %232 OriginUpperLeft
OpSource GLSL 450
OpName %31 "image_mipmapped_src"
OpName %33 "image_multisampled_src"
@@ -36,9 +39,12 @@ OpName %61 "sampler_cmp"
OpName %63 "image_2d_depth"
OpName %66 "local_id"
OpName %69 "main"
OpName %115 "queries"
OpName %182 "sample"
OpName %203 "sample_comparison"
OpName %105 "local_id"
OpName %107 "depth_load"
OpName %128 "queries"
OpName %176 "levels_queries"
OpName %205 "sample"
OpName %232 "sample_comparison"
OpDecorate %31 DescriptorSet 0
OpDecorate %31 Binding 0
OpDecorate %33 DescriptorSet 0
@@ -77,9 +83,11 @@ OpDecorate %61 Binding 1
OpDecorate %63 DescriptorSet 1
OpDecorate %63 Binding 2
OpDecorate %66 BuiltIn LocalInvocationId
OpDecorate %113 BuiltIn Position
OpDecorate %181 Location 0
OpDecorate %201 Location 0
OpDecorate %105 BuiltIn LocalInvocationId
OpDecorate %126 BuiltIn Position
OpDecorate %175 BuiltIn Position
OpDecorate %204 Location 0
OpDecorate %230 Location 0
%2 = OpTypeVoid
%4 = OpTypeInt 32 1
%3 = OpConstant %4 10
@@ -146,166 +154,203 @@ OpDecorate %201 Location 0
%67 = OpTypePointer Input %18
%66 = OpVariable %67 Input
%70 = OpTypeFunction %2
%79 = OpTypeVector %12 2
%87 = OpTypeVector %12 4
%102 = OpTypeVector %4 3
%114 = OpTypePointer Output %27
%113 = OpVariable %114 Output
%124 = OpConstant %12 0
%181 = OpVariable %114 Output
%186 = OpTypeVector %8 2
%188 = OpTypeSampledImage %21
%202 = OpTypePointer Output %8
%201 = OpVariable %202 Output
%208 = OpTypeSampledImage %29
%213 = OpConstant %8 0.0
%78 = OpTypeVector %12 2
%86 = OpTypeVector %12 4
%97 = OpTypeVector %4 3
%105 = OpVariable %67 Input
%127 = OpTypePointer Output %27
%126 = OpVariable %127 Output
%136 = OpConstant %12 0
%175 = OpVariable %127 Output
%204 = OpVariable %127 Output
%210 = OpTypeVector %8 2
%213 = OpTypeSampledImage %20
%216 = OpTypeSampledImage %21
%231 = OpTypePointer Output %8
%230 = OpVariable %231 Output
%237 = OpTypeSampledImage %29
%242 = OpConstant %8 0.0
%69 = OpFunction %2 None %70
%65 = OpLabel
%68 = OpLoad %18 %66
%71 = OpLoad %11 %31
%72 = OpLoad %13 %33
%73 = OpLoad %14 %35
%74 = OpLoad %15 %37
%75 = OpLoad %16 %39
%76 = OpLoad %17 %43
OpBranch %77
%77 = OpLabel
%78 = OpImageQuerySize %19 %74
%80 = OpVectorShuffle %79 %68 %68 0 1
%81 = OpBitcast %19 %80
%82 = OpIMul %19 %78 %81
%83 = OpCompositeConstruct %19 %3 %5
%84 = OpSMod %19 %82 %83
%85 = OpCompositeExtract %12 %68 2
%86 = OpBitcast %4 %85
%88 = OpImageFetch %87 %71 %84 Lod %86
%89 = OpCompositeExtract %12 %68 2
%90 = OpBitcast %4 %89
%91 = OpImageFetch %87 %72 %84 Sample %90
%73 = OpLoad %15 %37
%74 = OpLoad %16 %39
%75 = OpLoad %17 %43
OpBranch %76
%76 = OpLabel
%77 = OpImageQuerySize %19 %73
%79 = OpVectorShuffle %78 %68 %68 0 1
%80 = OpBitcast %19 %79
%81 = OpIMul %19 %77 %80
%82 = OpCompositeConstruct %19 %3 %5
%83 = OpSMod %19 %81 %82
%84 = OpCompositeExtract %12 %68 2
%85 = OpBitcast %4 %84
%87 = OpImageFetch %86 %71 %83 Lod %85
%88 = OpCompositeExtract %12 %68 2
%89 = OpBitcast %4 %88
%90 = OpImageFetch %86 %72 %83 Sample %89
%91 = OpImageRead %86 %73 %83
%92 = OpCompositeExtract %12 %68 2
%93 = OpBitcast %4 %92
%94 = OpImageFetch %27 %73 %84 Sample %93
%95 = OpCompositeExtract %8 %94 0
%96 = OpImageRead %87 %74 %84
%97 = OpCompositeExtract %12 %68 2
%98 = OpBitcast %4 %97
%99 = OpCompositeExtract %12 %68 2
%100 = OpBitcast %4 %99
%101 = OpIAdd %4 %100 %6
%103 = OpCompositeConstruct %102 %84 %98
%104 = OpImageFetch %87 %75 %103 Lod %101
%105 = OpCompositeExtract %4 %84 0
%106 = OpIAdd %87 %88 %91
%107 = OpConvertFToU %12 %95
%108 = OpCompositeConstruct %87 %107 %107 %107 %107
%109 = OpIAdd %87 %106 %108
%110 = OpIAdd %87 %109 %96
%111 = OpIAdd %87 %110 %104
OpImageWrite %76 %105 %111
%94 = OpCompositeExtract %12 %68 2
%95 = OpBitcast %4 %94
%96 = OpIAdd %4 %95 %6
%98 = OpCompositeConstruct %97 %83 %93
%99 = OpImageFetch %86 %74 %98 Lod %96
%100 = OpCompositeExtract %4 %83 0
%101 = OpIAdd %86 %87 %90
%102 = OpIAdd %86 %101 %91
%103 = OpIAdd %86 %102 %99
OpImageWrite %75 %100 %103
OpReturn
OpFunctionEnd
%115 = OpFunction %2 None %70
%112 = OpLabel
%116 = OpLoad %20 %45
%117 = OpLoad %21 %47
%118 = OpLoad %22 %49
%119 = OpLoad %23 %51
%120 = OpLoad %24 %53
%121 = OpLoad %25 %55
%122 = OpLoad %26 %57
OpBranch %123
%123 = OpLabel
%125 = OpImageQuerySizeLod %4 %116 %124
%126 = OpImageQuerySizeLod %19 %117 %124
%127 = OpImageQueryLevels %4 %117
%128 = OpImageQuerySizeLod %19 %117 %6
%129 = OpImageQuerySizeLod %102 %118 %124
%130 = OpVectorShuffle %19 %129 %129 0 1
%131 = OpImageQueryLevels %4 %118
%132 = OpImageQuerySizeLod %102 %118 %6
%133 = OpVectorShuffle %19 %132 %132 0 1
%134 = OpImageQuerySizeLod %102 %118 %124
%135 = OpCompositeExtract %4 %134 2
%136 = OpImageQuerySizeLod %19 %119 %124
%137 = OpImageQueryLevels %4 %119
%138 = OpImageQuerySizeLod %19 %119 %6
%139 = OpImageQuerySizeLod %102 %120 %124
%140 = OpVectorShuffle %19 %139 %139 0 0
%141 = OpImageQueryLevels %4 %120
%142 = OpImageQuerySizeLod %102 %120 %6
%143 = OpVectorShuffle %19 %142 %142 0 0
%144 = OpImageQuerySizeLod %102 %120 %124
%145 = OpCompositeExtract %4 %144 2
%146 = OpImageQuerySizeLod %102 %121 %124
%147 = OpImageQueryLevels %4 %121
%148 = OpImageQuerySizeLod %102 %121 %6
%149 = OpImageQuerySamples %4 %122
%150 = OpCompositeExtract %4 %126 1
%151 = OpIAdd %4 %125 %150
%152 = OpCompositeExtract %4 %128 1
%153 = OpIAdd %4 %151 %152
%154 = OpCompositeExtract %4 %130 1
%107 = OpFunction %2 None %70
%104 = OpLabel
%106 = OpLoad %18 %105
%108 = OpLoad %14 %35
%109 = OpLoad %15 %37
%110 = OpLoad %17 %43
OpBranch %111
%111 = OpLabel
%112 = OpImageQuerySize %19 %109
%113 = OpVectorShuffle %78 %106 %106 0 1
%114 = OpBitcast %19 %113
%115 = OpIMul %19 %112 %114
%116 = OpCompositeConstruct %19 %3 %5
%117 = OpSMod %19 %115 %116
%118 = OpCompositeExtract %12 %106 2
%119 = OpBitcast %4 %118
%120 = OpImageFetch %27 %108 %117 Sample %119
%121 = OpCompositeExtract %8 %120 0
%122 = OpCompositeExtract %4 %117 0
%123 = OpConvertFToU %12 %121
%124 = OpCompositeConstruct %86 %123 %123 %123 %123
OpImageWrite %110 %122 %124
OpReturn
OpFunctionEnd
%128 = OpFunction %2 None %70
%125 = OpLabel
%129 = OpLoad %20 %45
%130 = OpLoad %21 %47
%131 = OpLoad %22 %49
%132 = OpLoad %23 %51
%133 = OpLoad %24 %53
%134 = OpLoad %25 %55
OpBranch %135
%135 = OpLabel
%137 = OpImageQuerySizeLod %4 %129 %136
%138 = OpImageQuerySizeLod %19 %130 %136
%139 = OpImageQuerySizeLod %19 %130 %6
%140 = OpImageQuerySizeLod %97 %131 %136
%141 = OpVectorShuffle %19 %140 %140 0 1
%142 = OpImageQuerySizeLod %97 %131 %6
%143 = OpVectorShuffle %19 %142 %142 0 1
%144 = OpImageQuerySizeLod %19 %132 %136
%145 = OpImageQuerySizeLod %19 %132 %6
%146 = OpImageQuerySizeLod %97 %133 %136
%147 = OpVectorShuffle %19 %146 %146 0 0
%148 = OpImageQuerySizeLod %97 %133 %6
%149 = OpVectorShuffle %19 %148 %148 0 0
%150 = OpImageQuerySizeLod %97 %134 %136
%151 = OpImageQuerySizeLod %97 %134 %6
%152 = OpCompositeExtract %4 %138 1
%153 = OpIAdd %4 %137 %152
%154 = OpCompositeExtract %4 %139 1
%155 = OpIAdd %4 %153 %154
%156 = OpCompositeExtract %4 %133 1
%156 = OpCompositeExtract %4 %141 1
%157 = OpIAdd %4 %155 %156
%158 = OpIAdd %4 %157 %135
%159 = OpCompositeExtract %4 %136 1
%160 = OpIAdd %4 %158 %159
%161 = OpCompositeExtract %4 %138 1
%162 = OpIAdd %4 %160 %161
%163 = OpCompositeExtract %4 %140 1
%164 = OpIAdd %4 %162 %163
%165 = OpCompositeExtract %4 %143 1
%166 = OpIAdd %4 %164 %165
%167 = OpIAdd %4 %166 %145
%168 = OpCompositeExtract %4 %146 2
%158 = OpCompositeExtract %4 %143 1
%159 = OpIAdd %4 %157 %158
%160 = OpCompositeExtract %4 %144 1
%161 = OpIAdd %4 %159 %160
%162 = OpCompositeExtract %4 %145 1
%163 = OpIAdd %4 %161 %162
%164 = OpCompositeExtract %4 %147 1
%165 = OpIAdd %4 %163 %164
%166 = OpCompositeExtract %4 %149 1
%167 = OpIAdd %4 %165 %166
%168 = OpCompositeExtract %4 %150 2
%169 = OpIAdd %4 %167 %168
%170 = OpCompositeExtract %4 %148 2
%170 = OpCompositeExtract %4 %151 2
%171 = OpIAdd %4 %169 %170
%172 = OpIAdd %4 %171 %149
%173 = OpIAdd %4 %172 %127
%174 = OpIAdd %4 %173 %131
%175 = OpIAdd %4 %174 %147
%176 = OpIAdd %4 %175 %137
%177 = OpIAdd %4 %176 %141
%178 = OpConvertSToF %8 %177
%179 = OpCompositeConstruct %27 %178 %178 %178 %178
OpStore %113 %179
%172 = OpConvertSToF %8 %171
%173 = OpCompositeConstruct %27 %172 %172 %172 %172
OpStore %126 %173
OpReturn
OpFunctionEnd
%182 = OpFunction %2 None %70
%180 = OpLabel
%183 = OpLoad %21 %47
%184 = OpLoad %28 %59
OpBranch %185
%185 = OpLabel
%187 = OpCompositeConstruct %186 %7 %7
%189 = OpSampledImage %188 %183 %184
%190 = OpImageSampleImplicitLod %27 %189 %187
%191 = OpSampledImage %188 %183 %184
%192 = OpImageSampleImplicitLod %27 %191 %187 ConstOffset %30
%193 = OpSampledImage %188 %183 %184
%194 = OpImageSampleExplicitLod %27 %193 %187 Lod %9
%195 = OpSampledImage %188 %183 %184
%196 = OpImageSampleExplicitLod %27 %195 %187 Lod|ConstOffset %9 %30
%197 = OpFAdd %27 %190 %192
%198 = OpFAdd %27 %197 %194
%199 = OpFAdd %27 %198 %196
OpStore %181 %199
%176 = OpFunction %2 None %70
%174 = OpLabel
%177 = OpLoad %21 %47
%178 = OpLoad %22 %49
%179 = OpLoad %23 %51
%180 = OpLoad %24 %53
%181 = OpLoad %25 %55
%182 = OpLoad %26 %57
OpBranch %183
%183 = OpLabel
%184 = OpImageQueryLevels %4 %177
%185 = OpImageQueryLevels %4 %178
%186 = OpImageQuerySizeLod %97 %178 %136
%187 = OpCompositeExtract %4 %186 2
%188 = OpImageQueryLevels %4 %179
%189 = OpImageQueryLevels %4 %180
%190 = OpImageQuerySizeLod %97 %180 %136
%191 = OpCompositeExtract %4 %190 2
%192 = OpImageQueryLevels %4 %181
%193 = OpImageQuerySamples %4 %182
%194 = OpIAdd %4 %187 %191
%195 = OpIAdd %4 %194 %193
%196 = OpIAdd %4 %195 %184
%197 = OpIAdd %4 %196 %185
%198 = OpIAdd %4 %197 %192
%199 = OpIAdd %4 %198 %188
%200 = OpIAdd %4 %199 %189
%201 = OpConvertSToF %8 %200
%202 = OpCompositeConstruct %27 %201 %201 %201 %201
OpStore %175 %202
OpReturn
OpFunctionEnd
%203 = OpFunction %2 None %70
%200 = OpLabel
%204 = OpLoad %28 %61
%205 = OpLoad %29 %63
OpBranch %206
%206 = OpLabel
%207 = OpCompositeConstruct %186 %7 %7
%209 = OpSampledImage %208 %205 %204
%210 = OpImageSampleDrefImplicitLod %8 %209 %207 %7
%211 = OpSampledImage %208 %205 %204
%212 = OpImageSampleDrefExplicitLod %8 %211 %207 %7 Lod %213
%214 = OpFAdd %8 %210 %212
OpStore %201 %214
%205 = OpFunction %2 None %70
%203 = OpLabel
%206 = OpLoad %20 %45
%207 = OpLoad %21 %47
%208 = OpLoad %28 %59
OpBranch %209
%209 = OpLabel
%211 = OpCompositeConstruct %210 %7 %7
%212 = OpCompositeExtract %8 %211 0
%214 = OpSampledImage %213 %206 %208
%215 = OpImageSampleImplicitLod %27 %214 %212
%217 = OpSampledImage %216 %207 %208
%218 = OpImageSampleImplicitLod %27 %217 %211
%219 = OpSampledImage %216 %207 %208
%220 = OpImageSampleImplicitLod %27 %219 %211 ConstOffset %30
%221 = OpSampledImage %216 %207 %208
%222 = OpImageSampleExplicitLod %27 %221 %211 Lod %9
%223 = OpSampledImage %216 %207 %208
%224 = OpImageSampleExplicitLod %27 %223 %211 Lod|ConstOffset %9 %30
%225 = OpFAdd %27 %215 %218
%226 = OpFAdd %27 %225 %220
%227 = OpFAdd %27 %226 %222
%228 = OpFAdd %27 %227 %224
OpStore %204 %228
OpReturn
OpFunctionEnd
%232 = OpFunction %2 None %70
%229 = OpLabel
%233 = OpLoad %28 %61
%234 = OpLoad %29 %63
OpBranch %235
%235 = OpLabel
%236 = OpCompositeConstruct %210 %7 %7
%238 = OpSampledImage %237 %234 %233
%239 = OpImageSampleDrefImplicitLod %8 %238 %236 %7
%240 = OpSampledImage %237 %234 %233
%241 = OpImageSampleDrefExplicitLod %8 %240 %236 %7 Lod %242
%243 = OpFAdd %8 %239 %241
OpStore %230 %243
OpReturn
OpFunctionEnd

39
tests/out/wgsl/image.wgsl
View File

@@ -39,10 +39,18 @@ fn main([[builtin(local_invocation_id)]] local_id: vec3<u32>) {
let itc: vec2<i32> = ((dim * vec2<i32>(local_id.xy)) % vec2<i32>(10, 20));
let value1_: vec4<u32> = textureLoad(image_mipmapped_src, itc, i32(local_id.z));
let value2_: vec4<u32> = textureLoad(image_multisampled_src, itc, i32(local_id.z));
let value3_: f32 = textureLoad(image_depth_multisampled_src, itc, i32(local_id.z));
let value4_: vec4<u32> = textureLoad(image_storage_src, itc);
let value5_: vec4<u32> = textureLoad(image_array_src, itc, i32(local_id.z), (i32(local_id.z) + 1));
textureStore(image_dst, itc.x, ((((value1_ + value2_) + vec4<u32>(u32(value3_))) + value4_) + value5_));
textureStore(image_dst, itc.x, (((value1_ + value2_) + value4_) + value5_));
return;
}
[[stage(compute), workgroup_size(16, 1, 1)]]
fn depth_load([[builtin(local_invocation_id)]] local_id1: vec3<u32>) {
let dim1: vec2<i32> = textureDimensions(image_storage_src);
let itc1: vec2<i32> = ((dim1 * vec2<i32>(local_id1.xy)) % vec2<i32>(10, 20));
let val: f32 = textureLoad(image_depth_multisampled_src, itc1, i32(local_id1.z));
textureStore(image_dst, itc1.x, vec4<u32>(u32(val)));
return;
}
@@ -50,35 +58,42 @@ fn main([[builtin(local_invocation_id)]] local_id: vec3<u32>) {
fn queries() -> [[builtin(position)]] vec4<f32> {
let dim_1d: i32 = textureDimensions(image_1d);
let dim_2d: vec2<i32> = textureDimensions(image_2d);
let num_levels_2d: i32 = textureNumLevels(image_2d);
let dim_2d_lod: vec2<i32> = textureDimensions(image_2d, 1);
let dim_2d_array: vec2<i32> = textureDimensions(image_2d_array);
let num_levels_2d_array: i32 = textureNumLevels(image_2d_array);
let dim_2d_array_lod: vec2<i32> = textureDimensions(image_2d_array, 1);
let num_layers_2d: i32 = textureNumLayers(image_2d_array);
let dim_cube: vec2<i32> = textureDimensions(image_cube);
let num_levels_cube: i32 = textureNumLevels(image_cube);
let dim_cube_lod: vec2<i32> = textureDimensions(image_cube, 1);
let dim_cube_array: vec2<i32> = textureDimensions(image_cube_array);
let num_levels_cube_array: i32 = textureNumLevels(image_cube_array);
let dim_cube_array_lod: vec2<i32> = textureDimensions(image_cube_array, 1);
let num_layers_cube: i32 = textureNumLayers(image_cube_array);
let dim_3d: vec3<i32> = textureDimensions(image_3d);
let num_levels_3d: i32 = textureNumLevels(image_3d);
let dim_3d_lod: vec3<i32> = textureDimensions(image_3d, 1);
let num_samples_aa: i32 = textureNumSamples(image_aa);
let sum: i32 = ((((((((((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + num_layers_2d) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + num_layers_cube) + dim_3d.z) + dim_3d_lod.z) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array);
let sum: i32 = ((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z);
return vec4<f32>(f32(sum));
}
[[stage(vertex)]]
fn levels_queries() -> [[builtin(position)]] vec4<f32> {
let num_levels_2d: i32 = textureNumLevels(image_2d);
let num_levels_2d_array: i32 = textureNumLevels(image_2d_array);
let num_layers_2d: i32 = textureNumLayers(image_2d_array);
let num_levels_cube: i32 = textureNumLevels(image_cube);
let num_levels_cube_array: i32 = textureNumLevels(image_cube_array);
let num_layers_cube: i32 = textureNumLayers(image_cube_array);
let num_levels_3d: i32 = textureNumLevels(image_3d);
let num_samples_aa: i32 = textureNumSamples(image_aa);
let sum1: i32 = (((((((num_layers_2d + num_layers_cube) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array);
return vec4<f32>(f32(sum1));
}
[[stage(fragment)]]
fn sample() -> [[location(0)]] vec4<f32> {
let tc: vec2<f32> = vec2<f32>(0.5);
let s1d: vec4<f32> = textureSample(image_1d, sampler_reg, tc.x);
let s2d: vec4<f32> = textureSample(image_2d, sampler_reg, tc);
let s2d_offset: vec4<f32> = textureSample(image_2d, sampler_reg, tc, vec2<i32>(3, 1));
let s2d_level: vec4<f32> = textureSampleLevel(image_2d, sampler_reg, tc, 2.299999952316284);
let s2d_level_offset: vec4<f32> = textureSampleLevel(image_2d, sampler_reg, tc, 2.299999952316284, vec2<i32>(3, 1));
return (((s2d + s2d_offset) + s2d_level) + s2d_level_offset);
return ((((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset);
}
[[stage(fragment)]]

2
tests/snapshots.rs
View File

@@ -488,7 +488,7 @@ fn convert_wgsl() {
),
(
"image",
Targets::SPIRV | Targets::METAL | Targets::HLSL | Targets::WGSL,
Targets::SPIRV | Targets::METAL | Targets::HLSL | Targets::WGSL | Targets::GLSL,
),
("extra", Targets::SPIRV | Targets::METAL | Targets::WGSL),
(