Hello, in kernel function tessellation_main, i don’t see any boundary check on pid. But it looks like the thread grid size (threadgroupsPerGrid * threadsPerThreadgroup) doesn’t match the data size (tessellation factors buffer) . The program runs without any issue on my mac so I think my understanding about MTLComputeCommandEncoder.dispatchThreadgroups maybe wrong 
Compute is only introduced properly in chapter 16 where bounds checking is also covered. In this chapter compute was only used minimally.