diff options
| author | Bill Hollings <bill.hollings@brenwill.com> | 2022-09-14 22:19:15 +0300 |
|---|---|---|
| committer | Bill Hollings <bill.hollings@brenwill.com> | 2022-09-14 22:19:15 +0300 |
| commit | 5493b3030e4fe86d6f52a31ef97837f8f8d2736b (patch) | |
| tree | 68e281d891fa7efd1697fef3735ac7889db79ef2 | |
| parent | f6ca6178251c3c886d99781c5437df919fc21734 (diff) | |
MSL: Support OpPtrEqual, OpPtrNotEqual, and OpPtrDiff.
- Add CompilerMSL::emit_binary_ptr_op() and to_ptr_expression()
to emit binary pointer op. Compare matrix addresses without automatic
transpose() conversion, to avoid error taking address of temporary copy.
- Add Compiler::add_active_interface_variable() to also track active
interface vars in the entry point for SPIR-V 1.4 and above.
- For OpPtrAccessChain that ends in array element, use Element
as offset to existing index, otherwise it will access into
array dimension that doesn't exist.
- Dereference pointer function call arguments. Ultimately, this
dereferencing is actually backwards, and in future, we should aim
to properly support passing pointer variables between functions,
but such a refactoring was beyond the scope here.
- Use [] to declare array of pointers, as array<T*> is not supported in MSL.
- Add unit test shaders.
15 files changed, 749 insertions, 22 deletions
diff --git a/reference/shaders-msl-no-opt/asm/comp/opptrdiff-basic.spv14.asm.comp b/reference/shaders-msl-no-opt/asm/comp/opptrdiff-basic.spv14.asm.comp new file mode 100644 index 00000000..2a8b59f0 --- /dev/null +++ b/reference/shaders-msl-no-opt/asm/comp/opptrdiff-basic.spv14.asm.comp @@ -0,0 +1,51 @@ +#include <metal_stdlib> +#include <simd/simd.h> + +using namespace metal; + +struct _7 +{ + int _m0[1][4]; +}; + +struct _9 +{ + int _m0[1][17]; +}; + +struct _11 +{ + int _m0; +}; + +kernel void main0(device _7& _2 [[buffer(0)]], device _9& _3 [[buffer(1)]], constant _11& _4 [[buffer(2)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) +{ + if (int3(gl_WorkGroupID).x >= _4._m0) + { + return; + } + int _49; + if (int3(gl_LocalInvocationID).x == 1) + { + _3._m0[int3(gl_WorkGroupID).x][16] = &_2._m0[int3(gl_WorkGroupID).x] - &_2._m0[0]; + _49 = 0; + } + else + { + _49 = 0; + } + for (;;) + { + int _50 = _49 + 1; + _3._m0[int3(gl_WorkGroupID).x][(int3(gl_LocalInvocationID).x * 4) + _49] = &_2._m0[int3(gl_WorkGroupID).x][int3(gl_LocalInvocationID).x] - &_2._m0[int3(gl_WorkGroupID).x][_49]; + if (_50 == 4) + { + break; + } + else + { + _49 = _50; + } + } +} + diff --git a/reference/shaders-msl-no-opt/asm/comp/opptrdiff-opptraccesschain-elem-offset.spv14.asm.comp b/reference/shaders-msl-no-opt/asm/comp/opptrdiff-opptraccesschain-elem-offset.spv14.asm.comp new file mode 100644 index 00000000..69e76f3f --- /dev/null +++ b/reference/shaders-msl-no-opt/asm/comp/opptrdiff-opptraccesschain-elem-offset.spv14.asm.comp @@ -0,0 +1,45 @@ +#include <metal_stdlib> +#include <simd/simd.h> + +using namespace metal; + +struct _7 +{ + int _m0; + int _m1[1]; +}; + +struct _9 +{ + int2 _m0[1]; +}; + +kernel void main0(device _7& _2 [[buffer(0)]], device _9& _3 [[buffer(1)]]) +{ + int _28 = _2._m0; + device int* _4 = &_2._m1[0]; + device int* _5 = &_2._m1[0 + _28]; + int _34; + if (!(_28 <= 0)) + { + _34 = 0; + for (;;) + { + device int* _36 = _4; + device int* _37 = _5; + int _35 = _34 + 1; + _4 = &_36[1]; + _5 = &_37[-1]; + _3._m0[_34] = int2(_36 - _37, _37 - _36); + if (_34 >= _28) + { + break; + } + else + { + _34 = _35; + } + } + } +} + diff --git a/reference/shaders-msl-no-opt/asm/comp/opptrequal-basic.spv14.asm.comp b/reference/shaders-msl-no-opt/asm/comp/opptrequal-basic.spv14.asm.comp new file mode 100644 index 00000000..52916413 --- /dev/null +++ b/reference/shaders-msl-no-opt/asm/comp/opptrequal-basic.spv14.asm.comp @@ -0,0 +1,33 @@ +#include <metal_stdlib> +#include <simd/simd.h> + +using namespace metal; + +struct _7 +{ + uint _m0[1]; +}; + +kernel void main0(device _7& _2 [[buffer(0)]], device _7& _3 [[buffer(1)]], device _7& _4 [[buffer(2)]], device _7& _5 [[buffer(3)]]) +{ + uint _18 = 0u; + uint _28 = _18 + 1u; + _5._m0[_18] = uint(&_2 == &_3); + uint _32 = _28 + 1u; + _5._m0[_28] = uint(&_2._m0 == &_3._m0); + uint _36 = _32 + 1u; + _5._m0[_32] = uint(&_2._m0[0u] == &_3._m0[0u]); + uint _40 = _36 + 1u; + _5._m0[_36] = uint(&_2 == &_4); + uint _44 = _40 + 1u; + _5._m0[_40] = uint(&_2._m0 == &_4._m0); + uint _48 = _44 + 1u; + _5._m0[_44] = uint(&_2._m0[0u] == &_4._m0[0u]); + uint _52 = _48 + 1u; + _5._m0[_48] = uint(&_3 == &_4); + uint _56 = _52 + 1u; + _5._m0[_52] = uint(&_3._m0 == &_4._m0); + _5._m0[_56] = uint(&_3._m0[0u] == &_4._m0[0u]); + _5._m0[_56 + 1u] = uint(&_2 == &_2); +} + diff --git a/reference/shaders-msl-no-opt/asm/comp/opptrequal-row-maj-mtx-bypass-transpose.spv14.asm.comp b/reference/shaders-msl-no-opt/asm/comp/opptrequal-row-maj-mtx-bypass-transpose.spv14.asm.comp new file mode 100644 index 00000000..16d29c1d --- /dev/null +++ b/reference/shaders-msl-no-opt/asm/comp/opptrequal-row-maj-mtx-bypass-transpose.spv14.asm.comp @@ -0,0 +1,37 @@ +#include <metal_stdlib> +#include <simd/simd.h> + +using namespace metal; + +struct _6 +{ + float4x4 _m0; + float4x4 _m1; + float _m2; + float _m3; +}; + +struct _7 +{ + uint _m0[1]; +}; + +kernel void main0(device _6& _2 [[buffer(0)]], device _6& _3 [[buffer(1)]], device _7& _4 [[buffer(2)]]) +{ + uint _26 = 0u; + uint _39 = _26 + 1u; + _4._m0[_26] = (&_2._m2 == &_2._m3) ? 0u : 1u; + bool _40 = &_2._m2 == &_3._m2; + uint _43 = _39 + 1u; + _4._m0[_39] = _40 ? 0u : 1u; + bool _46 = (_40 ? &_2._m2 : &_2._m3) == (_40 ? &_3._m2 : &_3._m3); + uint _49 = _43 + 1u; + _4._m0[_43] = _46 ? 0u : 1u; + uint _54 = _49 + 1u; + _4._m0[_49] = ((_46 ? &_2._m2 : &_2._m3) == &((device float*)&_2._m0[0u])[0u]) ? 0u : 1u; + uint _56 = (&_2._m0 == &_2._m1) ? 0u : 1u; + uint _58 = _54 + 1u; + _4._m0[_54] = _56; + _4._m0[_58] = _56; +} + diff --git a/reference/shaders-msl-no-opt/asm/comp/opptrnotequal-basic.spv14.asm.comp b/reference/shaders-msl-no-opt/asm/comp/opptrnotequal-basic.spv14.asm.comp new file mode 100644 index 00000000..d9af2035 --- /dev/null +++ b/reference/shaders-msl-no-opt/asm/comp/opptrnotequal-basic.spv14.asm.comp @@ -0,0 +1,33 @@ +#include <metal_stdlib> +#include <simd/simd.h> + +using namespace metal; + +struct _7 +{ + uint _m0[1]; +}; + +kernel void main0(device _7& _2 [[buffer(0)]], device _7& _3 [[buffer(1)]], device _7& _4 [[buffer(2)]], device _7& _5 [[buffer(3)]]) +{ + uint _18 = 0u; + uint _28 = _18 + 1u; + _5._m0[_18] = uint(&_2 != &_3); + uint _32 = _28 + 1u; + _5._m0[_28] = uint(&_2._m0 != &_3._m0); + uint _36 = _32 + 1u; + _5._m0[_32] = uint(&_2._m0[0u] != &_3._m0[0u]); + uint _40 = _36 + 1u; + _5._m0[_36] = uint(&_2 != &_4); + uint _44 = _40 + 1u; + _5._m0[_40] = uint(&_2._m0 != &_4._m0); + uint _48 = _44 + 1u; + _5._m0[_44] = uint(&_2._m0[0u] != &_4._m0[0u]); + uint _52 = _48 + 1u; + _5._m0[_48] = uint(&_3 != &_4); + uint _56 = _52 + 1u; + _5._m0[_52] = uint(&_3._m0 != &_4._m0); + _5._m0[_56] = uint(&_3._m0[0u] != &_4._m0[0u]); + _5._m0[_56 + 1u] = uint(&_2 != &_2); +} + diff --git a/shaders-msl-no-opt/asm/comp/opptrdiff-basic.spv14.asm.comp b/shaders-msl-no-opt/asm/comp/opptrdiff-basic.spv14.asm.comp new file mode 100644 index 00000000..8319dfdb --- /dev/null +++ b/shaders-msl-no-opt/asm/comp/opptrdiff-basic.spv14.asm.comp @@ -0,0 +1,98 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 59 +; Schema: 0 + OpCapability Shader + OpCapability VariablePointersStorageBuffer + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %1 "main" %2 %3 %4 %5 %6 + OpExecutionMode %1 LocalSize 4 1 1 + OpDecorate %7 Block + OpMemberDecorate %7 0 Offset 0 + OpDecorate %8 ArrayStride 16 + OpDecorate %9 Block + OpMemberDecorate %9 0 Offset 0 + OpDecorate %10 ArrayStride 68 + OpDecorate %11 Block + OpMemberDecorate %11 0 Offset 0 + OpDecorate %12 ArrayStride 4 + OpDecorate %13 ArrayStride 4 + OpDecorate %2 DescriptorSet 0 + OpDecorate %2 Binding 0 + OpDecorate %3 DescriptorSet 0 + OpDecorate %3 Binding 1 + OpDecorate %4 DescriptorSet 0 + OpDecorate %4 Binding 2 + OpDecorate %5 BuiltIn LocalInvocationId + OpDecorate %6 BuiltIn WorkgroupId + %14 = OpTypeVoid + %15 = OpTypeBool + %16 = OpTypeInt 32 1 + %17 = OpConstant %16 0 + %18 = OpConstant %16 1 + %19 = OpConstant %16 4 + %20 = OpConstant %16 16 + %21 = OpConstant %16 17 + %22 = OpTypeVector %16 3 + %23 = OpTypePointer Input %22 + %12 = OpTypeArray %16 %19 + %8 = OpTypeRuntimeArray %12 + %7 = OpTypeStruct %8 + %24 = OpTypePointer StorageBuffer %7 + %25 = OpTypePointer StorageBuffer %12 + %13 = OpTypeArray %16 %21 + %10 = OpTypeRuntimeArray %13 + %9 = OpTypeStruct %10 + %26 = OpTypePointer StorageBuffer %9 + %27 = OpTypePointer StorageBuffer %13 + %28 = OpTypePointer StorageBuffer %16 + %11 = OpTypeStruct %16 + %29 = OpTypePointer Uniform %11 + %30 = OpTypePointer Uniform %16 + %2 = OpVariable %24 StorageBuffer + %3 = OpVariable %26 StorageBuffer + %4 = OpVariable %29 Uniform + %5 = OpVariable %23 Input + %6 = OpVariable %23 Input + %31 = OpTypeFunction %14 + %1 = OpFunction %14 None %31 + %32 = OpLabel + %33 = OpAccessChain %30 %4 %17 + %34 = OpLoad %16 %33 + %35 = OpLoad %22 %6 + %36 = OpCompositeExtract %16 %35 0 + %37 = OpLoad %22 %5 + %38 = OpCompositeExtract %16 %37 0 + %39 = OpAccessChain %25 %2 %17 %17 + %40 = OpAccessChain %25 %2 %17 %36 + %41 = OpSGreaterThanEqual %15 %36 %34 + OpSelectionMerge %42 None + OpBranchConditional %41 %43 %42 + %43 = OpLabel + OpReturn + %42 = OpLabel + %44 = OpIEqual %15 %38 %18 + OpSelectionMerge %45 None + OpBranchConditional %44 %46 %45 + %46 = OpLabel + %47 = OpPtrDiff %16 %40 %39 + %48 = OpAccessChain %28 %3 %17 %36 %20 + OpStore %48 %47 + OpBranch %45 + %45 = OpLabel + %49 = OpPhi %16 %17 %42 %17 %46 %50 %45 + %50 = OpIAdd %16 %49 %18 + %51 = OpIEqual %15 %50 %19 + %52 = OpIMul %16 %38 %19 + %53 = OpIAdd %16 %52 %49 + %54 = OpAccessChain %28 %40 %38 + %55 = OpAccessChain %28 %40 %49 + %56 = OpPtrDiff %16 %54 %55 + %57 = OpAccessChain %28 %3 %17 %36 %53 + OpStore %57 %56 + OpLoopMerge %58 %45 None + OpBranchConditional %51 %58 %45 + %58 = OpLabel + OpReturn + OpFunctionEnd diff --git a/shaders-msl-no-opt/asm/comp/opptrdiff-opptraccesschain-elem-offset.spv14.asm.comp b/shaders-msl-no-opt/asm/comp/opptrdiff-opptraccesschain-elem-offset.spv14.asm.comp new file mode 100644 index 00000000..85664919 --- /dev/null +++ b/shaders-msl-no-opt/asm/comp/opptrdiff-opptraccesschain-elem-offset.spv14.asm.comp @@ -0,0 +1,79 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 46 +; Schema: 0 + OpCapability Shader + OpCapability VariablePointersStorageBuffer + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %1 "main" %2 %3 %4 %5 + OpExecutionMode %1 LocalSize 1 1 1 + OpDecorate %6 ArrayStride 4 + OpDecorate %7 Block + OpMemberDecorate %7 0 Offset 0 + OpMemberDecorate %7 1 Offset 4 + OpDecorate %2 DescriptorSet 0 + OpDecorate %2 Binding 0 + OpDecorate %8 ArrayStride 8 + OpDecorate %9 Block + OpMemberDecorate %9 0 Offset 0 + OpDecorate %3 DescriptorSet 0 + OpDecorate %3 Binding 1 + OpDecorate %10 ArrayStride 4 + %11 = OpTypeVoid + %12 = OpTypeBool + %13 = OpTypeInt 32 1 + %14 = OpConstant %13 -1 + %15 = OpConstant %13 0 + %16 = OpConstant %13 1 + %17 = OpConstant %13 2 + %18 = OpConstant %13 3 + %19 = OpTypeVector %13 2 + %6 = OpTypeRuntimeArray %13 + %7 = OpTypeStruct %13 %6 + %20 = OpTypePointer StorageBuffer %7 + %2 = OpVariable %20 StorageBuffer + %8 = OpTypeRuntimeArray %19 + %9 = OpTypeStruct %8 + %21 = OpTypePointer StorageBuffer %9 + %3 = OpVariable %21 StorageBuffer + %10 = OpTypePointer StorageBuffer %13 + %22 = OpTypePointer Private %10 + %4 = OpVariable %22 Private + %5 = OpVariable %22 Private + %23 = OpTypePointer StorageBuffer %13 + %24 = OpTypePointer StorageBuffer %19 + %25 = OpTypeFunction %11 + %1 = OpFunction %11 None %25 + %26 = OpLabel + %27 = OpAccessChain %23 %2 %15 + %28 = OpLoad %13 %27 + %29 = OpAccessChain %10 %2 %16 %15 + OpStore %4 %29 + %30 = OpPtrAccessChain %10 %29 %28 + OpStore %5 %30 + %31 = OpSLessThanEqual %12 %28 %15 + OpSelectionMerge %32 None + OpBranchConditional %31 %32 %33 + %33 = OpLabel + %34 = OpPhi %13 %15 %26 %35 %33 + %36 = OpLoad %10 %4 + %37 = OpLoad %10 %5 + %38 = OpPtrAccessChain %10 %36 %16 + %39 = OpPtrAccessChain %10 %37 %14 + %35 = OpIAdd %13 %34 %16 + OpStore %4 %38 + OpStore %5 %39 + %40 = OpPtrDiff %13 %36 %37 + %41 = OpPtrDiff %13 %37 %36 + %42 = OpCompositeConstruct %19 %40 %41 + %43 = OpAccessChain %24 %3 %15 %34 + OpStore %43 %42 + %44 = OpSGreaterThanEqual %12 %34 %28 + OpLoopMerge %45 %33 None + OpBranchConditional %44 %45 %33 + %45 = OpLabel + OpBranch %32 + %32 = OpLabel + OpReturn + OpFunctionEnd diff --git a/shaders-msl-no-opt/asm/comp/opptrequal-basic.spv14.asm.comp b/shaders-msl-no-opt/asm/comp/opptrequal-basic.spv14.asm.comp new file mode 100644 index 00000000..5a97976c --- /dev/null +++ b/shaders-msl-no-opt/asm/comp/opptrequal-basic.spv14.asm.comp @@ -0,0 +1,96 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 64 +; Schema: 0 + OpCapability Shader + OpCapability VariablePointers + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %1 "main" %2 %3 %4 %5 + OpExecutionMode %1 LocalSize 1 1 1 + OpDecorate %6 ArrayStride 4 + OpDecorate %7 Block + OpMemberDecorate %7 0 Offset 0 + OpDecorate %2 DescriptorSet 0 + OpDecorate %2 Binding 0 + OpDecorate %3 DescriptorSet 0 + OpDecorate %3 Binding 1 + OpDecorate %4 DescriptorSet 0 + OpDecorate %4 Binding 2 + OpDecorate %5 DescriptorSet 0 + OpDecorate %5 Binding 3 + %8 = OpTypeVoid + %9 = OpTypeBool + %10 = OpTypeInt 32 0 + %11 = OpConstant %10 0 + %12 = OpConstant %10 1 + %6 = OpTypeRuntimeArray %10 + %7 = OpTypeStruct %6 + %13 = OpTypePointer StorageBuffer %7 + %14 = OpTypePointer StorageBuffer %6 + %15 = OpTypePointer StorageBuffer %10 + %2 = OpVariable %13 StorageBuffer + %3 = OpVariable %13 StorageBuffer + %4 = OpVariable %13 StorageBuffer + %5 = OpVariable %13 StorageBuffer + %16 = OpTypeFunction %8 + %1 = OpFunction %8 None %16 + %17 = OpLabel + %18 = OpCopyObject %10 %11 + %19 = OpAccessChain %14 %2 %11 + %20 = OpAccessChain %15 %2 %11 %11 + %21 = OpAccessChain %14 %3 %11 + %22 = OpAccessChain %15 %3 %11 %11 + %23 = OpAccessChain %14 %4 %11 + %24 = OpAccessChain %15 %4 %11 %11 + %25 = OpPtrEqual %9 %2 %3 + %26 = OpSelect %10 %25 %12 %11 + %27 = OpAccessChain %15 %5 %11 %18 + %28 = OpIAdd %10 %18 %12 + OpStore %27 %26 + %29 = OpPtrEqual %9 %19 %21 + %30 = OpSelect %10 %29 %12 %11 + %31 = OpAccessChain %15 %5 %11 %28 + %32 = OpIAdd %10 %28 %12 + OpStore %31 %30 + %33 = OpPtrEqual %9 %20 %22 + %34 = OpSelect %10 %33 %12 %11 + %35 = OpAccessChain %15 %5 %11 %32 + %36 = OpIAdd %10 %32 %12 + OpStore %35 %34 + %37 = OpPtrEqual %9 %2 %4 + %38 = OpSelect %10 %37 %12 %11 + %39 = OpAccessChain %15 %5 %11 %36 + %40 = OpIAdd %10 %36 %12 + OpStore %39 %38 + %41 = OpPtrEqual %9 %19 %23 + %42 = OpSelect %10 %41 %12 %11 + %43 = OpAccessChain %15 %5 %11 %40 + %44 = OpIAdd %10 %40 %12 + OpStore %43 %42 + %45 = OpPtrEqual %9 %20 %24 + %46 = OpSelect %10 %45 %12 %11 + %47 = OpAccessChain %15 %5 %11 %44 + %48 = OpIAdd %10 %44 %12 + OpStore %47 %46 + %49 = OpPtrEqual %9 %3 %4 + %50 = OpSelect %10 %49 %12 %11 + %51 = OpAccessChain %15 %5 %11 %48 + %52 = OpIAdd %10 %48 %12 + OpStore %51 %50 + %53 = OpPtrEqual %9 %21 %23 + %54 = OpSelect %10 %53 %12 %11 + %55 = OpAccessChain %15 %5 %11 %52 + %56 = OpIAdd %10 %52 %12 + OpStore %55 %54 + %57 = OpPtrEqual %9 %22 %24 + %58 = OpSelect %10 %57 %12 %11 + %59 = OpAccessChain %15 %5 %11 %56 + %60 = OpIAdd %10 %56 %12 + OpStore %59 %58 + %61 = OpPtrEqual %9 %2 %2 + %62 = OpSelect %10 %61 %12 %11 + %63 = OpAccessChain %15 %5 %11 %60 + OpStore %63 %62 + OpReturn + OpFunctionEnd diff --git a/shaders-msl-no-opt/asm/comp/opptrequal-row-maj-mtx-bypass-transpose.spv14.asm.comp b/shaders-msl-no-opt/asm/comp/opptrequal-row-maj-mtx-bypass-transpose.spv14.asm.comp new file mode 100644 index 00000000..89813b22 --- /dev/null +++ b/shaders-msl-no-opt/asm/comp/opptrequal-row-maj-mtx-bypass-transpose.spv14.asm.comp @@ -0,0 +1,98 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 63 +; Schema: 0 + OpCapability Shader + OpCapability VariablePointersStorageBuffer + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %1 "main" %2 %3 %4 + OpExecutionMode %1 LocalSize 1 1 1 + OpDecorate %5 ArrayStride 4 + OpDecorate %6 Block + OpDecorate %7 Block + OpMemberDecorate %6 0 ColMajor + OpMemberDecorate %6 0 Offset 0 + OpMemberDecorate %6 0 MatrixStride 16 + OpMemberDecorate %6 1 RowMajor + OpMemberDecorate %6 1 Offset 64 + OpMemberDecorate %6 1 MatrixStride 16 + OpMemberDecorate %6 2 Offset 128 + OpMemberDecorate %6 3 Offset 132 + OpMemberDecorate %7 0 Offset 0 + OpDecorate %2 DescriptorSet 0 + OpDecorate %2 Binding 0 + OpDecorate %3 DescriptorSet 0 + OpDecorate %3 Binding 1 + OpDecorate %4 DescriptorSet 0 + OpDecorate %4 Binding 2 + %8 = OpTypeVoid + %9 = OpTypeBool + %10 = OpTypeInt 32 0 + %11 = OpConstant %10 0 + %12 = OpConstant %10 1 + %13 = OpConstant %10 2 + %14 = OpConstant %10 3 + %15 = OpTypeFloat 32 + %5 = OpTypeRuntimeArray %10 + %16 = OpTypeVector %15 4 + %17 = OpTypeMatrix %16 4 + %6 = OpTypeStruct %17 %17 %15 %15 + %7 = OpTypeStruct %5 + %18 = OpTypePointer StorageBuffer %6 + %19 = OpTypePointer StorageBuffer %7 + %20 = OpTypePointer StorageBuffer %17 + %21 = OpTypePointer StorageBuffer %10 + %22 = OpTypePointer StorageBuffer %15 + %23 = OpTypePointer StorageBuffer %16 + %2 = OpVariable %18 StorageBuffer + %3 = OpVariable %18 StorageBuffer + %4 = OpVariable %19 StorageBuffer + %24 = OpTypeFunction %8 + %1 = OpFunction %8 None %24 + %25 = OpLabel + %26 = OpCopyObject %10 %11 + %27 = OpAccessChain %22 %2 %13 + %28 = OpAccessChain %22 %2 %14 + %29 = OpAccessChain %22 %3 %13 + %30 = OpAccessChain %22 %3 %14 + %31 = OpAccessChain %20 %2 %11 + %32 = OpAccessChain %20 %2 %12 + %33 = OpAccessChain %23 %2 %11 %11 + %34 = OpAccessChain %23 %2 %11 %12 + %35 = OpAccessChain %22 %2 %11 %11 %11 + %36 = OpPtrEqual %9 %27 %28 + %37 = OpSelect %10 %36 %11 %12 + %38 = OpAccessChain %21 %4 %11 %26 + %39 = OpIAdd %10 %26 %12 + OpStore %38 %37 + %40 = OpPtrEqual %9 %27 %29 + %41 = OpSelect %10 %40 %11 %12 + %42 = OpAccessChain %21 %4 %11 %39 + %43 = OpIAdd %10 %39 %12 + OpStore %42 %41 + %44 = OpSelect %22 %40 %27 %28 + %45 = OpSelect %22 %40 %29 %30 + %46 = OpPtrEqual %9 %44 %45 + %47 = OpSelect %10 %46 %11 %12 + %48 = OpAccessChain %21 %4 %11 %43 + %49 = OpIAdd %10 %43 %12 + OpStore %48 %47 + %50 = OpSelect %22 %46 %27 %28 + %51 = OpPtrEqual %9 %50 %35 + %52 = OpSelect %10 %51 %11 %12 + %53 = OpAccessChain %21 %4 %11 %49 + %54 = OpIAdd %10 %49 %12 + OpStore %53 %52 + %55 = OpPtrEqual %9 %31 %32 + %56 = OpSelect %10 %55 %11 %12 + %57 = OpAccessChain %21 %4 %11 %54 + %58 = OpIAdd %10 %54 %12 + OpStore %57 %56 + %59 = OpPtrEqual %9 %33 %34 + %60 = OpSelect %10 %59 %11 %12 + %61 = OpAccessChain %21 %4 %11 %58 + %62 = OpIAdd %10 %58 %12 + OpStore %61 %56 + OpReturn + OpFunctionEnd diff --git a/shaders-msl-no-opt/asm/comp/opptrnotequal-basic.spv14.asm.comp b/shaders-msl-no-opt/asm/comp/opptrnotequal-basic.spv14.asm.comp new file mode 100644 index 00000000..1cbf8045 --- /dev/null +++ b/shaders-msl-no-opt/asm/comp/opptrnotequal-basic.spv14.asm.comp @@ -0,0 +1,96 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 64 +; Schema: 0 + OpCapability Shader + OpCapability VariablePointers + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %1 "main" %2 %3 %4 %5 + OpExecutionMode %1 LocalSize 1 1 1 + OpDecorate %6 ArrayStride 4 + OpDecorate %7 Block + OpMemberDecorate %7 0 Offset 0 + OpDecorate %2 DescriptorSet 0 + OpDecorate %2 Binding 0 + OpDecorate %3 DescriptorSet 0 + OpDecorate %3 Binding 1 + OpDecorate %4 DescriptorSet 0 + OpDecorate %4 Binding 2 + OpDecorate %5 DescriptorSet 0 + OpDecorate %5 Binding 3 + %8 = OpTypeVoid + %9 = OpTypeBool + %10 = OpTypeInt 32 0 + %11 = OpConstant %10 0 + %12 = OpConstant %10 1 + %6 = OpTypeRuntimeArray %10 + %7 = OpTypeStruct %6 + %13 = OpTypePointer StorageBuffer %7 + %14 = OpTypePointer StorageBuffer %6 + %15 = OpTypePointer StorageBuffer %10 + %2 = OpVariable %13 StorageBuffer + %3 = OpVariable %13 StorageBuffer + %4 = OpVariable %13 StorageBuffer + %5 = OpVariable %13 StorageBuffer + %16 = OpTypeFunction %8 + %1 = OpFunction %8 None %16 + %17 = OpLabel + %18 = OpCopyObject %10 %11 + %19 = OpAccessChain %14 %2 %11 + %20 = OpAccessChain %15 %2 %11 %11 + %21 = OpAccessChain %14 %3 %11 + %22 = OpAccessChain %15 %3 %11 %11 + %23 = OpAccessChain %14 %4 %11 + %24 = OpAccessChain %15 %4 %11 %11 + %25 = OpPtrNotEqual %9 %2 %3 + %26 = OpSelect %10 %25 %12 %11 + %27 = OpAccessChain %15 %5 %11 %18 + %28 = OpIAdd %10 %18 %12 + OpStore %27 %26 + %29 = OpPtrNotEqual %9 %19 %21 + %30 = OpSelect %10 %29 %12 %11 + %31 = OpAccessChain %15 %5 %11 %28 + %32 = OpIAdd %10 %28 %12 + OpStore %31 %30 + %33 = OpPtrNotEqual %9 %20 %22 + %34 = OpSelect %10 %33 %12 %11 + %35 = OpAccessChain %15 %5 %11 %32 + %36 = OpIAdd %10 %32 %12 + OpStore %35 %34 + %37 = OpPtrNotEqual %9 %2 %4 + %38 = OpSelect %10 %37 %12 %11 + %39 = OpAccessChain %15 %5 %11 %36 + %40 = OpIAdd %10 %36 %12 + OpStore %39 %38 + %41 = OpPtrNotEqual %9 %19 %23 + %42 = OpSelect %10 %41 %12 %11 + %43 = OpAccessChain %15 %5 %11 %40 + %44 = OpIAdd %10 %40 %12 + OpStore %43 %42 + %45 = OpPtrNotEqual %9 %20 %24 + %46 = OpSelect %10 %45 %12 %11 + %47 = OpAccessChain %15 %5 %11 %44 + %48 = OpIAdd %10 %44 %12 + OpStore %47 %46 + %49 = OpPtrNotEqual %9 %3 %4 + %50 = OpSelect %10 %49 %12 %11 + %51 = OpAccessChain %15 %5 %11 %48 + %52 = OpIAdd %10 %48 %12 + OpStore %51 %50 + %53 = OpPtrNotEqual %9 %21 %23 + %54 = OpSelect %10 %53 %12 %11 + %55 = OpAccessChain %15 %5 %11 %52 + %56 = OpIAdd %10 %52 %12 + OpStore %55 %54 + %57 = OpPtrNotEqual %9 %22 %24 + %58 = OpSelect %10 %57 %12 %11 + %59 = OpAccessChain %15 %5 %11 %56 + %60 = OpIAdd %10 %56 %12 + OpStore %59 %58 + %61 = OpPtrNotEqual %9 %2 %2 + %62 = OpSelect %10 %61 %12 %11 + %63 = OpAccessChain %15 %5 %11 %60 + OpStore %63 %62 + OpReturn + OpFunctionEnd diff --git a/spirv_cross.cpp b/spirv_cross.cpp index 04ea35fa..599a3232 100644 --- a/spirv_cross.cpp +++ b/spirv_cross.cpp @@ -2375,6 +2375,19 @@ void Compiler::add_implied_read_expression(SPIRAccessChain &e, uint32_t source) e.implied_read_expressions.push_back(source); } +void Compiler::add_active_interface_variable(uint32_t var_id) +{ + active_interface_variables.insert(var_id); + + // In SPIR-V 1.4 and up we must also track the interface variable in the entry point. + if (ir.get_spirv_version() >= 0x10400) + { + auto &vars = get_entry_point().interface_variables; + if (find(begin(vars), end(vars), VariableID(var_id)) == end(vars)) + vars.push_back(var_id); + } +} + void Compiler::inherit_expression_dependencies(uint32_t dst, uint32_t source_expression) { // Don't inherit any expression dependencies if the expression in dst diff --git a/spirv_cross.hpp b/spirv_cross.hpp index 1d7e7c48..783af1e2 100644 --- a/spirv_cross.hpp +++ b/spirv_cross.hpp @@ -755,6 +755,7 @@ protected: void inherit_expression_dependencies(uint32_t dst, uint32_t source); void add_implied_read_expression(SPIRExpression &e, uint32_t source); void add_implied_read_expression(SPIRAccessChain &e, uint32_t source); + void add_active_interface_variable(uint32_t var_id); // For proper multiple entry point support, allow querying if an Input or Output // variable is part of that entry points interface. diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index bcd4f911..5d87cea5 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -8975,20 +8975,31 @@ const char *CompilerGLSL::index_to_swizzle(uint32_t index) } void CompilerGLSL::access_chain_internal_append_index(std::string &expr, uint32_t /*base*/, const SPIRType * /*type*/, - AccessChainFlags flags, bool & /*access_chain_is_arrayed*/, + AccessChainFlags flags, bool &access_chain_is_arrayed, uint32_t index) { bool index_is_literal = (flags & ACCESS_CHAIN_INDEX_IS_LITERAL_BIT) != 0; + bool ptr_chain = (flags & ACCESS_CHAIN_PTR_CHAIN_BIT) != 0; bool register_expression_read = (flags & ACCESS_CHAIN_SKIP_REGISTER_EXPRESSION_READ_BIT) == 0; - expr += "["; + string idx_expr = index_is_literal ? convert_to_string(index) : to_unpacked_expression(index, register_expression_read); - if (index_is_literal) - expr += convert_to_string(index); + // For the case where the base of an OpPtrAccessChain already ends in [n], + // we need to use the index as an offset to the existing index, otherwise, + // we can just use the index directly. + if (ptr_chain && access_chain_is_arrayed) + { + size_t split_pos = expr.find_last_of(']'); + string expr_front = expr.substr(0, split_pos); + string expr_back = expr.substr(split_pos); + expr = expr_front + " + " + enclose_expression(idx_expr) + expr_back; + } else - expr += to_unpacked_expression(index, register_expression_read); - - expr += "]"; + { + expr += "["; + expr += idx_expr; + expr += "]"; + } } bool CompilerGLSL::access_chain_needs_stage_io_builtin_translation(uint32_t) @@ -9049,10 +9060,12 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice bool pending_array_enclose = false; bool dimension_flatten = false; - const auto append_index = [&](uint32_t index, bool is_literal) { + const auto append_index = [&](uint32_t index, bool is_literal, bool is_ptr_chain = false) { AccessChainFlags mod_flags = flags; if (!is_literal) mod_flags &= ~ACCESS_CHAIN_INDEX_IS_LITERAL_BIT; + if (!is_ptr_chain) + mod_flags &= ~ACCESS_CHAIN_PTR_CHAIN_BIT; access_chain_internal_append_index(expr, base, type, mod_flags, access_chain_is_arrayed, index); check_physical_type_cast(expr, type, physical_type); }; @@ -9105,7 +9118,7 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice } else { - append_index(index, is_literal); + append_index(index, is_literal, true); } if (type->basetype == SPIRType::ControlPointArray) diff --git a/spirv_msl.cpp b/spirv_msl.cpp index 48e3e6a6..60b35021 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -1414,19 +1414,19 @@ string CompilerMSL::compile() activate_argument_buffer_resources(); if (swizzle_buffer_id) - active_interface_variables.insert(swizzle_buffer_id); + add_active_interface_variable(swizzle_buffer_id); if (buffer_size_buffer_id) - active_interface_variables.insert(buffer_size_buffer_id); + add_active_interface_variable(buffer_size_buffer_id); if (view_mask_buffer_id) - active_interface_variables.insert(view_mask_buffer_id); + add_active_interface_variable(view_mask_buffer_id); if (dynamic_offsets_buffer_id) - active_interface_variables.insert(dynamic_offsets_buffer_id); + add_active_interface_variable(dynamic_offsets_buffer_id); if (builtin_layer_id) - active_interface_variables.insert(builtin_layer_id); + add_active_interface_variable(builtin_layer_id); if (builtin_dispatch_base_id && !msl_options.supports_msl_version(1, 2)) - active_interface_variables.insert(builtin_dispatch_base_id); + add_active_interface_variable(builtin_dispatch_base_id); if (builtin_sample_mask_id) - active_interface_variables.insert(builtin_sample_mask_id); + add_active_interface_variable(builtin_sample_mask_id); // Create structs to hold input, output and uniform variables. // Do output first to ensure out. is declared at top of entry function. @@ -7239,6 +7239,23 @@ void CompilerMSL::emit_specialization_constants_and_structs() statement(""); } +void CompilerMSL::emit_binary_ptr_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op) +{ + bool forward = should_forward(op0) && should_forward(op1); + emit_op(result_type, result_id, join(to_ptr_expression(op0), " ", op, " ", to_ptr_expression(op1)), forward); + inherit_expression_dependencies(result_id, op0); + inherit_expression_dependencies(result_id, op1); +} + +string CompilerMSL::to_ptr_expression(uint32_t id, bool register_expression_read) +{ + auto *e = maybe_get<SPIRExpression>(id); + auto expr = enclose_expression(e && e->need_transpose ? e->expression : to_expression(id, register_expression_read)); + if (!should_dereference(id)) + expr = address_of_expression(expr); + return expr; +} + void CompilerMSL::emit_binary_unord_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op) { @@ -8027,6 +8044,7 @@ void CompilerMSL::check_physical_type_cast(std::string &expr, const SPIRType *ty void CompilerMSL::emit_instruction(const Instruction &instruction) { #define MSL_BOP(op) emit_binary_op(ops[0], ops[1], ops[2], ops[3], #op) +#define MSL_PTR_BOP(op) emit_binary_ptr_op(ops[0], ops[1], ops[2], ops[3], #op) #define MSL_BOP_CAST(op, type) \ emit_binary_op_cast(ops[0], ops[1], ops[2], ops[3], #op, type, opcode_is_sign_invariant(opcode)) #define MSL_UOP(op) emit_unary_op(ops[0], ops[1], ops[2], #op) @@ -8166,6 +8184,19 @@ void CompilerMSL::emit_instruction(const Instruction &instruction) MSL_UNORD_BOP(<=); break; + // Pointer math + case OpPtrEqual: + MSL_PTR_BOP(==); + break; + + case OpPtrNotEqual: + MSL_PTR_BOP(!=); + break; + + case OpPtrDiff: + MSL_PTR_BOP(-); + break; + // Derivatives case OpDPdx: case OpDPdxFine: @@ -10871,6 +10902,10 @@ string CompilerMSL::to_func_call_arg(const SPIRFunction::Parameter &arg, uint32_ constants.push_back(id); } } + // Dereference pointer variables where needed. + // FIXME: This dereference is actually backwards. We should really just support passing pointer variables between functions. + else if (should_dereference(id)) + arg_str += dereference_expression(type, CompilerGLSL::to_func_call_arg(arg, id)); else arg_str += CompilerGLSL::to_func_call_arg(arg, id); @@ -14170,17 +14205,14 @@ string CompilerMSL::type_to_array_glsl(const SPIRType &type) case SPIRType::AtomicCounter: case SPIRType::ControlPointArray: case SPIRType::RayQuery: - { return CompilerGLSL::type_to_array_glsl(type); - } + default: - { - if (using_builtin_array()) + if (type_is_array_of_pointers(type) || using_builtin_array()) return CompilerGLSL::type_to_array_glsl(type); else return ""; } - } } string CompilerMSL::constant_op_expression(const SPIRConstantOp &cop) @@ -16920,7 +16952,7 @@ void CompilerMSL::activate_argument_buffer_resources() uint32_t desc_set = get_decoration(self, DecorationDescriptorSet); if (descriptor_set_is_argument_buffer(desc_set)) - active_interface_variables.insert(self); + add_active_interface_variable(self); }); } diff --git a/spirv_msl.hpp b/spirv_msl.hpp index 4b9d88da..0e5c76db 100644 --- a/spirv_msl.hpp +++ b/spirv_msl.hpp @@ -739,6 +739,8 @@ protected: // If the underlying resource has been used for comparison then duplicate loads of that resource must be too // Use Metal's native frame-buffer fetch API for subpass inputs. void emit_texture_op(const Instruction &i, bool sparse) override; + void emit_binary_ptr_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op); + std::string to_ptr_expression(uint32_t id, bool register_expression_read = true); void emit_binary_unord_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op); void emit_instruction(const Instruction &instr) override; void emit_glsl_op(uint32_t result_type, uint32_t result_id, uint32_t op, const uint32_t *args, |