intern/cycles/device/metal/device_impl.mm

@@ -567,6 +567,10 @@ void MetalDevice::compile_and_load(int device_id, MetalPipelineType pso_type)
     thread_scoped_lock lock(existing_devices_mutex);
     if (MetalDevice *instance = get_device_by_ID(device_id, lock)) {
       if (mtlLibrary) {
+        if (error && [error localizedDescription]) {
+          VLOG_WARNING << "MSL compilation messages: " << [[error localizedDescription] UTF8String];
+        }
+
         instance->mtlLibrary[pso_type] = mtlLibrary;
 
         starttime = time_dt();

intern/cycles/device/metal/kernel.mm

@@ -715,7 +715,7 @@ void MetalKernelPipeline::compile()
           }
         }
       };
-      if (computePipelineStateDescriptor.linkedFunctions) {
+      if (linked_functions) {
         addComputePipelineFunctionsWithDescriptor();
       }
 
@@ -748,7 +748,7 @@ void MetalKernelPipeline::compile()
       }
 
       /* Add pipeline into the new archive (unless we did it earlier). */
-      if (pipeline && !computePipelineStateDescriptor.linkedFunctions) {
+      if (pipeline && !linked_functions) {
         addComputePipelineFunctionsWithDescriptor();
       }
     }

intern/cycles/kernel/device/gpu/parallel_sorted_index.h

@@ -38,7 +38,7 @@ ccl_device_inline void gpu_parallel_sort_bucket_pass(const uint num_states,
                                                      ccl_gpu_shared int *buckets,
                                                      const ushort local_id,
                                                      const ushort local_size,
-                                                     const ushort grid_id)
+                                                     const uint grid_id)
 {
   /* Zero the bucket sizes. */
   if (local_id < max_shaders) {
@@ -89,7 +89,7 @@ ccl_device_inline void gpu_parallel_sort_write_pass(const uint num_states,
                                                     ccl_gpu_shared int *local_offset,
                                                     const ushort local_id,
                                                     const ushort local_size,
-                                                    const ushort grid_id)
+                                                    const uint grid_id)
 {
   /* Calculate each partition's global offset from the prefix sum of the active state counts per
    * partition. */

intern/cycles/kernel/device/metal/compat.h

@@ -109,7 +109,7 @@ struct kernel_gpu_##name \
            const uint metal_global_id, \
            const ushort metal_local_id, \
            const ushort metal_local_size, \
-           const ushort metal_grid_id, \
+           const uint metal_grid_id, \
            uint simdgroup_size, \
            uint simd_lane_index, \
            uint simd_group_index, \
@@ -122,7 +122,7 @@ kernel void cycles_metal_##name(device const kernel_gpu_##name *params_struct, \
                                 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]], \
-                                const ushort metal_grid_id    [[threadgroup_position_in_grid]], \
+                                const uint metal_grid_id    [[threadgroup_position_in_grid]], \
                                 uint simdgroup_size [[threads_per_simdgroup]], \
                                 uint simd_lane_index [[thread_index_in_simdgroup]], \
                                 uint simd_group_index [[simdgroup_index_in_threadgroup]], \
@@ -135,7 +135,7 @@ void kernel_gpu_##name::run(thread MetalKernelContext& context, \
                   const uint metal_global_id, \
                   const ushort metal_local_id, \
                   const ushort metal_local_size, \
-                  const ushort metal_grid_id, \
+                  const uint metal_grid_id, \
                   uint simdgroup_size, \
                   uint simd_lane_index, \
                   uint simd_group_index, \