fix a bunch of case where SPIRType::op is wrong/corrupted

This commit is contained in:
Hugo Devillers 2023-11-30 15:42:46 +01:00
parent 950cad5913
commit af92037acb
18 changed files with 61 additions and 33 deletions

View File

@ -15,8 +15,8 @@ struct main0_patchOut
struct main0_in
{
uint3 m_86;
ushort2 m_90;
uint3 m_87;
ushort2 m_92;
float4 gl_Position;
};

View File

@ -68,7 +68,7 @@ struct main0_patchOut
struct main0_in
{
float3 in_tc_attr;
ushort2 m_196;
ushort2 m_197;
};
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], constant uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], device main0_in* spvIn [[buffer(22)]])

View File

@ -11,7 +11,7 @@ struct main0_out
struct main0_in
{
float4 vInputs;
ushort2 m_44;
ushort2 m_45;
};
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], device main0_in* spvIn [[buffer(22)]])

View File

@ -13,7 +13,7 @@ struct main0_out
struct main0_in
{
float3 in_tc_attr;
ushort2 m_104;
ushort2 m_105;
};
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], device main0_in* spvIn [[buffer(22)]])

View File

@ -10,8 +10,8 @@ struct main0_out
struct main0_in
{
uint3 m_82;
ushort2 m_86;
uint3 m_83;
ushort2 m_88;
float4 gl_Position;
};

View File

@ -20,7 +20,7 @@ struct main0_out
struct main0_in
{
float3 in_tc_attr;
ushort2 m_119;
ushort2 m_120;
};
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], device main0_in* spvIn [[buffer(22)]])

View File

@ -22,7 +22,7 @@ struct main0_patchOut
struct main0_in
{
float3 vPatchPosBase;
ushort2 m_996;
ushort2 m_997;
};
kernel void main0(constant UBO& _41 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], constant uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], device main0_in* spvIn [[buffer(22)]])

View File

@ -81,7 +81,7 @@ struct main0_out
struct main0_in
{
VertexOutput_1 p;
ushort2 m_173;
ushort2 m_174;
float4 gl_Position;
};

View File

@ -52,8 +52,8 @@ struct main0_out
struct main0_in
{
uint3 m_57;
ushort2 m_61;
uint3 m_58;
ushort2 m_63;
spvUnsafeArray<float, 2> gl_ClipDistance;
spvUnsafeArray<float, 1> gl_CullDistance;
};

View File

@ -17,8 +17,8 @@ struct main0_patchOut
struct main0_in
{
uint3 m_78;
ushort2 m_82;
uint3 m_79;
ushort2 m_84;
float4 gl_Position;
};

View File

@ -68,7 +68,7 @@ struct main0_patchOut
struct main0_in
{
float3 in_tc_attr;
ushort2 m_179;
ushort2 m_180;
};
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], constant uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], device main0_in* spvIn [[buffer(22)]])

View File

@ -52,7 +52,7 @@ struct main0_out
struct main0_in
{
float4 vInputs;
ushort2 m_43;
ushort2 m_44;
};
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], device main0_in* spvIn [[buffer(22)]])

View File

@ -13,7 +13,7 @@ struct main0_out
struct main0_in
{
float3 in_tc_attr;
ushort2 m_103;
ushort2 m_104;
};
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], device main0_in* spvIn [[buffer(22)]])

View File

@ -10,8 +10,8 @@ struct main0_out
struct main0_in
{
uint3 m_82;
ushort2 m_86;
uint3 m_83;
ushort2 m_88;
float4 gl_Position;
};

View File

@ -20,7 +20,7 @@ struct main0_out
struct main0_in
{
float3 in_tc_attr;
ushort2 m_107;
ushort2 m_108;
};
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], device main0_in* spvIn [[buffer(22)]])

View File

@ -24,7 +24,7 @@ struct main0_patchOut
struct main0_in
{
float3 vPatchPosBase;
ushort2 m_430;
ushort2 m_431;
};
static inline __attribute__((always_inline))

View File

