Handle edge cases in OpCopyMemory.
Implement this by synthesizing an OpLoad/OpStore pair instead.
This commit is contained in:
Hans-Kristian Arntzen
parent
d57ab68a21
commit
4ca06c7278
@ -0,0 +1,125 @@
|
||||
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||
#pragma clang diagnostic ignored "-Wmissing-braces"
|
||||
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, size_t Num>
|
||||
struct spvUnsafeArray
|
||||
{
|
||||
T elements[Num ? Num : 1];
|
||||
|
||||
thread T& operator [] (size_t pos) thread
|
||||
{
|
||||
return elements[pos];
|
||||
}
|
||||
constexpr const thread T& operator [] (size_t pos) const thread
|
||||
{
|
||||
return elements[pos];
|
||||
}
|
||||
|
||||
device T& operator [] (size_t pos) device
|
||||
{
|
||||
return elements[pos];
|
||||
}
|
||||
constexpr const device T& operator [] (size_t pos) const device
|
||||
{
|
||||
return elements[pos];
|
||||
}
|
||||
|
||||
constexpr const constant T& operator [] (size_t pos) const constant
|
||||
{
|
||||
return elements[pos];
|
||||
}
|
||||
|
||||
threadgroup T& operator [] (size_t pos) threadgroup
|
||||
{
|
||||
return elements[pos];
|
||||
}
|
||||
constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
|
||||
{
|
||||
return elements[pos];
|
||||
}
|
||||
};
|
||||
|
||||
struct cb1_struct
|
||||
{
|
||||
float4 _m0[1];
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float3 vocp0;
|
||||
float4 vocp1;
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 v0 [[attribute(0)]];
|
||||
float4 v1 [[attribute(1)]];
|
||||
float3 vicp0 [[attribute(2)]];
|
||||
float4 vicp1 [[attribute(4)]];
|
||||
};
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
void fork0_epilogue(thread const float4& _87, thread const float4& _88, thread const float4& _89, device half (&gl_TessLevelOuter)[3])
|
||||
{
|
||||
gl_TessLevelOuter[0u] = half(_87.x);
|
||||
gl_TessLevelOuter[1u] = half(_88.x);
|
||||
gl_TessLevelOuter[2u] = half(_89.x);
|
||||
}
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
void fork0(uint vForkInstanceId, device half (&gl_TessLevelOuter)[3], thread spvUnsafeArray<float4, 4> (&opc), constant cb1_struct& cb0_0, thread float4& v_48, thread float4& v_49, thread float4& v_50)
|
||||
{
|
||||
float4 r0;
|
||||
r0.x = as_type<float>(vForkInstanceId);
|
||||
opc[as_type<int>(r0.x)].x = cb0_0._m0[0u].x;
|
||||
v_48 = opc[0u];
|
||||
v_49 = opc[1u];
|
||||
v_50 = opc[2u];
|
||||
fork0_epilogue(v_48, v_49, v_50, gl_TessLevelOuter);
|
||||
}
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
void fork1_epilogue(thread const float4& _109, device half &gl_TessLevelInner)
|
||||
{
|
||||
gl_TessLevelInner = half(_109.x);
|
||||
}
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
void fork1(device half &gl_TessLevelInner, thread spvUnsafeArray<float4, 4> (&opc), constant cb1_struct& cb0_0, thread float4& v_56)
|
||||
{
|
||||
opc[3u].x = cb0_0._m0[0u].x;
|
||||
v_56 = opc[3u];
|
||||
fork1_epilogue(v_56, gl_TessLevelInner);
|
||||
}
|
||||
|
||||
kernel void main0(main0_in in [[stage_in]], constant cb1_struct& cb0_0 [[buffer(0)]], uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device MTLTriangleTessellationFactorsHalf* spvTessLevel [[buffer(26)]], threadgroup main0_in* gl_in [[threadgroup(0)]])
|
||||
{
|
||||
device main0_out* gl_out = &spvOut[gl_PrimitiveID * 3];
|
||||
if (gl_InvocationID < spvIndirectParams[0])
|
||||
gl_in[gl_InvocationID] = in;
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
if (gl_InvocationID >= 3)
|
||||
return;
|
||||
spvUnsafeArray<spvUnsafeArray<float4, 3>, 2> vicp;
|
||||
spvUnsafeArray<float4, 3> _153 = spvUnsafeArray<float4, 3>({ gl_in[0].v0, gl_in[1].v0, gl_in[2].v0 });
|
||||
vicp[0u] = _153;
|
||||
spvUnsafeArray<float4, 3> _154 = spvUnsafeArray<float4, 3>({ gl_in[0].v1, gl_in[1].v1, gl_in[2].v1 });
|
||||
vicp[1u] = _154;
|
||||
gl_out[gl_InvocationID].vocp0 = gl_in[gl_InvocationID].vicp0;
|
||||
gl_out[gl_InvocationID].vocp1 = gl_in[gl_InvocationID].vicp1;
|
||||
spvUnsafeArray<float4, 4> opc;
|
||||
float4 v_48;
|
||||
float4 v_49;
|
||||
float4 v_50;
|
||||
fork0(0u, spvTessLevel[gl_PrimitiveID].edgeTessellationFactor, opc, cb0_0, v_48, v_49, v_50);
|
||||
fork0(1u, spvTessLevel[gl_PrimitiveID].edgeTessellationFactor, opc, cb0_0, v_48, v_49, v_50);
|
||||
fork0(2u, spvTessLevel[gl_PrimitiveID].edgeTessellationFactor, opc, cb0_0, v_48, v_49, v_50);
|
||||
float4 v_56;
|
||||
fork1(spvTessLevel[gl_PrimitiveID].insideTessellationFactor, opc, cb0_0, v_56);
|
||||
}
|
||||
|
||||
@ -0,0 +1,73 @@
|
||||
#version 450
|
||||
layout(vertices = 3) out;
|
||||
|
||||
layout(binding = 0, std140) uniform cb1_struct
|
||||
{
|
||||
vec4 _m0[1];
|
||||
} cb0_0;
|
||||
|
||||
layout(location = 0) in vec4 v0[];
|
||||
layout(location = 1) in vec4 v1[];
|
||||
layout(location = 2) in vec3 vicp0[];
|
||||
layout(location = 3) out vec3 vocp0[3];
|
||||
layout(location = 4) in vec4 vicp1[];
|
||||
layout(location = 5) out vec4 vocp1[3];
|
||||
vec4 opc[4];
|
||||
vec4 vicp[2][3];
|
||||
vec4 _48;
|
||||
vec4 _49;
|
||||
vec4 _50;
|
||||
vec4 _56;
|
||||
|
||||
void fork0_epilogue(vec4 _87, vec4 _88, vec4 _89)
|
||||
{
|
||||
gl_TessLevelOuter[0u] = _87.x;
|
||||
gl_TessLevelOuter[1u] = _88.x;
|
||||
gl_TessLevelOuter[2u] = _89.x;
|
||||
}
|
||||
|
||||
void fork0(uint vForkInstanceId)
|
||||
{
|
||||
vec4 r0;
|
||||
r0.x = uintBitsToFloat(vForkInstanceId);
|
||||
opc[floatBitsToInt(r0.x)].x = cb0_0._m0[0u].x;
|
||||
_48 = opc[0u];
|
||||
_49 = opc[1u];
|
||||
_50 = opc[2u];
|
||||
fork0_epilogue(_48, _49, _50);
|
||||
}
|
||||
|
||||
void fork1_epilogue(vec4 _109)
|
||||
{
|
||||
gl_TessLevelInner[0u] = _109.x;
|
||||
}
|
||||
|
||||
void fork1()
|
||||
{
|
||||
opc[3u].x = cb0_0._m0[0u].x;
|
||||
_56 = opc[3u];
|
||||
fork1_epilogue(_56);
|
||||
}
|
||||
|
||||
void main()
|
||||
{
|
||||
vec4 _126_unrolled[3];
|
||||
for (int i = 0; i < int(3); i++)
|
||||
{
|
||||
_126_unrolled[i] = v0[i];
|
||||
}
|
||||
vicp[0u] = _126_unrolled;
|
||||
vec4 _127_unrolled[3];
|
||||
for (int i = 0; i < int(3); i++)
|
||||
{
|
||||
_127_unrolled[i] = v1[i];
|
||||
}
|
||||
vicp[1u] = _127_unrolled;
|
||||
vocp0[gl_InvocationID] = vicp0[gl_InvocationID];
|
||||
vocp1[gl_InvocationID] = vicp1[gl_InvocationID];
|
||||
fork0(0u);
|
||||
fork0(1u);
|
||||
fork0(2u);
|
||||
fork1();
|
||||
}
|
||||
|
||||
199
shaders-msl-no-opt/asm/tesc/copy-memory-control-point.asm.tesc
Normal file
199
shaders-msl-no-opt/asm/tesc/copy-memory-control-point.asm.tesc
Normal file
@ -0,0 +1,199 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Wine VKD3D Shader Compiler; 2
|
||||
; Bound: 126
|
||||
; Schema: 0
|
||||
OpCapability Tessellation
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint TessellationControl %1 "main" %4 %30 %80 %101 %103 %108 %110 %115 %117
|
||||
OpExecutionMode %1 OutputVertices 3
|
||||
OpExecutionMode %1 Triangles
|
||||
OpExecutionMode %1 SpacingEqual
|
||||
OpExecutionMode %1 VertexOrderCw
|
||||
OpName %1 "main"
|
||||
OpName %11 "opc"
|
||||
OpName %14 "cb1_struct"
|
||||
OpName %16 "cb0_0"
|
||||
OpName %22 "vicp"
|
||||
OpName %23 "fork0"
|
||||
OpName %26 "vForkInstanceId"
|
||||
OpName %34 "r0"
|
||||
OpName %32 "fork0_epilogue"
|
||||
OpName %75 "fork1"
|
||||
OpName %81 "fork1_epilogue"
|
||||
OpName %101 "v0"
|
||||
OpName %103 "v1"
|
||||
OpName %108 "vicp0"
|
||||
OpName %110 "vocp0"
|
||||
OpName %115 "vicp1"
|
||||
OpName %117 "vocp1"
|
||||
OpDecorate %4 BuiltIn InvocationId
|
||||
OpDecorate %13 ArrayStride 16
|
||||
OpDecorate %14 Block
|
||||
OpMemberDecorate %14 0 Offset 0
|
||||
OpDecorate %16 DescriptorSet 0
|
||||
OpDecorate %16 Binding 0
|
||||
OpDecorate %30 BuiltIn TessLevelOuter
|
||||
OpDecorate %30 Patch
|
||||
OpDecorate %30 Patch
|
||||
OpDecorate %30 Patch
|
||||
OpDecorate %30 Patch
|
||||
OpDecorate %80 BuiltIn TessLevelInner
|
||||
OpDecorate %80 Patch
|
||||
OpDecorate %80 Patch
|
||||
OpDecorate %101 Location 0
|
||||
OpDecorate %103 Location 1
|
||||
OpDecorate %108 Location 2
|
||||
OpDecorate %110 Location 3
|
||||
OpDecorate %115 Location 4
|
||||
OpDecorate %117 Location 5
|
||||
%2 = OpTypeInt 32 1
|
||||
%3 = OpTypePointer Input %2
|
||||
%4 = OpVariable %3 Input
|
||||
%5 = OpTypeFloat 32
|
||||
%6 = OpTypeVector %5 4
|
||||
%7 = OpTypeInt 32 0
|
||||
%8 = OpConstant %7 4
|
||||
%9 = OpTypeArray %6 %8
|
||||
%10 = OpTypePointer Private %9
|
||||
%11 = OpVariable %10 Private
|
||||
%12 = OpConstant %7 1
|
||||
%13 = OpTypeArray %6 %12
|
||||
%14 = OpTypeStruct %13
|
||||
%15 = OpTypePointer Uniform %14
|
||||
%16 = OpVariable %15 Uniform
|
||||
%17 = OpConstant %7 3
|
||||
%18 = OpTypeArray %6 %17
|
||||
%19 = OpConstant %7 2
|
||||
%20 = OpTypeArray %18 %19
|
||||
%21 = OpTypePointer Private %20
|
||||
%22 = OpVariable %21 Private
|
||||
%24 = OpTypeVoid
|
||||
%25 = OpTypeFunction %24 %7
|
||||
%28 = OpTypeArray %5 %8
|
||||
%29 = OpTypePointer Output %28
|
||||
%30 = OpVariable %29 Output
|
||||
%31 = OpConstant %7 0
|
||||
%33 = OpTypePointer Function %6
|
||||
%36 = OpTypePointer Function %5
|
||||
%38 = OpTypePointer Uniform %6
|
||||
%40 = OpTypePointer Uniform %5
|
||||
%46 = OpTypePointer Private %6
|
||||
%48 = OpTypePointer Private %5
|
||||
%52 = OpVariable %46 Private
|
||||
%55 = OpVariable %46 Private
|
||||
%58 = OpVariable %46 Private
|
||||
%60 = OpTypeFunction %24 %46 %46 %46
|
||||
%69 = OpTypePointer Output %5
|
||||
%76 = OpTypeFunction %24
|
||||
%78 = OpTypeArray %5 %19
|
||||
%79 = OpTypePointer Output %78
|
||||
%80 = OpVariable %79 Output
|
||||
%89 = OpVariable %46 Private
|
||||
%91 = OpTypeFunction %24 %46
|
||||
%98 = OpTypePointer Private %18
|
||||
%100 = OpTypePointer Input %18
|
||||
%101 = OpVariable %100 Input
|
||||
%103 = OpVariable %100 Input
|
||||
%105 = OpTypeVector %5 3
|
||||
%106 = OpTypeArray %105 %17
|
||||
%107 = OpTypePointer Input %106
|
||||
%108 = OpVariable %107 Input
|
||||
%109 = OpTypePointer Output %106
|
||||
%110 = OpVariable %109 Output
|
||||
%111 = OpTypePointer Output %105
|
||||
%112 = OpTypePointer Input %105
|
||||
%115 = OpVariable %100 Input
|
||||
%116 = OpTypePointer Output %18
|
||||
%117 = OpVariable %116 Output
|
||||
%118 = OpTypePointer Output %6
|
||||
%119 = OpTypePointer Input %6
|
||||
%23 = OpFunction %24 None %25
|
||||
%26 = OpFunctionParameter %7
|
||||
%27 = OpLabel
|
||||
%34 = OpVariable %33 Function
|
||||
%35 = OpBitcast %5 %26
|
||||
%37 = OpInBoundsAccessChain %36 %34 %31
|
||||
OpStore %37 %35
|
||||
%39 = OpAccessChain %38 %16 %31 %31
|
||||
%41 = OpInBoundsAccessChain %40 %39 %31
|
||||
%42 = OpLoad %5 %41
|
||||
%43 = OpInBoundsAccessChain %36 %34 %31
|
||||
%44 = OpLoad %5 %43
|
||||
%45 = OpBitcast %2 %44
|
||||
%47 = OpAccessChain %46 %11 %45
|
||||
%49 = OpInBoundsAccessChain %48 %47 %31
|
||||
OpStore %49 %42
|
||||
%50 = OpAccessChain %46 %11 %31
|
||||
%51 = OpLoad %6 %50
|
||||
OpStore %52 %51
|
||||
%53 = OpAccessChain %46 %11 %12
|
||||
%54 = OpLoad %6 %53
|
||||
OpStore %55 %54
|
||||
%56 = OpAccessChain %46 %11 %19
|
||||
%57 = OpLoad %6 %56
|
||||
OpStore %58 %57
|
||||
%59 = OpFunctionCall %24 %32 %52 %55 %58
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%32 = OpFunction %24 None %60
|
||||
%61 = OpFunctionParameter %46
|
||||
%62 = OpFunctionParameter %46
|
||||
%63 = OpFunctionParameter %46
|
||||
%64 = OpLabel
|
||||
%65 = OpLoad %6 %61
|
||||
%66 = OpLoad %6 %62
|
||||
%67 = OpLoad %6 %63
|
||||
%68 = OpCompositeExtract %5 %65 0
|
||||
%70 = OpAccessChain %69 %30 %31
|
||||
OpStore %70 %68
|
||||
%71 = OpCompositeExtract %5 %66 0
|
||||
%72 = OpAccessChain %69 %30 %12
|
||||
OpStore %72 %71
|
||||
%73 = OpCompositeExtract %5 %67 0
|
||||
%74 = OpAccessChain %69 %30 %19
|
||||
OpStore %74 %73
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%75 = OpFunction %24 None %76
|
||||
%77 = OpLabel
|
||||
%82 = OpAccessChain %38 %16 %31 %31
|
||||
%83 = OpInBoundsAccessChain %40 %82 %31
|
||||
%84 = OpLoad %5 %83
|
||||
%85 = OpAccessChain %46 %11 %17
|
||||
%86 = OpInBoundsAccessChain %48 %85 %31
|
||||
OpStore %86 %84
|
||||
%87 = OpAccessChain %46 %11 %17
|
||||
%88 = OpLoad %6 %87
|
||||
OpStore %89 %88
|
||||
%90 = OpFunctionCall %24 %81 %89
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%81 = OpFunction %24 None %91
|
||||
%92 = OpFunctionParameter %46
|
||||
%93 = OpLabel
|
||||
%94 = OpLoad %6 %92
|
||||
%95 = OpCompositeExtract %5 %94 0
|
||||
%96 = OpAccessChain %69 %80 %31
|
||||
OpStore %96 %95
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%1 = OpFunction %24 None %76
|
||||
%97 = OpLabel
|
||||
%99 = OpInBoundsAccessChain %98 %22 %31
|
||||
OpCopyMemory %99 %101
|
||||
%102 = OpInBoundsAccessChain %98 %22 %12
|
||||
OpCopyMemory %102 %103
|
||||
%104 = OpLoad %2 %4
|
||||
%113 = OpAccessChain %111 %110 %104
|
||||
%114 = OpAccessChain %112 %108 %104
|
||||
OpCopyMemory %113 %114
|
||||
%120 = OpAccessChain %118 %117 %104
|
||||
%121 = OpAccessChain %119 %115 %104
|
||||
OpCopyMemory %120 %121
|
||||
%122 = OpFunctionCall %24 %23 %31
|
||||
%123 = OpFunctionCall %24 %23 %12
|
||||
%124 = OpFunctionCall %24 %23 %19
|
||||
%125 = OpFunctionCall %24 %75
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
199
shaders-no-opt/asm/tesc/copy-memory-control-point.asm.tesc
Normal file
199
shaders-no-opt/asm/tesc/copy-memory-control-point.asm.tesc
Normal file
@ -0,0 +1,199 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Wine VKD3D Shader Compiler; 2
|
||||
; Bound: 126
|
||||
; Schema: 0
|
||||
OpCapability Tessellation
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint TessellationControl %1 "main" %4 %30 %80 %101 %103 %108 %110 %115 %117
|
||||
OpExecutionMode %1 OutputVertices 3
|
||||
OpExecutionMode %1 Triangles
|
||||
OpExecutionMode %1 SpacingEqual
|
||||
OpExecutionMode %1 VertexOrderCw
|
||||
OpName %1 "main"
|
||||
OpName %11 "opc"
|
||||
OpName %14 "cb1_struct"
|
||||
OpName %16 "cb0_0"
|
||||
OpName %22 "vicp"
|
||||
OpName %23 "fork0"
|
||||
OpName %26 "vForkInstanceId"
|
||||
OpName %34 "r0"
|
||||
OpName %32 "fork0_epilogue"
|
||||
OpName %75 "fork1"
|
||||
OpName %81 "fork1_epilogue"
|
||||
OpName %101 "v0"
|
||||
OpName %103 "v1"
|
||||
OpName %108 "vicp0"
|
||||
OpName %110 "vocp0"
|
||||
OpName %115 "vicp1"
|
||||
OpName %117 "vocp1"
|
||||
OpDecorate %4 BuiltIn InvocationId
|
||||
OpDecorate %13 ArrayStride 16
|
||||
OpDecorate %14 Block
|
||||
OpMemberDecorate %14 0 Offset 0
|
||||
OpDecorate %16 DescriptorSet 0
|
||||
OpDecorate %16 Binding 0
|
||||
OpDecorate %30 BuiltIn TessLevelOuter
|
||||
OpDecorate %30 Patch
|
||||
OpDecorate %30 Patch
|
||||
OpDecorate %30 Patch
|
||||
OpDecorate %30 Patch
|
||||
OpDecorate %80 BuiltIn TessLevelInner
|
||||
OpDecorate %80 Patch
|
||||
OpDecorate %80 Patch
|
||||
OpDecorate %101 Location 0
|
||||
OpDecorate %103 Location 1
|
||||
OpDecorate %108 Location 2
|
||||
OpDecorate %110 Location 3
|
||||
OpDecorate %115 Location 4
|
||||
OpDecorate %117 Location 5
|
||||
%2 = OpTypeInt 32 1
|
||||
%3 = OpTypePointer Input %2
|
||||
%4 = OpVariable %3 Input
|
||||
%5 = OpTypeFloat 32
|
||||
%6 = OpTypeVector %5 4
|
||||
%7 = OpTypeInt 32 0
|
||||
%8 = OpConstant %7 4
|
||||
%9 = OpTypeArray %6 %8
|
||||
%10 = OpTypePointer Private %9
|
||||
%11 = OpVariable %10 Private
|
||||
%12 = OpConstant %7 1
|
||||
%13 = OpTypeArray %6 %12
|
||||
%14 = OpTypeStruct %13
|
||||
%15 = OpTypePointer Uniform %14
|
||||
%16 = OpVariable %15 Uniform
|
||||
%17 = OpConstant %7 3
|
||||
%18 = OpTypeArray %6 %17
|
||||
%19 = OpConstant %7 2
|
||||
%20 = OpTypeArray %18 %19
|
||||
%21 = OpTypePointer Private %20
|
||||
%22 = OpVariable %21 Private
|
||||
%24 = OpTypeVoid
|
||||
%25 = OpTypeFunction %24 %7
|
||||
%28 = OpTypeArray %5 %8
|
||||
%29 = OpTypePointer Output %28
|
||||
%30 = OpVariable %29 Output
|
||||
%31 = OpConstant %7 0
|
||||
%33 = OpTypePointer Function %6
|
||||
%36 = OpTypePointer Function %5
|
||||
%38 = OpTypePointer Uniform %6
|
||||
%40 = OpTypePointer Uniform %5
|
||||
%46 = OpTypePointer Private %6
|
||||
%48 = OpTypePointer Private %5
|
||||
%52 = OpVariable %46 Private
|
||||
%55 = OpVariable %46 Private
|
||||
%58 = OpVariable %46 Private
|
||||
%60 = OpTypeFunction %24 %46 %46 %46
|
||||
%69 = OpTypePointer Output %5
|
||||
%76 = OpTypeFunction %24
|
||||
%78 = OpTypeArray %5 %19
|
||||
%79 = OpTypePointer Output %78
|
||||
%80 = OpVariable %79 Output
|
||||
%89 = OpVariable %46 Private
|
||||
%91 = OpTypeFunction %24 %46
|
||||
%98 = OpTypePointer Private %18
|
||||
%100 = OpTypePointer Input %18
|
||||
%101 = OpVariable %100 Input
|
||||
%103 = OpVariable %100 Input
|
||||
%105 = OpTypeVector %5 3
|
||||
%106 = OpTypeArray %105 %17
|
||||
%107 = OpTypePointer Input %106
|
||||
%108 = OpVariable %107 Input
|
||||
%109 = OpTypePointer Output %106
|
||||
%110 = OpVariable %109 Output
|
||||
%111 = OpTypePointer Output %105
|
||||
%112 = OpTypePointer Input %105
|
||||
%115 = OpVariable %100 Input
|
||||
%116 = OpTypePointer Output %18
|
||||
%117 = OpVariable %116 Output
|
||||
%118 = OpTypePointer Output %6
|
||||
%119 = OpTypePointer Input %6
|
||||
%23 = OpFunction %24 None %25
|
||||
%26 = OpFunctionParameter %7
|
||||
%27 = OpLabel
|
||||
%34 = OpVariable %33 Function
|
||||
%35 = OpBitcast %5 %26
|
||||
%37 = OpInBoundsAccessChain %36 %34 %31
|
||||
OpStore %37 %35
|
||||
%39 = OpAccessChain %38 %16 %31 %31
|
||||
%41 = OpInBoundsAccessChain %40 %39 %31
|
||||
%42 = OpLoad %5 %41
|
||||
%43 = OpInBoundsAccessChain %36 %34 %31
|
||||
%44 = OpLoad %5 %43
|
||||
%45 = OpBitcast %2 %44
|
||||
%47 = OpAccessChain %46 %11 %45
|
||||
%49 = OpInBoundsAccessChain %48 %47 %31
|
||||
OpStore %49 %42
|
||||
%50 = OpAccessChain %46 %11 %31
|
||||
%51 = OpLoad %6 %50
|
||||
OpStore %52 %51
|
||||
%53 = OpAccessChain %46 %11 %12
|
||||
%54 = OpLoad %6 %53
|
||||
OpStore %55 %54
|
||||
%56 = OpAccessChain %46 %11 %19
|
||||
%57 = OpLoad %6 %56
|
||||
OpStore %58 %57
|
||||
%59 = OpFunctionCall %24 %32 %52 %55 %58
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%32 = OpFunction %24 None %60
|
||||
%61 = OpFunctionParameter %46
|
||||
%62 = OpFunctionParameter %46
|
||||
%63 = OpFunctionParameter %46
|
||||
%64 = OpLabel
|
||||
%65 = OpLoad %6 %61
|
||||
%66 = OpLoad %6 %62
|
||||
%67 = OpLoad %6 %63
|
||||
%68 = OpCompositeExtract %5 %65 0
|
||||
%70 = OpAccessChain %69 %30 %31
|
||||
OpStore %70 %68
|
||||
%71 = OpCompositeExtract %5 %66 0
|
||||
%72 = OpAccessChain %69 %30 %12
|
||||
OpStore %72 %71
|
||||
%73 = OpCompositeExtract %5 %67 0
|
||||
%74 = OpAccessChain %69 %30 %19
|
||||
OpStore %74 %73
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%75 = OpFunction %24 None %76
|
||||
%77 = OpLabel
|
||||
%82 = OpAccessChain %38 %16 %31 %31
|
||||
%83 = OpInBoundsAccessChain %40 %82 %31
|
||||
%84 = OpLoad %5 %83
|
||||
%85 = OpAccessChain %46 %11 %17
|
||||
%86 = OpInBoundsAccessChain %48 %85 %31
|
||||
OpStore %86 %84
|
||||
%87 = OpAccessChain %46 %11 %17
|
||||
%88 = OpLoad %6 %87
|
||||
OpStore %89 %88
|
||||
%90 = OpFunctionCall %24 %81 %89
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%81 = OpFunction %24 None %91
|
||||
%92 = OpFunctionParameter %46
|
||||
%93 = OpLabel
|
||||
%94 = OpLoad %6 %92
|
||||
%95 = OpCompositeExtract %5 %94 0
|
||||
%96 = OpAccessChain %69 %80 %31
|
||||
OpStore %96 %95
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%1 = OpFunction %24 None %76
|
||||
%97 = OpLabel
|
||||
%99 = OpInBoundsAccessChain %98 %22 %31
|
||||
OpCopyMemory %99 %101
|
||||
%102 = OpInBoundsAccessChain %98 %22 %12
|
||||
OpCopyMemory %102 %103
|
||||
%104 = OpLoad %2 %4
|
||||
%113 = OpAccessChain %111 %110 %104
|
||||
%114 = OpAccessChain %112 %108 %104
|
||||
OpCopyMemory %113 %114
|
||||
%120 = OpAccessChain %118 %117 %104
|
||||
%121 = OpAccessChain %119 %115 %104
|
||||
OpCopyMemory %120 %121
|
||||
%122 = OpFunctionCall %24 %23 %31
|
||||
%123 = OpFunctionCall %24 %23 %12
|
||||
%124 = OpFunctionCall %24 %23 %19
|
||||
%125 = OpFunctionCall %24 %75
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -302,8 +302,20 @@ struct Instruction
|
||||
{
|
||||
uint16_t op = 0;
|
||||
uint16_t count = 0;
|
||||
// If offset is 0 (not a valid offset into the instruction stream),
|
||||
// we have an instruction stream which is embedded in the object.
|
||||
uint32_t offset = 0;
|
||||
uint32_t length = 0;
|
||||
|
||||
inline bool is_embedded() const
|
||||
{
|
||||
return offset == 0;
|
||||
}
|
||||
};
|
||||
|
||||
struct EmbeddedInstruction : Instruction
|
||||
{
|
||||
SmallVector<uint32_t> ops;
|
||||
};
|
||||
|
||||
enum Types
|
||||
|
||||
@ -513,10 +513,19 @@ protected:
|
||||
if (!instr.length)
|
||||
return nullptr;
|
||||
|
||||
if (instr.is_embedded())
|
||||
{
|
||||
auto &embedded = static_cast<const EmbeddedInstruction &>(instr);
|
||||
assert(embedded.ops.size() == instr.length);
|
||||
return embedded.ops.data();
|
||||
}
|
||||
else
|
||||
{
|
||||
if (instr.offset + instr.length > ir.spirv.size())
|
||||
SPIRV_CROSS_THROW("Compiler::stream() out of range.");
|
||||
return &ir.spirv[instr.offset];
|
||||
}
|
||||
}
|
||||
|
||||
ParsedIR ir;
|
||||
// Marks variables which have global scope and variables which can alias with other variables
|
||||
|
||||
@ -40,6 +40,13 @@ using namespace spv;
|
||||
using namespace SPIRV_CROSS_NAMESPACE;
|
||||
using namespace std;
|
||||
|
||||
enum ExtraSubExpressionType
|
||||
{
|
||||
// Create masks above any legal ID range to allow multiple address spaces into the extra_sub_expressions map.
|
||||
EXTRA_SUB_EXPRESSION_TYPE_STREAM_OFFSET = 0x10000000,
|
||||
EXTRA_SUB_EXPRESSION_TYPE_AUX = 0x20000000
|
||||
};
|
||||
|
||||
static bool is_unsigned_opcode(Op op)
|
||||
{
|
||||
// Don't have to be exhaustive, only relevant for legacy target checking ...
|
||||
@ -7364,7 +7371,7 @@ void CompilerGLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
|
||||
{
|
||||
// Make sure we have a unique ID here to avoid aliasing the extra sub-expressions between clamp and NMin sub-op.
|
||||
// IDs cannot exceed 24 bits, so we can make use of the higher bits for some unique flags.
|
||||
uint32_t &max_id = extra_sub_expressions[id | 0x80000000u];
|
||||
uint32_t &max_id = extra_sub_expressions[id | EXTRA_SUB_EXPRESSION_TYPE_AUX];
|
||||
if (!max_id)
|
||||
max_id = ir.increase_bound_by(1);
|
||||
|
||||
@ -10314,10 +10321,27 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
||||
uint32_t rhs = ops[1];
|
||||
if (lhs != rhs)
|
||||
{
|
||||
flush_variable_declaration(lhs);
|
||||
flush_variable_declaration(rhs);
|
||||
statement(to_expression(lhs), " = ", to_unpacked_expression(rhs), ";");
|
||||
register_write(lhs);
|
||||
uint32_t &tmp_id = extra_sub_expressions[instruction.offset | EXTRA_SUB_EXPRESSION_TYPE_STREAM_OFFSET];
|
||||
if (!tmp_id)
|
||||
tmp_id = ir.increase_bound_by(1);
|
||||
uint32_t tmp_type_id = expression_type(rhs).parent_type;
|
||||
|
||||
EmbeddedInstruction fake_load, fake_store;
|
||||
fake_load.op = OpLoad;
|
||||
fake_load.length = 3;
|
||||
fake_load.ops.push_back(tmp_type_id);
|
||||
fake_load.ops.push_back(tmp_id);
|
||||
fake_load.ops.push_back(rhs);
|
||||
|
||||
fake_store.op = OpStore;
|
||||
fake_store.length = 2;
|
||||
fake_store.ops.push_back(lhs);
|
||||
fake_store.ops.push_back(tmp_id);
|
||||
|
||||
// Load and Store do a *lot* of workarounds, and we'd like to reuse them as much as possible.
|
||||
// Synthesize a fake Load and Store pair for CopyMemory.
|
||||
emit_instruction(fake_load);
|
||||
emit_instruction(fake_store);
|
||||
}
|
||||
break;
|
||||
}
|
||||
|
||||
Loading…
Reference in New Issue
Block a user