{"record":{"id":"8bf2980ecdb7b2ec","repo":"jax-ml/jax","slug":"requested-too-many-cuda-threads-per-block-each-mo","errorCode":null,"errorMessage":"Requested too many CUDA threads per block. Each Mosaic thread corresponds to 128 CUDA threads. At most {max_mosaic_threads} are supported, got {self}","messagePattern":"Requested too many CUDA threads per block\\. Each Mosaic thread corresponds to 128 CUDA threads\\. At most (.+?) are supported, got (.+?)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/core.py","lineNumber":1656,"sourceCode":"  def __post_init__(self):\n    if len(self.cluster) > 3:\n      raise ValueError(f\"cluster= must be at most 3D, got {self}.\")\n    if len(self.grid_names) != len(self.grid):\n      raise ValueError(\n          f\"grid_names must have the same length as grid, got {self}.\"\n      )\n    if len(self.cluster_names) != len(self.cluster):\n      raise ValueError(\n          f\"cluster_names must have the same length as cluster, got {self}.\"\n      )\n    if (self.thread_name is None) != (self.num_threads is None):\n      raise ValueError(\n          \"num_threads and thread_name must be either both set or both None,\"\n          f\" got {self}\"\n      )\n    max_mosaic_threads = 2048 // 128\n    if self.num_threads is not None and self.num_threads > max_mosaic_threads:\n      raise ValueError(\n          \"Requested too many CUDA threads per block. Each Mosaic thread\"\n          f\" corresponds to 128 CUDA threads. At most {max_mosaic_threads}\"\n          f\" are supported, got {self}\"\n      )\n    object.__setattr__(self, \"grid\", tuple(self.grid))\n    object.__setattr__(self, \"grid_names\", tuple(self.grid_names))\n    object.__setattr__(self, \"cluster\", tuple(self.cluster))\n    object.__setattr__(self, \"cluster_names\", tuple(self.cluster_names))\n\n  @property\n  def default_memory_space(self) -> MemorySpace:\n    return MemorySpace.GMEM\n\n  @property\n  def shape(self) -> collections.OrderedDict[object, int]:\n    pairs: Iterable[tuple[object, int]]\n    if self.num_threads is not None:\n      pairs = zip(","sourceCodeStart":1638,"sourceCodeEnd":1674,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/core.py#L1638-L1674","documentation":"Launch-config validation: each Mosaic 'thread' maps to 128 CUDA threads and a CUDA block supports at most 2048 threads, so num_threads may not exceed 2048 // 128 = 16. Requesting more raises this error before launch.","triggerScenarios":"Passing num_threads > 16 (e.g. num_threads=32) to the Mosaic GPU kernel launch config while also setting thread_name.","commonSituations":"Tuning warp-specialized kernels and assuming num_threads maps 1:1 to CUDA warps; porting kernels between Mosaic versions where the thread model changed to 128-CUDA-thread units.","solutions":["Reduce num_threads to at most 16","If more parallelism is needed, restructure across warps/cluster instead of more Mosaic threads"],"exampleFix":"// before\nnum_threads=32\n\n// after\nnum_threads=16  # max: 2048 CUDA threads // 128","handlingStrategy":"validation","validationCode":"MAX_MOSAIC_THREADS = 2048 // 128  # 16\nassert num_threads is None or num_threads <= MAX_MOSAIC_THREADS","typeGuard":"def threads_supported(num_threads: int | None) -> bool:\n    return num_threads is None or 1 <= num_threads <= 16","tryCatchPattern":"try:\n    cfg = LaunchConfig(num_threads=n, thread_name='warp')\nexcept ValueError:\n    cfg = LaunchConfig(num_threads=min(n, 16), thread_name='warp')","preventionTips":["Remember each Mosaic thread = 128 CUDA threads; hard cap is 16","Clamp num_threads at config creation","Expose num_threads as a tunable bounded to 16"],"tags":["jax","pallas","mosaic-gpu","thread-limits","resource-limit","validation"],"backgroundTag":"exceeds-thread-limit","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}