AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Merge pull request #1489 from cdavis5e/msl-fix-atomic-image-coord

Browse Source 
MSL: Fix calculation of atomic image buffer address.
This commit is contained in:
Hans-Kristian Arntzen 2020-10-14 10:44:59 +02:00 committed by GitHub
commit 7332b44c3c
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
14 changed files with 78 additions and 13 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
CMakeLists.txt
View File

@ -323,7 +323,7 @@ if (SPIRV_CROSS_STATIC)
endif()
set(spirv-cross-abi-major 0)
set(spirv-cross-abi-minor 39)
set(spirv-cross-abi-minor 40)
set(spirv-cross-abi-patch 0)
if (SPIRV_CROSS_SHARED)

14
main.cpp
View File

@ -563,6 +563,8 @@ struct CLIArguments
bool msl_vertex_for_tessellation = false;
uint32_t msl_additional_fixed_sample_mask = 0xffffffff;
bool msl_arrayed_subpass_input = false;
uint32_t msl_r32ui_linear_texture_alignment = 4;
uint32_t msl_r32ui_alignment_constant_id = 65535;
bool glsl_emit_push_constant_as_ubo = false;
bool glsl_emit_ubo_as_plain_uniforms = false;
bool glsl_force_flattened_io_blocks = false;
@ -768,7 +770,11 @@ static void print_help_msl()
"\t[--msl-additional-fixed-sample-mask <mask>]:\n"
"\t\tSet an additional fixed sample mask. If the shader outputs a sample mask, then the final sample mask will be a bitwise AND of the two.\n"
"\t[--msl-arrayed-subpass-input]:\n\t\tAssume that images of dimension SubpassData have multiple layers. Layered input attachments are accessed relative to BuiltInLayer.\n"
"\t\tThis option has no effect if multiview is also enabled.\n");
"\t\tThis option has no effect if multiview is also enabled.\n"
"\t[--msl-r32ui-linear-texture-align <alignment>]:\n\t\tThe required alignment of linear textures of format MTLPixelFormatR32Uint.\n"
"\t\tThis is used to align the row stride for atomic accesses to such images.\n"
"\t[--msl-r32ui-linear-texture-align-constant-id <id>]:\n\t\tThe function constant ID to use for the linear texture alignment.\n"
"\t\tOn MSL 1.2 or later, you can override the alignment by setting this function constant.\n");
// clang-format on
}
@ -1007,6 +1013,8 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
msl_opts.vertex_for_tessellation = args.msl_vertex_for_tessellation;
msl_opts.additional_fixed_sample_mask = args.msl_additional_fixed_sample_mask;
msl_opts.arrayed_subpass_input = args.msl_arrayed_subpass_input;
msl_opts.r32ui_linear_texture_alignment = args.msl_r32ui_linear_texture_alignment;
msl_opts.r32ui_alignment_constant_id = args.msl_r32ui_alignment_constant_id;
msl_comp->set_msl_options(msl_opts);
for (auto &v : args.msl_discrete_descriptor_sets)
msl_comp->add_discrete_descriptor_set(v);
@ -1427,6 +1435,10 @@ static int main_inner(int argc, char *argv[])
cbs.add("--msl-additional-fixed-sample-mask",
[&args](CLIParser &parser) { args.msl_additional_fixed_sample_mask = parser.next_hex_uint(); });
cbs.add("--msl-arrayed-subpass-input", [&args](CLIParser &) { args.msl_arrayed_subpass_input = true; });
cbs.add("--msl-r32ui-linear-texture-align",
[&args](CLIParser &parser) { args.msl_r32ui_linear_texture_alignment = parser.next_uint(); });
cbs.add("--msl-r32ui-linear-texture-align-constant-id",
[&args](CLIParser &parser) { args.msl_r32ui_alignment_constant_id = parser.next_uint(); });
cbs.add("--extension", [&args](CLIParser &parser) { args.extensions.push_back(parser.next_string()); });
cbs.add("--rename-entry-point", [&args](CLIParser &parser) {
auto old_name = parser.next_string();

5
reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.argument.msl2.comp
View File

@ -23,8 +23,11 @@ struct spvDescriptorSetBuffer0
sampler uTextureSmplr [[id(4)]];
};
// The required alignment of a linear texture of R32Uint format.
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{

5
reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.comp
View File

@ -14,8 +14,11 @@ struct SSBO
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
// The required alignment of a linear texture of R32Uint format.
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
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]])
{

5
reference/opt/shaders-msl/frag/pixel-interlock-ordered.msl2.argument.frag
View File

@ -35,8 +35,11 @@ struct spvDescriptorSetBuffer0
device Buffer2* m_52 [[id(7), raster_order_group(0)]];
};
// The required alignment of a linear texture of R32Uint format.
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
fragment void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]])
{

5
reference/opt/shaders-msl/frag/pixel-interlock-ordered.msl2.frag
View File

@ -23,8 +23,11 @@ struct Buffer2
uint quux;
};
// The required alignment of a linear texture of R32Uint format.
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
fragment void main0(device Buffer3& _9 [[buffer(0)]], volatile device Buffer& _42 [[buffer(2), raster_order_group(0)]], device Buffer2& _52 [[buffer(3), raster_order_group(0)]], texture2d<float, access::write> img4 [[texture(0)]], texture2d<float, access::write> img [[texture(1), raster_order_group(0)]], texture2d<float> img3 [[texture(2), raster_order_group(0)]], texture2d<uint> img2 [[texture(3), raster_order_group(0)]], device atomic_uint* img2_atomic [[buffer(1), raster_order_group(0)]])
{

5
reference/shaders-msl/comp/image-atomic-automatic-bindings.argument.msl2.comp
View File

@ -23,8 +23,11 @@ struct spvDescriptorSetBuffer0
sampler uTextureSmplr [[id(4)]];
};
// The required alignment of a linear texture of R32Uint format.
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{

5
reference/shaders-msl/comp/image-atomic-automatic-bindings.comp
View File

@ -14,8 +14,11 @@ struct SSBO
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
// The required alignment of a linear texture of R32Uint format.
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
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]])
{

5
reference/shaders-msl/frag/pixel-interlock-ordered.msl2.argument.frag
View File

@ -35,8 +35,11 @@ struct spvDescriptorSetBuffer0
device Buffer2* m_52 [[id(7), raster_order_group(0)]];
};
// The required alignment of a linear texture of R32Uint format.
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
fragment void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]])
{

5
reference/shaders-msl/frag/pixel-interlock-ordered.msl2.frag
View File

@ -23,8 +23,11 @@ struct Buffer2
uint quux;
};
// The required alignment of a linear texture of R32Uint format.
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
fragment void main0(device Buffer3& _9 [[buffer(0)]], volatile device Buffer& _42 [[buffer(2), raster_order_group(0)]], device Buffer2& _52 [[buffer(3), raster_order_group(0)]], texture2d<float, access::write> img4 [[texture(0)]], texture2d<float, access::write> img [[texture(1), raster_order_group(0)]], texture2d<float> img3 [[texture(2), raster_order_group(0)]], texture2d<uint> img2 [[texture(3), raster_order_group(0)]], device atomic_uint* img2_atomic [[buffer(1), raster_order_group(0)]])
{

8
spirv_cross_c.cpp
View File

@ -666,6 +666,14 @@ spvc_result spvc_compiler_options_set_uint(spvc_compiler_options options, spvc_c
case SPVC_COMPILER_OPTION_MSL_ARRAYED_SUBPASS_INPUT:
options->msl.arrayed_subpass_input = value != 0;
break;
case SPVC_COMPILER_OPTION_MSL_R32UI_LINEAR_TEXTURE_ALIGNMENT:
options->msl.r32ui_linear_texture_alignment = value;
break;
case SPVC_COMPILER_OPTION_MSL_R32UI_ALIGNMENT_CONSTANT_ID:
options->msl.r32ui_alignment_constant_id = value;
break;
#endif
default:

4
spirv_cross_c.h
View File

@ -33,7 +33,7 @@ extern "C" {
/* Bumped if ABI or API breaks backwards compatibility. */
#define SPVC_C_API_VERSION_MAJOR 0
/* Bumped if APIs or enumerations are added in a backwards compatible way. */
#define SPVC_C_API_VERSION_MINOR 39
#define SPVC_C_API_VERSION_MINOR 40
/* Bumped if internal implementation details change. */
#define SPVC_C_API_VERSION_PATCH 0
@ -638,6 +638,8 @@ typedef enum spvc_compiler_option
SPVC_COMPILER_OPTION_MSL_MULTIVIEW_LAYERED_RENDERING = 67 | SPVC_COMPILER_OPTION_MSL_BIT,
SPVC_COMPILER_OPTION_MSL_ARRAYED_SUBPASS_INPUT = 68 | SPVC_COMPILER_OPTION_MSL_BIT,
SPVC_COMPILER_OPTION_MSL_R32UI_LINEAR_TEXTURE_ALIGNMENT = 69 | SPVC_COMPILER_OPTION_MSL_BIT,
SPVC_COMPILER_OPTION_MSL_R32UI_ALIGNMENT_CONSTANT_ID = 70 | SPVC_COMPILER_OPTION_MSL_BIT,
SPVC_COMPILER_OPTION_INT_MAX = 0x7fffffff
} spvc_compiler_option;

21
spirv_msl.cpp
View File

@ -4194,8 +4194,25 @@ void CompilerMSL::emit_custom_functions()
// Emulate texture2D atomic operations
case SPVFuncImplImage2DAtomicCoords:
{
if (msl_options.supports_msl_version(1, 2))
{
statement("// The required alignment of a linear texture of R32Uint format.");
statement("constant uint spvLinearTextureAlignmentOverride [[function_constant(",
msl_options.r32ui_alignment_constant_id, ")]];");
statement("constant uint spvLinearTextureAlignment = ",
"is_function_constant_defined(spvLinearTextureAlignmentOverride) ? ",
"spvLinearTextureAlignmentOverride : ", msl_options.r32ui_linear_texture_alignment, ";");
}
else
{
statement("// The required alignment of a linear texture of R32Uint format.");
statement("constant uint spvLinearTextureAlignment = ", msl_options.r32ui_linear_texture_alignment,
";");
}
statement("// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics");
statement("#define spvImage2DAtomicCoord(tc, tex) (((tex).get_width() * (tc).x) + (tc).y)");
statement("#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + ",
" spvLinearTextureAlignment / 4 - 1) & ~(",
" spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)");
statement("");
break;
}
@ -13516,7 +13533,7 @@ void CompilerMSL::analyze_argument_buffers()
{
uint32_t buffer_resource_index = get_metal_resource_index(var, SPIRType::AtomicCounter, 0);
resources_in_set[desc_set].push_back(
{ &var, to_name(var_id) + "_atomic", SPIRType::Struct, buffer_resource_index, 0 });
{ &var, to_name(var_id) + "_atomic", SPIRType::Struct, buffer_resource_index, 0 });
}
}
}

2
spirv_msl.hpp
View File

@ -268,6 +268,8 @@ public:
Platform platform = macOS;
uint32_t msl_version = make_msl_version(1, 2);
uint32_t texel_buffer_texture_width = 4096; // Width of 2D Metal textures used as 1D texel buffers
uint32_t r32ui_linear_texture_alignment = 4;
uint32_t r32ui_alignment_constant_id = 65535;
uint32_t swizzle_buffer_index = 30;
uint32_t indirect_params_buffer_index = 29;
uint32_t shader_output_buffer_index = 28;