From 4b5ab528ac1203a9201230d1c910d1c8f735d906 Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Fri, 28 May 2021 22:39:18 -0400 Subject: [PATCH] Comparison sampling tests --- tests/in/image.wgsl | 14 ++ tests/out/image.msl | 14 ++ tests/out/image.spvasm | 351 +++++++++++++++++++++++------------------ tests/out/image.wgsl | 12 ++ 4 files changed, 234 insertions(+), 157 deletions(-) diff --git a/tests/in/image.wgsl b/tests/in/image.wgsl index 440cd351d6..05cfe75059 100644 --- a/tests/in/image.wgsl +++ b/tests/in/image.wgsl @@ -58,3 +58,17 @@ fn queries() -> [[builtin(position)]] vec4 { num_levels_2d + num_levels_2d_array + num_levels_3d + num_levels_cube + num_levels_cube_array; return vec4(f32(sum)); } + +[[group(1), binding(0)]] +var sampler_cmp: sampler_comparison; +[[group(1), binding(1)]] +var image_2d_depth: texture_depth_2d; + +[[stage(fragment)]] +fn sample_comparison() -> [[location(0)]] f32 { + let tc = vec2(0.5); + let dref = 0.5; + let s2d_depth = textureSampleCompare(image_2d_depth, sampler_cmp, tc, dref); + let s2d_depth_level = textureSampleCompareLevel(image_2d_depth, sampler_cmp, tc, dref); + return s2d_depth + s2d_depth_level; +} diff --git a/tests/out/image.msl b/tests/out/image.msl index a8d37578eb..e569e45044 100644 --- a/tests/out/image.msl +++ b/tests/out/image.msl @@ -30,3 +30,17 @@ vertex queriesOutput queries( ) { return queriesOutput { metal::float4(static_cast((((((((((((((((((int(image_1d.get_width()) + int2(image_2d.get_width(), image_2d.get_height()).y) + int2(image_2d.get_width(1), image_2d.get_height(1)).y) + int2(image_2d_array.get_width(), image_2d_array.get_height()).y) + int2(image_2d_array.get_width(1), image_2d_array.get_height(1)).y) + int(image_2d_array.get_array_size())) + int3(image_cube.get_width()).y) + int3(image_cube.get_width(1)).y) + int3(image_cube_array.get_width()).y) + int3(image_cube_array.get_width(1)).y) + int(image_cube_array.get_array_size())) + int3(image_3d.get_width(), image_3d.get_height(), image_3d.get_depth()).z) + int3(image_3d.get_width(1), image_3d.get_height(1), image_3d.get_depth(1)).z) + int(image_aa.get_num_samples())) + int(image_2d.get_num_mip_levels())) + int(image_2d_array.get_num_mip_levels())) + int(image_3d.get_num_mip_levels())) + int(image_cube.get_num_mip_levels())) + int(image_cube_array.get_num_mip_levels()))) }; } + + +struct sample_comparisonOutput { + float member2 [[color(0)]]; +}; +fragment sample_comparisonOutput sample_comparison( + metal::sampler sampler_cmp [[user(fake0)]] +, metal::depth2d image_2d_depth [[user(fake0)]] +) { + metal::float2 _e12 = metal::float2(0.5); + float _e14 = image_2d_depth.sample_compare(sampler_cmp, _e12, 0.5); + float _e15 = image_2d_depth.sample_compare(sampler_cmp, _e12, 0.5); + return sample_comparisonOutput { _e14 + _e15 }; +} diff --git a/tests/out/image.spvasm b/tests/out/image.spvasm index 08d1f5a8ed..d3d3b7fde5 100644 --- a/tests/out/image.spvasm +++ b/tests/out/image.spvasm @@ -1,180 +1,217 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 127 +; Bound: 150 OpCapability Image1D OpCapability Shader OpCapability ImageQuery %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %43 "main" %40 -OpEntryPoint Vertex %61 "queries" %59 -OpExecutionMode %43 LocalSize 16 1 1 +OpEntryPoint GLCompute %50 "main" %47 +OpEntryPoint Vertex %68 "queries" %66 +OpEntryPoint Fragment %137 "sample_comparison" %135 +OpExecutionMode %50 LocalSize 16 1 1 +OpExecutionMode %137 OriginUpperLeft OpSource GLSL 450 -OpName %21 "image_src" -OpName %23 "image_dst" -OpName %25 "image_1d" -OpName %27 "image_2d" -OpName %29 "image_2d_array" -OpName %31 "image_cube" -OpName %33 "image_cube_array" -OpName %35 "image_3d" -OpName %37 "image_aa" -OpName %40 "local_id" -OpName %43 "main" -OpName %61 "queries" -OpDecorate %21 NonWritable -OpDecorate %21 DescriptorSet 0 -OpDecorate %21 Binding 1 -OpDecorate %23 NonReadable -OpDecorate %23 DescriptorSet 0 -OpDecorate %23 Binding 2 -OpDecorate %25 DescriptorSet 0 -OpDecorate %25 Binding 0 -OpDecorate %27 DescriptorSet 0 -OpDecorate %27 Binding 1 -OpDecorate %29 DescriptorSet 0 -OpDecorate %29 Binding 2 -OpDecorate %31 DescriptorSet 0 -OpDecorate %31 Binding 3 -OpDecorate %33 DescriptorSet 0 -OpDecorate %33 Binding 4 -OpDecorate %35 DescriptorSet 0 -OpDecorate %35 Binding 5 -OpDecorate %37 DescriptorSet 0 -OpDecorate %37 Binding 6 -OpDecorate %40 BuiltIn LocalInvocationId -OpDecorate %59 BuiltIn Position +OpName %24 "image_src" +OpName %26 "image_dst" +OpName %28 "image_1d" +OpName %30 "image_2d" +OpName %32 "image_2d_array" +OpName %34 "image_cube" +OpName %36 "image_cube_array" +OpName %38 "image_3d" +OpName %40 "image_aa" +OpName %42 "sampler_cmp" +OpName %44 "image_2d_depth" +OpName %47 "local_id" +OpName %50 "main" +OpName %68 "queries" +OpName %137 "sample_comparison" +OpDecorate %24 NonWritable +OpDecorate %24 DescriptorSet 0 +OpDecorate %24 Binding 1 +OpDecorate %26 NonReadable +OpDecorate %26 DescriptorSet 0 +OpDecorate %26 Binding 2 +OpDecorate %28 DescriptorSet 0 +OpDecorate %28 Binding 0 +OpDecorate %30 DescriptorSet 0 +OpDecorate %30 Binding 1 +OpDecorate %32 DescriptorSet 0 +OpDecorate %32 Binding 2 +OpDecorate %34 DescriptorSet 0 +OpDecorate %34 Binding 3 +OpDecorate %36 DescriptorSet 0 +OpDecorate %36 Binding 4 +OpDecorate %38 DescriptorSet 0 +OpDecorate %38 Binding 5 +OpDecorate %40 DescriptorSet 0 +OpDecorate %40 Binding 6 +OpDecorate %42 DescriptorSet 1 +OpDecorate %42 Binding 0 +OpDecorate %44 DescriptorSet 1 +OpDecorate %44 Binding 1 +OpDecorate %47 BuiltIn LocalInvocationId +OpDecorate %66 BuiltIn Position +OpDecorate %135 Location 0 %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 10 %5 = OpConstant %4 20 %6 = OpConstant %4 1 -%8 = OpTypeInt 32 0 -%7 = OpTypeImage %8 2D 0 0 0 2 Rgba8ui -%9 = OpTypeImage %8 1D 0 0 0 2 R32ui -%10 = OpTypeVector %8 3 -%11 = OpTypeVector %4 2 -%13 = OpTypeFloat 32 -%12 = OpTypeImage %13 1D 0 0 0 1 Unknown -%14 = OpTypeImage %13 2D 0 0 0 1 Unknown -%15 = OpTypeImage %13 2D 0 1 0 1 Unknown -%16 = OpTypeImage %13 Cube 0 0 0 1 Unknown -%17 = OpTypeImage %13 Cube 0 1 0 1 Unknown -%18 = OpTypeImage %13 3D 0 0 0 1 Unknown -%19 = OpTypeImage %13 2D 0 0 1 1 Unknown -%20 = OpTypeVector %13 4 -%22 = OpTypePointer UniformConstant %7 -%21 = OpVariable %22 UniformConstant -%24 = OpTypePointer UniformConstant %9 -%23 = OpVariable %24 UniformConstant -%26 = OpTypePointer UniformConstant %12 -%25 = OpVariable %26 UniformConstant -%28 = OpTypePointer UniformConstant %14 -%27 = OpVariable %28 UniformConstant -%30 = OpTypePointer UniformConstant %15 -%29 = OpVariable %30 UniformConstant -%32 = OpTypePointer UniformConstant %16 -%31 = OpVariable %32 UniformConstant -%34 = OpTypePointer UniformConstant %17 -%33 = OpVariable %34 UniformConstant -%36 = OpTypePointer UniformConstant %18 -%35 = OpVariable %36 UniformConstant -%38 = OpTypePointer UniformConstant %19 -%37 = OpVariable %38 UniformConstant -%41 = OpTypePointer Input %10 -%40 = OpVariable %41 Input -%44 = OpTypeFunction %2 -%49 = OpTypeVector %8 2 -%55 = OpTypeVector %8 4 -%60 = OpTypePointer Output %20 -%59 = OpVariable %60 Output -%70 = OpConstant %8 0 -%75 = OpTypeVector %4 3 -%43 = OpFunction %2 None %44 -%39 = OpLabel -%42 = OpLoad %10 %40 -%45 = OpLoad %7 %21 -%46 = OpLoad %9 %23 -OpBranch %47 -%47 = OpLabel -%48 = OpImageQuerySize %11 %45 -%50 = OpVectorShuffle %49 %42 %42 0 1 -%51 = OpBitcast %11 %50 -%52 = OpIMul %11 %48 %51 -%53 = OpCompositeConstruct %11 %3 %5 -%54 = OpSMod %11 %52 %53 -%56 = OpImageRead %55 %45 %54 -%57 = OpCompositeExtract %4 %54 0 -OpImageWrite %46 %57 %56 +%8 = OpTypeFloat 32 +%7 = OpConstant %8 0.5 +%10 = OpTypeInt 32 0 +%9 = OpTypeImage %10 2D 0 0 0 2 Rgba8ui +%11 = OpTypeImage %10 1D 0 0 0 2 R32ui +%12 = OpTypeVector %10 3 +%13 = OpTypeVector %4 2 +%14 = OpTypeImage %8 1D 0 0 0 1 Unknown +%15 = OpTypeImage %8 2D 0 0 0 1 Unknown +%16 = OpTypeImage %8 2D 0 1 0 1 Unknown +%17 = OpTypeImage %8 Cube 0 0 0 1 Unknown +%18 = OpTypeImage %8 Cube 0 1 0 1 Unknown +%19 = OpTypeImage %8 3D 0 0 0 1 Unknown +%20 = OpTypeImage %8 2D 0 0 1 1 Unknown +%21 = OpTypeVector %8 4 +%22 = OpTypeSampler +%23 = OpTypeImage %8 2D 1 0 0 1 Unknown +%25 = OpTypePointer UniformConstant %9 +%24 = OpVariable %25 UniformConstant +%27 = OpTypePointer UniformConstant %11 +%26 = OpVariable %27 UniformConstant +%29 = OpTypePointer UniformConstant %14 +%28 = OpVariable %29 UniformConstant +%31 = OpTypePointer UniformConstant %15 +%30 = OpVariable %31 UniformConstant +%33 = OpTypePointer UniformConstant %16 +%32 = OpVariable %33 UniformConstant +%35 = OpTypePointer UniformConstant %17 +%34 = OpVariable %35 UniformConstant +%37 = OpTypePointer UniformConstant %18 +%36 = OpVariable %37 UniformConstant +%39 = OpTypePointer UniformConstant %19 +%38 = OpVariable %39 UniformConstant +%41 = OpTypePointer UniformConstant %20 +%40 = OpVariable %41 UniformConstant +%43 = OpTypePointer UniformConstant %22 +%42 = OpVariable %43 UniformConstant +%45 = OpTypePointer UniformConstant %23 +%44 = OpVariable %45 UniformConstant +%48 = OpTypePointer Input %12 +%47 = OpVariable %48 Input +%51 = OpTypeFunction %2 +%56 = OpTypeVector %10 2 +%62 = OpTypeVector %10 4 +%67 = OpTypePointer Output %21 +%66 = OpVariable %67 Output +%77 = OpConstant %10 0 +%82 = OpTypeVector %4 3 +%136 = OpTypePointer Output %8 +%135 = OpVariable %136 Output +%141 = OpTypeVector %8 2 +%143 = OpTypeSampledImage %23 +%148 = OpConstant %8 0.0 +%50 = OpFunction %2 None %51 +%46 = OpLabel +%49 = OpLoad %12 %47 +%52 = OpLoad %9 %24 +%53 = OpLoad %11 %26 +OpBranch %54 +%54 = OpLabel +%55 = OpImageQuerySize %13 %52 +%57 = OpVectorShuffle %56 %49 %49 0 1 +%58 = OpBitcast %13 %57 +%59 = OpIMul %13 %55 %58 +%60 = OpCompositeConstruct %13 %3 %5 +%61 = OpSMod %13 %59 %60 +%63 = OpImageRead %62 %52 %61 +%64 = OpCompositeExtract %4 %61 0 +OpImageWrite %53 %64 %63 OpReturn OpFunctionEnd -%61 = OpFunction %2 None %44 -%58 = OpLabel -%62 = OpLoad %12 %25 -%63 = OpLoad %14 %27 -%64 = OpLoad %15 %29 -%65 = OpLoad %16 %31 -%66 = OpLoad %17 %33 -%67 = OpLoad %18 %35 -%68 = OpLoad %19 %37 -OpBranch %69 -%69 = OpLabel -%71 = OpImageQuerySizeLod %4 %62 %70 -%72 = OpImageQuerySizeLod %11 %63 %70 -%73 = OpImageQueryLevels %4 %63 -%74 = OpImageQuerySizeLod %11 %63 %6 -%76 = OpImageQuerySizeLod %75 %64 %70 -%77 = OpVectorShuffle %11 %76 %76 0 1 -%78 = OpImageQueryLevels %4 %64 -%79 = OpImageQuerySizeLod %75 %64 %6 -%80 = OpVectorShuffle %11 %79 %79 0 1 -%81 = OpImageQuerySizeLod %75 %64 %70 -%82 = OpCompositeExtract %4 %81 2 -%83 = OpImageQuerySizeLod %11 %65 %70 -%84 = OpVectorShuffle %75 %83 %83 0 0 0 -%85 = OpImageQueryLevels %4 %65 -%86 = OpImageQuerySizeLod %11 %65 %6 -%87 = OpVectorShuffle %75 %86 %86 0 0 0 -%88 = OpImageQuerySizeLod %75 %66 %70 -%89 = OpImageQueryLevels %4 %66 -%90 = OpImageQuerySizeLod %75 %66 %6 -%91 = OpImageQuerySizeLod %75 %66 %70 -%92 = OpCompositeExtract %4 %91 2 -%93 = OpImageQuerySizeLod %75 %67 %70 -%94 = OpImageQueryLevels %4 %67 -%95 = OpImageQuerySizeLod %75 %67 %6 -%96 = OpImageQuerySamples %4 %68 -%97 = OpCompositeExtract %4 %72 1 -%98 = OpIAdd %4 %71 %97 -%99 = OpCompositeExtract %4 %74 1 -%100 = OpIAdd %4 %98 %99 -%101 = OpCompositeExtract %4 %77 1 -%102 = OpIAdd %4 %100 %101 -%103 = OpCompositeExtract %4 %80 1 -%104 = OpIAdd %4 %102 %103 -%105 = OpIAdd %4 %104 %82 -%106 = OpCompositeExtract %4 %84 1 +%68 = OpFunction %2 None %51 +%65 = OpLabel +%69 = OpLoad %14 %28 +%70 = OpLoad %15 %30 +%71 = OpLoad %16 %32 +%72 = OpLoad %17 %34 +%73 = OpLoad %18 %36 +%74 = OpLoad %19 %38 +%75 = OpLoad %20 %40 +OpBranch %76 +%76 = OpLabel +%78 = OpImageQuerySizeLod %4 %69 %77 +%79 = OpImageQuerySizeLod %13 %70 %77 +%80 = OpImageQueryLevels %4 %70 +%81 = OpImageQuerySizeLod %13 %70 %6 +%83 = OpImageQuerySizeLod %82 %71 %77 +%84 = OpVectorShuffle %13 %83 %83 0 1 +%85 = OpImageQueryLevels %4 %71 +%86 = OpImageQuerySizeLod %82 %71 %6 +%87 = OpVectorShuffle %13 %86 %86 0 1 +%88 = OpImageQuerySizeLod %82 %71 %77 +%89 = OpCompositeExtract %4 %88 2 +%90 = OpImageQuerySizeLod %13 %72 %77 +%91 = OpVectorShuffle %82 %90 %90 0 0 0 +%92 = OpImageQueryLevels %4 %72 +%93 = OpImageQuerySizeLod %13 %72 %6 +%94 = OpVectorShuffle %82 %93 %93 0 0 0 +%95 = OpImageQuerySizeLod %82 %73 %77 +%96 = OpImageQueryLevels %4 %73 +%97 = OpImageQuerySizeLod %82 %73 %6 +%98 = OpImageQuerySizeLod %82 %73 %77 +%99 = OpCompositeExtract %4 %98 2 +%100 = OpImageQuerySizeLod %82 %74 %77 +%101 = OpImageQueryLevels %4 %74 +%102 = OpImageQuerySizeLod %82 %74 %6 +%103 = OpImageQuerySamples %4 %75 +%104 = OpCompositeExtract %4 %79 1 +%105 = OpIAdd %4 %78 %104 +%106 = OpCompositeExtract %4 %81 1 %107 = OpIAdd %4 %105 %106 -%108 = OpCompositeExtract %4 %87 1 +%108 = OpCompositeExtract %4 %84 1 %109 = OpIAdd %4 %107 %108 -%110 = OpCompositeExtract %4 %88 1 +%110 = OpCompositeExtract %4 %87 1 %111 = OpIAdd %4 %109 %110 -%112 = OpCompositeExtract %4 %90 1 -%113 = OpIAdd %4 %111 %112 -%114 = OpIAdd %4 %113 %92 -%115 = OpCompositeExtract %4 %93 2 +%112 = OpIAdd %4 %111 %89 +%113 = OpCompositeExtract %4 %91 1 +%114 = OpIAdd %4 %112 %113 +%115 = OpCompositeExtract %4 %94 1 %116 = OpIAdd %4 %114 %115 -%117 = OpCompositeExtract %4 %95 2 +%117 = OpCompositeExtract %4 %95 1 %118 = OpIAdd %4 %116 %117 -%119 = OpIAdd %4 %118 %96 -%120 = OpIAdd %4 %119 %73 -%121 = OpIAdd %4 %120 %78 -%122 = OpIAdd %4 %121 %94 -%123 = OpIAdd %4 %122 %85 -%124 = OpIAdd %4 %123 %89 -%125 = OpConvertSToF %13 %124 -%126 = OpCompositeConstruct %20 %125 %125 %125 %125 -OpStore %59 %126 +%119 = OpCompositeExtract %4 %97 1 +%120 = OpIAdd %4 %118 %119 +%121 = OpIAdd %4 %120 %99 +%122 = OpCompositeExtract %4 %100 2 +%123 = OpIAdd %4 %121 %122 +%124 = OpCompositeExtract %4 %102 2 +%125 = OpIAdd %4 %123 %124 +%126 = OpIAdd %4 %125 %103 +%127 = OpIAdd %4 %126 %80 +%128 = OpIAdd %4 %127 %85 +%129 = OpIAdd %4 %128 %101 +%130 = OpIAdd %4 %129 %92 +%131 = OpIAdd %4 %130 %96 +%132 = OpConvertSToF %8 %131 +%133 = OpCompositeConstruct %21 %132 %132 %132 %132 +OpStore %66 %133 +OpReturn +OpFunctionEnd +%137 = OpFunction %2 None %51 +%134 = OpLabel +%138 = OpLoad %22 %42 +%139 = OpLoad %23 %44 +OpBranch %140 +%140 = OpLabel +%142 = OpCompositeConstruct %141 %7 %7 +%144 = OpSampledImage %143 %139 %138 +%145 = OpImageSampleDrefImplicitLod %8 %144 %142 %7 +%146 = OpSampledImage %143 %139 %138 +%147 = OpImageSampleDrefExplicitLod %8 %146 %142 %7 Lod %148 +%149 = OpFAdd %8 %145 %147 +OpStore %135 %149 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/image.wgsl b/tests/out/image.wgsl index f92a356c4d..4d1ead8261 100644 --- a/tests/out/image.wgsl +++ b/tests/out/image.wgsl @@ -16,6 +16,10 @@ var image_cube_array: texture_cube_array; var image_3d: texture_3d; [[group(0), binding(6)]] var image_aa: texture_multisampled_2d; +[[group(1), binding(0)]] +var sampler_cmp: sampler_comparison; +[[group(1), binding(1)]] +var image_2d_depth: texture_depth_2d; [[stage(compute), workgroup_size(16, 1, 1)]] fn main([[builtin(local_invocation_id)]] local_id: vec3) { @@ -49,3 +53,11 @@ fn queries() -> [[builtin(position)]] vec4 { let _e32: i32 = textureNumSamples(image_aa); return vec4(f32(((((((((((((((((((_e9 + _e10.y) + _e13.y) + _e14.y) + _e17.y) + _e18) + _e19.y) + _e22.y) + _e23.y) + _e26.y) + _e27) + _e28.z) + _e31.z) + _e32) + _e11) + _e15) + _e29) + _e20) + _e24))); } + +[[stage(fragment)]] +fn sample_comparison() -> [[location(0)]] f32 { + let _e12: vec2 = vec2(0.5); + let _e14: f32 = textureSampleCompare(image_2d_depth, sampler_cmp, _e12, 0.5); + let _e15: f32 = textureSampleCompareLevel(image_2d_depth, sampler_cmp, _e12, 0.5); + return (_e14 + _e15); +}