MSL: Handle non-thread storage class in Modf/Frexp pointer versions.

This commit is contained in:
Hans-Kristian Arntzen 2021-11-07 10:49:26 +01:00
parent 04293e03fd
commit e40d19bdbf
5 changed files with 204 additions and 2 deletions

View File

@ -0,0 +1,44 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct _17
{
float2 _m0;
float2 _m1;
};
struct _4
{
uint2 _m0[324];
};
struct _7
{
float2 _m0[648];
};
struct _10
{
float2 _m0[648];
};
kernel void main0(const device _4& _5 [[buffer(0)]], device _7& _8 [[buffer(1)]], device _10& _11 [[buffer(2)]])
{
for (uint _39 = 0u; _39 < 648u; _39 += 2u)
{
uint2 _40 = _5._m0[_39 / 2u];
float2 _41 = as_type<float2>(_40);
float2 _76;
float2 _61 = modf(_41, _76);
_8._m0[_39] = _76;
_8._m0[_39 + 1u] = _61;
_17 _64;
_64._m0 = modf(_41, _64._m1);
_17 _42 = _64;
_11._m0[_39] = _42._m1;
_11._m0[_39 + 1u] = _42._m0;
}
}

View File

@ -0,0 +1,25 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct main0_out
{
float4 f [[user(locn0)]];
float4 gl_Position [[position]];
};
struct main0_in
{
float4 f2 [[attribute(0)]];
};
vertex void main0(main0_in in [[stage_in]], uint gl_VertexIndex [[vertex_id]], uint gl_BaseVertex [[base_vertex]], uint gl_InstanceIndex [[instance_id]], uint gl_BaseInstance [[base_instance]], device main0_out* spvOut [[buffer(28)]], device uint* spvIndirectParams [[buffer(29)]])
{
device main0_out& out = spvOut[(gl_InstanceIndex - gl_BaseInstance) * spvIndirectParams[0] + gl_VertexIndex - gl_BaseVertex];
float4 _35;
float4 _21 = modf(in.f2, _35);
out.f = _35;
out.gl_Position = _21;
}

View File

