Updated spirv-cross.

This commit is contained in:
Бранимир Караџић 2019-11-28 20:20:59 -08:00
parent 84f36f2def
commit 97b4fec1ff
13 changed files with 119 additions and 14 deletions

View File

@ -0,0 +1,25 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#pragma clang diagnostic ignored "-Wunused-variable"
#include <metal_stdlib>
#include <simd/simd.h>
#include <metal_atomic>
using namespace metal;
struct SSBO
{
float4 outdata;
};
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)
kernel void main0(device SSBO& _31 [[buffer(1)]], texture2d<uint> uImage [[texture(0)]], device atomic_uint* uImage_atomic [[buffer(0)]], texture2d<float> uTexture [[texture(1)]], sampler uTextureSmplr [[sampler(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
uint _26 = atomic_fetch_add_explicit((device atomic_uint*)&uImage_atomic[spvImage2DAtomicCoord(int2(gl_GlobalInvocationID.xy), uImage)], 10u, memory_order_relaxed);
_31.outdata = uTexture.sample(uTextureSmplr, float2(gl_GlobalInvocationID.xy), level(0.0)) + float4(float(_26));
}

View File

@ -68,7 +68,7 @@ struct main0_in
uint in_var_TEXCOORD0 [[user(locn0)]]; uint in_var_TEXCOORD0 [[user(locn0)]];
}; };
fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d<uint> RWShadowTileNumCulledObjects [[texture(2)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]]) fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d<uint> RWShadowTileNumCulledObjects [[texture(0)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]])
{ {
main0_out out = {}; main0_out out = {};
uint2 _77 = uint2(gl_FragCoord.xy); uint2 _77 = uint2(gl_FragCoord.xy);

View File

@ -68,7 +68,7 @@ struct main0_in
uint in_var_TEXCOORD0 [[user(locn0)]]; uint in_var_TEXCOORD0 [[user(locn0)]];
}; };
fragment main0_out main0(main0_in in [[stage_in]], constant uint* spvBufferSizeConstants [[buffer(25)]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d<uint> RWShadowTileNumCulledObjects [[texture(2)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]]) fragment main0_out main0(main0_in in [[stage_in]], constant uint* spvBufferSizeConstants [[buffer(25)]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d<uint> RWShadowTileNumCulledObjects [[texture(0)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]])
{ {
main0_out out = {}; main0_out out = {};
constant uint& CulledObjectBoxBoundsBufferSize = spvBufferSizeConstants[0]; constant uint& CulledObjectBoxBoundsBufferSize = spvBufferSizeConstants[0];

View File

@ -0,0 +1,26 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#pragma clang diagnostic ignored "-Wunused-variable"
#include <metal_stdlib>
#include <simd/simd.h>
#include <metal_atomic>
using namespace metal;
struct SSBO
{
float4 outdata;
};
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)
kernel void main0(device SSBO& _31 [[buffer(1)]], texture2d<uint> uImage [[texture(0)]], device atomic_uint* uImage_atomic [[buffer(0)]], texture2d<float> uTexture [[texture(1)]], sampler uTextureSmplr [[sampler(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
uint _26 = atomic_fetch_add_explicit((device atomic_uint*)&uImage_atomic[spvImage2DAtomicCoord(int2(gl_GlobalInvocationID.xy), uImage)], 10u, memory_order_relaxed);
uint ret = _26;
_31.outdata = uTexture.sample(uTextureSmplr, float2(gl_GlobalInvocationID.xy), level(0.0)) + float4(float(ret));
}

View File

@ -33,8 +33,12 @@ void main()
_4.u16 = uint16_t(_30); _4.u16 = uint16_t(_30);
_4.f32 = float(_31); _4.f32 = float(_31);
_4.f32 = float(int16_t(_32)); _4.f32 = float(int16_t(_32));
_4.f32 = float(_29);
_4.f32 = float(int(_30));
_4.f32 = float(uint16_t(_31)); _4.f32 = float(uint16_t(_31));
_4.f32 = float(_32); _4.f32 = float(_32);
_4.f32 = float(uint(_29));
_4.f32 = float(_30);
_4.s16 = int16_t(_33); _4.s16 = int16_t(_33);
_4.u16 = uint16_t(int16_t(_33)); _4.u16 = uint16_t(int16_t(_33));
_4.u16 = uint16_t(_33); _4.u16 = uint16_t(_33);

View File

@ -68,7 +68,7 @@ struct main0_in
uint in_var_TEXCOORD0 [[user(locn0)]]; uint in_var_TEXCOORD0 [[user(locn0)]];
}; };
fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d<uint> RWShadowTileNumCulledObjects [[texture(2)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]]) fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d<uint> RWShadowTileNumCulledObjects [[texture(0)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]])
{ {
main0_out out = {}; main0_out out = {};
uint2 _77 = uint2(gl_FragCoord.xy); uint2 _77 = uint2(gl_FragCoord.xy);

View File

@ -68,7 +68,7 @@ struct main0_in
uint in_var_TEXCOORD0 [[user(locn0)]]; uint in_var_TEXCOORD0 [[user(locn0)]];
}; };
fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d<uint> RWShadowTileNumCulledObjects [[texture(2)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]]) fragment main0_out main0(main0_in in [[stage_in]], const device type_StructuredBuffer_v4float& CulledObjectBoxBounds [[buffer(0)]], constant type_Globals& _Globals [[buffer(1)]], texture2d<uint> RWShadowTileNumCulledObjects [[texture(0)]], device atomic_uint* RWShadowTileNumCulledObjects_atomic [[buffer(2)]], float4 gl_FragCoord [[position]])
{ {
main0_out out = {}; main0_out out = {};
uint2 _77 = uint2(gl_FragCoord.xy); uint2 _77 = uint2(gl_FragCoord.xy);

View File

@ -0,0 +1,16 @@
#version 450
layout(local_size_x = 1) in;
layout(set = 0, binding = 0, r32ui) uniform uimage2D uImage;
layout(set = 0, binding = 1) uniform sampler2D uTexture;
layout(set = 0, binding = 2) buffer SSBO
{
vec4 outdata;
};
void main()
{
uint ret = imageAtomicAdd(uImage, ivec2(gl_GlobalInvocationID.xy), 10u);
outdata = textureLod(uTexture, vec2(gl_GlobalInvocationID.xy), 0.0) + float(ret);
}

View File

@ -111,11 +111,23 @@
OpStore %ptr_f32 %s16_to_f32_signed OpStore %ptr_f32 %s16_to_f32_signed
%u16_to_f32_signed = OpConvertSToF %float %u16 %u16_to_f32_signed = OpConvertSToF %float %u16
OpStore %ptr_f32 %u16_to_f32_signed OpStore %ptr_f32 %u16_to_f32_signed
%s32_to_f32_signed = OpConvertSToF %float %s32
OpStore %ptr_f32 %s32_to_f32_signed
%u32_to_f32_signed = OpConvertSToF %float %u32
OpStore %ptr_f32 %u32_to_f32_signed
; UToF
%s16_to_f32_unsigned = OpConvertUToF %float %s16 %s16_to_f32_unsigned = OpConvertUToF %float %s16
OpStore %ptr_f32 %s16_to_f32_unsigned OpStore %ptr_f32 %s16_to_f32_unsigned
%u16_to_f32_unsigned = OpConvertUToF %float %u16 %u16_to_f32_unsigned = OpConvertUToF %float %u16
OpStore %ptr_f32 %u16_to_f32_unsigned OpStore %ptr_f32 %u16_to_f32_unsigned
%s32_to_f32_unsigned = OpConvertUToF %float %s32
OpStore %ptr_f32 %s32_to_f32_unsigned
%u32_to_f32_unsigned = OpConvertUToF %float %u32
OpStore %ptr_f32 %u32_to_f32_unsigned
; FToS ; FToS
%f32_to_s16_signed = OpConvertFToS %short %f32 %f32_to_s16_signed = OpConvertFToS %short %f32
OpStore %ptr_s16 %f32_to_s16_signed OpStore %ptr_s16 %f32_to_s16_signed

View File

@ -315,6 +315,10 @@ public:
const std::string &get_cleansed_entry_point_name(const std::string &name, const std::string &get_cleansed_entry_point_name(const std::string &name,
spv::ExecutionModel execution_model) const; spv::ExecutionModel execution_model) const;
// Traverses all reachable opcodes and sets active_builtins to a bitmask of all builtin variables which are accessed in the shader.
void update_active_builtins();
bool has_active_builtin(spv::BuiltIn builtin, spv::StorageClass storage);
// Query and modify OpExecutionMode. // Query and modify OpExecutionMode.
const Bitset &get_execution_mode_bitset() const; const Bitset &get_execution_mode_bitset() const;
@ -833,10 +837,6 @@ protected:
uint32_t cull_distance_count = 0; uint32_t cull_distance_count = 0;
bool position_invariant = false; bool position_invariant = false;
// Traverses all reachable opcodes and sets active_builtins to a bitmask of all builtin variables which are accessed in the shader.
void update_active_builtins();
bool has_active_builtin(spv::BuiltIn builtin, spv::StorageClass storage);
void analyze_parameter_preservation( void analyze_parameter_preservation(
SPIRFunction &entry, const CFG &cfg, SPIRFunction &entry, const CFG &cfg,
const std::unordered_map<uint32_t, std::unordered_set<uint32_t>> &variable_to_blocks, const std::unordered_map<uint32_t, std::unordered_set<uint32_t>> &variable_to_blocks,

View File

@ -9104,9 +9104,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
auto &arg_type = expression_type(ops[2]); auto &arg_type = expression_type(ops[2]);
auto func = type_to_glsl_constructor(type); auto func = type_to_glsl_constructor(type);
// If we're sign-extending or zero-extending, we need to make sure we cast from the correct type. if (arg_type.width < type.width || type_is_floating_point(type))
// For truncation, it does not matter, so don't emit useless casts.
if (arg_type.width < type.width)
emit_unary_func_op_cast(result_type, id, ops[2], func.c_str(), input_type, type.basetype); emit_unary_func_op_cast(result_type, id, ops[2], func.c_str(), input_type, type.basetype);
else else
emit_unary_func_op(result_type, id, ops[2], func.c_str()); emit_unary_func_op(result_type, id, ops[2], func.c_str());
@ -12119,7 +12117,12 @@ void CompilerGLSL::emit_block_chain(SPIRBlock &block)
SPIRBlock::ContinueBlockType continue_type = SPIRBlock::ContinueNone; SPIRBlock::ContinueBlockType continue_type = SPIRBlock::ContinueNone;
if (block.continue_block) if (block.continue_block)
{
continue_type = continue_block_type(get<SPIRBlock>(block.continue_block)); continue_type = continue_block_type(get<SPIRBlock>(block.continue_block));
// If we know we cannot emit a loop, mark the block early as a complex loop so we don't force unnecessary recompiles.
if (continue_type == SPIRBlock::ComplexLoop)
block.complex_continue = true;
}
// If we have loop variables, stop masking out access to the variable now. // If we have loop variables, stop masking out access to the variable now.
for (auto var_id : block.loop_variables) for (auto var_id : block.loop_variables)

View File

@ -9627,9 +9627,14 @@ uint32_t CompilerMSL::get_metal_resource_index(SPIRVariable &var, SPIRType::Base
// If a matching binding has been specified, find and use it. // If a matching binding has been specified, find and use it.
auto itr = resource_bindings.find({ execution.model, var_desc_set, var_binding }); auto itr = resource_bindings.find({ execution.model, var_desc_set, var_binding });
auto resource_decoration = var_type.basetype == SPIRType::SampledImage && basetype == SPIRType::Sampler ? // Atomic helper buffers for image atomics need to use secondary bindings as well.
SPIRVCrossDecorationResourceIndexSecondary : bool use_secondary_binding = (var_type.basetype == SPIRType::SampledImage && basetype == SPIRType::Sampler) ||
SPIRVCrossDecorationResourceIndexPrimary; basetype == SPIRType::AtomicCounter;
auto resource_decoration = use_secondary_binding ?
SPIRVCrossDecorationResourceIndexSecondary :
SPIRVCrossDecorationResourceIndexPrimary;
if (plane == 1) if (plane == 1)
resource_decoration = SPIRVCrossDecorationResourceIndexTertiary; resource_decoration = SPIRVCrossDecorationResourceIndexTertiary;
if (plane == 2) if (plane == 2)

View File

@ -1,5 +1,19 @@
#!/usr/bin/env python3 #!/usr/bin/env python3
# Copyright 2015-2019 Arm Limited
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
import sys import sys
import os import os
import os.path import os.path