Skip to content

Commit

Permalink
Merge branch 'main' into xfb
Browse files Browse the repository at this point in the history
  • Loading branch information
cdavis5e authored Sep 19, 2023
2 parents 109959e + 5e963d6 commit 3bd855f
Show file tree
Hide file tree
Showing 29 changed files with 1,308 additions and 125 deletions.
13 changes: 8 additions & 5 deletions main.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -676,6 +676,7 @@ struct CLIArguments
bool msl_manual_helper_invocation_updates = true;
bool msl_check_discarded_frag_stores = false;
bool msl_sample_dref_lod_array_as_grad = false;
bool msl_runtime_array_rich_descriptor = false;
const char *msl_combined_sampler_suffix = nullptr;
CompilerMSL::Options::PrimitiveType msl_xfb_primitive_type = CompilerMSL::Options::PrimitiveType::Dynamic;
bool glsl_emit_push_constant_as_ubo = false;
Expand Down Expand Up @@ -865,7 +866,8 @@ static void print_help_msl()
"\t\tRequires MSL 2.0 to be enabled.\n"
"\t[--msl-argument-buffer-tier]:\n\t\tWhen using Metal argument buffers, indicate the Metal argument buffer tier level supported by the Metal platform.\n"
"\t\tUses same values as Metal MTLArgumentBuffersTier enumeration (0 = Tier1, 1 = Tier2).\n"
"\t\tSetting this value also enables msl-argument-buffers.\n"
"\t\tNOTE: Setting this value no longer enables msl-argument-buffers implicitly.\n"
"\t[--msl-runtime-array-rich-descriptor]:\n\t\tWhen declaring a runtime array of SSBOs, declare an array of {ptr, len} pairs to support OpArrayLength.\n"
"\t[--msl-texture-buffer-native]:\n\t\tEnable native support for texel buffers. Otherwise, it is emulated as a normal texture.\n"
"\t[--msl-framebuffer-fetch]:\n\t\tImplement subpass inputs with frame buffer fetch.\n"
"\t\tEmits [[color(N)]] inputs in fragment stage.\n"
Expand Down Expand Up @@ -1236,6 +1238,7 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
msl_opts.sample_dref_lod_array_as_grad = args.msl_sample_dref_lod_array_as_grad;
msl_opts.xfb_primitive_type = args.msl_xfb_primitive_type;
msl_opts.ios_support_base_vertex_instance = true;
msl_opts.runtime_array_rich_descriptor = args.msl_runtime_array_rich_descriptor;
msl_comp->set_msl_options(msl_opts);
for (auto &v : args.msl_discrete_descriptor_sets)
msl_comp->add_discrete_descriptor_set(v);
Expand Down Expand Up @@ -1642,10 +1645,8 @@ static int main_inner(int argc, char *argv[])
cbs.add("--msl-pad-fragment-output", [&args](CLIParser &) { args.msl_pad_fragment_output = true; });
cbs.add("--msl-domain-lower-left", [&args](CLIParser &) { args.msl_domain_lower_left = true; });
cbs.add("--msl-argument-buffers", [&args](CLIParser &) { args.msl_argument_buffers = true; });
cbs.add("--msl-argument-buffer-tier", [&args](CLIParser &parser) {
args.msl_argument_buffers_tier = parser.next_uint();
args.msl_argument_buffers = true;
});
cbs.add("--msl-argument-buffer-tier",
[&args](CLIParser &parser) { args.msl_argument_buffers_tier = parser.next_uint(); });
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",
Expand Down Expand Up @@ -1795,6 +1796,8 @@ static int main_inner(int argc, char *argv[])
cbs.add("--msl-combined-sampler-suffix", [&args](CLIParser &parser) {
args.msl_combined_sampler_suffix = parser.next_string();
});
cbs.add("--msl-runtime-array-rich-descriptor",
[&args](CLIParser &) { args.msl_runtime_array_rich_descriptor = true; });
cbs.add("--msl-xfb-primitive-type",
[&args](CLIParser &parser)
{
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,7 @@ using namespace metal;
static inline __attribute__((always_inline))
void _main(thread const uint3& id, texture2d<float, access::read_write> TargetTexture)
{
TargetTexture.fence();
float2 loaded = TargetTexture.read(uint2(id.xy)).xy;
float2 storeTemp = loaded + float2(1.0);
TargetTexture.write(storeTemp.xyyy, uint2((id.xy + uint2(1u))));
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,7 @@ using namespace metal;

fragment void main0(texture2d_ms<float> uImageMS [[texture(0)]], texture2d_array<float, access::read_write> uImageArray [[texture(1)]], texture2d<float, access::write> uImage [[texture(2)]])
{
uImageArray.fence();
uImage.write(uImageMS.read(uint2(int2(1, 2)), 2), uint2(int2(2, 3)));
uImageArray.write(uImageArray.read(uint2(int3(1, 2, 4).xy), uint(int3(1, 2, 4).z)), uint2(int3(2, 3, 7).xy), uint(int3(2, 3, 7).z));
}
Expand Down
Original file line number Diff line number Diff line change
@@ -0,0 +1,159 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"

#include <metal_stdlib>
#include <simd/simd.h>
#if __METAL_VERSION__ >= 230
#include <metal_raytracing>
using namespace metal::raytracing;
#endif

using namespace metal;

intersection_params spvMakeIntersectionParams(uint flags)
{
intersection_params ip;
if ((flags & 1) != 0)
ip.force_opacity(forced_opacity::opaque);
if ((flags & 2) != 0)
ip.force_opacity(forced_opacity::non_opaque);
if ((flags & 4) != 0)
ip.accept_any_intersection(true);
if ((flags & 16) != 0)
ip.set_triangle_cull_mode(triangle_cull_mode::back);
if ((flags & 32) != 0)
ip.set_triangle_cull_mode(triangle_cull_mode::front);
if ((flags & 64) != 0)
ip.set_opacity_cull_mode(opacity_cull_mode::opaque);
if ((flags & 128) != 0)
ip.set_opacity_cull_mode(opacity_cull_mode::non_opaque);
if ((flags & 256) != 0)
ip.set_geometry_cull_mode(geometry_cull_mode::triangle);
if ((flags & 512) != 0)
ip.set_geometry_cull_mode(geometry_cull_mode::bounding_box);
return ip;
}

template<typename T>
struct spvDescriptor
{
T value;
};

template<typename T>
struct spvBufferDescriptor
{
T value;
int length;
const device T& operator -> () const device
{
return value;
}
const device T& operator * () const device
{
return value;
}
};

template<typename T>
struct spvDescriptorArray
{
spvDescriptorArray(const device spvDescriptor<T>* ptr) : ptr(ptr)
{
}
const device T& operator [] (size_t i) const
{
return ptr[i].value;
}
const device spvDescriptor<T>* ptr;
};

template<typename T>
struct spvDescriptorArray<device T*>
{
spvDescriptorArray(const device spvBufferDescriptor<device T*>* ptr) : ptr(ptr)
{
}
const device T* operator [] (size_t i) const
{
return ptr[i].value;
}
const int length(int i) const
{
return ptr[i].length;
}
const device spvBufferDescriptor<device T*>* ptr;
};

struct Ssbo
{
uint val;
uint data[1];
};

struct Ubo
{
uint val;
};

struct main0_in
{
uint inputId [[user(locn0)]];
};

fragment void main0(main0_in in [[stage_in]], const device spvBufferDescriptor<const device Ssbo*>* ssbo_ [[buffer(4)]], const device spvDescriptor<constant Ubo*>* ubo_ [[buffer(5)]], const device spvDescriptor<texture2d<float>>* smp_textures_ [[buffer(0)]], const device spvDescriptor<texture2d<float>>* textures_ [[buffer(2)]], const device spvDescriptor<texture2d<float>>* images_ [[buffer(6)]], const device spvDescriptor<sampler>* smp_texturesSmplr_ [[buffer(1)]], const device spvDescriptor<sampler>* smp_ [[buffer(3)]], const device spvDescriptor<raytracing::acceleration_structure<raytracing::instancing>>* tlas_ [[buffer(7)]])
{
spvDescriptorArray<texture2d<float>> smp_textures {smp_textures_};
spvDescriptorArray<sampler> smp_texturesSmplr {smp_texturesSmplr_};
spvDescriptorArray<texture2d<float>> textures {textures_};
spvDescriptorArray<sampler> smp {smp_};
spvDescriptorArray<const device Ssbo*> ssbo {ssbo_};
spvDescriptorArray<constant Ubo*> ubo {ubo_};
spvDescriptorArray<texture2d<float>> images {images_};
spvDescriptorArray<raytracing::acceleration_structure<raytracing::instancing>> tlas {tlas_};

uint _231 = in.inputId;
raytracing::intersection_query<raytracing::instancing, raytracing::triangle_data> rayQuery;
raytracing::intersection_query<raytracing::instancing, raytracing::triangle_data> rayQuery_1;
if (smp_textures[_231].sample(smp_texturesSmplr[_231], float2(0.0), level(0.0)).w > 0.5)
{
discard_fragment();
}
uint _249 = in.inputId + 8u;
if (textures[_231].sample(smp[_249], float2(0.0), level(0.0)).w > 0.5)
{
discard_fragment();
}
if (ssbo[_231]->val == 2u)
{
discard_fragment();
}
if (int((ssbo.length(123) - 4) / 4) == 25)
{
discard_fragment();
}
if (ubo[_231]->val == 2u)
{
discard_fragment();
}
if (images[_231].read(uint2(int2(0))).w > 0.5)
{
discard_fragment();
}
rayQuery.reset(ray(float3(0.0), float3(1.0), 0.00999999977648258209228515625, 1.0), tlas[in.inputId], spvMakeIntersectionParams(0u));
bool _301 = rayQuery.next();
if (smp_textures[_231].sample(smp_texturesSmplr[_231], float2(0.0), level(0.0)).w > 0.5)
{
discard_fragment();
}
if (textures[_231].sample(smp[_231], float2(0.0), level(0.0)).w > 0.5)
{
discard_fragment();
}
if (images[_231].read(uint2(int2(0))).w > 0.5)
{
discard_fragment();
}
rayQuery_1.reset(ray(float3(0.0), float3(1.0), 0.00999999977648258209228515625, 1.0), tlas[in.inputId], spvMakeIntersectionParams(0u));
bool _336 = rayQuery_1.next();
}

4 changes: 1 addition & 3 deletions reference/opt/shaders/asm/vert/invariant-block.sso.asm.vert
Original file line number Diff line number Diff line change
Expand Up @@ -2,14 +2,12 @@

out gl_PerVertex
{
vec4 gl_Position;
invariant vec4 gl_Position;
float gl_PointSize;
float gl_ClipDistance[1];
float gl_CullDistance[1];
};

invariant gl_Position;

void main()
{
gl_Position = vec4(1.0);
Expand Down
4 changes: 1 addition & 3 deletions reference/opt/shaders/asm/vert/invariant.sso.asm.vert
Original file line number Diff line number Diff line change
Expand Up @@ -2,11 +2,9 @@

out gl_PerVertex
{
vec4 gl_Position;
invariant vec4 gl_Position;
};

invariant gl_Position;

void main()
{
gl_Position = vec4(1.0);
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,7 @@ using namespace metal;
static inline __attribute__((always_inline))
void _main(thread const uint3& id, texture2d<float, access::read_write> TargetTexture)
{
TargetTexture.fence();
float2 loaded = TargetTexture.read(uint2(id.xy)).xy;
float2 storeTemp = loaded + float2(1.0);
TargetTexture.write(storeTemp.xyyy, uint2((id.xy + uint2(1u))));
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,7 @@ using namespace metal;
fragment void main0(texture2d_ms<float> uImageMS [[texture(0)]], texture2d_array<float, access::read_write> uImageArray [[texture(1)]], texture2d<float, access::write> uImage [[texture(2)]])
{
float4 a = uImageMS.read(uint2(int2(1, 2)), 2);
uImageArray.fence();
float4 b = uImageArray.read(uint2(int3(1, 2, 4).xy), uint(int3(1, 2, 4).z));
uImage.write(a, uint2(int2(2, 3)));
uImageArray.write(b, uint2(int3(2, 3, 7).xy), uint(int3(2, 3, 7).z));
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -44,7 +44,7 @@ struct main0_in
};

static inline __attribute__((always_inline))
uint doRay(thread const float3& rayOrigin, thread const float3& rayDirection, thread const float& rayDistance, thread raytracing::intersection_query<raytracing::instancing, raytracing::triangle_data>& rayQuery, thread const raytracing::acceleration_structure<raytracing::instancing>& topLevelAS)
uint doRay(thread const float3& rayOrigin, thread const float3& rayDirection, thread const float& rayDistance, thread raytracing::intersection_query<raytracing::instancing, raytracing::triangle_data>& rayQuery, const raytracing::acceleration_structure<raytracing::instancing> topLevelAS)
{
rayQuery.reset(ray(rayOrigin, rayDirection, 0.001000000047497451305389404296875, rayDistance), topLevelAS, spvMakeIntersectionParams(4u));
for (;;)
Expand Down
Loading

0 comments on commit 3bd855f

Please sign in to comment.