てらモス🌹

PTX 9.5 just landed with some interesting new additions - sin/cos .2pi - Angles can now be supplied in cycles tensormap.replace .alloc_class - TMA descriptors can explicitly track small/large tensor allocations and query that state trace.mark + .metadata_section - Kernel code can emit timestamped performance fabric publish and sync instructions like that of TMA commit_group and wait_group cp.async.bulk.wait_group.sync_restrict::proxy::async - A narrower async copy wait, that orders prior async proxy accesses against subsequent async proxy without unnecessarily ordering the generic proxy. A weaker ordering in a sense (so it could be more async)? fence.alias.sync_restrict::read.release.sys - A system scope release fence specifically for prior reads through aliased generic addresses. %locality_domain_id and %locality_domain_supplemental_sm - Lets a kernel discover which locality domain its cluster was launched on and lets persistent kernels/schedulers distinguish guaranteed locality, domain SMs from supplemental SMs. This means we can now query the number of domains and SMs per domain, allocate memory inside a particular locality domain, and create a green context whose SMs come from that same domain. I remeber what @pranjalssh is grouping SMs into those two sets, then assigning an A-row to one requester side so its repeated reads build useful residency in only one L2 half. CUDA locality domains now let you request something closer to - allocate memory in locality domain X run the kernel on SMs from locality domain X docs.nvidia.com/cuda/paralle…