a kŗ”h‰ć@s4dZddlmZejddd„ƒZejd dd„ƒZdS) a­ Grid Dependency Control (GDC) is a mechanism used when enabling programmatic dependent launch to launch and synchronize grids. These APIs expose GDC to the programmer. Programmatic dependent launch is supported on SM90 (Hopper) and beyond. For PTX reference on grid dependency control see https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-griddepcontrol. é)ŚcoreNc Cstjddgtjdd|ddS)aU GDC wait is a blocking instruction that waits for all instructions in a prior kernel to complete before continuing. This ensures all memory operations happening before the wait is visible to instructions after it, e.g. if the prior kernel writes to address "x" the new values will be visible in this kernel after the wait. This instruction is also safe to execute when programatic dependent launch is disabled. See https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-griddepcontrol for more details. z griddepcontrol.wait; // dummy $0ś=rFé©ZdtypeŚis_pureŚpackŚ_builderN©rZinline_asm_elementwiseŚint32©r©r śL/var/www/auris/lib/python3.9/site-packages/triton/language/extra/cuda/gdc.pyŚgdc_wait s ’rc Cstjddgtjdd|ddS)a‚ This operation when launched with programmatic dependent launch signals that the next program may launch once all programs in the current kernel call this function or complete. Repeated calls to this function have no effect past the first call, and the first call should be treated by the programmer as a hint to the runtime system to launch the next kernel. This instruction is also safe to execute when programatic dependent launch is disabled. See https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-griddepcontrol for more details. z-griddepcontrol.launch_dependents; // dummy $0rFrrNr r r r r Śgdc_launch_dependentss’r)N)N)Ś__doc__Ztriton.languagerŚexternrrr r r r Śs