diff options
| author | Michael Jones <michael_jones> | 2022-01-07 18:28:43 +0300 |
|---|---|---|
| committer | Michael Jones <michael_p_jones@apple.com> | 2022-01-07 19:19:31 +0300 |
| commit | efe3d60a2c8306aefd41bc304548da35b67c252c (patch) | |
| tree | 1b789ceaf37bbee5e4237f12febb69211a95c33c | |
| parent | b7ea6e9d0ef910fe668d0fa3454ac94400471804 (diff) | |
Cycles: Fix Metal build
This patch fixes a couple of new Metal kernel compilation errors: 1) a kernel parameter count overflow, and 2) missing address space qualifiers.
Reviewed By: brecht
Differential Revision: https://developer.blender.org/D13763
| -rw-r--r-- | intern/cycles/kernel/device/gpu/kernel.h | 4 | ||||
| -rw-r--r-- | intern/cycles/kernel/device/metal/compat.h | 8 |
2 files changed, 8 insertions, 4 deletions
diff --git a/intern/cycles/kernel/device/gpu/kernel.h b/intern/cycles/kernel/device/gpu/kernel.h index 027b2a7a8c7..00c727b48cb 100644 --- a/intern/cycles/kernel/device/gpu/kernel.h +++ b/intern/cycles/kernel/device/gpu/kernel.h @@ -821,8 +821,8 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS) if (guiding_pass_flow != PASS_UNUSED) { kernel_assert(render_pass_motion != PASS_UNUSED); - const float *motion_in = buffer + render_pass_motion; - float *flow_out = guiding_pixel + guiding_pass_flow; + ccl_global const float *motion_in = buffer + render_pass_motion; + ccl_global float *flow_out = guiding_pixel + guiding_pass_flow; flow_out[0] = -motion_in[0] * pixel_scale; flow_out[1] = -motion_in[1] * pixel_scale; diff --git a/intern/cycles/kernel/device/metal/compat.h b/intern/cycles/kernel/device/metal/compat.h index a51afc37fc0..1222b68f0ee 100644 --- a/intern/cycles/kernel/device/metal/compat.h +++ b/intern/cycles/kernel/device/metal/compat.h @@ -98,8 +98,12 @@ using namespace metal::raytracing; #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) +#define FN17(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, p17) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15; p16; p17; +#define FN18(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, p17, p18) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15; p16; p17; p18; +#define FN19(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, p17, p18, p19) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15; p16; p17; p18; p19; +#define FN20(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, p17, p18, p19, p20) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15; p16; p17; p18; p19; p20; +#define GET_LAST_ARG(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, p17, p18, p19, p20, ...) p20 +#define PARAMS_MAKER(...) GET_LAST_ARG(__VA_ARGS__, FN20, FN19, FN18, FN17, FN16, FN15, FN14, FN13, FN12, FN11, FN10, FN9, FN8, FN7, FN6, FN5, FN4, FN3, FN2, FN1, FN0) /* Generate a struct containing the entry-point parameters and a "run" * method which can access them implicitly via this-> */ |