GPU Metal: bound the threadgroup size when creating pipelines

Pass the tuned thread count as maxTotalThreadsPerThreadgroup, the Metal
equivalent of __launch_bounds__. Without it the register allocation can
leave room for fewer threads than the tuning table assumes, and the
launch was clamped: on an M1 Max (applegpu_g13s) the two TRD tracker
kernels got 384 instead of 512, and the TPC compression gather kernel
832 instead of 1024, which breaks its block-wide scans.

Author: ktf

GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm

@@ -64,17 +64,24 @@
     return 1;
   }
 
+  // The kernel sizes its threadgroup memory from the tuning table at compile
+  // time, so a launch must not exceed what it was built for either way.
+  const uint32_t tunedThreads = this->template getKernelProperties<T, I>().nThreads;
+
+  // Like __launch_bounds__: without it the register allocation may leave room
+  // for fewer threads than the kernel is tuned for.
+  MTLComputePipelineDescriptor* desc = [MTLComputePipelineDescriptor new];
+  desc.computeFunction = krnl;
+  desc.maxTotalThreadsPerThreadgroup = tunedThreads;
   NSError* error = nil;
-  id<MTLComputePipelineState> pso = [mInternals->device newComputePipelineStateWithFunction:krnl error:&error];
+  id<MTLComputePipelineState> pso = [mInternals->device newComputePipelineStateWithDescriptor:desc options:MTLPipelineOptionNone reflection:nil error:&error];
+  [desc release];
   if (pso == nil) {
     NSLog(@"%@", error);
     GPUError("Error creating Metal pipeline state: %s", [kname cStringUsingEncoding:NSUTF8StringEncoding]);
     return 1;
   }
 
-  // The kernel sizes its threadgroup memory from the tuning table at compile
-  // time, so a launch must not exceed what it was built for either way.
-  const uint32_t tunedThreads = this->template getKernelProperties<T, I>().nThreads;
   if (tunedThreads > pso.maxTotalThreadsPerThreadgroup) {
     // Fewer threads than the kernel was built for: safe for the shared arrays,
     // but the scan and compaction kernels assume every slot has an owner.