Using VK_KHR_multiview with input attachment image of arrayLayers=1 causes type mismatch between shader and resource. #2617
Description
I'm trying to use this WBOIT composition pipeline with the render pass, which uses VK_KHR_multiview extension with viewMask=0b1.
#version 450
#extension GL_AMD_shader_trinary_minmax : enable
const float EPSILON = 1e-5f;
layout (location = 0) out vec4 outColor;
layout (input_attachment_index = 0, set = 0, binding = 0) uniform subpassInput inputAccumulation;
layout (input_attachment_index = 1, set = 0, binding = 1) uniform subpassInput inputRevealage;
bool isApproximatelyEqual(float a, float b) {
return abs(a - b) <= (abs(a) < abs(b) ? abs(b) : abs(a)) * EPSILON;
}
float trinaryMax(vec3 v) {
return max3(v.x, v.y, v.z);
}
void main(){
float revealage = subpassLoad(inputRevealage).r;
if (isApproximatelyEqual(revealage, 1.0f)){
discard;
}
vec4 accumulation = subpassLoad(inputAccumulation);
if (isinf(trinaryMax(abs(accumulation.rgb)))) {
accumulation.rgb = accumulation.aaa;
}
vec3 averageColor = accumulation.rgb / max(accumulation.a, EPSILON);
outColor = vec4(averageColor, 1.0 - revealage);
}which is compiled to MSL by:
[mvk-info] Compiling Metal shader with MathMode Fast, MathFloatingPointFunctions Fast, and PreserveInvariance disabled.
[mvk-info] Converting SPIR-V:
; SPIR-V
; Version: 1.5
; Generator: Google Shaderc over Glslang; 11
; Bound: 116
; Schema: 0
OpCapability Shader
OpCapability InputAttachment
%1 = OpExtInstImport "GLSL.std.450"
%53 = OpExtInstImport "SPV_AMD_shader_trinary_minmax"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %4 "main" %60 %79 %108
OpExecutionMode %4 OriginUpperLeft
OpSource GLSL 450
OpSourceExtension "GL_AMD_shader_trinary_minmax"
OpSourceExtension "GL_GOOGLE_cpp_style_line_directive"
OpSourceExtension "GL_GOOGLE_include_directive"
OpName %4 "main"
OpName %12 "isApproximatelyEqual(f1;f1;"
OpName %10 "a"
OpName %11 "b"
OpName %18 "trinaryMax(vf3;"
OpName %17 "v"
OpName %57 "revealage"
OpName %60 "inputRevealage"
OpName %70 "param"
OpName %72 "param"
OpName %78 "accumulation"
OpName %79 "inputAccumulation"
OpName %85 "param"
OpName %98 "averageColor"
OpName %108 "outColor"
OpDecorate %60 Binding 1
OpDecorate %60 DescriptorSet 0
OpDecorate %60 InputAttachmentIndex 1
OpDecorate %79 Binding 0
OpDecorate %79 DescriptorSet 0
OpDecorate %79 InputAttachmentIndex 0
OpDecorate %108 Location 0
%2 = OpTypeVoid
%3 = OpTypeFunction %2
%6 = OpTypeFloat 32
%7 = OpTypePointer Function %6
%8 = OpTypeBool
%9 = OpTypeFunction %8 %7 %7
%14 = OpTypeVector %6 3
%15 = OpTypePointer Function %14
%16 = OpTypeFunction %6 %15
%38 = OpConstant %6 9.99999975e-06
%43 = OpTypeInt 32 0
%44 = OpConstant %43 0
%47 = OpConstant %43 1
%50 = OpConstant %43 2
%58 = OpTypeImage %6 SubpassData 0 0 0 2 Unknown
%59 = OpTypePointer UniformConstant %58
%60 = OpVariable %59 UniformConstant
%62 = OpTypeInt 32 1
%63 = OpConstant %62 0
%64 = OpTypeVector %62 2
%65 = OpConstantComposite %64 %63 %63
%66 = OpTypeVector %6 4
%69 = OpConstant %6 1
%77 = OpTypePointer Function %66
%79 = OpVariable %59 UniformConstant
%101 = OpConstant %43 3
%107 = OpTypePointer Output %66
%108 = OpVariable %107 Output
%4 = OpFunction %2 None %3
%5 = OpLabel
%57 = OpVariable %7 Function
%70 = OpVariable %7 Function
%72 = OpVariable %7 Function
%78 = OpVariable %77 Function
%85 = OpVariable %15 Function
%98 = OpVariable %15 Function
%61 = OpLoad %58 %60
%67 = OpImageRead %66 %61 %65
%68 = OpCompositeExtract %6 %67 0
OpStore %57 %68
%71 = OpLoad %6 %57
OpStore %70 %71
OpStore %72 %69
%73 = OpFunctionCall %8 %12 %70 %72
OpSelectionMerge %75 None
OpBranchConditional %73 %74 %75
%74 = OpLabel
OpKill
%75 = OpLabel
%80 = OpLoad %58 %79
%81 = OpImageRead %66 %80 %65
OpStore %78 %81
%82 = OpLoad %66 %78
%83 = OpVectorShuffle %14 %82 %82 0 1 2
%84 = OpExtInst %14 %1 FAbs %83
OpStore %85 %84
%86 = OpFunctionCall %6 %18 %85
%87 = OpIsInf %8 %86
OpSelectionMerge %89 None
OpBranchConditional %87 %88 %89
%88 = OpLabel
%90 = OpLoad %66 %78
%91 = OpVectorShuffle %14 %90 %90 3 3 3
%92 = OpAccessChain %7 %78 %44
%93 = OpCompositeExtract %6 %91 0
OpStore %92 %93
%94 = OpAccessChain %7 %78 %47
%95 = OpCompositeExtract %6 %91 1
OpStore %94 %95
%96 = OpAccessChain %7 %78 %50
%97 = OpCompositeExtract %6 %91 2
OpStore %96 %97
OpBranch %89
%89 = OpLabel
%99 = OpLoad %66 %78
%100 = OpVectorShuffle %14 %99 %99 0 1 2
%102 = OpAccessChain %7 %78 %101
%103 = OpLoad %6 %102
%104 = OpExtInst %6 %1 FMax %103 %38
%105 = OpCompositeConstruct %14 %104 %104 %104
%106 = OpFDiv %14 %100 %105
OpStore %98 %106
%109 = OpLoad %14 %98
%110 = OpLoad %6 %57
%111 = OpFSub %6 %69 %110
%112 = OpCompositeExtract %6 %109 0
%113 = OpCompositeExtract %6 %109 1
%114 = OpCompositeExtract %6 %109 2
%115 = OpCompositeConstruct %66 %112 %113 %114 %111
OpStore %108 %115
OpReturn
OpFunctionEnd
%12 = OpFunction %8 None %9
%10 = OpFunctionParameter %7
%11 = OpFunctionParameter %7
%13 = OpLabel
%29 = OpVariable %7 Function
%20 = OpLoad %6 %10
%21 = OpLoad %6 %11
%22 = OpFSub %6 %20 %21
%23 = OpExtInst %6 %1 FAbs %22
%24 = OpLoad %6 %10
%25 = OpExtInst %6 %1 FAbs %24
%26 = OpLoad %6 %11
%27 = OpExtInst %6 %1 FAbs %26
%28 = OpFOrdLessThan %8 %25 %27
OpSelectionMerge %31 None
OpBranchConditional %28 %30 %34
%30 = OpLabel
%32 = OpLoad %6 %11
%33 = OpExtInst %6 %1 FAbs %32
OpStore %29 %33
OpBranch %31
%34 = OpLabel
%35 = OpLoad %6 %10
%36 = OpExtInst %6 %1 FAbs %35
OpStore %29 %36
OpBranch %31
%31 = OpLabel
%37 = OpLoad %6 %29
%39 = OpFMul %6 %37 %38
%40 = OpFOrdLessThanEqual %8 %23 %39
OpReturnValue %40
OpFunctionEnd
%18 = OpFunction %6 None %16
%17 = OpFunctionParameter %15
%19 = OpLabel
%45 = OpAccessChain %7 %17 %44
%46 = OpLoad %6 %45
%48 = OpAccessChain %7 %17 %47
%49 = OpLoad %6 %48
%51 = OpAccessChain %7 %17 %50
%52 = OpLoad %6 %51
%54 = OpExtInst %6 %53 FMax3AMD %46 %49 %52
OpReturnValue %54
OpFunctionEnd
End SPIR-V
Converted MSL:
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct spvDescriptorSetBuffer0
{
texture2d_array<float> inputAccumulation [[id(0)]];
texture2d_array<float> inputRevealage [[id(1)]];
};
struct main0_out
{
float4 outColor [[color(0)]];
};
static inline __attribute__((always_inline))
bool isApproximatelyEqual(thread const float& a, thread const float& b)
{
float _29;
if (abs(a) < abs(b))
{
_29 = abs(b);
}
else
{
_29 = abs(a);
}
return abs(a - b) <= (_29 * 9.9999997473787516355514526367188e-06);
}
static inline __attribute__((always_inline))
float trinaryMax(thread const float3& v)
{
return max3(v.x, v.y, v.z);
}
fragment main0_out main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant uint* spvViewMask [[buffer(27)]], float4 gl_FragCoord [[position]], uint gl_ViewIndex [[render_target_array_index]])
{
main0_out out = {};
gl_ViewIndex += spvViewMask[0];
float revealage = spvDescriptorSet0.inputRevealage.read(uint2(gl_FragCoord.xy), gl_ViewIndex).x;
float param = revealage;
float param_1 = 1.0;
if (isApproximatelyEqual(param, param_1))
{
discard_fragment();
}
float4 accumulation = spvDescriptorSet0.inputAccumulation.read(uint2(gl_FragCoord.xy), gl_ViewIndex);
float3 param_2 = abs(accumulation.xyz);
if (isinf(trinaryMax(param_2)))
{
float4 _90 = accumulation;
accumulation.x = _90.www.x;
accumulation.y = _90.www.y;
accumulation.z = _90.www.z;
}
float3 averageColor = accumulation.xyz / float3(fast::max(accumulation.w, 9.9999997473787516355514526367188e-06));
out.outColor = float4(averageColor, 1.0 - revealage);
return out;
}
End MSL
As the render pass uses the multiview, Vulkan input attachments (inputAccumulation and inputRevealage) are converted to metal::texture2d_array, and read by gl_ViewIndex, as expected.
However, when I'm using those input attachment images with arrayLayers=1 and usage specified as only attachment usage, the issue arise.
MoltenVK/MoltenVK/MoltenVK/GPUObjects/MVKImage.mm
Lines 2329 to 2345 in 48d5cfb
MoltenVK forcibly disable the attempt to create VkImageView with type=VK_IMAGE_VIEW_TYPE_2D_ARRAY from single-layer image of attachment-usage only, instead it creates MTLTexture whose slice is 1. This causes type mismatch between the compiled MSL shader (metal::texture2d_array) and resource (metal::texture2d).
If I add VK_IMAGE_USAGE_STORAGE_BIT usage at the image creation, this issue can be avoided.