This commit is contained in:
Hans-Kristian Arntzen 2017-12-05 09:58:12 +01:00
Родитель 12988801c4
Коммит aa2557c7df
6 изменённых файлов: 24 добавлений и 35 удалений

Просмотреть файл

@ -3,15 +3,7 @@
using namespace metal;
constant uint3 gl_WorkGroupSize = uint3(8u, 4u, 2u);
kernel void main0(uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], uint3 gl_NumWorkGroups [[threadgroups_per_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
{
uint3 local_id = gl_LocalInvocationID;
uint3 global_id = gl_GlobalInvocationID;
uint local_index = gl_LocalInvocationIndex;
uint3 work_group_size = gl_WorkGroupSize;
uint3 num_work_groups = gl_NumWorkGroups;
uint3 work_group_id = gl_WorkGroupID;
}

Просмотреть файл

@ -12,12 +12,10 @@ struct SSBO
kernel void main0(device SSBO& _13 [[buffer(0)]])
{
float4 _17 = _13.data;
float2 _27 = float2(10.0);
float2 _28 = _17.yz + _27;
_13.data = float4(_17.x, _28, _17.w);
_13.data = float4(_17.x, _17.yz + float2(10.0), _17.w);
_13.data = (_17 + _17) + _17;
_13.data = _28.xxyy;
_13.data = float4(_28.y);
_13.data = float4((_17.zw + _27)[_13.index]);
_13.data = (_17.yz + float2(10.0)).xxyy;
_13.data = float4((_17.yz + float2(10.0)).y);
_13.data = float4((_17.zw + float2(10.0))[_13.index]);
}

Просмотреть файл

@ -5,7 +5,6 @@ using namespace metal;
kernel void main0(texture2d<float> uImageIn [[texture(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], texture2d<float, access::write> uImageOut [[texture(1)]])
{
int2 _23 = int2(gl_GlobalInvocationID.xy);
uImageOut.write(uImageIn.read(uint2((_23 + int2(uImageIn.get_width(), uImageIn.get_height())))), uint2(_23));
uImageOut.write(uImageIn.read(uint2((int2(gl_GlobalInvocationID.xy) + int2(uImageIn.get_width(), uImageIn.get_height())))), uint2(int2(gl_GlobalInvocationID.xy)));
}

Просмотреть файл

@ -28,10 +28,6 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic
}
for (int _68 = 0; _68 < 20; _65 = _69 + 1, _68 = _65)
{
if (_68 == 10)
{
break;
}
return;
}
_27.out_data[gl_GlobalInvocationID.x] = float4(10.0);

Просмотреть файл

@ -679,7 +679,8 @@ ShaderResources Compiler::get_shader_resources(const unordered_set<uint32_t> *ac
if (var.storage == StorageClassInput && interface_variable_exists_in_entry_point(var.self))
{
if (meta[type.self].decoration.decoration_flags & (1ull << DecorationBlock))
res.stage_inputs.push_back({ var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
res.stage_inputs.push_back(
{ var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
else
res.stage_inputs.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
}
@ -692,7 +693,8 @@ ShaderResources Compiler::get_shader_resources(const unordered_set<uint32_t> *ac
else if (var.storage == StorageClassOutput && interface_variable_exists_in_entry_point(var.self))
{
if (meta[type.self].decoration.decoration_flags & (1ull << DecorationBlock))
res.stage_outputs.push_back({ var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
res.stage_outputs.push_back(
{ var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
else
res.stage_outputs.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
}
@ -700,21 +702,21 @@ ShaderResources Compiler::get_shader_resources(const unordered_set<uint32_t> *ac
else if (type.storage == StorageClassUniform &&
(meta[type.self].decoration.decoration_flags & (1ull << DecorationBlock)))
{
res.uniform_buffers.push_back({ var.self, var.basetype, type.self,
get_remapped_declared_block_name(var.self) });
res.uniform_buffers.push_back(
{ var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
}
// Old way to declare SSBOs.
else if (type.storage == StorageClassUniform &&
(meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)))
{
res.storage_buffers.push_back({ var.self, var.basetype, type.self,
get_remapped_declared_block_name(var.self) });
res.storage_buffers.push_back(
{ var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
}
// Modern way to declare SSBOs.
else if (type.storage == StorageClassStorageBuffer)
{
res.storage_buffers.push_back({ var.self, var.basetype, type.self,
get_remapped_declared_block_name(var.self) });
res.storage_buffers.push_back(
{ var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
}
// Push constant blocks
else if (type.storage == StorageClassPushConstant)

Просмотреть файл

@ -1532,7 +1532,7 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
emit_barrier(ops[0], ops[1], ops[2]);
break;
// OpOuterProduct
// OpOuterProduct
default:
CompilerGLSL::emit_instruction(instruction);
@ -1837,10 +1837,10 @@ void CompilerMSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
break;
}
// TODO:
// GLSLstd450InterpolateAtCentroid (centroid_no_perspective qualifier)
// GLSLstd450InterpolateAtSample (sample_no_perspective qualifier)
// GLSLstd450InterpolateAtOffset
// TODO:
// GLSLstd450InterpolateAtCentroid (centroid_no_perspective qualifier)
// GLSLstd450InterpolateAtSample (sample_no_perspective qualifier)
// GLSLstd450InterpolateAtOffset
default:
CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count);
@ -2725,11 +2725,13 @@ string CompilerMSL::ensure_valid_name(string name, string pfx)
void CompilerMSL::replace_illegal_names()
{
static const unordered_set<string> keywords = {
"kernel", "bias",
"kernel",
"bias",
};
static const unordered_set<string> illegal_func_names = {
"main", "saturate",
"main",
"saturate",
};
for (auto &id : ids)