@ -0,0 +1,116 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 91
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %2 "main"
OpExecutionMode %2 LocalSize 1 1 1
OpDecorate %_arr_v2uint_uint_324 ArrayStride 8
OpMemberDecorate %_struct_6 0 NonWritable
OpMemberDecorate %_struct_6 0 Offset 0
OpDecorate %_struct_6 BufferBlock
OpDecorate %7 DescriptorSet 0
OpDecorate %7 Binding 0
OpDecorate %_arr_v2float_uint_648 ArrayStride 8
OpMemberDecorate %_struct_9 0 Offset 0
OpDecorate %_struct_9 BufferBlock
OpDecorate %11 DescriptorSet 0
OpDecorate %11 Binding 1
OpDecorate %_arr_v2float_uint_648_0 ArrayStride 8
OpMemberDecorate %_struct_13 0 Offset 0
OpDecorate %_struct_13 BufferBlock
OpDecorate %14 DescriptorSet 0
OpDecorate %14 Binding 2
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%_ptr_Function_v2float = OpTypePointer Function %v2float
%_struct_19 = OpTypeStruct %v2float %v2float
%10 = OpTypeFunction %_struct_19 %_ptr_Function_v2float
%_ptr_Function__struct_19 = OpTypePointer Function %_struct_19
%uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint
%uint_0 = OpConstant %uint 0
%uint_648 = OpConstant %uint 648
%bool = OpTypeBool
%v2uint = OpTypeVector %uint 2
%_ptr_Function_v2uint = OpTypePointer Function %v2uint
%uint_324 = OpConstant %uint 324
%_arr_v2uint_uint_324 = OpTypeArray %v2uint %uint_324
%_struct_6 = OpTypeStruct %_arr_v2uint_uint_324
%_ptr_Uniform__struct_6 = OpTypePointer Uniform %_struct_6
%7 = OpVariable %_ptr_Uniform__struct_6 Uniform
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%uint_2 = OpConstant %uint 2
%_ptr_Uniform_v2uint = OpTypePointer Uniform %v2uint
%_arr_v2float_uint_648 = OpTypeArray %v2float %uint_648
%_struct_9 = OpTypeStruct %_arr_v2float_uint_648
%_ptr_Uniform__struct_9 = OpTypePointer Uniform %_struct_9
%11 = OpVariable %_ptr_Uniform__struct_9 Uniform
%uint_1 = OpConstant %uint 1
%_ptr_Uniform_v2float = OpTypePointer Uniform %v2float
%_arr_v2float_uint_648_0 = OpTypeArray %v2float %uint_648
%_struct_13 = OpTypeStruct %_arr_v2float_uint_648_0
%_ptr_Uniform__struct_13 = OpTypePointer Uniform %_struct_13
%14 = OpVariable %_ptr_Uniform__struct_13 Uniform
%int_1 = OpConstant %int 1
%2 = OpFunction %void None %3
%5 = OpLabel
%46 = OpVariable %_ptr_Function_uint Function
%47 = OpVariable %_ptr_Function_v2uint Function
%48 = OpVariable %_ptr_Function_v2float Function
%50 = OpVariable %_ptr_Function__struct_19 Function
OpStore %46 %uint_0
OpBranch %30
%30 = OpLabel
OpLoopMerge %32 %33 None
OpBranch %34
%34 = OpLabel
%35 = OpLoad %uint %46
%38 = OpULessThan %bool %35 %uint_648
OpBranchConditional %38 %31 %32
%31 = OpLabel
%49 = OpLoad %uint %46
%51 = OpUDiv %uint %49 %uint_2
%53 = OpAccessChain %_ptr_Uniform_v2uint %7 %int_0 %51
%54 = OpLoad %v2uint %53
OpStore %47 %54
%56 = OpLoad %v2uint %47
%57 = OpBitcast %v2float %56
OpStore %48 %57
%62 = OpLoad %uint %46
%64 = OpIAdd %uint %62 %uint_1
%65 = OpLoad %v2float %48
%66 = OpLoad %uint %46
%68 = OpAccessChain %_ptr_Uniform_v2float %11 %int_0 %66
%69 = OpExtInst %v2float %1 Modf %65 %68
%70 = OpAccessChain %_ptr_Uniform_v2float %11 %int_0 %64
OpStore %70 %69
%73 = OpLoad %v2float %48
%74 = OpExtInst %_struct_19 %1 ModfStruct %73
OpStore %50 %74
%79 = OpLoad %uint %46
%81 = OpAccessChain %_ptr_Function_v2float %50 %int_1
%82 = OpLoad %v2float %81
%83 = OpAccessChain %_ptr_Uniform_v2float %14 %int_0 %79
OpStore %83 %82
%84 = OpLoad %uint %46
%85 = OpIAdd %uint %84 %uint_1
%86 = OpAccessChain %_ptr_Function_v2float %50 %int_0
%87 = OpLoad %v2float %86
%88 = OpAccessChain %_ptr_Uniform_v2float %14 %int_0 %85
OpStore %88 %87
OpBranch %33
%33 = OpLabel
%89 = OpLoad %uint %46
%90 = OpIAdd %uint %89 %uint_2
OpStore %46 %90
OpBranch %30
%32 = OpLabel
OpReturn
OpFunctionEnd

View File

@ -0,0 +1,9 @@
#version 450
layout(location = 0) out vec4 f;
layout(location = 0) in vec4 f2;
void main()
{
gl_Position = modf(f2, f);
}

View File

@ -9163,8 +9163,16 @@ void CompilerMSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
case GLSLstd450Frexp:
{
// Special case. If the variable is a scalar access chain, we cannot use it directly. We have to emit a temporary.
// Another special case is if the variable is in a storage class which is not thread.
auto *ptr = maybe_get<SPIRExpression>(args[1]);
if (ptr && ptr->access_chain && is_scalar(expression_type(args[1])))
auto &type = expression_type(args[1]);
bool is_thread_storage = storage_class_array_is_thread(type.storage);
if (type.storage == StorageClassOutput && capture_output_to_buffer)
is_thread_storage = false;
if (!is_thread_storage ||
(ptr && ptr->access_chain && is_scalar(expression_type(args[1]))))
{
register_call_out_argument(args[1]);
forced_temporaries.insert(id);
@ -9175,7 +9183,7 @@ void CompilerMSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
if (!tmp_id)
tmp_id = ir.increase_bound_by(1);
uint32_t tmp_type_id = get_pointee_type_id(ptr->expression_type);
uint32_t tmp_type_id = get_pointee_type_id(expression_type_id(args[1]));
emit_uninitialized_temporary_expression(tmp_type_id, tmp_id);
emit_binary_func_op(result_type, id, args[0], tmp_id, eop == GLSLstd450Modf ? "modf" : "frexp");
statement(to_expression(args[1]), " = ", to_expression(tmp_id), ";");