When cycles is activated using the M1 Metal GPU blender quits with the error below. The changes made to Cycles code to introduce OptiX temporal denoising support seems to have caused it.
commit: 8393ccd07634b3152b18d4d527b1460dab9dbe06
Failed to compile library:
program_source:123688:5: error: expected ';' at end of declaration list
ccl_gpu_kernel_signature(filter_guiding_preprocess,
^
program_source:133:3: note: expanded from macro 'ccl_gpu_kernel_signature'
void run(thread MetalKernelContext& context, \
^
program_source:123688:5: error: no member named 'run' in 'kernel_gpu_filter_guiding_preprocess'
ccl_gpu_kernel_signature(filter_guiding_preprocess,
^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
program_source:155:18: note: expanded from macro 'ccl_gpu_kernel_signature'
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); \
~~~~~~~~~~~~~ ^
program_source:123688:5: error: out-of-line definition of 'run' does not match any declaration in 'kernel_gpu_filter_guiding_preprocess'
ccl_gpu_kernel_signature(filter_guiding_preprocess,
^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
program_source:157:25: note: expanded from macro 'ccl_gpu_kernel_signature'
void kernel_gpu_##name::run(thread MetalKernelContext& context, \
^~~
program_source:123709:30: error: use of undeclared identifier 'width'
const int y = work_index / width;
^
program_source:123710:34: error: use of undeclared identifier 'width'
const int x = work_index - y * width;
^
program_source:123712:12: error: use of undeclared identifier 'width'
if (x >= width || y >= height) {
^
program_source:123712:26: error: reference to non-static member function must be called
if (x >= width || y >= height) {
^~~~~~
program_source:123716:48: error: use of undeclared identifier 'width'
const uint64_t guiding_pixel_index = x + y * width;
^
program_source:123717:37: error: use of undeclared identifier 'guiding_buffer'
ccl_global float *guiding_pixel = guiding_buffer + guiding_pixel_index * guiding_pass_stride;
^
program_source:123717:76: error: use of undeclared identifier 'guiding_pass_stride'
ccl_global float *guiding_pixel = guiding_buffer + guiding_pixel_index * guiding_pass_stride;
^
program_source:123719:39: error: use of undeclared identifier 'render_offset'
const uint64_t render_pixel_index = render_offset + (x + full_x) + (y + full_y) * render_stride;
^
program_source:123719:60: error: use of undeclared identifier 'full_x'
const uint64_t render_pixel_index = render_offset + (x + full_x) + (y + full_y) * render_stride;
^
program_source:123719:75: error: use of undeclared identifier 'full_y'
const uint64_t render_pixel_index = render_offset + (x + full_x) + (y + full_y) * render_stride;
^
program_source:123719:85: error: use of undeclared identifier 'render_stride'
const uint64_t render_pixel_index = render_offset + (x + full_x) + (y + full_y) * render_stride;
^
program_source:123720:36: error: use of undeclared identifier 'render_buffer'
ccl_global const float *buffer = render_buffer + render_pixel_index * render_pass_stride;
^
program_source:123720:73: error: use of undeclared identifier 'render_pass_stride'
ccl_global const float *buffer = render_buffer + render_pixel_index * render_pass_stride;
^
program_source:123723:7: error: use of undeclared identifier 'render_pass_sample_count'
if (render_pass_sample_count == PASS_UNUSED) {
^
program_source:123724:26: error: use of undeclared identifier 'num_samples'
pixel_scale = 1.0f / num_samples;
^
program_source:123727:49: error: use of undeclared identifier 'render_pass_sample_count'
pixel_scale = 1.0f / __float_as_uint(buffer[render_pass_sample_count]);
^
program_source:123731:7: error: use of undeclared identifier 'guiding_pass_albedo'
if (guiding_pass_albedo != PASS_UNUSED) {
^
program_source:123734:49: error: use of undeclared identifier 'render_pass_denoising_albedo'
ccl_global const float *aledo_in = buffer + render_pass_denoising_albedo;
^
program_source:123735:52: error: use of undeclared identifier 'guiding_pass_albedo'
ccl_global float *albedo_out = guiding_pixel + guiding_pass_albedo;
^
program_source:123743:7: error: use of undeclared identifier 'guiding_pass_normal'
if (guiding_pass_normal != PASS_UNUSED) {
^
program_source:123746:50: error: use of undeclared identifier 'render_pass_denoising_normal'
ccl_global const float *normal_in = buffer + render_pass_denoising_normal;
^
program_source:123747:52: error: use of undeclared identifier 'guiding_pass_normal'
ccl_global float *normal_out = guiding_pixel + guiding_pass_normal;
^
program_source:123755:7: error: use of undeclared identifier 'guiding_pass_flow'
if (guiding_pass_flow != PASS_UNUSED) {
^
program_source:123758:17: error: pointer type must have explicit address space qualifier
const float *motion_in = buffer + render_pass_motion;
^
program_source:123758:39: error: use of undeclared identifier 'render_pass_motion'
const float *motion_in = buffer + render_pass_motion;
^
program_source:123759:11: error: pointer type must have explicit address space qualifier
float *flow_out = guiding_pixel + guiding_pass_flow;
^
program_source:123759:39: error: use of undeclared identifier 'guiding_pass_flow'
float *flow_out = guiding_pixel + guiding_pass_flow;