@ -4160,11 +4160,12 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch)
continue;
// Create a fake variable to put at the location.
uint32_t offset = ir.increase_bound_by(4);
uint32_t offset = ir.increase_bound_by(5);
uint32_t type_id = offset;
uint32_t array_type_id = offset + 1;
uint32_t ptr_type_id = offset + 2;
uint32_t var_id = offset + 3;
uint32_t vec_type_id = offset + 1;
uint32_t array_type_id = offset + 2;
uint32_t ptr_type_id = offset + 3;
uint32_t var_id = offset + 4;
SPIRType type { spv::Op::OpTypeInt };
switch (input.second.format)
@ -4180,14 +4181,22 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch)
type.width = 32;
break;
}
type.vecsize = input.second.vecsize;
set<SPIRType>(type_id, type);
if (input.second.vecsize > 1)
{
type.op = spv::Op::OpTypeVector;
type.vecsize = input.second.vecsize;
set<SPIRType>(vec_type_id, type);
type_id = vec_type_id;
}
type.op = spv::Op::OpTypeArray;
type.array.push_back(0);
type.array_size_literal.push_back(true);
type.parent_type = type_id;
set<SPIRType>(array_type_id, type);
type.op = spv::Op::OpTypePointer;
type.pointer = true;
type.pointer_depth++;
type.parent_type = array_type_id;
@ -4218,11 +4227,12 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch)
continue;
// Create a fake variable to put at the location.
uint32_t offset = ir.increase_bound_by(4);
uint32_t offset = ir.increase_bound_by(5);
uint32_t type_id = offset;
uint32_t array_type_id = offset + 1;
uint32_t ptr_type_id = offset + 2;
uint32_t var_id = offset + 3;
uint32_t vec_type_id = offset + 1;
uint32_t array_type_id = offset + 2;
uint32_t ptr_type_id = offset + 3;
uint32_t var_id = offset + 4;
SPIRType type { spv::Op::OpTypeInt };
switch (output.second.format)
@ -4238,17 +4248,25 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch)
type.width = 32;
break;
}
type.vecsize = output.second.vecsize;
set<SPIRType>(type_id, type);
if (output.second.vecsize > 1)
{
type.op = spv::Op::OpTypeVector;
type.vecsize = output.second.vecsize;
set<SPIRType>(vec_type_id, type);
type_id = vec_type_id;
}
if (is_tesc_shader())
{
type.op = spv::Op::OpTypeArray;
type.array.push_back(0);
type.array_size_literal.push_back(true);
type.parent_type = type_id;
type.parent_type = vec_type_id;
set<SPIRType>(array_type_id, type);
}
type.op = spv::Op::OpTypePointer;
type.pointer = true;
type.pointer_depth++;
type.parent_type = is_tesc_shader() ? array_type_id : type_id;
@ -4895,6 +4913,7 @@ void CompilerMSL::ensure_member_packing_rules_msl(SPIRType &ib_type, uint32_t in
{
type.columns = 1;
assert(type.array.empty());
type.op = spv::Op::OpTypeArray;
type.array.push_back(1);
type.array_size_literal.push_back(true);
}
@ -4911,6 +4930,7 @@ void CompilerMSL::ensure_member_packing_rules_msl(SPIRType &ib_type, uint32_t in
type.vecsize = type.columns;
type.columns = 1;
assert(type.array.empty());
type.op = spv::Op::OpTypeArray;
type.array.push_back(1);
type.array_size_literal.push_back(true);
}
@ -10809,7 +10829,7 @@ string CompilerMSL::to_function_name(const TextureFunctionNameArguments &args)
string CompilerMSL::convert_to_f32(const string &expr, uint32_t components)
{
SPIRType t { spv::Op::OpTypeFloat };
SPIRType t { components > 1 ? spv::Op::OpTypeVector : spv::Op::OpTypeFloat };
t.basetype = SPIRType::Float;
t.vecsize = components;
t.columns = 1;
@ -18063,6 +18083,7 @@ void CompilerMSL::add_argument_buffer_padding_type(uint32_t mbr_type_id, SPIRTyp
uint32_t ary_type_id = ir.increase_bound_by(1);
auto &ary_type = set<SPIRType>(ary_type_id, spv::Op::OpTypeArray);
ary_type = get<SPIRType>(type_id);
ary_type.op = spv::Op::OpTypeArray;
ary_type.array.push_back(count);
ary_type.array_size_literal.push_back(true);
ary_type.parent_type = type_id;

View File

@ -571,6 +571,7 @@ void Parser::parse(const Instruction &instruction)
auto &vecbase = set<SPIRType>(id, op);
vecbase = base;
vecbase.op = op;
vecbase.vecsize = vecsize;
vecbase.self = id;
vecbase.parent_type = ops[1];
@ -586,6 +587,7 @@ void Parser::parse(const Instruction &instruction)
auto &matrixbase = set<SPIRType>(id, op);
matrixbase = base;
matrixbase.op = op;
matrixbase.columns = colcount;
matrixbase.self = id;
matrixbase.parent_type = ops[1];
@ -601,6 +603,7 @@ void Parser::parse(const Instruction &instruction)
auto &base = get<SPIRType>(tid);
arraybase = base;
arraybase.op = op;
arraybase.parent_type = tid;
uint32_t cid = ops[2];
@ -632,6 +635,7 @@ void Parser::parse(const Instruction &instruction)
forward_pointer_fixups.push_back({ id, ops[1] });
arraybase = base;
arraybase.op = op;
arraybase.array.push_back(0);
arraybase.array_size_literal.push_back(true);
arraybase.parent_type = ops[1];
@ -685,7 +689,10 @@ void Parser::parse(const Instruction &instruction)
auto &ptrbase = set<SPIRType>(id, op);
if (base)
{
ptrbase = *base;
ptrbase.op = op;
}
ptrbase.pointer = true;
ptrbase.pointer_depth++;