mirror of
https://github.com/KhronosGroup/SPIRV-Cross.git
synced 2024-11-10 06:00:07 +00:00
Merge remote-tracking branch 'upstream/master'
This commit is contained in:
Lukas Hermanns
commit
84351d3aed
@ -27,14 +27,14 @@ matrix:
|
||||
- os: windows
|
||||
before_install:
|
||||
- choco install python3
|
||||
- export PATH="/c/Python37:/c/Python37/Scripts:$PATH"
|
||||
- export PATH="/c/Python38:/c/Python38/Scripts:$PATH"
|
||||
env:
|
||||
- GENERATOR="Visual Studio 15 2017"
|
||||
- ARTIFACT=vs2017-32bit
|
||||
- os: windows
|
||||
before_install:
|
||||
- choco install python3
|
||||
- export PATH="/c/Python37:/c/Python37/Scripts:$PATH"
|
||||
- export PATH="/c/Python38:/c/Python38/Scripts:$PATH"
|
||||
env:
|
||||
- GENERATOR="Visual Studio 15 2017 Win64"
|
||||
- ARTIFACT=vs2017-64bit
|
||||
|
||||
@ -308,7 +308,7 @@ if (SPIRV_CROSS_STATIC)
|
||||
endif()
|
||||
|
||||
set(spirv-cross-abi-major 0)
|
||||
set(spirv-cross-abi-minor 18)
|
||||
set(spirv-cross-abi-minor 19)
|
||||
set(spirv-cross-abi-patch 0)
|
||||
|
||||
if (SPIRV_CROSS_SHARED)
|
||||
|
||||
6
main.cpp
6
main.cpp
@ -525,6 +525,7 @@ struct CLIArguments
|
||||
bool vulkan_glsl_disable_ext_samplerless_texture_functions = false;
|
||||
bool emit_line_directives = false;
|
||||
SmallVector<uint32_t> msl_discrete_descriptor_sets;
|
||||
SmallVector<uint32_t> msl_device_argument_buffers;
|
||||
SmallVector<pair<uint32_t, uint32_t>> msl_dynamic_buffers;
|
||||
SmallVector<PLSArg> pls_in;
|
||||
SmallVector<PLSArg> pls_out;
|
||||
@ -603,6 +604,7 @@ static void print_help()
|
||||
"\t[--msl-framebuffer-fetch]\n"
|
||||
"\t[--msl-emulate-cube-array]\n"
|
||||
"\t[--msl-discrete-descriptor-set <index>]\n"
|
||||
"\t[--msl-device-argument-buffer <index>]\n"
|
||||
"\t[--msl-multiview]\n"
|
||||
"\t[--msl-view-index-from-device-index]\n"
|
||||
"\t[--msl-dispatch-base]\n"
|
||||
@ -776,6 +778,8 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
|
||||
msl_comp->set_msl_options(msl_opts);
|
||||
for (auto &v : args.msl_discrete_descriptor_sets)
|
||||
msl_comp->add_discrete_descriptor_set(v);
|
||||
for (auto &v : args.msl_device_argument_buffers)
|
||||
msl_comp->set_argument_buffer_device_address_space(v, true);
|
||||
uint32_t i = 0;
|
||||
for (auto &v : args.msl_dynamic_buffers)
|
||||
msl_comp->add_dynamic_buffer(v.first, v.second, i++);
|
||||
@ -1096,6 +1100,8 @@ static int main_inner(int argc, char *argv[])
|
||||
cbs.add("--msl-argument-buffers", [&args](CLIParser &) { args.msl_argument_buffers = true; });
|
||||
cbs.add("--msl-discrete-descriptor-set",
|
||||
[&args](CLIParser &parser) { args.msl_discrete_descriptor_sets.push_back(parser.next_uint()); });
|
||||
cbs.add("--msl-device-argument-buffer",
|
||||
[&args](CLIParser &parser) { args.msl_device_argument_buffers.push_back(parser.next_uint()); });
|
||||
cbs.add("--msl-texture-buffer-native", [&args](CLIParser &) { args.msl_texture_buffer_native = true; });
|
||||
cbs.add("--msl-framebuffer-fetch", [&args](CLIParser &) { args.msl_framebuffer_fetch = true; });
|
||||
cbs.add("--msl-invariant-float-math", [&args](CLIParser &) { args.msl_invariant_float_math = true; });
|
||||
|
||||
28
reference/opt/shaders-hlsl/asm/comp/bitcast_icmp.asm.comp
Normal file
28
reference/opt/shaders-hlsl/asm/comp/bitcast_icmp.asm.comp
Normal file
@ -0,0 +1,28 @@
|
||||
RWByteAddressBuffer _5 : register(u0);
|
||||
RWByteAddressBuffer _6 : register(u1);
|
||||
|
||||
void comp_main()
|
||||
{
|
||||
bool4 _31 = bool4(int(_5.Load4(16).x) < int4(_5.Load4(0)).x, int(_5.Load4(16).y) < int4(_5.Load4(0)).y, int(_5.Load4(16).z) < int4(_5.Load4(0)).z, int(_5.Load4(16).w) < int4(_5.Load4(0)).w);
|
||||
bool4 _32 = bool4(int(_5.Load4(16).x) <= int4(_5.Load4(0)).x, int(_5.Load4(16).y) <= int4(_5.Load4(0)).y, int(_5.Load4(16).z) <= int4(_5.Load4(0)).z, int(_5.Load4(16).w) <= int4(_5.Load4(0)).w);
|
||||
bool4 _33 = bool4(_5.Load4(16).x < uint(int4(_5.Load4(0)).x), _5.Load4(16).y < uint(int4(_5.Load4(0)).y), _5.Load4(16).z < uint(int4(_5.Load4(0)).z), _5.Load4(16).w < uint(int4(_5.Load4(0)).w));
|
||||
bool4 _34 = bool4(_5.Load4(16).x <= uint(int4(_5.Load4(0)).x), _5.Load4(16).y <= uint(int4(_5.Load4(0)).y), _5.Load4(16).z <= uint(int4(_5.Load4(0)).z), _5.Load4(16).w <= uint(int4(_5.Load4(0)).w));
|
||||
bool4 _35 = bool4(int(_5.Load4(16).x) > int4(_5.Load4(0)).x, int(_5.Load4(16).y) > int4(_5.Load4(0)).y, int(_5.Load4(16).z) > int4(_5.Load4(0)).z, int(_5.Load4(16).w) > int4(_5.Load4(0)).w);
|
||||
bool4 _36 = bool4(int(_5.Load4(16).x) >= int4(_5.Load4(0)).x, int(_5.Load4(16).y) >= int4(_5.Load4(0)).y, int(_5.Load4(16).z) >= int4(_5.Load4(0)).z, int(_5.Load4(16).w) >= int4(_5.Load4(0)).w);
|
||||
bool4 _37 = bool4(_5.Load4(16).x > uint(int4(_5.Load4(0)).x), _5.Load4(16).y > uint(int4(_5.Load4(0)).y), _5.Load4(16).z > uint(int4(_5.Load4(0)).z), _5.Load4(16).w > uint(int4(_5.Load4(0)).w));
|
||||
bool4 _38 = bool4(_5.Load4(16).x >= uint(int4(_5.Load4(0)).x), _5.Load4(16).y >= uint(int4(_5.Load4(0)).y), _5.Load4(16).z >= uint(int4(_5.Load4(0)).z), _5.Load4(16).w >= uint(int4(_5.Load4(0)).w));
|
||||
_6.Store4(0, uint4(_31.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _31.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _31.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _31.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
|
||||
_6.Store4(0, uint4(_32.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _32.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _32.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _32.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
|
||||
_6.Store4(0, uint4(_33.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _33.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _33.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _33.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
|
||||
_6.Store4(0, uint4(_34.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _34.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _34.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _34.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
|
||||
_6.Store4(0, uint4(_35.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _35.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _35.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _35.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
|
||||
_6.Store4(0, uint4(_36.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _36.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _36.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _36.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
|
||||
_6.Store4(0, uint4(_37.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _37.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _37.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _37.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
|
||||
_6.Store4(0, uint4(_38.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _38.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _38.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _38.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void main()
|
||||
{
|
||||
comp_main();
|
||||
}
|
||||
29
reference/opt/shaders-msl/asm/comp/bitcast_icmp.asm.comp
Normal file
29
reference/opt/shaders-msl/asm/comp/bitcast_icmp.asm.comp
Normal file
@ -0,0 +1,29 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct _3
|
||||
{
|
||||
int4 _m0;
|
||||
uint4 _m1;
|
||||
};
|
||||
|
||||
struct _4
|
||||
{
|
||||
uint4 _m0;
|
||||
int4 _m1;
|
||||
};
|
||||
|
||||
kernel void main0(device _3& restrict _5 [[buffer(0)]], device _4& restrict _6 [[buffer(1)]])
|
||||
{
|
||||
_6._m0 = select(uint4(0u), uint4(1u), int4(_5._m1) < _5._m0);
|
||||
_6._m0 = select(uint4(0u), uint4(1u), int4(_5._m1) <= _5._m0);
|
||||
_6._m0 = select(uint4(0u), uint4(1u), _5._m1 < uint4(_5._m0));
|
||||
_6._m0 = select(uint4(0u), uint4(1u), _5._m1 <= uint4(_5._m0));
|
||||
_6._m0 = select(uint4(0u), uint4(1u), int4(_5._m1) > _5._m0);
|
||||
_6._m0 = select(uint4(0u), uint4(1u), int4(_5._m1) >= _5._m0);
|
||||
_6._m0 = select(uint4(0u), uint4(1u), _5._m1 > uint4(_5._m0));
|
||||
_6._m0 = select(uint4(0u), uint4(1u), _5._m1 >= uint4(_5._m0));
|
||||
}
|
||||
|
||||
@ -0,0 +1,85 @@
|
||||
#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 UBO
|
||||
{
|
||||
float4 v;
|
||||
};
|
||||
|
||||
struct spvDescriptorSetBuffer0
|
||||
{
|
||||
array<texture2d<float>, 10000> uSamplers [[id(0)]];
|
||||
array<sampler, 10000> uSamplersSmplr [[id(10000)]];
|
||||
};
|
||||
|
||||
struct spvDescriptorSetBuffer1
|
||||
{
|
||||
spvUnsafeArray<constant thread UBO*, 10000> vs [[id(0)]];
|
||||
};
|
||||
|
||||
struct spvDescriptorSetBuffer2
|
||||
{
|
||||
texture2d<float> uSampler [[id(0)]];
|
||||
sampler uSamplerSmplr [[id(1)]];
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0(main0_in in [[stage_in]], const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], constant spvDescriptorSetBuffer2& spvDescriptorSet2 [[buffer(2)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
out.FragColor = (spvDescriptorSet0.uSamplers[9999].sample(spvDescriptorSet0.uSamplersSmplr[9999], in.vUV) + spvDescriptorSet1.vs[5000]->v) + spvDescriptorSet2.uSampler.sample(spvDescriptorSet2.uSamplerSmplr, in.vUV);
|
||||
return out;
|
||||
}
|
||||
|
||||
27
reference/opt/shaders/asm/comp/bitcast_icmp.asm.comp
Normal file
27
reference/opt/shaders/asm/comp/bitcast_icmp.asm.comp
Normal file
@ -0,0 +1,27 @@
|
||||
#version 310 es
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
layout(binding = 0, std430) restrict buffer _3_5
|
||||
{
|
||||
ivec4 _m0;
|
||||
uvec4 _m1;
|
||||
} _5;
|
||||
|
||||
layout(binding = 1, std430) restrict buffer _4_6
|
||||
{
|
||||
uvec4 _m0;
|
||||
ivec4 _m1;
|
||||
} _6;
|
||||
|
||||
void main()
|
||||
{
|
||||
_6._m0 = mix(uvec4(0u), uvec4(1u), lessThan(ivec4(_5._m1), _5._m0));
|
||||
_6._m0 = mix(uvec4(0u), uvec4(1u), lessThanEqual(ivec4(_5._m1), _5._m0));
|
||||
_6._m0 = mix(uvec4(0u), uvec4(1u), lessThan(_5._m1, uvec4(_5._m0)));
|
||||
_6._m0 = mix(uvec4(0u), uvec4(1u), lessThanEqual(_5._m1, uvec4(_5._m0)));
|
||||
_6._m0 = mix(uvec4(0u), uvec4(1u), greaterThan(ivec4(_5._m1), _5._m0));
|
||||
_6._m0 = mix(uvec4(0u), uvec4(1u), greaterThanEqual(ivec4(_5._m1), _5._m0));
|
||||
_6._m0 = mix(uvec4(0u), uvec4(1u), greaterThan(_5._m1, uvec4(_5._m0)));
|
||||
_6._m0 = mix(uvec4(0u), uvec4(1u), greaterThanEqual(_5._m1, uvec4(_5._m0)));
|
||||
}
|
||||
|
||||
@ -0,0 +1,51 @@
|
||||
static float4 A;
|
||||
static float4 B;
|
||||
static float4 FragColor;
|
||||
|
||||
struct SPIRV_Cross_Input
|
||||
{
|
||||
float4 A : TEXCOORD0;
|
||||
float4 B : TEXCOORD1;
|
||||
};
|
||||
|
||||
struct SPIRV_Cross_Output
|
||||
{
|
||||
float4 FragColor : SV_Target0;
|
||||
};
|
||||
|
||||
float4 test_vector()
|
||||
{
|
||||
bool4 le = bool4(!(A.x >= B.x), !(A.y >= B.y), !(A.z >= B.z), !(A.w >= B.w));
|
||||
bool4 leq = bool4(!(A.x > B.x), !(A.y > B.y), !(A.z > B.z), !(A.w > B.w));
|
||||
bool4 ge = bool4(!(A.x <= B.x), !(A.y <= B.y), !(A.z <= B.z), !(A.w <= B.w));
|
||||
bool4 geq = bool4(!(A.x < B.x), !(A.y < B.y), !(A.z < B.z), !(A.w < B.w));
|
||||
bool4 eq = bool4(A.x == B.x, A.y == B.y, A.z == B.z, A.w == B.w);
|
||||
bool4 neq = bool4(A.x != B.x, A.y != B.y, A.z != B.z, A.w != B.w);
|
||||
return ((((float4(le.x ? 1.0f.xxxx.x : 0.0f.xxxx.x, le.y ? 1.0f.xxxx.y : 0.0f.xxxx.y, le.z ? 1.0f.xxxx.z : 0.0f.xxxx.z, le.w ? 1.0f.xxxx.w : 0.0f.xxxx.w) + float4(leq.x ? 1.0f.xxxx.x : 0.0f.xxxx.x, leq.y ? 1.0f.xxxx.y : 0.0f.xxxx.y, leq.z ? 1.0f.xxxx.z : 0.0f.xxxx.z, leq.w ? 1.0f.xxxx.w : 0.0f.xxxx.w)) + float4(ge.x ? 1.0f.xxxx.x : 0.0f.xxxx.x, ge.y ? 1.0f.xxxx.y : 0.0f.xxxx.y, ge.z ? 1.0f.xxxx.z : 0.0f.xxxx.z, ge.w ? 1.0f.xxxx.w : 0.0f.xxxx.w)) + float4(geq.x ? 1.0f.xxxx.x : 0.0f.xxxx.x, geq.y ? 1.0f.xxxx.y : 0.0f.xxxx.y, geq.z ? 1.0f.xxxx.z : 0.0f.xxxx.z, geq.w ? 1.0f.xxxx.w : 0.0f.xxxx.w)) + float4(eq.x ? 1.0f.xxxx.x : 0.0f.xxxx.x, eq.y ? 1.0f.xxxx.y : 0.0f.xxxx.y, eq.z ? 1.0f.xxxx.z : 0.0f.xxxx.z, eq.w ? 1.0f.xxxx.w : 0.0f.xxxx.w)) + float4(neq.x ? 1.0f.xxxx.x : 0.0f.xxxx.x, neq.y ? 1.0f.xxxx.y : 0.0f.xxxx.y, neq.z ? 1.0f.xxxx.z : 0.0f.xxxx.z, neq.w ? 1.0f.xxxx.w : 0.0f.xxxx.w);
|
||||
}
|
||||
|
||||
float test_scalar()
|
||||
{
|
||||
bool le = !(A.x >= B.x);
|
||||
bool leq = !(A.x > B.x);
|
||||
bool ge = !(A.x <= B.x);
|
||||
bool geq = !(A.x < B.x);
|
||||
bool eq = A.x == B.x;
|
||||
bool neq = A.x != B.x;
|
||||
return ((((float(le) + float(leq)) + float(ge)) + float(geq)) + float(eq)) + float(neq);
|
||||
}
|
||||
|
||||
void frag_main()
|
||||
{
|
||||
FragColor = test_vector() + test_scalar().xxxx;
|
||||
}
|
||||
|
||||
SPIRV_Cross_Output main(SPIRV_Cross_Input stage_input)
|
||||
{
|
||||
A = stage_input.A;
|
||||
B = stage_input.B;
|
||||
frag_main();
|
||||
SPIRV_Cross_Output stage_output;
|
||||
stage_output.FragColor = FragColor;
|
||||
return stage_output;
|
||||
}
|
||||
28
reference/shaders-hlsl/asm/comp/bitcast_icmp.asm.comp
Normal file
28
reference/shaders-hlsl/asm/comp/bitcast_icmp.asm.comp
Normal file
@ -0,0 +1,28 @@
|
||||
RWByteAddressBuffer _5 : register(u0);
|
||||
RWByteAddressBuffer _6 : register(u1);
|
||||
|
||||
void comp_main()
|
||||
{
|
||||
bool4 _31 = bool4(int(_5.Load4(16).x) < int4(_5.Load4(0)).x, int(_5.Load4(16).y) < int4(_5.Load4(0)).y, int(_5.Load4(16).z) < int4(_5.Load4(0)).z, int(_5.Load4(16).w) < int4(_5.Load4(0)).w);
|
||||
bool4 _32 = bool4(int(_5.Load4(16).x) <= int4(_5.Load4(0)).x, int(_5.Load4(16).y) <= int4(_5.Load4(0)).y, int(_5.Load4(16).z) <= int4(_5.Load4(0)).z, int(_5.Load4(16).w) <= int4(_5.Load4(0)).w);
|
||||
bool4 _33 = bool4(_5.Load4(16).x < uint(int4(_5.Load4(0)).x), _5.Load4(16).y < uint(int4(_5.Load4(0)).y), _5.Load4(16).z < uint(int4(_5.Load4(0)).z), _5.Load4(16).w < uint(int4(_5.Load4(0)).w));
|
||||
bool4 _34 = bool4(_5.Load4(16).x <= uint(int4(_5.Load4(0)).x), _5.Load4(16).y <= uint(int4(_5.Load4(0)).y), _5.Load4(16).z <= uint(int4(_5.Load4(0)).z), _5.Load4(16).w <= uint(int4(_5.Load4(0)).w));
|
||||
bool4 _35 = bool4(int(_5.Load4(16).x) > int4(_5.Load4(0)).x, int(_5.Load4(16).y) > int4(_5.Load4(0)).y, int(_5.Load4(16).z) > int4(_5.Load4(0)).z, int(_5.Load4(16).w) > int4(_5.Load4(0)).w);
|
||||
bool4 _36 = bool4(int(_5.Load4(16).x) >= int4(_5.Load4(0)).x, int(_5.Load4(16).y) >= int4(_5.Load4(0)).y, int(_5.Load4(16).z) >= int4(_5.Load4(0)).z, int(_5.Load4(16).w) >= int4(_5.Load4(0)).w);
|
||||
bool4 _37 = bool4(_5.Load4(16).x > uint(int4(_5.Load4(0)).x), _5.Load4(16).y > uint(int4(_5.Load4(0)).y), _5.Load4(16).z > uint(int4(_5.Load4(0)).z), _5.Load4(16).w > uint(int4(_5.Load4(0)).w));
|
||||
bool4 _38 = bool4(_5.Load4(16).x >= uint(int4(_5.Load4(0)).x), _5.Load4(16).y >= uint(int4(_5.Load4(0)).y), _5.Load4(16).z >= uint(int4(_5.Load4(0)).z), _5.Load4(16).w >= uint(int4(_5.Load4(0)).w));
|
||||
_6.Store4(0, uint4(_31.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _31.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _31.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _31.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
|
||||
_6.Store4(0, uint4(_32.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _32.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _32.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _32.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
|
||||
_6.Store4(0, uint4(_33.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _33.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _33.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _33.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
|
||||
_6.Store4(0, uint4(_34.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _34.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _34.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _34.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
|
||||
_6.Store4(0, uint4(_35.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _35.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _35.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _35.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
|
||||
_6.Store4(0, uint4(_36.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _36.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _36.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _36.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
|
||||
_6.Store4(0, uint4(_37.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _37.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _37.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _37.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
|
||||
_6.Store4(0, uint4(_38.x ? uint4(1u, 1u, 1u, 1u).x : uint4(0u, 0u, 0u, 0u).x, _38.y ? uint4(1u, 1u, 1u, 1u).y : uint4(0u, 0u, 0u, 0u).y, _38.z ? uint4(1u, 1u, 1u, 1u).z : uint4(0u, 0u, 0u, 0u).z, _38.w ? uint4(1u, 1u, 1u, 1u).w : uint4(0u, 0u, 0u, 0u).w));
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void main()
|
||||
{
|
||||
comp_main();
|
||||
}
|
||||
29
reference/shaders-msl/asm/comp/bitcast_icmp.asm.comp
Normal file
29
reference/shaders-msl/asm/comp/bitcast_icmp.asm.comp
Normal file
@ -0,0 +1,29 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct _3
|
||||
{
|
||||
int4 _m0;
|
||||
uint4 _m1;
|
||||
};
|
||||
|
||||
struct _4
|
||||
{
|
||||
uint4 _m0;
|
||||
int4 _m1;
|
||||
};
|
||||
|
||||
kernel void main0(device _3& restrict _5 [[buffer(0)]], device _4& restrict _6 [[buffer(1)]])
|
||||
{
|
||||
_6._m0 = select(uint4(0u), uint4(1u), int4(_5._m1) < _5._m0);
|
||||
_6._m0 = select(uint4(0u), uint4(1u), int4(_5._m1) <= _5._m0);
|
||||
_6._m0 = select(uint4(0u), uint4(1u), _5._m1 < uint4(_5._m0));
|
||||
_6._m0 = select(uint4(0u), uint4(1u), _5._m1 <= uint4(_5._m0));
|
||||
_6._m0 = select(uint4(0u), uint4(1u), int4(_5._m1) > _5._m0);
|
||||
_6._m0 = select(uint4(0u), uint4(1u), int4(_5._m1) >= _5._m0);
|
||||
_6._m0 = select(uint4(0u), uint4(1u), _5._m1 > uint4(_5._m0));
|
||||
_6._m0 = select(uint4(0u), uint4(1u), _5._m1 >= uint4(_5._m0));
|
||||
}
|
||||
|
||||
@ -0,0 +1,97 @@
|
||||
#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 UBO
|
||||
{
|
||||
float4 v;
|
||||
};
|
||||
|
||||
struct spvDescriptorSetBuffer0
|
||||
{
|
||||
array<texture2d<float>, 10000> uSamplers [[id(0)]];
|
||||
array<sampler, 10000> uSamplersSmplr [[id(10000)]];
|
||||
};
|
||||
|
||||
struct spvDescriptorSetBuffer1
|
||||
{
|
||||
spvUnsafeArray<constant thread UBO*, 10000> vs [[id(0)]];
|
||||
};
|
||||
|
||||
struct spvDescriptorSetBuffer2
|
||||
{
|
||||
texture2d<float> uSampler [[id(0)]];
|
||||
sampler uSamplerSmplr [[id(1)]];
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
float4 samp_array(thread const array<texture2d<float>, 10000> uSamplers, thread const array<sampler, 10000> uSamplersSmplr, thread float2& vUV, const device spvUnsafeArray<constant UBO*, 10000> (&vs))
|
||||
{
|
||||
return uSamplers[9999].sample(uSamplersSmplr[9999], vUV) + vs[5000]->v;
|
||||
}
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
float4 samp_single(thread float2& vUV, thread texture2d<float> uSampler, thread const sampler uSamplerSmplr)
|
||||
{
|
||||
return uSampler.sample(uSamplerSmplr, vUV);
|
||||
}
|
||||
|
||||
fragment main0_out main0(main0_in in [[stage_in]], const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], constant spvDescriptorSetBuffer2& spvDescriptorSet2 [[buffer(2)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
out.FragColor = samp_array(spvDescriptorSet0.uSamplers, spvDescriptorSet0.uSamplersSmplr, in.vUV, spvDescriptorSet1.vs) + samp_single(in.vUV, spvDescriptorSet2.uSampler, spvDescriptorSet2.uSamplerSmplr);
|
||||
return out;
|
||||
}
|
||||
|
||||
33
reference/shaders-no-opt/asm/frag/unordered-compare.asm.frag
Normal file
33
reference/shaders-no-opt/asm/frag/unordered-compare.asm.frag
Normal file
@ -0,0 +1,33 @@
|
||||
#version 450
|
||||
|
||||
layout(location = 0) in vec4 A;
|
||||
layout(location = 1) in vec4 B;
|
||||
layout(location = 0) out vec4 FragColor;
|
||||
|
||||
vec4 test_vector()
|
||||
{
|
||||
bvec4 le = not(greaterThanEqual(A, B));
|
||||
bvec4 leq = not(greaterThan(A, B));
|
||||
bvec4 ge = not(lessThanEqual(A, B));
|
||||
bvec4 geq = not(lessThan(A, B));
|
||||
bvec4 eq = not(notEqual(A, B));
|
||||
bvec4 neq = not(equal(A, B));
|
||||
return ((((mix(vec4(0.0), vec4(1.0), le) + mix(vec4(0.0), vec4(1.0), leq)) + mix(vec4(0.0), vec4(1.0), ge)) + mix(vec4(0.0), vec4(1.0), geq)) + mix(vec4(0.0), vec4(1.0), eq)) + mix(vec4(0.0), vec4(1.0), neq);
|
||||
}
|
||||
|
||||
float test_scalar()
|
||||
{
|
||||
bool le = !(A.x >= B.x);
|
||||
bool leq = !(A.x > B.x);
|
||||
bool ge = !(A.x <= B.x);
|
||||
bool geq = !(A.x < B.x);
|
||||
bool eq = !(A.x != B.x);
|
||||
bool neq = !(A.x == B.x);
|
||||
return ((((float(le) + float(leq)) + float(ge)) + float(geq)) + float(eq)) + float(neq);
|
||||
}
|
||||
|
||||
void main()
|
||||
{
|
||||
FragColor = test_vector() + vec4(test_scalar());
|
||||
}
|
||||
|
||||
@ -0,0 +1,27 @@
|
||||
#version 450
|
||||
|
||||
#ifndef SPIRV_CROSS_CONSTANT_ID_0
|
||||
#define SPIRV_CROSS_CONSTANT_ID_0 0
|
||||
#endif
|
||||
const int omap_r = SPIRV_CROSS_CONSTANT_ID_0;
|
||||
#ifndef SPIRV_CROSS_CONSTANT_ID_1
|
||||
#define SPIRV_CROSS_CONSTANT_ID_1 1
|
||||
#endif
|
||||
const int omap_g = SPIRV_CROSS_CONSTANT_ID_1;
|
||||
#ifndef SPIRV_CROSS_CONSTANT_ID_2
|
||||
#define SPIRV_CROSS_CONSTANT_ID_2 2
|
||||
#endif
|
||||
const int omap_b = SPIRV_CROSS_CONSTANT_ID_2;
|
||||
#ifndef SPIRV_CROSS_CONSTANT_ID_3
|
||||
#define SPIRV_CROSS_CONSTANT_ID_3 3
|
||||
#endif
|
||||
const int omap_a = SPIRV_CROSS_CONSTANT_ID_3;
|
||||
|
||||
layout(location = 0) out vec4 FragColor;
|
||||
layout(location = 0) in vec4 vColor;
|
||||
|
||||
void main()
|
||||
{
|
||||
FragColor = vec4(vColor[omap_r], vColor[omap_g], vColor[omap_b], vColor[omap_a]);
|
||||
}
|
||||
|
||||
27
reference/shaders/asm/comp/bitcast_icmp.asm.comp
Normal file
27
reference/shaders/asm/comp/bitcast_icmp.asm.comp
Normal file
@ -0,0 +1,27 @@
|
||||
#version 310 es
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
layout(binding = 0, std430) restrict buffer _3_5
|
||||
{
|
||||
ivec4 _m0;
|
||||
uvec4 _m1;
|
||||
} _5;
|
||||
|
||||
layout(binding = 1, std430) restrict buffer _4_6
|
||||
{
|
||||
uvec4 _m0;
|
||||
ivec4 _m1;
|
||||
} _6;
|
||||
|
||||
void main()
|
||||
{
|
||||
_6._m0 = mix(uvec4(0u), uvec4(1u), lessThan(ivec4(_5._m1), _5._m0));
|
||||
_6._m0 = mix(uvec4(0u), uvec4(1u), lessThanEqual(ivec4(_5._m1), _5._m0));
|
||||
_6._m0 = mix(uvec4(0u), uvec4(1u), lessThan(_5._m1, uvec4(_5._m0)));
|
||||
_6._m0 = mix(uvec4(0u), uvec4(1u), lessThanEqual(_5._m1, uvec4(_5._m0)));
|
||||
_6._m0 = mix(uvec4(0u), uvec4(1u), greaterThan(ivec4(_5._m1), _5._m0));
|
||||
_6._m0 = mix(uvec4(0u), uvec4(1u), greaterThanEqual(ivec4(_5._m1), _5._m0));
|
||||
_6._m0 = mix(uvec4(0u), uvec4(1u), greaterThan(_5._m1, uvec4(_5._m0)));
|
||||
_6._m0 = mix(uvec4(0u), uvec4(1u), greaterThanEqual(_5._m1, uvec4(_5._m0)));
|
||||
}
|
||||
|
||||
177
shaders-hlsl-no-opt/asm/frag/unordered-compare.asm.frag
Normal file
177
shaders-hlsl-no-opt/asm/frag/unordered-compare.asm.frag
Normal file
@ -0,0 +1,177 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 7
|
||||
; Bound: 132
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Fragment %main "main" %A %B %FragColor
|
||||
OpExecutionMode %main OriginUpperLeft
|
||||
OpSource GLSL 450
|
||||
OpName %main "main"
|
||||
OpName %test_vector_ "test_vector("
|
||||
OpName %test_scalar_ "test_scalar("
|
||||
OpName %le "le"
|
||||
OpName %A "A"
|
||||
OpName %B "B"
|
||||
OpName %leq "leq"
|
||||
OpName %ge "ge"
|
||||
OpName %geq "geq"
|
||||
OpName %eq "eq"
|
||||
OpName %neq "neq"
|
||||
OpName %le_0 "le"
|
||||
OpName %leq_0 "leq"
|
||||
OpName %ge_0 "ge"
|
||||
OpName %geq_0 "geq"
|
||||
OpName %eq_0 "eq"
|
||||
OpName %neq_0 "neq"
|
||||
OpName %FragColor "FragColor"
|
||||
OpDecorate %A Location 0
|
||||
OpDecorate %B Location 1
|
||||
OpDecorate %FragColor Location 0
|
||||
%void = OpTypeVoid
|
||||
%3 = OpTypeFunction %void
|
||||
%float = OpTypeFloat 32
|
||||
%v4float = OpTypeVector %float 4
|
||||
%8 = OpTypeFunction %v4float
|
||||
%11 = OpTypeFunction %float
|
||||
%bool = OpTypeBool
|
||||
%v4bool = OpTypeVector %bool 4
|
||||
%_ptr_Function_v4bool = OpTypePointer Function %v4bool
|
||||
%_ptr_Input_v4float = OpTypePointer Input %v4float
|
||||
%A = OpVariable %_ptr_Input_v4float Input
|
||||
%B = OpVariable %_ptr_Input_v4float Input
|
||||
%float_0 = OpConstant %float 0
|
||||
%float_1 = OpConstant %float 1
|
||||
%47 = OpConstantComposite %v4float %float_0 %float_0 %float_0 %float_0
|
||||
%48 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
|
||||
%_ptr_Function_bool = OpTypePointer Function %bool
|
||||
%uint = OpTypeInt 32 0
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Input_float = OpTypePointer Input %float
|
||||
%_ptr_Output_v4float = OpTypePointer Output %v4float
|
||||
%FragColor = OpVariable %_ptr_Output_v4float Output
|
||||
%main = OpFunction %void None %3
|
||||
%5 = OpLabel
|
||||
%128 = OpFunctionCall %v4float %test_vector_
|
||||
%129 = OpFunctionCall %float %test_scalar_
|
||||
%130 = OpCompositeConstruct %v4float %129 %129 %129 %129
|
||||
%131 = OpFAdd %v4float %128 %130
|
||||
OpStore %FragColor %131
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%test_vector_ = OpFunction %v4float None %8
|
||||
%10 = OpLabel
|
||||
%le = OpVariable %_ptr_Function_v4bool Function
|
||||
%leq = OpVariable %_ptr_Function_v4bool Function
|
||||
%ge = OpVariable %_ptr_Function_v4bool Function
|
||||
%geq = OpVariable %_ptr_Function_v4bool Function
|
||||
%eq = OpVariable %_ptr_Function_v4bool Function
|
||||
%neq = OpVariable %_ptr_Function_v4bool Function
|
||||
%20 = OpLoad %v4float %A
|
||||
%22 = OpLoad %v4float %B
|
||||
%23 = OpFUnordLessThan %v4bool %20 %22
|
||||
OpStore %le %23
|
||||
%25 = OpLoad %v4float %A
|
||||
%26 = OpLoad %v4float %B
|
||||
%27 = OpFUnordLessThanEqual %v4bool %25 %26
|
||||
OpStore %leq %27
|
||||
%29 = OpLoad %v4float %A
|
||||
%30 = OpLoad %v4float %B
|
||||
%31 = OpFUnordGreaterThan %v4bool %29 %30
|
||||
OpStore %ge %31
|
||||
%33 = OpLoad %v4float %A
|
||||
%34 = OpLoad %v4float %B
|
||||
%35 = OpFUnordGreaterThanEqual %v4bool %33 %34
|
||||
OpStore %geq %35
|
||||
%37 = OpLoad %v4float %A
|
||||
%38 = OpLoad %v4float %B
|
||||
%39 = OpFUnordEqual %v4bool %37 %38
|
||||
OpStore %eq %39
|
||||
%41 = OpLoad %v4float %A
|
||||
%42 = OpLoad %v4float %B
|
||||
%43 = OpFUnordNotEqual %v4bool %41 %42
|
||||
OpStore %neq %43
|
||||
%44 = OpLoad %v4bool %le
|
||||
%49 = OpSelect %v4float %44 %48 %47
|
||||
%50 = OpLoad %v4bool %leq
|
||||
%51 = OpSelect %v4float %50 %48 %47
|
||||
%52 = OpFAdd %v4float %49 %51
|
||||
%53 = OpLoad %v4bool %ge
|
||||
%54 = OpSelect %v4float %53 %48 %47
|
||||
%55 = OpFAdd %v4float %52 %54
|
||||
%56 = OpLoad %v4bool %geq
|
||||
%57 = OpSelect %v4float %56 %48 %47
|
||||
%58 = OpFAdd %v4float %55 %57
|
||||
%59 = OpLoad %v4bool %eq
|
||||
%60 = OpSelect %v4float %59 %48 %47
|
||||
%61 = OpFAdd %v4float %58 %60
|
||||
%62 = OpLoad %v4bool %neq
|
||||
%63 = OpSelect %v4float %62 %48 %47
|
||||
%64 = OpFAdd %v4float %61 %63
|
||||
OpReturnValue %64
|
||||
OpFunctionEnd
|
||||
%test_scalar_ = OpFunction %float None %11
|
||||
%13 = OpLabel
|
||||
%le_0 = OpVariable %_ptr_Function_bool Function
|
||||
%leq_0 = OpVariable %_ptr_Function_bool Function
|
||||
%ge_0 = OpVariable %_ptr_Function_bool Function
|
||||
%geq_0 = OpVariable %_ptr_Function_bool Function
|
||||
%eq_0 = OpVariable %_ptr_Function_bool Function
|
||||
%neq_0 = OpVariable %_ptr_Function_bool Function
|
||||
%72 = OpAccessChain %_ptr_Input_float %A %uint_0
|
||||
%73 = OpLoad %float %72
|
||||
%74 = OpAccessChain %_ptr_Input_float %B %uint_0
|
||||
%75 = OpLoad %float %74
|
||||
%76 = OpFUnordLessThan %bool %73 %75
|
||||
OpStore %le_0 %76
|
||||
%78 = OpAccessChain %_ptr_Input_float %A %uint_0
|
||||
%79 = OpLoad %float %78
|
||||
%80 = OpAccessChain %_ptr_Input_float %B %uint_0
|
||||
%81 = OpLoad %float %80
|
||||
%82 = OpFUnordLessThanEqual %bool %79 %81
|
||||
OpStore %leq_0 %82
|
||||
%84 = OpAccessChain %_ptr_Input_float %A %uint_0
|
||||
%85 = OpLoad %float %84
|
||||
%86 = OpAccessChain %_ptr_Input_float %B %uint_0
|
||||
%87 = OpLoad %float %86
|
||||
%88 = OpFUnordGreaterThan %bool %85 %87
|
||||
OpStore %ge_0 %88
|
||||
%90 = OpAccessChain %_ptr_Input_float %A %uint_0
|
||||
%91 = OpLoad %float %90
|
||||
%92 = OpAccessChain %_ptr_Input_float %B %uint_0
|
||||
%93 = OpLoad %float %92
|
||||
%94 = OpFUnordGreaterThanEqual %bool %91 %93
|
||||
OpStore %geq_0 %94
|
||||
%96 = OpAccessChain %_ptr_Input_float %A %uint_0
|
||||
%97 = OpLoad %float %96
|
||||
%98 = OpAccessChain %_ptr_Input_float %B %uint_0
|
||||
%99 = OpLoad %float %98
|
||||
%100 = OpFUnordEqual %bool %97 %99
|
||||
OpStore %eq_0 %100
|
||||
%102 = OpAccessChain %_ptr_Input_float %A %uint_0
|
||||
%103 = OpLoad %float %102
|
||||
%104 = OpAccessChain %_ptr_Input_float %B %uint_0
|
||||
%105 = OpLoad %float %104
|
||||
%106 = OpFUnordNotEqual %bool %103 %105
|
||||
OpStore %neq_0 %106
|
||||
%107 = OpLoad %bool %le_0
|
||||
%108 = OpSelect %float %107 %float_1 %float_0
|
||||
%109 = OpLoad %bool %leq_0
|
||||
%110 = OpSelect %float %109 %float_1 %float_0
|
||||
%111 = OpFAdd %float %108 %110
|
||||
%112 = OpLoad %bool %ge_0
|
||||
%113 = OpSelect %float %112 %float_1 %float_0
|
||||
%114 = OpFAdd %float %111 %113
|
||||
%115 = OpLoad %bool %geq_0
|
||||
%116 = OpSelect %float %115 %float_1 %float_0
|
||||
%117 = OpFAdd %float %114 %116
|
||||
%118 = OpLoad %bool %eq_0
|
||||
%119 = OpSelect %float %118 %float_1 %float_0
|
||||
%120 = OpFAdd %float %117 %119
|
||||
%121 = OpLoad %bool %neq_0
|
||||
%122 = OpSelect %float %121 %float_1 %float_0
|
||||
%123 = OpFAdd %float %120 %122
|
||||
OpReturnValue %123
|
||||
OpFunctionEnd
|
||||
101
shaders-hlsl/asm/comp/bitcast_icmp.asm.comp
Normal file
101
shaders-hlsl/asm/comp/bitcast_icmp.asm.comp
Normal file
@ -0,0 +1,101 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 1
|
||||
; Bound: 30
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %func "main"
|
||||
OpExecutionMode %func LocalSize 1 1 1
|
||||
OpSource ESSL 310
|
||||
OpSourceExtension "GL_GOOGLE_cpp_style_line_directive"
|
||||
OpSourceExtension "GL_GOOGLE_include_directive"
|
||||
OpMemberDecorate %input_struct 0 Offset 0
|
||||
OpMemberDecorate %input_struct 1 Offset 16
|
||||
OpMemberDecorate %output_struct 0 Offset 0
|
||||
OpMemberDecorate %output_struct 1 Offset 16
|
||||
OpDecorate %input_struct BufferBlock
|
||||
OpDecorate %inputs DescriptorSet 0
|
||||
OpDecorate %inputs Binding 0
|
||||
OpDecorate %inputs Restrict
|
||||
OpDecorate %output_struct BufferBlock
|
||||
OpDecorate %outputs DescriptorSet 0
|
||||
OpDecorate %outputs Binding 1
|
||||
OpDecorate %outputs Restrict
|
||||
|
||||
%void = OpTypeVoid
|
||||
%main_func = OpTypeFunction %void
|
||||
|
||||
%bool = OpTypeBool
|
||||
%bvec4 = OpTypeVector %bool 4
|
||||
|
||||
%uint = OpTypeInt 32 0
|
||||
%uvec4 = OpTypeVector %uint 4
|
||||
|
||||
%int = OpTypeInt 32 1
|
||||
%ivec4 = OpTypeVector %int 4
|
||||
|
||||
%ivec4_ptr = OpTypePointer Uniform %ivec4
|
||||
%uvec4_ptr = OpTypePointer Uniform %uvec4
|
||||
|
||||
%zero = OpConstant %int 0
|
||||
%one = OpConstant %int 1
|
||||
%uzero = OpConstant %uint 0
|
||||
%uone = OpConstant %uint 1
|
||||
%utrue = OpConstantComposite %uvec4 %uone %uone %uone %uone
|
||||
%ufalse = OpConstantComposite %uvec4 %uzero %uzero %uzero %uzero
|
||||
|
||||
%input_struct = OpTypeStruct %ivec4 %uvec4
|
||||
%input_struct_ptr = OpTypePointer Uniform %input_struct
|
||||
%inputs = OpVariable %input_struct_ptr Uniform
|
||||
%output_struct = OpTypeStruct %uvec4 %ivec4
|
||||
%output_struct_ptr = OpTypePointer Uniform %output_struct
|
||||
%outputs = OpVariable %output_struct_ptr Uniform
|
||||
|
||||
%func = OpFunction %void None %main_func
|
||||
%block = OpLabel
|
||||
|
||||
%input1_ptr = OpAccessChain %ivec4_ptr %inputs %zero
|
||||
%input0_ptr = OpAccessChain %uvec4_ptr %inputs %one
|
||||
%input1 = OpLoad %ivec4 %input1_ptr
|
||||
%input0 = OpLoad %uvec4 %input0_ptr
|
||||
|
||||
%output_ptr_uvec4 = OpAccessChain %uvec4_ptr %outputs %zero
|
||||
|
||||
%result_slt = OpSLessThan %bvec4 %input0 %input1
|
||||
%result_sle = OpSLessThanEqual %bvec4 %input0 %input1
|
||||
%result_ult = OpULessThan %bvec4 %input0 %input1
|
||||
%result_ule = OpULessThanEqual %bvec4 %input0 %input1
|
||||
%result_sgt = OpSGreaterThan %bvec4 %input0 %input1
|
||||
%result_sge = OpSGreaterThanEqual %bvec4 %input0 %input1
|
||||
%result_ugt = OpUGreaterThan %bvec4 %input0 %input1
|
||||
%result_uge = OpUGreaterThanEqual %bvec4 %input0 %input1
|
||||
|
||||
%int_slt = OpSelect %uvec4 %result_slt %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_slt
|
||||
|
||||
%int_sle = OpSelect %uvec4 %result_sle %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_sle
|
||||
|
||||
%int_ult = OpSelect %uvec4 %result_ult %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_ult
|
||||
|
||||
%int_ule = OpSelect %uvec4 %result_ule %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_ule
|
||||
|
||||
%int_sgt = OpSelect %uvec4 %result_sgt %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_sgt
|
||||
|
||||
%int_sge = OpSelect %uvec4 %result_sge %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_sge
|
||||
|
||||
%int_ugt = OpSelect %uvec4 %result_ugt %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_ugt
|
||||
|
||||
%int_uge = OpSelect %uvec4 %result_uge %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_uge
|
||||
|
||||
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
101
shaders-msl/asm/comp/bitcast_icmp.asm.comp
Normal file
101
shaders-msl/asm/comp/bitcast_icmp.asm.comp
Normal file
@ -0,0 +1,101 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 1
|
||||
; Bound: 30
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %func "main"
|
||||
OpExecutionMode %func LocalSize 1 1 1
|
||||
OpSource ESSL 310
|
||||
OpSourceExtension "GL_GOOGLE_cpp_style_line_directive"
|
||||
OpSourceExtension "GL_GOOGLE_include_directive"
|
||||
OpMemberDecorate %input_struct 0 Offset 0
|
||||
OpMemberDecorate %input_struct 1 Offset 16
|
||||
OpMemberDecorate %output_struct 0 Offset 0
|
||||
OpMemberDecorate %output_struct 1 Offset 16
|
||||
OpDecorate %input_struct BufferBlock
|
||||
OpDecorate %inputs DescriptorSet 0
|
||||
OpDecorate %inputs Binding 0
|
||||
OpDecorate %inputs Restrict
|
||||
OpDecorate %output_struct BufferBlock
|
||||
OpDecorate %outputs DescriptorSet 0
|
||||
OpDecorate %outputs Binding 1
|
||||
OpDecorate %outputs Restrict
|
||||
|
||||
%void = OpTypeVoid
|
||||
%main_func = OpTypeFunction %void
|
||||
|
||||
%bool = OpTypeBool
|
||||
%bvec4 = OpTypeVector %bool 4
|
||||
|
||||
%uint = OpTypeInt 32 0
|
||||
%uvec4 = OpTypeVector %uint 4
|
||||
|
||||
%int = OpTypeInt 32 1
|
||||
%ivec4 = OpTypeVector %int 4
|
||||
|
||||
%ivec4_ptr = OpTypePointer Uniform %ivec4
|
||||
%uvec4_ptr = OpTypePointer Uniform %uvec4
|
||||
|
||||
%zero = OpConstant %int 0
|
||||
%one = OpConstant %int 1
|
||||
%uzero = OpConstant %uint 0
|
||||
%uone = OpConstant %uint 1
|
||||
%utrue = OpConstantComposite %uvec4 %uone %uone %uone %uone
|
||||
%ufalse = OpConstantComposite %uvec4 %uzero %uzero %uzero %uzero
|
||||
|
||||
%input_struct = OpTypeStruct %ivec4 %uvec4
|
||||
%input_struct_ptr = OpTypePointer Uniform %input_struct
|
||||
%inputs = OpVariable %input_struct_ptr Uniform
|
||||
%output_struct = OpTypeStruct %uvec4 %ivec4
|
||||
%output_struct_ptr = OpTypePointer Uniform %output_struct
|
||||
%outputs = OpVariable %output_struct_ptr Uniform
|
||||
|
||||
%func = OpFunction %void None %main_func
|
||||
%block = OpLabel
|
||||
|
||||
%input1_ptr = OpAccessChain %ivec4_ptr %inputs %zero
|
||||
%input0_ptr = OpAccessChain %uvec4_ptr %inputs %one
|
||||
%input1 = OpLoad %ivec4 %input1_ptr
|
||||
%input0 = OpLoad %uvec4 %input0_ptr
|
||||
|
||||
%output_ptr_uvec4 = OpAccessChain %uvec4_ptr %outputs %zero
|
||||
|
||||
%result_slt = OpSLessThan %bvec4 %input0 %input1
|
||||
%result_sle = OpSLessThanEqual %bvec4 %input0 %input1
|
||||
%result_ult = OpULessThan %bvec4 %input0 %input1
|
||||
%result_ule = OpULessThanEqual %bvec4 %input0 %input1
|
||||
%result_sgt = OpSGreaterThan %bvec4 %input0 %input1
|
||||
%result_sge = OpSGreaterThanEqual %bvec4 %input0 %input1
|
||||
%result_ugt = OpUGreaterThan %bvec4 %input0 %input1
|
||||
%result_uge = OpUGreaterThanEqual %bvec4 %input0 %input1
|
||||
|
||||
%int_slt = OpSelect %uvec4 %result_slt %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_slt
|
||||
|
||||
%int_sle = OpSelect %uvec4 %result_sle %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_sle
|
||||
|
||||
%int_ult = OpSelect %uvec4 %result_ult %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_ult
|
||||
|
||||
%int_ule = OpSelect %uvec4 %result_ule %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_ule
|
||||
|
||||
%int_sgt = OpSelect %uvec4 %result_sgt %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_sgt
|
||||
|
||||
%int_sge = OpSelect %uvec4 %result_sge %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_sge
|
||||
|
||||
%int_ugt = OpSelect %uvec4 %result_ugt %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_ugt
|
||||
|
||||
%int_uge = OpSelect %uvec4 %result_uge %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_uge
|
||||
|
||||
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -0,0 +1,26 @@
|
||||
#version 450
|
||||
|
||||
layout(location = 0) out vec4 FragColor;
|
||||
layout(location = 0) in vec2 vUV;
|
||||
layout(set = 0, binding = 0) uniform sampler2D uSamplers[10000];
|
||||
layout(set = 2, binding = 0) uniform sampler2D uSampler;
|
||||
|
||||
layout(set = 1, binding = 0) uniform UBO
|
||||
{
|
||||
vec4 v;
|
||||
} vs[10000];
|
||||
|
||||
vec4 samp_array()
|
||||
{
|
||||
return texture(uSamplers[9999], vUV) + vs[5000].v;
|
||||
}
|
||||
|
||||
vec4 samp_single()
|
||||
{
|
||||
return texture(uSampler, vUV);
|
||||
}
|
||||
|
||||
void main()
|
||||
{
|
||||
FragColor = samp_array() + samp_single();
|
||||
}
|
||||
177
shaders-no-opt/asm/frag/unordered-compare.asm.frag
Normal file
177
shaders-no-opt/asm/frag/unordered-compare.asm.frag
Normal file
@ -0,0 +1,177 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 7
|
||||
; Bound: 132
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Fragment %main "main" %A %B %FragColor
|
||||
OpExecutionMode %main OriginUpperLeft
|
||||
OpSource GLSL 450
|
||||
OpName %main "main"
|
||||
OpName %test_vector_ "test_vector("
|
||||
OpName %test_scalar_ "test_scalar("
|
||||
OpName %le "le"
|
||||
OpName %A "A"
|
||||
OpName %B "B"
|
||||
OpName %leq "leq"
|
||||
OpName %ge "ge"
|
||||
OpName %geq "geq"
|
||||
OpName %eq "eq"
|
||||
OpName %neq "neq"
|
||||
OpName %le_0 "le"
|
||||
OpName %leq_0 "leq"
|
||||
OpName %ge_0 "ge"
|
||||
OpName %geq_0 "geq"
|
||||
OpName %eq_0 "eq"
|
||||
OpName %neq_0 "neq"
|
||||
OpName %FragColor "FragColor"
|
||||
OpDecorate %A Location 0
|
||||
OpDecorate %B Location 1
|
||||
OpDecorate %FragColor Location 0
|
||||
%void = OpTypeVoid
|
||||
%3 = OpTypeFunction %void
|
||||
%float = OpTypeFloat 32
|
||||
%v4float = OpTypeVector %float 4
|
||||
%8 = OpTypeFunction %v4float
|
||||
%11 = OpTypeFunction %float
|
||||
%bool = OpTypeBool
|
||||
%v4bool = OpTypeVector %bool 4
|
||||
%_ptr_Function_v4bool = OpTypePointer Function %v4bool
|
||||
%_ptr_Input_v4float = OpTypePointer Input %v4float
|
||||
%A = OpVariable %_ptr_Input_v4float Input
|
||||
%B = OpVariable %_ptr_Input_v4float Input
|
||||
%float_0 = OpConstant %float 0
|
||||
%float_1 = OpConstant %float 1
|
||||
%47 = OpConstantComposite %v4float %float_0 %float_0 %float_0 %float_0
|
||||
%48 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
|
||||
%_ptr_Function_bool = OpTypePointer Function %bool
|
||||
%uint = OpTypeInt 32 0
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Input_float = OpTypePointer Input %float
|
||||
%_ptr_Output_v4float = OpTypePointer Output %v4float
|
||||
%FragColor = OpVariable %_ptr_Output_v4float Output
|
||||
%main = OpFunction %void None %3
|
||||
%5 = OpLabel
|
||||
%128 = OpFunctionCall %v4float %test_vector_
|
||||
%129 = OpFunctionCall %float %test_scalar_
|
||||
%130 = OpCompositeConstruct %v4float %129 %129 %129 %129
|
||||
%131 = OpFAdd %v4float %128 %130
|
||||
OpStore %FragColor %131
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%test_vector_ = OpFunction %v4float None %8
|
||||
%10 = OpLabel
|
||||
%le = OpVariable %_ptr_Function_v4bool Function
|
||||
%leq = OpVariable %_ptr_Function_v4bool Function
|
||||
%ge = OpVariable %_ptr_Function_v4bool Function
|
||||
%geq = OpVariable %_ptr_Function_v4bool Function
|
||||
%eq = OpVariable %_ptr_Function_v4bool Function
|
||||
%neq = OpVariable %_ptr_Function_v4bool Function
|
||||
%20 = OpLoad %v4float %A
|
||||
%22 = OpLoad %v4float %B
|
||||
%23 = OpFUnordLessThan %v4bool %20 %22
|
||||
OpStore %le %23
|
||||
%25 = OpLoad %v4float %A
|
||||
%26 = OpLoad %v4float %B
|
||||
%27 = OpFUnordLessThanEqual %v4bool %25 %26
|
||||
OpStore %leq %27
|
||||
%29 = OpLoad %v4float %A
|
||||
%30 = OpLoad %v4float %B
|
||||
%31 = OpFUnordGreaterThan %v4bool %29 %30
|
||||
OpStore %ge %31
|
||||
%33 = OpLoad %v4float %A
|
||||
%34 = OpLoad %v4float %B
|
||||
%35 = OpFUnordGreaterThanEqual %v4bool %33 %34
|
||||
OpStore %geq %35
|
||||
%37 = OpLoad %v4float %A
|
||||
%38 = OpLoad %v4float %B
|
||||
%39 = OpFUnordEqual %v4bool %37 %38
|
||||
OpStore %eq %39
|
||||
%41 = OpLoad %v4float %A
|
||||
%42 = OpLoad %v4float %B
|
||||
%43 = OpFUnordNotEqual %v4bool %41 %42
|
||||
OpStore %neq %43
|
||||
%44 = OpLoad %v4bool %le
|
||||
%49 = OpSelect %v4float %44 %48 %47
|
||||
%50 = OpLoad %v4bool %leq
|
||||
%51 = OpSelect %v4float %50 %48 %47
|
||||
%52 = OpFAdd %v4float %49 %51
|
||||
%53 = OpLoad %v4bool %ge
|
||||
%54 = OpSelect %v4float %53 %48 %47
|
||||
%55 = OpFAdd %v4float %52 %54
|
||||
%56 = OpLoad %v4bool %geq
|
||||
%57 = OpSelect %v4float %56 %48 %47
|
||||
%58 = OpFAdd %v4float %55 %57
|
||||
%59 = OpLoad %v4bool %eq
|
||||
%60 = OpSelect %v4float %59 %48 %47
|
||||
%61 = OpFAdd %v4float %58 %60
|
||||
%62 = OpLoad %v4bool %neq
|
||||
%63 = OpSelect %v4float %62 %48 %47
|
||||
%64 = OpFAdd %v4float %61 %63
|
||||
OpReturnValue %64
|
||||
OpFunctionEnd
|
||||
%test_scalar_ = OpFunction %float None %11
|
||||
%13 = OpLabel
|
||||
%le_0 = OpVariable %_ptr_Function_bool Function
|
||||
%leq_0 = OpVariable %_ptr_Function_bool Function
|
||||
%ge_0 = OpVariable %_ptr_Function_bool Function
|
||||
%geq_0 = OpVariable %_ptr_Function_bool Function
|
||||
%eq_0 = OpVariable %_ptr_Function_bool Function
|
||||
%neq_0 = OpVariable %_ptr_Function_bool Function
|
||||
%72 = OpAccessChain %_ptr_Input_float %A %uint_0
|
||||
%73 = OpLoad %float %72
|
||||
%74 = OpAccessChain %_ptr_Input_float %B %uint_0
|
||||
%75 = OpLoad %float %74
|
||||
%76 = OpFUnordLessThan %bool %73 %75
|
||||
OpStore %le_0 %76
|
||||
%78 = OpAccessChain %_ptr_Input_float %A %uint_0
|
||||
%79 = OpLoad %float %78
|
||||
%80 = OpAccessChain %_ptr_Input_float %B %uint_0
|
||||
%81 = OpLoad %float %80
|
||||
%82 = OpFUnordLessThanEqual %bool %79 %81
|
||||
OpStore %leq_0 %82
|
||||
%84 = OpAccessChain %_ptr_Input_float %A %uint_0
|
||||
%85 = OpLoad %float %84
|
||||
%86 = OpAccessChain %_ptr_Input_float %B %uint_0
|
||||
%87 = OpLoad %float %86
|
||||
%88 = OpFUnordGreaterThan %bool %85 %87
|
||||
OpStore %ge_0 %88
|
||||
%90 = OpAccessChain %_ptr_Input_float %A %uint_0
|
||||
%91 = OpLoad %float %90
|
||||
%92 = OpAccessChain %_ptr_Input_float %B %uint_0
|
||||
%93 = OpLoad %float %92
|
||||
%94 = OpFUnordGreaterThanEqual %bool %91 %93
|
||||
OpStore %geq_0 %94
|
||||
%96 = OpAccessChain %_ptr_Input_float %A %uint_0
|
||||
%97 = OpLoad %float %96
|
||||
%98 = OpAccessChain %_ptr_Input_float %B %uint_0
|
||||
%99 = OpLoad %float %98
|
||||
%100 = OpFUnordEqual %bool %97 %99
|
||||
OpStore %eq_0 %100
|
||||
%102 = OpAccessChain %_ptr_Input_float %A %uint_0
|
||||
%103 = OpLoad %float %102
|
||||
%104 = OpAccessChain %_ptr_Input_float %B %uint_0
|
||||
%105 = OpLoad %float %104
|
||||
%106 = OpFUnordNotEqual %bool %103 %105
|
||||
OpStore %neq_0 %106
|
||||
%107 = OpLoad %bool %le_0
|
||||
%108 = OpSelect %float %107 %float_1 %float_0
|
||||
%109 = OpLoad %bool %leq_0
|
||||
%110 = OpSelect %float %109 %float_1 %float_0
|
||||
%111 = OpFAdd %float %108 %110
|
||||
%112 = OpLoad %bool %ge_0
|
||||
%113 = OpSelect %float %112 %float_1 %float_0
|
||||
%114 = OpFAdd %float %111 %113
|
||||
%115 = OpLoad %bool %geq_0
|
||||
%116 = OpSelect %float %115 %float_1 %float_0
|
||||
%117 = OpFAdd %float %114 %116
|
||||
%118 = OpLoad %bool %eq_0
|
||||
%119 = OpSelect %float %118 %float_1 %float_0
|
||||
%120 = OpFAdd %float %117 %119
|
||||
%121 = OpLoad %bool %neq_0
|
||||
%122 = OpSelect %float %121 %float_1 %float_0
|
||||
%123 = OpFAdd %float %120 %122
|
||||
OpReturnValue %123
|
||||
OpFunctionEnd
|
||||
@ -0,0 +1,49 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 7
|
||||
; Bound: 27
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Fragment %main "main" %FragColor %vColor
|
||||
OpExecutionMode %main OriginUpperLeft
|
||||
OpSource GLSL 450
|
||||
OpName %main "main"
|
||||
OpName %FragColor "FragColor"
|
||||
OpName %vColor "vColor"
|
||||
OpName %omap_r "omap_r"
|
||||
OpName %omap_g "omap_g"
|
||||
OpName %omap_b "omap_b"
|
||||
OpName %omap_a "omap_a"
|
||||
OpDecorate %FragColor Location 0
|
||||
OpDecorate %vColor Location 0
|
||||
OpDecorate %omap_r SpecId 0
|
||||
OpDecorate %omap_g SpecId 1
|
||||
OpDecorate %omap_b SpecId 2
|
||||
OpDecorate %omap_a SpecId 3
|
||||
%void = OpTypeVoid
|
||||
%3 = OpTypeFunction %void
|
||||
%float = OpTypeFloat 32
|
||||
%v4float = OpTypeVector %float 4
|
||||
%_ptr_Output_v4float = OpTypePointer Output %v4float
|
||||
%FragColor = OpVariable %_ptr_Output_v4float Output
|
||||
%_ptr_Input_v4float = OpTypePointer Input %v4float
|
||||
%vColor = OpVariable %_ptr_Input_v4float Input
|
||||
%int = OpTypeInt 32 1
|
||||
%omap_r = OpSpecConstant %int 0
|
||||
%_ptr_Input_float = OpTypePointer Input %float
|
||||
%omap_g = OpSpecConstant %int 1
|
||||
%omap_b = OpSpecConstant %int 2
|
||||
%omap_a = OpSpecConstant %int 3
|
||||
%main = OpFunction %void None %3
|
||||
%5 = OpLabel
|
||||
%loaded = OpLoad %v4float %vColor
|
||||
%r = OpVectorExtractDynamic %float %loaded %omap_r
|
||||
%g = OpVectorExtractDynamic %float %loaded %omap_g
|
||||
%b = OpVectorExtractDynamic %float %loaded %omap_b
|
||||
%a = OpVectorExtractDynamic %float %loaded %omap_a
|
||||
%rgba = OpCompositeConstruct %v4float %r %g %b %a
|
||||
OpStore %FragColor %rgba
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
101
shaders/asm/comp/bitcast_icmp.asm.comp
Normal file
101
shaders/asm/comp/bitcast_icmp.asm.comp
Normal file
@ -0,0 +1,101 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 1
|
||||
; Bound: 30
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %func "main"
|
||||
OpExecutionMode %func LocalSize 1 1 1
|
||||
OpSource ESSL 310
|
||||
OpSourceExtension "GL_GOOGLE_cpp_style_line_directive"
|
||||
OpSourceExtension "GL_GOOGLE_include_directive"
|
||||
OpMemberDecorate %input_struct 0 Offset 0
|
||||
OpMemberDecorate %input_struct 1 Offset 16
|
||||
OpMemberDecorate %output_struct 0 Offset 0
|
||||
OpMemberDecorate %output_struct 1 Offset 16
|
||||
OpDecorate %input_struct BufferBlock
|
||||
OpDecorate %inputs DescriptorSet 0
|
||||
OpDecorate %inputs Binding 0
|
||||
OpDecorate %inputs Restrict
|
||||
OpDecorate %output_struct BufferBlock
|
||||
OpDecorate %outputs DescriptorSet 0
|
||||
OpDecorate %outputs Binding 1
|
||||
OpDecorate %outputs Restrict
|
||||
|
||||
%void = OpTypeVoid
|
||||
%main_func = OpTypeFunction %void
|
||||
|
||||
%bool = OpTypeBool
|
||||
%bvec4 = OpTypeVector %bool 4
|
||||
|
||||
%uint = OpTypeInt 32 0
|
||||
%uvec4 = OpTypeVector %uint 4
|
||||
|
||||
%int = OpTypeInt 32 1
|
||||
%ivec4 = OpTypeVector %int 4
|
||||
|
||||
%ivec4_ptr = OpTypePointer Uniform %ivec4
|
||||
%uvec4_ptr = OpTypePointer Uniform %uvec4
|
||||
|
||||
%zero = OpConstant %int 0
|
||||
%one = OpConstant %int 1
|
||||
%uzero = OpConstant %uint 0
|
||||
%uone = OpConstant %uint 1
|
||||
%utrue = OpConstantComposite %uvec4 %uone %uone %uone %uone
|
||||
%ufalse = OpConstantComposite %uvec4 %uzero %uzero %uzero %uzero
|
||||
|
||||
%input_struct = OpTypeStruct %ivec4 %uvec4
|
||||
%input_struct_ptr = OpTypePointer Uniform %input_struct
|
||||
%inputs = OpVariable %input_struct_ptr Uniform
|
||||
%output_struct = OpTypeStruct %uvec4 %ivec4
|
||||
%output_struct_ptr = OpTypePointer Uniform %output_struct
|
||||
%outputs = OpVariable %output_struct_ptr Uniform
|
||||
|
||||
%func = OpFunction %void None %main_func
|
||||
%block = OpLabel
|
||||
|
||||
%input1_ptr = OpAccessChain %ivec4_ptr %inputs %zero
|
||||
%input0_ptr = OpAccessChain %uvec4_ptr %inputs %one
|
||||
%input1 = OpLoad %ivec4 %input1_ptr
|
||||
%input0 = OpLoad %uvec4 %input0_ptr
|
||||
|
||||
%output_ptr_uvec4 = OpAccessChain %uvec4_ptr %outputs %zero
|
||||
|
||||
%result_slt = OpSLessThan %bvec4 %input0 %input1
|
||||
%result_sle = OpSLessThanEqual %bvec4 %input0 %input1
|
||||
%result_ult = OpULessThan %bvec4 %input0 %input1
|
||||
%result_ule = OpULessThanEqual %bvec4 %input0 %input1
|
||||
%result_sgt = OpSGreaterThan %bvec4 %input0 %input1
|
||||
%result_sge = OpSGreaterThanEqual %bvec4 %input0 %input1
|
||||
%result_ugt = OpUGreaterThan %bvec4 %input0 %input1
|
||||
%result_uge = OpUGreaterThanEqual %bvec4 %input0 %input1
|
||||
|
||||
%int_slt = OpSelect %uvec4 %result_slt %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_slt
|
||||
|
||||
%int_sle = OpSelect %uvec4 %result_sle %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_sle
|
||||
|
||||
%int_ult = OpSelect %uvec4 %result_ult %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_ult
|
||||
|
||||
%int_ule = OpSelect %uvec4 %result_ule %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_ule
|
||||
|
||||
%int_sgt = OpSelect %uvec4 %result_sgt %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_sgt
|
||||
|
||||
%int_sge = OpSelect %uvec4 %result_sge %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_sge
|
||||
|
||||
%int_ugt = OpSelect %uvec4 %result_ugt %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_ugt
|
||||
|
||||
%int_uge = OpSelect %uvec4 %result_uge %utrue %ufalse
|
||||
OpStore %output_ptr_uvec4 %int_uge
|
||||
|
||||
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -61,7 +61,7 @@ void CFG::build_immediate_dominators()
|
||||
if (immediate_dominators[block])
|
||||
{
|
||||
assert(immediate_dominators[edge]);
|
||||
immediate_dominators[block] = find_common_dominator(block, edge);
|
||||
immediate_dominators[block] = find_common_dominator(immediate_dominators[block], edge);
|
||||
}
|
||||
else
|
||||
immediate_dominators[block] = edge;
|
||||
|
||||
@ -946,6 +946,26 @@ spvc_result spvc_compiler_msl_add_discrete_descriptor_set(spvc_compiler compiler
|
||||
#endif
|
||||
}
|
||||
|
||||
spvc_result spvc_compiler_msl_set_argument_buffer_device_address_space(spvc_compiler compiler, unsigned desc_set, spvc_bool device_address)
|
||||
{
|
||||
#if SPIRV_CROSS_C_API_MSL
|
||||
if (compiler->backend != SPVC_BACKEND_MSL)
|
||||
{
|
||||
compiler->context->report_error("MSL function used on a non-MSL backend.");
|
||||
return SPVC_ERROR_INVALID_ARGUMENT;
|
||||
}
|
||||
|
||||
auto &msl = *static_cast<CompilerMSL *>(compiler->compiler.get());
|
||||
msl.set_argument_buffer_device_address_space(desc_set, bool(device_address));
|
||||
return SPVC_SUCCESS;
|
||||
#else
|
||||
(void)desc_set;
|
||||
(void)device_address;
|
||||
compiler->context->report_error("MSL function used on a non-MSL backend.");
|
||||
return SPVC_ERROR_INVALID_ARGUMENT;
|
||||
#endif
|
||||
}
|
||||
|
||||
spvc_bool spvc_compiler_msl_is_vertex_attribute_used(spvc_compiler compiler, unsigned location)
|
||||
{
|
||||
#if SPIRV_CROSS_C_API_MSL
|
||||
|
||||
@ -33,7 +33,7 @@ extern "C" {
|
||||
/* Bumped if ABI or API breaks backwards compatibility. */
|
||||
#define SPVC_C_API_VERSION_MAJOR 0
|
||||
/* Bumped if APIs or enumerations are added in a backwards compatible way. */
|
||||
#define SPVC_C_API_VERSION_MINOR 18
|
||||
#define SPVC_C_API_VERSION_MINOR 19
|
||||
/* Bumped if internal implementation details change. */
|
||||
#define SPVC_C_API_VERSION_PATCH 0
|
||||
|
||||
@ -619,6 +619,7 @@ SPVC_PUBLIC_API spvc_result spvc_compiler_msl_add_vertex_attribute(spvc_compiler
|
||||
SPVC_PUBLIC_API spvc_result spvc_compiler_msl_add_resource_binding(spvc_compiler compiler,
|
||||
const spvc_msl_resource_binding *binding);
|
||||
SPVC_PUBLIC_API spvc_result spvc_compiler_msl_add_discrete_descriptor_set(spvc_compiler compiler, unsigned desc_set);
|
||||
SPVC_PUBLIC_API spvc_result spvc_compiler_msl_set_argument_buffer_device_address_space(spvc_compiler compiler, unsigned desc_set, spvc_bool device_address);
|
||||
SPVC_PUBLIC_API spvc_bool spvc_compiler_msl_is_vertex_attribute_used(spvc_compiler compiler, unsigned location);
|
||||
SPVC_PUBLIC_API spvc_bool spvc_compiler_msl_is_resource_used(spvc_compiler compiler,
|
||||
SpvExecutionModel model,
|
||||
|
||||
160
spirv_glsl.cpp
160
spirv_glsl.cpp
@ -4178,8 +4178,18 @@ void CompilerGLSL::emit_unrolled_unary_op(uint32_t result_type, uint32_t result_
|
||||
}
|
||||
|
||||
void CompilerGLSL::emit_unrolled_binary_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1,
|
||||
const char *op)
|
||||
const char *op, bool negate, SPIRType::BaseType expected_type)
|
||||
{
|
||||
auto &type0 = expression_type(op0);
|
||||
auto &type1 = expression_type(op1);
|
||||
|
||||
SPIRType target_type0 = type0;
|
||||
SPIRType target_type1 = type1;
|
||||
target_type0.basetype = expected_type;
|
||||
target_type1.basetype = expected_type;
|
||||
target_type0.vecsize = 1;
|
||||
target_type1.vecsize = 1;
|
||||
|
||||
auto &type = get<SPIRType>(result_type);
|
||||
auto expr = type_to_glsl_constructor(type);
|
||||
expr += '(';
|
||||
@ -4187,11 +4197,25 @@ void CompilerGLSL::emit_unrolled_binary_op(uint32_t result_type, uint32_t result
|
||||
{
|
||||
// Make sure to call to_expression multiple times to ensure
|
||||
// that these expressions are properly flushed to temporaries if needed.
|
||||
expr += to_extract_component_expression(op0, i);
|
||||
if (negate)
|
||||
expr += "!(";
|
||||
|
||||
if (expected_type != SPIRType::Unknown && type0.basetype != expected_type)
|
||||
expr += bitcast_expression(target_type0, type0.basetype, to_extract_component_expression(op0, i));
|
||||
else
|
||||
expr += to_extract_component_expression(op0, i);
|
||||
|
||||
expr += ' ';
|
||||
expr += op;
|
||||
expr += ' ';
|
||||
expr += to_extract_component_expression(op1, i);
|
||||
|
||||
if (expected_type != SPIRType::Unknown && type1.basetype != expected_type)
|
||||
expr += bitcast_expression(target_type1, type1.basetype, to_extract_component_expression(op1, i));
|
||||
else
|
||||
expr += to_extract_component_expression(op1, i);
|
||||
|
||||
if (negate)
|
||||
expr += ")";
|
||||
|
||||
if (i + 1 < type.vecsize)
|
||||
expr += ", ";
|
||||
@ -6868,8 +6892,16 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
|
||||
else if (ir.ids[index].get_type() == TypeConstant && !is_packed && !row_major_matrix_needs_conversion)
|
||||
{
|
||||
auto &c = get<SPIRConstant>(index);
|
||||
expr += ".";
|
||||
expr += index_to_swizzle(c.scalar());
|
||||
if (c.specialization)
|
||||
{
|
||||
// If the index is a spec constant, we cannot turn extract into a swizzle.
|
||||
expr += join("[", to_expression(index), "]");
|
||||
}
|
||||
else
|
||||
{
|
||||
expr += ".";
|
||||
expr += index_to_swizzle(c.scalar());
|
||||
}
|
||||
}
|
||||
else if (index_is_literal)
|
||||
{
|
||||
@ -7841,6 +7873,10 @@ uint32_t CompilerGLSL::get_integer_width_for_instruction(const Instruction &inst
|
||||
case OpSLessThanEqual:
|
||||
case OpSGreaterThan:
|
||||
case OpSGreaterThanEqual:
|
||||
case OpULessThan:
|
||||
case OpULessThanEqual:
|
||||
case OpUGreaterThan:
|
||||
case OpUGreaterThanEqual:
|
||||
return expression_type(ops[2]).width;
|
||||
|
||||
default:
|
||||
@ -8863,7 +8899,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
||||
auto &type = get<SPIRType>(result_type);
|
||||
|
||||
if (type.vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "||");
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "||", false, SPIRType::Unknown);
|
||||
else
|
||||
GLSL_BOP(||);
|
||||
break;
|
||||
@ -8877,7 +8913,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
||||
auto &type = get<SPIRType>(result_type);
|
||||
|
||||
if (type.vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "&&");
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "&&", false, SPIRType::Unknown);
|
||||
else
|
||||
GLSL_BOP(&&);
|
||||
break;
|
||||
@ -8934,7 +8970,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
||||
case OpUGreaterThan:
|
||||
case OpSGreaterThan:
|
||||
{
|
||||
auto type = opcode == OpUGreaterThan ? SPIRType::UInt : SPIRType::Int;
|
||||
auto type = opcode == OpUGreaterThan ? uint_type : int_type;
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
GLSL_BFOP_CAST(greaterThan, type);
|
||||
else
|
||||
@ -8954,7 +8990,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
||||
case OpUGreaterThanEqual:
|
||||
case OpSGreaterThanEqual:
|
||||
{
|
||||
auto type = opcode == OpUGreaterThanEqual ? SPIRType::UInt : SPIRType::Int;
|
||||
auto type = opcode == OpUGreaterThanEqual ? uint_type : int_type;
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
GLSL_BFOP_CAST(greaterThanEqual, type);
|
||||
else
|
||||
@ -8974,7 +9010,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
||||
case OpULessThan:
|
||||
case OpSLessThan:
|
||||
{
|
||||
auto type = opcode == OpULessThan ? SPIRType::UInt : SPIRType::Int;
|
||||
auto type = opcode == OpULessThan ? uint_type : int_type;
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
GLSL_BFOP_CAST(lessThan, type);
|
||||
else
|
||||
@ -8994,7 +9030,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
||||
case OpULessThanEqual:
|
||||
case OpSLessThanEqual:
|
||||
{
|
||||
auto type = opcode == OpULessThanEqual ? SPIRType::UInt : SPIRType::Int;
|
||||
auto type = opcode == OpULessThanEqual ? uint_type : int_type;
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
GLSL_BFOP_CAST(lessThanEqual, type);
|
||||
else
|
||||
@ -10104,28 +10140,98 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
||||
break;
|
||||
|
||||
case OpFUnordEqual:
|
||||
GLSL_BFOP(unsupported_FUnordEqual);
|
||||
break;
|
||||
|
||||
case OpFUnordNotEqual:
|
||||
GLSL_BFOP(unsupported_FUnordNotEqual);
|
||||
break;
|
||||
|
||||
case OpFUnordLessThan:
|
||||
GLSL_BFOP(unsupported_FUnordLessThan);
|
||||
break;
|
||||
|
||||
case OpFUnordGreaterThan:
|
||||
GLSL_BFOP(unsupported_FUnordGreaterThan);
|
||||
break;
|
||||
|
||||
case OpFUnordLessThanEqual:
|
||||
GLSL_BFOP(unsupported_FUnordLessThanEqual);
|
||||
break;
|
||||
|
||||
case OpFUnordGreaterThanEqual:
|
||||
GLSL_BFOP(unsupported_FUnordGreaterThanEqual);
|
||||
{
|
||||
// GLSL doesn't specify if floating point comparisons are ordered or unordered,
|
||||
// but glslang always emits ordered floating point compares for GLSL.
|
||||
// To get unordered compares, we can test the opposite thing and invert the result.
|
||||
// This way, we force true when there is any NaN present.
|
||||
uint32_t op0 = ops[2];
|
||||
uint32_t op1 = ops[3];
|
||||
|
||||
string expr;
|
||||
if (expression_type(op0).vecsize > 1)
|
||||
{
|
||||
const char *comp_op = nullptr;
|
||||
switch (opcode)
|
||||
{
|
||||
case OpFUnordEqual:
|
||||
comp_op = "notEqual";
|
||||
break;
|
||||
|
||||
case OpFUnordNotEqual:
|
||||
comp_op = "equal";
|
||||
break;
|
||||
|
||||
case OpFUnordLessThan:
|
||||
comp_op = "greaterThanEqual";
|
||||
break;
|
||||
|
||||
case OpFUnordLessThanEqual:
|
||||
comp_op = "greaterThan";
|
||||
break;
|
||||
|
||||
case OpFUnordGreaterThan:
|
||||
comp_op = "lessThanEqual";
|
||||
break;
|
||||
|
||||
case OpFUnordGreaterThanEqual:
|
||||
comp_op = "lessThan";
|
||||
break;
|
||||
|
||||
default:
|
||||
assert(0);
|
||||
break;
|
||||
}
|
||||
|
||||
expr = join("not(", comp_op, "(", to_unpacked_expression(op0), ", ", to_unpacked_expression(op1), "))");
|
||||
}
|
||||
else
|
||||
{
|
||||
const char *comp_op = nullptr;
|
||||
switch (opcode)
|
||||
{
|
||||
case OpFUnordEqual:
|
||||
comp_op = " != ";
|
||||
break;
|
||||
|
||||
case OpFUnordNotEqual:
|
||||
comp_op = " == ";
|
||||
break;
|
||||
|
||||
case OpFUnordLessThan:
|
||||
comp_op = " >= ";
|
||||
break;
|
||||
|
||||
case OpFUnordLessThanEqual:
|
||||
comp_op = " > ";
|
||||
break;
|
||||
|
||||
case OpFUnordGreaterThan:
|
||||
comp_op = " <= ";
|
||||
break;
|
||||
|
||||
case OpFUnordGreaterThanEqual:
|
||||
comp_op = " < ";
|
||||
break;
|
||||
|
||||
default:
|
||||
assert(0);
|
||||
break;
|
||||
}
|
||||
|
||||
expr = join("!(", to_enclosed_unpacked_expression(op0), comp_op, to_enclosed_unpacked_expression(op1), ")");
|
||||
}
|
||||
|
||||
emit_op(ops[0], ops[1], expr, should_forward(op0) && should_forward(op1));
|
||||
inherit_expression_dependencies(ops[1], op0);
|
||||
inherit_expression_dependencies(ops[1], op1);
|
||||
break;
|
||||
}
|
||||
|
||||
case OpReportIntersectionNV:
|
||||
statement("reportIntersectionNV(", to_expression(ops[0]), ", ", to_expression(ops[1]), ");");
|
||||
|
||||
@ -475,7 +475,8 @@ protected:
|
||||
void emit_unary_func_op(uint32_t result_type, uint32_t result_id, uint32_t op0, const char *op);
|
||||
void emit_unrolled_unary_op(uint32_t result_type, uint32_t result_id, uint32_t operand, const char *op);
|
||||
void emit_binary_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op);
|
||||
void emit_unrolled_binary_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op);
|
||||
void emit_unrolled_binary_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op,
|
||||
bool negate, SPIRType::BaseType expected_type);
|
||||
void emit_binary_op_cast(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op,
|
||||
SPIRType::BaseType input_type, bool skip_cast_if_equal_type);
|
||||
|
||||
|
||||
@ -4032,6 +4032,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
||||
// If we need to do implicit bitcasts, make sure we do it with the correct type.
|
||||
uint32_t integer_width = get_integer_width_for_instruction(instruction);
|
||||
auto int_type = to_signed_basetype(integer_width);
|
||||
auto uint_type = to_unsigned_basetype(integer_width);
|
||||
|
||||
switch (opcode)
|
||||
{
|
||||
@ -4194,7 +4195,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
||||
auto id = ops[1];
|
||||
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "==");
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "==", false, SPIRType::Unknown);
|
||||
else
|
||||
HLSL_BOP_CAST(==, int_type);
|
||||
break;
|
||||
@ -4202,12 +4203,19 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
||||
|
||||
case OpLogicalEqual:
|
||||
case OpFOrdEqual:
|
||||
case OpFUnordEqual:
|
||||
{
|
||||
// HLSL != operator is unordered.
|
||||
// https://docs.microsoft.com/en-us/windows/win32/direct3d10/d3d10-graphics-programming-guide-resources-float-rules.
|
||||
// isnan() is apparently implemented as x != x as well.
|
||||
// We cannot implement UnordEqual as !(OrdNotEqual), as HLSL cannot express OrdNotEqual.
|
||||
// HACK: FUnordEqual will be implemented as FOrdEqual.
|
||||
|
||||
auto result_type = ops[0];
|
||||
auto id = ops[1];
|
||||
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "==");
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "==", false, SPIRType::Unknown);
|
||||
else
|
||||
HLSL_BOP(==);
|
||||
break;
|
||||
@ -4219,7 +4227,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
||||
auto id = ops[1];
|
||||
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "!=");
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "!=", false, SPIRType::Unknown);
|
||||
else
|
||||
HLSL_BOP_CAST(!=, int_type);
|
||||
break;
|
||||
@ -4227,12 +4235,23 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
||||
|
||||
case OpLogicalNotEqual:
|
||||
case OpFOrdNotEqual:
|
||||
case OpFUnordNotEqual:
|
||||
{
|
||||
// HLSL != operator is unordered.
|
||||
// https://docs.microsoft.com/en-us/windows/win32/direct3d10/d3d10-graphics-programming-guide-resources-float-rules.
|
||||
// isnan() is apparently implemented as x != x as well.
|
||||
|
||||
// FIXME: FOrdNotEqual cannot be implemented in a crisp and simple way here.
|
||||
// We would need to do something like not(UnordEqual), but that cannot be expressed either.
|
||||
// Adding a lot of NaN checks would be a breaking change from perspective of performance.
|
||||
// SPIR-V will generally use isnan() checks when this even matters.
|
||||
// HACK: FOrdNotEqual will be implemented as FUnordEqual.
|
||||
|
||||
auto result_type = ops[0];
|
||||
auto id = ops[1];
|
||||
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "!=");
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "!=", false, SPIRType::Unknown);
|
||||
else
|
||||
HLSL_BOP(!=);
|
||||
break;
|
||||
@ -4243,10 +4262,10 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
||||
{
|
||||
auto result_type = ops[0];
|
||||
auto id = ops[1];
|
||||
auto type = opcode == OpUGreaterThan ? SPIRType::UInt : SPIRType::Int;
|
||||
auto type = opcode == OpUGreaterThan ? uint_type : int_type;
|
||||
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">");
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">", false, type);
|
||||
else
|
||||
HLSL_BOP_CAST(>, type);
|
||||
break;
|
||||
@ -4258,21 +4277,33 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
||||
auto id = ops[1];
|
||||
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">");
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">", false, SPIRType::Unknown);
|
||||
else
|
||||
HLSL_BOP(>);
|
||||
break;
|
||||
}
|
||||
|
||||
case OpFUnordGreaterThan:
|
||||
{
|
||||
auto result_type = ops[0];
|
||||
auto id = ops[1];
|
||||
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<=", true, SPIRType::Unknown);
|
||||
else
|
||||
CompilerGLSL::emit_instruction(instruction);
|
||||
break;
|
||||
}
|
||||
|
||||
case OpUGreaterThanEqual:
|
||||
case OpSGreaterThanEqual:
|
||||
{
|
||||
auto result_type = ops[0];
|
||||
auto id = ops[1];
|
||||
|
||||
auto type = opcode == OpUGreaterThanEqual ? SPIRType::UInt : SPIRType::Int;
|
||||
auto type = opcode == OpUGreaterThanEqual ? uint_type : int_type;
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">=");
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">=", false, type);
|
||||
else
|
||||
HLSL_BOP_CAST(>=, type);
|
||||
break;
|
||||
@ -4284,21 +4315,33 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
||||
auto id = ops[1];
|
||||
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">=");
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">=", false, SPIRType::Unknown);
|
||||
else
|
||||
HLSL_BOP(>=);
|
||||
break;
|
||||
}
|
||||
|
||||
case OpFUnordGreaterThanEqual:
|
||||
{
|
||||
auto result_type = ops[0];
|
||||
auto id = ops[1];
|
||||
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<", true, SPIRType::Unknown);
|
||||
else
|
||||
CompilerGLSL::emit_instruction(instruction);
|
||||
break;
|
||||
}
|
||||
|
||||
case OpULessThan:
|
||||
case OpSLessThan:
|
||||
{
|
||||
auto result_type = ops[0];
|
||||
auto id = ops[1];
|
||||
|
||||
auto type = opcode == OpULessThan ? SPIRType::UInt : SPIRType::Int;
|
||||
auto type = opcode == OpULessThan ? uint_type : int_type;
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<");
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<", false, type);
|
||||
else
|
||||
HLSL_BOP_CAST(<, type);
|
||||
break;
|
||||
@ -4310,21 +4353,33 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
||||
auto id = ops[1];
|
||||
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<");
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<", false, SPIRType::Unknown);
|
||||
else
|
||||
HLSL_BOP(<);
|
||||
break;
|
||||
}
|
||||
|
||||
case OpFUnordLessThan:
|
||||
{
|
||||
auto result_type = ops[0];
|
||||
auto id = ops[1];
|
||||
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">=", true, SPIRType::Unknown);
|
||||
else
|
||||
CompilerGLSL::emit_instruction(instruction);
|
||||
break;
|
||||
}
|
||||
|
||||
case OpULessThanEqual:
|
||||
case OpSLessThanEqual:
|
||||
{
|
||||
auto result_type = ops[0];
|
||||
auto id = ops[1];
|
||||
|
||||
auto type = opcode == OpULessThanEqual ? SPIRType::UInt : SPIRType::Int;
|
||||
auto type = opcode == OpULessThanEqual ? uint_type : int_type;
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<=");
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<=", false, type);
|
||||
else
|
||||
HLSL_BOP_CAST(<=, type);
|
||||
break;
|
||||
@ -4336,12 +4391,24 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
||||
auto id = ops[1];
|
||||
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<=");
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], "<=", false, SPIRType::Unknown);
|
||||
else
|
||||
HLSL_BOP(<=);
|
||||
break;
|
||||
}
|
||||
|
||||
case OpFUnordLessThanEqual:
|
||||
{
|
||||
auto result_type = ops[0];
|
||||
auto id = ops[1];
|
||||
|
||||
if (expression_type(ops[2]).vecsize > 1)
|
||||
emit_unrolled_binary_op(result_type, id, ops[2], ops[3], ">", true, SPIRType::Unknown);
|
||||
else
|
||||
CompilerGLSL::emit_instruction(instruction);
|
||||
break;
|
||||
}
|
||||
|
||||
case OpImageQueryLod:
|
||||
emit_texture_op(instruction);
|
||||
break;
|
||||
|
||||
@ -74,6 +74,17 @@ void CompilerMSL::add_discrete_descriptor_set(uint32_t desc_set)
|
||||
argument_buffer_discrete_mask |= 1u << desc_set;
|
||||
}
|
||||
|
||||
void CompilerMSL::set_argument_buffer_device_address_space(uint32_t desc_set, bool device_storage)
|
||||
{
|
||||
if (desc_set < kMaxArgumentBuffers)
|
||||
{
|
||||
if (device_storage)
|
||||
argument_buffer_device_storage_mask |= 1u << desc_set;
|
||||
else
|
||||
argument_buffer_device_storage_mask &= ~(1u << desc_set);
|
||||
}
|
||||
}
|
||||
|
||||
bool CompilerMSL::is_msl_vertex_attribute_used(uint32_t location)
|
||||
{
|
||||
return vtx_attrs_in_use.count(location) != 0;
|
||||
@ -9577,7 +9588,13 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
|
||||
// constant SSBO * constant (&array)[N].
|
||||
// However, this only matters for argument buffers, since for MSL 1.0 style codegen,
|
||||
// we emit the buffer array on stack instead, and that seems to work just fine apparently.
|
||||
address_space = "constant";
|
||||
|
||||
// If the argument was marked as being in device address space, any pointer to member would
|
||||
// be const device, not constant.
|
||||
if (argument_buffer_device_storage_mask & (1u << desc_set))
|
||||
address_space = "const device";
|
||||
else
|
||||
address_space = "constant";
|
||||
}
|
||||
}
|
||||
}
|
||||
@ -12667,8 +12684,20 @@ void CompilerMSL::analyze_argument_buffers()
|
||||
argument_buffer_ids[desc_set] = next_id;
|
||||
|
||||
auto &buffer_type = set<SPIRType>(type_id);
|
||||
buffer_type.storage = StorageClassUniform;
|
||||
|
||||
buffer_type.basetype = SPIRType::Struct;
|
||||
|
||||
if ((argument_buffer_device_storage_mask & (1u << desc_set)) != 0)
|
||||
{
|
||||
buffer_type.storage = StorageClassStorageBuffer;
|
||||
// Make sure the argument buffer gets marked as const device.
|
||||
set_decoration(next_id, DecorationNonWritable);
|
||||
// Need to mark the type as a Block to enable this.
|
||||
set_decoration(type_id, DecorationBlock);
|
||||
}
|
||||
else
|
||||
buffer_type.storage = StorageClassUniform;
|
||||
|
||||
set_name(type_id, join("spvDescriptorSetBuffer", desc_set));
|
||||
|
||||
auto &ptr_type = set<SPIRType>(ptr_type_id);
|
||||
|
||||
@ -429,6 +429,10 @@ public:
|
||||
// This corresponds to VK_KHR_push_descriptor in Vulkan.
|
||||
void add_discrete_descriptor_set(uint32_t desc_set);
|
||||
|
||||
// If an argument buffer is large enough, it may need to be in the device storage space rather than
|
||||
// constant. Opt-in to this behavior here on a per set basis.
|
||||
void set_argument_buffer_device_address_space(uint32_t desc_set, bool device_storage);
|
||||
|
||||
// Query after compilation is done. This allows you to check if a location or set/binding combination was used by the shader.
|
||||
bool is_msl_vertex_attribute_used(uint32_t location);
|
||||
|
||||
@ -852,6 +856,8 @@ protected:
|
||||
|
||||
uint32_t argument_buffer_ids[kMaxArgumentBuffers];
|
||||
uint32_t argument_buffer_discrete_mask = 0;
|
||||
uint32_t argument_buffer_device_storage_mask = 0;
|
||||
|
||||
void analyze_argument_buffers();
|
||||
bool descriptor_set_is_argument_buffer(uint32_t desc_set) const;
|
||||
|
||||
|
||||
@ -226,6 +226,11 @@ def cross_compile_msl(shader, spirv, opt, iterations, paths):
|
||||
msl_args.append('--msl-dynamic-buffer')
|
||||
msl_args.append('1')
|
||||
msl_args.append('2')
|
||||
if '.device-argument-buffer.' in shader:
|
||||
msl_args.append('--msl-device-argument-buffer')
|
||||
msl_args.append('0')
|
||||
msl_args.append('--msl-device-argument-buffer')
|
||||
msl_args.append('1')
|
||||
|
||||
subprocess.check_call(msl_args)
|
||||
|
||||
|
||||
Loading…
Reference in New Issue
Block a user