From 3a4c8f406a3a3bf0627477c6183a594fa707a6e2 Mon Sep 17 00:00:00 2001 From: Michael Jones Date: Tue, 9 Nov 2021 21:30:46 +0000 Subject: Cycles: Adapt shared kernel/device/gpu layer for MSL This patch adapts the shared kernel entrypoints so that they can be compiled as MSL (Metal Shading Language). Where possible, the adaptations avoid changes in common code. In MSL, kernel function inputs are explicitly bound to resources. In the case of argument buffers, we declare a struct containing the kernel arguments, accessible via device pointer. This differs from CUDA and HIP where kernel function arguments are declared as traditional C-style function parameters. This patch adapts the entrypoints declared in kernel.h so that they can be translated via a new `ccl_gpu_kernel_signature` macro into the required parameter struct + kernel entrypoint pairing for MSL. MSL buffer attribution must be applied to function parameters or non-static class data members. To allow universal access to the integrator state, kernel data, and texture fetch adapters, we wrap all of the shared kernel code in a `MetalKernelContext` class. This is achieved by bracketing the appropriate kernel headers with "context_begin.h" and "context_end.h" on Metal. When calling deeper into the kernel code, we must reference the context class (e.g. `context.integrator_init_from_camera`). This extra prefixing is performed by a set of defines in "context_end.h". These will require explicit maintenance if entrypoints change. We invite discussion on more maintainable ways to enforce correctness. Lambda expressions are not supported on MSL, so a new `ccl_gpu_kernel_lambda` macro generates an inline function object and optionally capturing any required state. This yields the same behaviour. This approach is applied to all parallel_... implementations which are templated by operation. The lambda expressions in the film_convert... kernels don't adapt cleanly to use function objects. However, these entrypoints can be macro-generated more concisely to avoid lambda expressions entirely, instead relying on constant folding to handle the pixel/channel conversions. A separate implementation of `gpu_parallel_active_index_array` is provided for Metal to workaround some subtle differences in SIMD width, and also to encapsulate some required thread parameters which must be declared as explicit entrypoint function parameters. Ref T92212 Reviewed By: brecht Maniphest Tasks: T92212 Differential Revision: https://developer.blender.org/D13109 --- intern/cycles/kernel/device/metal/compat.h | 124 +++++++++++++++++++++++++++++ 1 file changed, 124 insertions(+) (limited to 'intern/cycles/kernel/device/metal/compat.h') diff --git a/intern/cycles/kernel/device/metal/compat.h b/intern/cycles/kernel/device/metal/compat.h index 77cea30914c..57c6845e508 100644 --- a/intern/cycles/kernel/device/metal/compat.h +++ b/intern/cycles/kernel/device/metal/compat.h @@ -58,6 +58,95 @@ using namespace metal; #define kernel_assert(cond) +#define ccl_gpu_global_id_x() metal_global_id +#define ccl_gpu_warp_size simdgroup_size +#define ccl_gpu_thread_idx_x simd_group_index +#define ccl_gpu_thread_mask(thread_warp) uint64_t((1ull << thread_warp) - 1) + +#define ccl_gpu_ballot(predicate) ((uint64_t)((simd_vote::vote_t)simd_ballot(predicate))) +#define ccl_gpu_popc(x) popcount(x) + +// clang-format off + +/* kernel.h adapters */ + +#define ccl_gpu_kernel(...) + +/* convert a comma-separated list into a semicolon-separated list (so that we can generate a struct based on kernel entrypoint parameters) */ +#define FN0() +#define FN1(p1) p1; +#define FN2(p1, p2) p1; p2; +#define FN3(p1, p2, p3) p1; p2; p3; +#define FN4(p1, p2, p3, p4) p1; p2; p3; p4; +#define FN5(p1, p2, p3, p4, p5) p1; p2; p3; p4; p5; +#define FN6(p1, p2, p3, p4, p5, p6) p1; p2; p3; p4; p5; p6; +#define FN7(p1, p2, p3, p4, p5, p6, p7) p1; p2; p3; p4; p5; p6; p7; +#define FN8(p1, p2, p3, p4, p5, p6, p7, p8) p1; p2; p3; p4; p5; p6; p7; p8; +#define FN9(p1, p2, p3, p4, p5, p6, p7, p8, p9) p1; p2; p3; p4; p5; p6; p7; p8; p9; +#define FN10(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; +#define FN11(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; +#define FN12(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; +#define FN13(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; +#define FN14(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; +#define FN15(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15; +#define FN16(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15; p16; +#define GET_LAST_ARG(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, ...) p16 +#define PARAMS_MAKER(...) GET_LAST_ARG(__VA_ARGS__, FN16, FN15, FN14, FN13, FN12, FN11, FN10, FN9, FN8, FN7, FN6, FN5, FN4, FN3, FN2, FN1, FN0) + +/* generate a struct containing the entrypoint parameters and a "run" method which can access them implicitly via this-> */ +#define ccl_gpu_kernel_signature(name, ...) \ +struct kernel_gpu_##name \ +{ \ + PARAMS_MAKER(__VA_ARGS__)(__VA_ARGS__) \ + void run(thread MetalKernelContext& context, \ + threadgroup int *simdgroup_offset, \ + const uint metal_global_id, \ + const ushort metal_local_id, \ + const ushort metal_local_size, \ + uint simdgroup_size, \ + uint simd_lane_index, \ + uint simd_group_index, \ + uint num_simd_groups) ccl_global const; \ +}; \ +kernel void kernel_metal_##name(device const kernel_gpu_##name *params_struct, \ + constant KernelParamsMetal &ccl_restrict _launch_params_metal, \ + constant MetalAncillaries *_metal_ancillaries, \ + threadgroup int *simdgroup_offset[[ threadgroup(0) ]], \ + const uint metal_global_id [[thread_position_in_grid]], \ + const ushort metal_local_id [[thread_position_in_threadgroup]], \ + const ushort metal_local_size [[threads_per_threadgroup]], \ + uint simdgroup_size [[threads_per_simdgroup]], \ + uint simd_lane_index [[thread_index_in_simdgroup]], \ + uint simd_group_index [[simdgroup_index_in_threadgroup]], \ + uint num_simd_groups [[simdgroups_per_threadgroup]]) { \ + MetalKernelContext context(_launch_params_metal, _metal_ancillaries); \ + INIT_DEBUG_BUFFER \ + params_struct->run(context, simdgroup_offset, metal_global_id, metal_local_id, metal_local_size, simdgroup_size, simd_lane_index, simd_group_index, num_simd_groups); \ +} \ +void kernel_gpu_##name::run(thread MetalKernelContext& context, \ + threadgroup int *simdgroup_offset, \ + const uint metal_global_id, \ + const ushort metal_local_id, \ + const ushort metal_local_size, \ + uint simdgroup_size, \ + uint simd_lane_index, \ + uint simd_group_index, \ + uint num_simd_groups) ccl_global const + +#define ccl_gpu_kernel_call(x) context.x + +/* define a function object where "func" is the lambda body, and additional parameters are used to specify captured state */ +#define ccl_gpu_kernel_lambda(func, ...) \ + struct KernelLambda \ + { \ + KernelLambda(ccl_private MetalKernelContext &_context) : context(_context) {} \ + ccl_private MetalKernelContext &context; \ + __VA_ARGS__; \ + int operator()(const int state) const { return (func); } \ + }ccl_gpu_kernel_lambda_pass(context); ccl_gpu_kernel_lambda_pass + +// clang-format on + /* make_type definitions with Metal style element initializers */ #ifdef make_float2 # undef make_float2 @@ -124,3 +213,38 @@ using namespace metal; #define logf(x) trigmode::log(float(x)) #define NULL 0 + +/* texture bindings and sampler setup */ + +struct Texture2DParamsMetal { + texture2d tex; +}; +struct Texture3DParamsMetal { + texture3d tex; +}; + +struct MetalAncillaries { + device Texture2DParamsMetal *textures_2d; + device Texture3DParamsMetal *textures_3d; +}; + +enum SamplerType { + SamplerFilterNearest_AddressRepeat, + SamplerFilterNearest_AddressClampEdge, + SamplerFilterNearest_AddressClampZero, + + SamplerFilterLinear_AddressRepeat, + SamplerFilterLinear_AddressClampEdge, + SamplerFilterLinear_AddressClampZero, + + SamplerCount +}; + +constant constexpr array metal_samplers = { + sampler(address::repeat, filter::nearest), + sampler(address::clamp_to_edge, filter::nearest), + sampler(address::clamp_to_zero, filter::nearest), + sampler(address::repeat, filter::linear), + sampler(address::clamp_to_edge, filter::linear), + sampler(address::clamp_to_zero, filter::linear), +}; \ No newline at end of file -- cgit v1.2.3 From f56562043521a5c160585aea3f28167b4d3bc77d Mon Sep 17 00:00:00 2001 From: Patrick Mours Date: Wed, 10 Nov 2021 14:37:15 +0100 Subject: Fix T92985: CUDA errors with Cycles film convert kernels rB3a4c8f406a3a3bf0627477c6183a594fa707a6e2 changed the macros that create the film convert kernel entry points, but in the process accidentally changed the parameter definition to one of those (which caused CUDA launch and misaligned address errors) and changed the implementation as well. This restores the correct implementation from before. In addition, the `ccl_gpu_kernel_threads` macro did not work as intended and caused the generated launch bounds to end up with an incorrect input for the second parameter (it was set to "thread_num_registers", rather than the result of the block number calculation). I'm not entirely sure why, as the macro definition looked sound to me. Decided to simply go with two separate macros instead, to simplify and solve this. Also changed how state is captured with the `ccl_gpu_kernel_lambda` macro slightly, to avoid a compiler warning (expression has no effect) that otherwise occurred. Maniphest Tasks: T92985 Differential Revision: https://developer.blender.org/D13175 --- intern/cycles/kernel/device/metal/compat.h | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) (limited to 'intern/cycles/kernel/device/metal/compat.h') diff --git a/intern/cycles/kernel/device/metal/compat.h b/intern/cycles/kernel/device/metal/compat.h index 57c6845e508..2fa9b7fed44 100644 --- a/intern/cycles/kernel/device/metal/compat.h +++ b/intern/cycles/kernel/device/metal/compat.h @@ -70,7 +70,8 @@ using namespace metal; /* kernel.h adapters */ -#define ccl_gpu_kernel(...) +#define ccl_gpu_kernel(block_num_threads, thread_num_registers) +#define ccl_gpu_kernel_threads(block_num_threads) /* convert a comma-separated list into a semicolon-separated list (so that we can generate a struct based on kernel entrypoint parameters) */ #define FN0() @@ -143,7 +144,7 @@ void kernel_gpu_##name::run(thread MetalKernelContext& context, \ ccl_private MetalKernelContext &context; \ __VA_ARGS__; \ int operator()(const int state) const { return (func); } \ - }ccl_gpu_kernel_lambda_pass(context); ccl_gpu_kernel_lambda_pass + } ccl_gpu_kernel_lambda_pass(context) // clang-format on @@ -247,4 +248,4 @@ constant constexpr array metal_samplers = { sampler(address::repeat, filter::linear), sampler(address::clamp_to_edge, filter::linear), sampler(address::clamp_to_zero, filter::linear), -}; \ No newline at end of file +}; -- cgit v1.2.3 From 1143bf281afc69b931f7d0eb1daa4b800dcc513d Mon Sep 17 00:00:00 2001 From: Campbell Barton Date: Sat, 13 Nov 2021 13:07:13 +1100 Subject: Cleanup: spelling in comments, comment block formatting --- intern/cycles/kernel/device/metal/compat.h | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) (limited to 'intern/cycles/kernel/device/metal/compat.h') diff --git a/intern/cycles/kernel/device/metal/compat.h b/intern/cycles/kernel/device/metal/compat.h index 2fa9b7fed44..a839917a907 100644 --- a/intern/cycles/kernel/device/metal/compat.h +++ b/intern/cycles/kernel/device/metal/compat.h @@ -73,7 +73,8 @@ using namespace metal; #define ccl_gpu_kernel(block_num_threads, thread_num_registers) #define ccl_gpu_kernel_threads(block_num_threads) -/* convert a comma-separated list into a semicolon-separated list (so that we can generate a struct based on kernel entrypoint parameters) */ +/* Convert a comma-separated list into a semicolon-separated list + * (so that we can generate a struct based on kernel entry-point parameters). */ #define FN0() #define FN1(p1) p1; #define FN2(p1, p2) p1; p2; @@ -94,7 +95,8 @@ using namespace metal; #define GET_LAST_ARG(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, ...) p16 #define PARAMS_MAKER(...) GET_LAST_ARG(__VA_ARGS__, FN16, FN15, FN14, FN13, FN12, FN11, FN10, FN9, FN8, FN7, FN6, FN5, FN4, FN3, FN2, FN1, FN0) -/* generate a struct containing the entrypoint parameters and a "run" method which can access them implicitly via this-> */ +/* Generate a struct containing the entry-point parameters and a "run" + * method which can access them implicitly via this-> */ #define ccl_gpu_kernel_signature(name, ...) \ struct kernel_gpu_##name \ { \ -- cgit v1.2.3 From 64003fa4b0b1699998a5b048d980eb775d547d8d Mon Sep 17 00:00:00 2001 From: Michael Jones Date: Tue, 16 Nov 2021 13:41:29 +0000 Subject: Cycles: Adapt volumetric lambda functions to work on MSL This patch adapts the existing volumetric read/write lambda functions for Metal. Lambda expressions are not supported on MSL, so two new macros `VOLUME_READ_LAMBDA` and `VOLUME_WRITE_LAMBDA` have been defined with a default implementation which, on Metal, is overridden to use inline function objects. This patch also removes the last remaining mention of the now-unused `ccl_addr_space`. Ref T92212 Reviewed By: leesonw Maniphest Tasks: T92212 Differential Revision: https://developer.blender.org/D13234 --- intern/cycles/kernel/device/metal/compat.h | 25 +++++++++++++++++++++++++ 1 file changed, 25 insertions(+) (limited to 'intern/cycles/kernel/device/metal/compat.h') diff --git a/intern/cycles/kernel/device/metal/compat.h b/intern/cycles/kernel/device/metal/compat.h index a839917a907..4a2c39d90fd 100644 --- a/intern/cycles/kernel/device/metal/compat.h +++ b/intern/cycles/kernel/device/metal/compat.h @@ -150,6 +150,31 @@ void kernel_gpu_##name::run(thread MetalKernelContext& context, \ // clang-format on +/* volumetric lambda functions - use function objects for lambda-like functionality */ +#define VOLUME_READ_LAMBDA(function_call) \ + struct FnObjectRead { \ + KernelGlobals kg; \ + ccl_private MetalKernelContext *context; \ + int state; \ +\ + VolumeStack operator()(const int i) const \ + { \ + return context->function_call; \ + } \ + } volume_read_lambda_pass{kg, this, state}; + +#define VOLUME_WRITE_LAMBDA(function_call) \ + struct FnObjectWrite { \ + KernelGlobals kg; \ + ccl_private MetalKernelContext *context; \ + int state; \ +\ + void operator()(const int i, VolumeStack entry) const \ + { \ + context->function_call; \ + } \ + } volume_write_lambda_pass{kg, this, state}; + /* make_type definitions with Metal style element initializers */ #ifdef make_float2 # undef make_float2 -- cgit v1.2.3 From 9937d5379ca936b4ba93534185477fa7e529181c Mon Sep 17 00:00:00 2001 From: Brecht Van Lommel Date: Tue, 16 Nov 2021 14:03:59 +0100 Subject: Cycles: add packed_float3 type for storage Introduce a packed_float3 type for smaller storage that is exactly 3 floats, instead of 4. For computation float3 is still used since it can use SIMD instructions. Ref T92212 Differential Revision: https://developer.blender.org/D13243 --- intern/cycles/kernel/device/metal/compat.h | 1 + 1 file changed, 1 insertion(+) (limited to 'intern/cycles/kernel/device/metal/compat.h') diff --git a/intern/cycles/kernel/device/metal/compat.h b/intern/cycles/kernel/device/metal/compat.h index 4a2c39d90fd..19358e063d8 100644 --- a/intern/cycles/kernel/device/metal/compat.h +++ b/intern/cycles/kernel/device/metal/compat.h @@ -42,6 +42,7 @@ using namespace metal; #define ccl_device_forceinline ccl_device #define ccl_device_noinline ccl_device __attribute__((noinline)) #define ccl_device_noinline_cpu ccl_device +#define ccl_device_inline_method ccl_device #define ccl_global device #define ccl_static_constant static constant constexpr #define ccl_device_constant constant -- cgit v1.2.3 From d19e35873f67c90b251ca38e007a83aa1eada211 Mon Sep 17 00:00:00 2001 From: Michael Jones Date: Thu, 18 Nov 2021 14:25:05 +0100 Subject: Cycles: several small fixes and additions for MSL This patch contains many small leftover fixes and additions that are required for Metal-enablement: - Address space fixes and a few other small compile fixes - Addition of missing functionality to the Metal adapter headers - Addition of various scattered `__KERNEL_METAL__` blocks (e.g. for atomic support & maths functions) Ref T92212 Differential Revision: https://developer.blender.org/D13263 --- intern/cycles/kernel/device/metal/compat.h | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) (limited to 'intern/cycles/kernel/device/metal/compat.h') diff --git a/intern/cycles/kernel/device/metal/compat.h b/intern/cycles/kernel/device/metal/compat.h index 19358e063d8..080109e3b83 100644 --- a/intern/cycles/kernel/device/metal/compat.h +++ b/intern/cycles/kernel/device/metal/compat.h @@ -34,6 +34,7 @@ using namespace metal; #pragma clang diagnostic ignored "-Wunused-variable" #pragma clang diagnostic ignored "-Wsign-compare" +#pragma clang diagnostic ignored "-Wuninitialized" /* Qualifiers */ @@ -65,7 +66,7 @@ using namespace metal; #define ccl_gpu_thread_mask(thread_warp) uint64_t((1ull << thread_warp) - 1) #define ccl_gpu_ballot(predicate) ((uint64_t)((simd_vote::vote_t)simd_ballot(predicate))) -#define ccl_gpu_popc(x) popcount(x) +#define ccl_gpu_syncthreads() threadgroup_barrier(mem_flags::mem_threadgroup); // clang-format off @@ -124,7 +125,6 @@ kernel void kernel_metal_##name(device const kernel_gpu_##name *params_struct, \ uint simd_group_index [[simdgroup_index_in_threadgroup]], \ uint num_simd_groups [[simdgroups_per_threadgroup]]) { \ MetalKernelContext context(_launch_params_metal, _metal_ancillaries); \ - INIT_DEBUG_BUFFER \ params_struct->run(context, simdgroup_offset, metal_global_id, metal_local_id, metal_local_size, simdgroup_size, simd_lane_index, simd_group_index, num_simd_groups); \ } \ void kernel_gpu_##name::run(thread MetalKernelContext& context, \ @@ -230,6 +230,7 @@ void kernel_gpu_##name::run(thread MetalKernelContext& context, \ #define sinhf(x) sinh(float(x)) #define coshf(x) cosh(float(x)) #define tanhf(x) tanh(float(x)) +#define saturatef(x) saturate(float(x)) /* Use native functions with possibly lower precision for performance, * no issues found so far. */ @@ -243,6 +244,8 @@ void kernel_gpu_##name::run(thread MetalKernelContext& context, \ #define NULL 0 +#define __device__ + /* texture bindings and sampler setup */ struct Texture2DParamsMetal { @@ -257,6 +260,9 @@ struct MetalAncillaries { device Texture3DParamsMetal *textures_3d; }; +#include "util/half.h" +#include "util/types.h" + enum SamplerType { SamplerFilterNearest_AddressRepeat, SamplerFilterNearest_AddressClampEdge, -- cgit v1.2.3 From d1f944c18634f215c3da0484ac3b80e994118680 Mon Sep 17 00:00:00 2001 From: Michael Jones Date: Thu, 18 Nov 2021 14:25:30 +0100 Subject: Cycles: declare constants at program scope on Metal MSL requires that constant address space literals be declared at program scope. This patch moves the `blackbody_table_r/g/b` and `cie_colour_match` constants into separate files so they can be declared at the appropriate scope. Ref T92212 Differential Revision: https://developer.blender.org/D13241 --- intern/cycles/kernel/device/metal/compat.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'intern/cycles/kernel/device/metal/compat.h') diff --git a/intern/cycles/kernel/device/metal/compat.h b/intern/cycles/kernel/device/metal/compat.h index 080109e3b83..61597a4acfc 100644 --- a/intern/cycles/kernel/device/metal/compat.h +++ b/intern/cycles/kernel/device/metal/compat.h @@ -45,7 +45,7 @@ using namespace metal; #define ccl_device_noinline_cpu ccl_device #define ccl_device_inline_method ccl_device #define ccl_global device -#define ccl_static_constant static constant constexpr +#define ccl_inline_constant static constant constexpr #define ccl_device_constant constant #define ccl_constant const device #define ccl_gpu_shared threadgroup -- cgit v1.2.3