Pith. sign in

REVIEW 2 major objections 7 minor 50 references

CuAsmRL: Optimizing GPU SASS Schedules via Deep Reinforcement Learning

T0 review · 2 major / 7 minor · reviewed 2026-08-10 · deepseek-v4-flash

Pith's one-line read An RL agent reorders GPU assembly to speed up LLM kernels by up to 26 percent.

desk verdict First RL-based SASS reorderer with a real artifact and credible measured gains; the headline speedup needs error bars and the 'transparent' claim needs scoping, but this deserves a serious referee. read the letter →

arxiv 2501.08071 v1 pith:JLOBCHJI submitted 2025-01-14 cs.AR cs.LG

classification cs.ARcs.LG
keywords GPUinstructionschedulingSASSreinforcementlearningCUDAkernelsTritoncompilerlargelanguagemodelsPPOAmpere
verification ladder T0 review T1 audit T2 compute T3 formal

The pith

A machine-rendered reading of the paper's core claim, the machinery that carries it, and where it could break.

The reading

CuAsmRL claims that a deep reinforcement-learning agent can find instruction schedules for NVIDIA GPUs that beat the -O3 schedule produced by NVIDIA's own compiler, without touching kernel source code. The agent plays an assembly game: starting from an already-optimized SASS binary, it repeatedly swaps a memory load/store instruction with its neighbor, runs the mutated kernel on an A100 GPU, and uses the measured runtime as reward. Integrated into the Triton compiler as a SASS-to-SASS pass, CuAsmRL improved throughput on six representative LLM kernels by up to 26 percent and on average about 9 percent. The significance is that specialized CUDA kernels are already heavily hand-tuned, so an automatic layer that still finds headroom suggests instruction scheduling is a general, under-exploited optimization surface.

What carries the argument

The engine is an assembly game, a Markov decision process whose states are embeddings of SASS instruction streams, whose actions are swaps of a memory instruction with the instruction above or below, and whose reward is the percentage runtime improvement measured against the initial -O3 schedule. Correctness is enforced by action masking: a candidate swap is killed if it violates a register dependency, a barrier dependency, or a stall-count dependency. Stall counts come from a microbenchmarked table for common fixed-latency integer instructions and from a static analysis pass that infers them from the already-valid original schedule; memory instructions whose dependencies cannot be resolved are denylisted and never moved. The policy is a CNN encoder over the embedded instruction matrix followed by an MLP, trained with proximal policy optimization.

What would settle it

Take a kernel from outside the LLM domain, apply the one-line CuAsmRL integration on an A100, and compare its output against the original on randomized inputs; if any optimized cubin returns wrong results, the completeness of the dependency masks is refuted, and if training never yields a faster schedule on such kernels, the claim of general applicability is weakened.

Watch

Extended reading notes

Core claim

On its own terms, the paper's central claim is that SASS instruction scheduling is a substantial and still-open performance surface for GPU kernels, and that a reinforcement-learning agent can automate the rescheduling human experts currently do by hand. Starting from the best schedule ptxas produces at -O3, CuAsmRL restricts its actions to moving memory load/store instructions such as LDG, LDGSTS, and STG one position up or down inside a basic block, with legal moves filtered by dependency masks for registers, barrier codes, and stall counts. The reward is the measured reduction in kernel runtime, and the agent is trained with PPO to maximize cumulative reward. In evaluation on six kernels used by LLMs, including fused feed-forward, batch matrix multiplication, flash attention, fused GEMM with LeakyReLU, softmax, and RMSNorm, the optimized schedules outperform Triton on every kernel, up to 26 percent and on average 9 percent. The paper also claims the learned moves are interpretable: one reordering of HMMA before LDGSTS in fused GEMM contributes 7 percent and is tied to operand-cache reuse.

Load-bearing premise

The reordering is safe only if the hand-coded dependency rules and the measured stall-count table cover every SASS hazard that exists; if any hazard is missing, an allowed swap can silently change the kernel's results, and the paper's own Section 7 asks users to verify optimized kernels manually.

Editorial extensions

If this is right

  • If the measured 9 percent average speedup holds, users of Triton-compiled LLM kernels can gain a transparent speedup on A100 simply by switching to CuAsmRL's decorator, with no source-level changes.
  • The discovered moves, such as placing HMMA before LDGSTS and moving LDGSTS ahead of a predicated-off LDS, can be read as concrete scheduling rules for ptxas or other SASS-level passes.
  • Because optimization happens after code generation, the approach applies to any cubin that can be disassembled and reassembled, not only to kernels written in Triton.
  • The comparison showing CuAsmRL matching or beating CuBLAS and FlashAttention-2 baselines on several kernels implies that schedule-level optimization can close part of the gap between compiler-generated and hand-written GPU code.

Reading between the lines

Editorial extensions of the paper, not claims the author makes directly.

  • The paper does not claim the speedups transfer to other GPU architectures; the stall-count table is A100-specific, so porting would require re-benchmarking before the same gains can be expected on newer hardware.
  • The authors leave implicit that the roughly five-hour per-kernel training cost could be amortized by pre-training a policy across many kernels, turning the optimizer into a default compiler pass rather than an offline per-kernel search.
  • A practical deployment consequence the paper does not spell out: transparent applies to source code, not to correctness, so an automated equivalence check would be needed before the optimized cubin could ship without human review.
Share X Bluesky LinkedIn Reddit HN

Editorial analysis

A structured set of objections, weighed in public.

Desk editor's note, referee report, and a circularity audit.

Referee Report

2 major / 7 minor

Summary. CuAsmRL formulates GPU SASS instruction scheduling as a reinforcement-learning 'assembly game' in which an agent repeatedly swaps memory instructions with adjacent instructions, subject to dependency masks derived from register, barrier, stall-count, and hand-coded 'additional' dependencies. The reward is measured kernel throughput on an A100 GPU. The approach is integrated into OpenAI Triton: it disassembles the compiled cubin, trains an agent to reorder the SASS schedule, and deploys the best schedule found during training. On six LLM-relevant kernels (fused feed-forward, GEMM with LeakyReLU, batch matrix multiply, flash-attention, softmax, and rmsnorm), the paper reports up to 26% and on average 9% throughput improvement over Triton, and it analyzes the discovered reorderings and their effect on memory throughput.

Significance. If the measured speedups hold, CuAsmRL is a practical drop-in SASS-to-SASS optimizer that extends a widely used compiler framework (Triton) with low-level scheduling improvements. The artifact is publicly available on Zenodo/GitHub, the paper reports training statistics, and the method for inferring stall counts from existing schedules is a useful contribution. The paper also provides a concrete case study (HMMA/LDGSTS reordering and the .reuse-flag interaction) that could inform future compiler optimizations. However, the significance is tempered by the absence of statistical validation for the headline speedups and by the manual-verification caveat for correctness, which together make the strength of the central claim uncertain.

major comments (2)
  1. [Section 5.3, Figure 6, Section 4.2] The central speedup claim lacks statistical validation. The reported numbers come from selecting the 'best optimized cubin found throughout the assembly game' (Section 4.2) using a reward signal that is itself a noisy execution-time measurement (Section 3.6: average of 100 iterations, with standard deviation typically within 1% of two measurements). Selecting an argmax over thousands of noisy evaluations biases the apparent improvement upward. The final benchmark in Section 5.1 uses the same 100-iteration measurement procedure and is not independent of the search. The paper should report per-kernel mean and standard deviation across multiple RL training seeds, include error bars or confidence intervals in Figure 6, and validate the chosen schedule with an independent measurement protocol (e.g., different iteration counts, clock sources, or re-measurement after the fact). For kernels with reported gains near 2%, the effect may be entirely within run-to-run noise; without these numbers the headline 'up to 26%, on average 9%' is not yet established.
  2. [Section 7, Section 3.5, Section 4.3] The claim that CuAsmRL optimizes kernels 'transparently' (Abstract, Section 1) is contradicted by the paper's own limitation statement that 'users are required to manually verify the optimized kernels as in §5.7.' The correctness of every reordering depends on the completeness of the hand-coded 'additional dependencies' (Section 3.5) and the stall-count table (Section 4.3), and the paper provides no formal verification or exhaustive testing; probabilistic testing with randomized inputs is a sanity check, not a proof. The authors should either (a) weaken the transparency claim to 'source-transparent' with a clear statement that correctness requires per-kernel verification, or (b) provide stronger evidence of correctness, such as systematic output comparison on many random inputs for each of the six kernels, a description of the verification performed in §5.7, and an explicit discussion of which SASS hazard classes are covered by the additional-dependency heuristics.
minor comments (7)
  1. [Section 5.7.1] There is a missing space in 'we observe7% improvement'; it should read 'we observe 7% improvement'.
  2. [References] Reference [46] appears corrupted: 'Rico Zhang, Biao an/cudad Sennrich' is not a valid citation; please correct the author and title.
  3. [Section 3.5, Algorithm 1] The pseudocode for masking stall counts is difficult to evaluate: the semantics of 'is_user(inst_to_check, cur)' are not defined, and the loop breaks at the first encountered user without considering multiple users or producers in the opposite direction. Please clarify why this is sufficient for safe reordering.
  4. [Section 5.1] The Cutlass baseline is described as using the 'default configuration' and is reported as 10x slower than Triton. This comparison is not informative unless the exact configuration is specified; consider reporting results with a tuned configuration or removing the baseline.
  5. [Section 5.5] The hyperparameter sensitivity analysis (Figure 8) is shown only for one kernel (fused GEMM with LeakyReLU). Please state whether similar behavior was observed for the other kernels or restrict the robustness claim accordingly.
  6. [Section 3.6] The statement that the standard deviation of two measurements is typically within 1% would be more useful if the actual per-kernel variances were reported alongside Figure 6.
  7. [Section 2.3] The example SASS instruction appears to be missing from the rendered text: after 'A typical SASS instruction is shown below' there is a blank line rather than a code listing.

Circularity Check

0 steps flagged · score 0.0 of 10

No significant circularity: the reported speedups are direct GPU timing measurements, and the dependency/stall machinery that constrains reordering is independent calibration, not a fitted predictor of the headline results.

full rationale

CuAsmRL's central claim (up to 26%, mean 9% over Triton, Figure 6) is an empirical optimization result: the RL agent mutates SASS schedules and receives a reward equal to measured kernel throughput (Eq. 3), and the reported numbers are measured execution times of the best cubin found. No equation in the paper defines the target speedup in terms of an input parameter, and no fitted parameter is renamed as a prediction. The stall-count table (Section 4.3) is obtained by dependency-based microbenchmarks on A100, and the static-analysis pass (Section 3.2) infers stall counts conservatively from the original schedule ('because the original schedule is always valid, the inferred value would be either overestimated or exact'); this is a correctness constraint on the search, not a model that generates the speedup. Heuristic dependency rules are manually identified and used only to mask invalid reorderings. The only self-citations ([8], [9]) are artifact and prior-work references that are not load-bearing for the throughput claim. The paper itself flags in Section 7 that users must manually verify optimized kernels and that a cost model is absent; these are correctness and statistical limitations (along with the absence of error bars and selection over the best cubin), but they do not make the derivation circular.

Assumptions & free parameters 0 free parameters · 5 assumptions · 0 invented entities

The central claim rests on empirically measured speedups rather than a derivation. The dependency masks and stall counts are hand-derived from reverse-engineered SASS semantics; these are the main axioms. No free parameters are fitted to make the results work: RL hyperparameters come from a prior study, and the stall counts are hardware measurements. No new physical or conceptual entities are introduced beyond the 'assembly game' formulation, which is a framing device, not a postulated entity.

assumptions (5)
  • domain assumption SASS instruction semantics and control-code behavior as reverse-engineered in prior works are accurate for Ampere.
    Sections 2.3 and 3.2 rely on decoded barrier, stall, and yield semantics to build the dependency masks.
  • domain assumption The microbenchmarked stall counts (Table 1) reflect true fixed-latency instruction latencies on A100.
    Section 4.3 derives the stall count table via dependency-based microbenchmarks; masking correctness depends on these values.
  • ad hoc to paper The hand-coded 'additional dependencies' (e.g., LDGSTS sequences) are the only extra hazards beyond register/barrier/stall dependencies.
    Section 3.5 states 'we have to identify them manually because of the lack of publically available data'; completeness is assumed without proof.
  • domain assumption The analysis pass's inferred stall counts are always safe (overestimated or exact).
    Section 3.2 argues that because the original schedule is valid, inferred values are either overestimated or exact, but this relies on the microbenchmarked table being correct.
  • domain assumption Performance measurement using 100 warmup plus 100 timed iterations with L2 cache clearing is stable and representative.
    Sections 3.6 and 5.1 use this protocol; the paper reports standard deviation within 1% but no confidence intervals.

how reviews work

0 comments
Cite this review

Pith. "Pith review of CuAsmRL: Optimizing GPU SASS Schedules via Deep Reinforcement Learning." pith.science (2026). https://pith.science/paper/JLOBCHJI

@misc{pith2026250108071,
  author       = {Pith},
  title        = {Pith review of: CuAsmRL: Optimizing GPU SASS Schedules via Deep Reinforcement Learning},
  year         = {2026},
  howpublished = {\url{https://pith.science/paper/JLOBCHJI}},
  note         = {Machine review of arXiv:2501.08071}
}
abstract

Large language models (LLMs) are remarked by their substantial computational requirements. To mitigate the cost, researchers develop specialized CUDA kernels, which often fuse several tensor operations to maximize the utilization of GPUs as much as possible. However, those specialized kernels may still leave performance on the table as CUDA assembly experts show that manual optimization of GPU SASS schedules can lead to better performance, and trial-and-error is largely employed to manually find the best GPU SASS schedules. In this work, we employ an automatic approach to optimize GPU SASS schedules, which thus can be integrated into existing compiler frameworks. The key to automatic optimization is training an RL agent to mimic how human experts perform manual scheduling. To this end, we formulate an assembly game, where RL agents can play to find the best GPU SASS schedules. The assembly game starts from a \textit{-O3} optimized SASS schedule, and the RL agents can iteratively apply actions to mutate the current schedules. Positive rewards are generated if the mutated schedules get higher throughput by executing on GPUs. Experiments show that CuAsmRL can further improve the performance of existing specialized CUDA kernels transparently by up to $26\%$, and on average $9\%$. Moreover, it is used as a tool to reveal potential optimization moves learned automatically.

Figures

Figures reproduced from arXiv: 2501.08071 by the authors.

Figure 1
Figure 1. CUDA compilation as documented by NVIDIA [25]. C++/Python and PTX are highlighted in green, indicating they are the common programming interfaces. SASS is a GPUs-native assembly and is highlighted in red, meaning it is undocumented. Cubin is an executable binary and is in gray. while for Python, Triton [40] can be used. The compilation process has several stages: first, the code is compiled to PTX, which is an inter… view at source ↗
Figure 2
Figure 2. Overall workflow of CuAsmRL. CuAsmRL takes as input the source code targeting Triton’s programming interface. Then it uses an autotuner to enumerate and find the optimal kernel configurations. Then the code is compiled with the optimal kernel configurations via Triton’s compi￾lation pipeline. Finally, an RL agent is trained to play the assembly game to optimize the SASS schedules, which out￾puts an optimized cubin. … view at source ↗
Figure 3
Figure 3. Assembly Game. At each iteration 𝑖, the SASS file is embedded, and the embedding is fed to the RL agent as state 𝑆𝑖 . The RL agent is represented by a deep neural network. The output of the RL agent is an action 𝐴𝑖 that changes the SASS file. Then the mutated SASS file is assembled and sent to execution on the target GPU. A reward 𝑅𝑖 is sent back to the agent and the mutated SASS file is transitioned to the next sta… view at source ↗
Figures from the paper (10 more)
Figure 4
Figure 4. Figure 4: Embedding. Different fields of SASS instruction such as control code, opcode, and operands are embedded individually and then concatenated to a vector. Dummy val￾ues (-1) are used for the absent fields and operand padding. Different vectors are concatenated in a row-wi…
Figure 5
Figure 5. Figure 5: An example of an action, which reorders the SASS instructions. However, allowing each instruction to be reordered intro￾duces a massive action space, as a kernel can have thousands of lines of SASS instructions. Considering the latency hid￾ing process is mostly about p…
Figure 6
Figure 6. Figure 6: Overall kernel throughput comparison. The throughput of Triton is normalized to 1, and the others are normalized accordingly. A high value indicates a better performance compared to Triton. bmm: batch matrix multiplication, fused_ff : fused feed-forward, rmsnorm: root-…
Figure 7
Figure 7. Figure 7: Percentages of stall count of fixed-latency instruc￾tions that are resolved by the built-in stall count table (db), inferred by the analysis pass (infer-only), and deny-listed (not resolved) on average for kernels listed in [PITH_FULL_IMAGE:figures/full_fig_p010_7.png]
Figure 8
Figure 8. Figure 8: Episodic returns for different hyperparameter settings. The green line is the default setting. default hyperparameters setting come from a work which performs large-scale case study across various domains [11] [PITH_FULL_IMAGE:figures/full_fig_p011_8.png]
Figure 9
Figure 9. Figure 9: A reordering for fused GEMM and the epilogue. Scheduling the HMMA instruction before the LDGSTS in￾struction achieves better performance. 5.7.1 Fused GEMM with LeakyReLU [PITH_FULL_IMAGE:figures/full_fig_p011_9.png]
Figure 10
Figure 10. Figure 10: Memory chart for fused GEMM LeakyReLU optimized by CuAsmRL [PITH_FULL_IMAGE:figures/full_fig_p016_10.png]
Figure 11
Figure 11. Figure 11: Memory chart for fused GEMM LeakyReLU optimized by Triton [PITH_FULL_IMAGE:figures/full_fig_p016_11.png]
Figure 12
Figure 12. Figure 12: Time series plots of training statistics during the assembly game. Appendix D Comparison of PTX and SASS 1 add.s32 %r121, %r204, 18432; 2 add.s32 %r123, %r204, 20480; 3 add.s32 %r125, %r204, 22528; 4 selp.b32 %r120, 16, 0, %p10; 5 cp.async.cg.shared.global [ %r119 + 0…
Figure 13
Figure 13. Figure 13: A reordering that schedule the LDGSTS instruction earlier than the predicated LDS instruction. This optimization move is applied to multiple kernels [PITH_FULL_IMAGE:figures/full_fig_p018_13.png]

Discussion (0). Continue with ORCID to comment.

Reference graph

Works this paper leans on

50 extracted references · 29 canonical work pages

  1. [1]

    Hamdy Abdelkhalik, Yehia Arafa, Nandakishore Santhi, and Abdel- Hameed Badawy. 2022. Demystifying the Nvidia Ampere Archi- tecture through Microbenchmarking and Instruction-level Analysis. arXiv:2208.11174 [cs.AR] https://arxiv.org/abs/2208.11174

  2. [2]

    Ravichandra Addanki, Shaileshh Bojja Venkatakrishnan, Shreyan Gupta, Hongzi Mao, and Mohammad Alizadeh. 2019. Placeto: Learning Generalizable Device Placement Algorithms for Distributed Machine Learning. In Proceedings of the 33rd International Conference on Neural Information Processing Systems. Curran Associates Inc., Red Hook, NY, USA, Article 358, 11 pages

  3. [3]

    Greg Brockman, Vicki Cheung, Ludwig Pettersson, Jonas Schneider, John Schulman, Jie Tang, and Wojciech Zaremba. 2016. OpenAI Gym. arXiv:1606.01540 [cs.LG]

  4. [4]

    Cloudcores. 2024. Cuasm. https : / / github. com / cloudcores / CuAssembler

  5. [5]

    Tri Dao, Dan Fu, Stefano Ermon, Atri Rudra, and Christopher Ré. 2022. FlashAttention: Fast and Memory-Efficient Exact Atten- tion with IO-Awareness. In Advances in Neural Information Pro- cessing Systems , S. Koyejo, S. Mohamed, A. Agarwal, D. Belgrave, K. Cho, and A. Oh (Eds.), Vol. 35. Curran Associates, Inc., 16344– 16359. https://proceedings.neurips.c...

  6. [6]

    Alhussein Fawzi, Matej Balog, Anderson Huang, Randrianarisoa Ramo- lairisoa, Arthur Guez, Demis Hassabis, Pushmeet Kohli, and Timo- thy P Lillicrap. 2022. Discovering faster matrix multiplication algo- rithms with reinforcement learning. Nature 610, 7930 (2022), 47–53. https://doi.org/10.1038/s41586-022-05172-4

  7. [7]

    Ameer Haj-Ali, Qijing (Jenny) Huang, John Xiang, William Moses, Krste Asanovic, John Wawrzynek, and Ion Stoica. 2020. AutoPhase: Juggling HLS Phase Orderings in Random Forests with Deep Rein- forcement Learning. In Proceedings of Machine Learning and Sys- tems, I. Dhillon, D. Papailiopoulos, and V. Sze (Eds.), Vol. 2. 70–

  8. [8]

    Guoliang He. 2025. Reproduction. https://doi.org/10.5281/zenodo. 14058861 CGO ’25 Artifact. CGO ’25, March 01–05, 2025, Las Vegas, NV, USA Guoliang He and Eiko Yoneki

Show all 50 references
  1. [9]

    Guoliang He, Sean Parker, and Eiko Yoneki. 2023. X-RLflow: Graph Re- inforcement Learning for Neural Network Subgraphs Transformation. arXiv:2304.14698 [cs.LG] https://arxiv.org/abs/2304.14698

  2. [10]

    Pieter Hijma, Stijn Heldens, Alessio Sclocco, Ben van Werkhoven, and Henri E. Bal. 2023. Optimization Techniques for GPU Programming. ACM Comput. Surv. 55, 11, Article 239 (mar 2023), 81 pages. https: //doi.org/10.1145/3570638

  3. [11]

    Shengyi Huang, Rousslan Fernand Julien Dossa, Antonin Raffin, Anssi Kanervisto, and Weixun Wang. 2022. The 37 Implementation Details of Proximal Policy Optimization. In ICLR Blog Track. https://iclr-blog- track.github.io/2022/03/25/ppo-implementation-details/ https://iclr- blo...

  4. [12]

    Intel. 2024. MaxAs. https://github.com/NervanaSystems/maxas

  5. [13]

    Zhe Jia, Marco Maggioni, Jeffrey Smith, and Daniele Paolo Scarpazza

  6. [14]

    Scarpazza

    Zhe Jia, Marco Maggioni, Benjamin Staiger, and Daniele P. Scarpazza

  7. [15]

    Albert Q. Jiang, Alexandre Sablayrolles, Antoine Roux, Arthur Mensch, Blanche Savary, Chris Bamford, Devendra Singh Chaplot, Diego de las Casas, Emma Bou Hanna, Florian Bressand, Gianna Lengyel, Guil- laume Bour, Guillaume Lample, Lélio Renard Lavaud, Lucile Saulnier, Marie-An...

  8. [16]

    kernl.ai. 2024. kernl. https://github.com/ELS-RD/kernl

  9. [17]

    Amy McGovern, Eliot Moss, and Andrew G. Barto. 2002. Basic-block Instruction Scheduling Using Reinforcement Learning and Rollouts. https://api.semanticscholar.org/CorpusID:1231595

  10. [18]

    Azalia Mirhoseini, Anna Goldie, Hieu Pham, Benoit Steiner, Quoc V Le, and Jeff Dean. 2018. A hierarchical model for device placement. In International Conference on Learning Representations

  11. [19]

    Rusu, Joel Veness, Marc G

    Volodymyr Mnih, Koray Kavukcuoglu, David Silver, Andrei A. Rusu, Joel Veness, Marc G. Bellemare, Alex Graves, Martin A. Riedmiller, Andreas Fidjeland, Georg Ostrovski, Stig Petersen, Charles Beattie, Amir Sadik, Ioannis Antonoglou, Helen King, Dharshan Kumaran, Daan Wierstra, ...

  12. [20]

    NVIDIA. 2024. CUDA c++ programming guide . https://docs.nvidia. com/cuda/cuda-c-programming-guide/index.html

  13. [21]

    NVIDIA. 2024. CUDA performance metrics. https://developer.nvidia. com/blog/how-implement-performance-metrics-cuda-cc/

  14. [22]

    NVIDIA. 2024. CUDA performance metrics. https://docs.nvidia.com/ cuda/cublas/

  15. [23]

    NVIDIA. 2024. NVIDIA. https://developer.nvidia.com/nsight-compute

  16. [24]

    NVIDIA. 2024. NVIDIA. https://docs.nvidia.com/nsight-visual- studio- edition/4.6/Content/Analysis/Report/CudaExperiments/ KernelLevel/PerformanceCounters.htm

  17. [25]

    NVIDIA. 2024. NVIDIA CUDA compiler. https://docs.nvidia.com/cuda/ cuda-compiler-driver-nvcc/index.html

  18. [26]

    NVIDIA. 2024. NVIDIA kepler GPU . https://www.nvidia.com/content/ dam/en-zz/Solutions/Data-Center/tesla-product-literature/NVIDIA- Kepler-GK110-GK210-Architecture-Whitepaper.pdf

  19. [27]

    NVIDIA. 2024. NVIDIA ptx. https://docs.nvidia.com/cuda/parallel- thread-execution/index.html

  20. [28]

    NVIDIA. 2024. NVIDIA sass . https://docs.nvidia.com/cuda/cuda- binary-utilities/index.html

  21. [29]

    OpenAI. 2024. OpenAI. https://openai.com/research/video-generation- models-as-world-simulators

  22. [30]

    OpenAI. 2024. OpenAI. https : / / twitter. com / sama / status / 1756089361609981993

  23. [31]

    OpenAI. 2024. OpenAI. https://openai.com/blog/chatgpt

  24. [32]

    OpenAI, Ilge Akkaya, Marcin Andrychowicz, Maciek Chociej, Ma- teusz Litwin, Bob McGrew, Arthur Petron, Alex Paino, Matthias Plap- pert, Glenn Powell, Raphael Ribas, Jonas Schneider, Nikolas Tezak, Jerry Tworek, Peter Welinder, Lilian Weng, Qiming Yuan, Wojciech Zaremba, and Le...

  25. [33]

    Pytorch2. 2024. Pytorch2. https://pytorch.org/blog/pytorch-2-paper- tutorial/

  26. [34]

    Rodrigues, Sam S

    Shane Ryoo, Christopher I. Rodrigues, Sam S. Stone, John A. Stratton, Sain-Zee Ueng, Sara S. Baghsorkhi, and Wen mei W. Hwu. 2008. Pro- gram optimization carving for GPU computing. J. Parallel and Distrib. Comput. 68, 10 (2008), 1389–1401. https://doi.org/10.1016/j.jpdc.2008. ...

  27. [35]

    John Schulman, Filip Wolski, Prafulla Dhariwal, Alec Radford, and Oleg Klimov. 2017. Proximal Policy Optimization Algorithms. arXiv:1707.06347 [cs.LG]

  28. [36]

    Ghassan Shobaki, Austin Kerbow, and Stanislav Mekhanoshin. 2020. Optimizing occupancy and ILP on the GPU using a combinatorial approach. InProceedings of the 18th ACM/IEEE International Symposium on Code Generation and Optimization (San Diego, CA, USA) (CGO ’20). Association f...

  29. [37]

    Ghassan Shobaki, Pınar Muyan-Özçelik, Josh Hutton, Bruce Linck, Vladislav Malyshenko, Austin Kerbow, Ronaldo Ramirez-Ortega, and Vahl Scott Gordon. 2024. Instruction Scheduling for the GPU on the GPU. In 2024 IEEE/ACM International Symposium on Code Generation and Optimization...

  30. [38]

    David Silver, Thomas Hubert, Julian Schrittwieser, Ioannis Antonoglou, Matthew Lai, Arthur Guez, Marc Lanctot, Laurent Sifre, Dharshan Kumaran, Thore Graepel, Timothy Lillicrap, Karen Simonyan, and Demis Hassabis. 2018. A general reinforcement learning algorithm that masters c...

  31. [39]

    Sutton and Andrew G

    Richard S. Sutton and Andrew G. Barto. 2018. Reinforcement Learning: An Introduction. A Bradford Book, Cambridge, MA, USA

  32. [40]

    Philippe Tillet, H. T. Kung, and David Cox. 2019. Triton: an interme- diate language and compiler for tiled neural network computations. In Proceedings of the 3rd ACM SIGPLAN International Workshop on Machine Learning and Programming Languages (Phoenix, AZ, USA) (MAPL 2019). A...

  33. [41]

    Hugo Touvron, Louis Martin, Kevin Stone, Peter Albert, Amjad Alma- hairi, Yasmine Babaei, Nikolay Bashlykov, Soumya Batra, Prajjwal Bhargava, Shruti Bhosale, Dan Bikel, Lukas Blecher, Cristian Canton Ferrer, Moya Chen, Guillem Cucurull, David Esiobu, Jude Fernandes, Jeremy Fu,...

  34. [42]

    Triton. 2024. Trion-repo. https://github.com/triton-lang/triton CuAsmRL: Optimizing GPU SASS Schedules via Deep Reinforcement Learning CGO ’25, March 01–05, 2025, Las Vegas, NV, USA

  35. [43]

    Ashish Vaswani, Noam Shazeer, Niki Parmar, Jakob Uszkoreit, Llion Jones, Aidan N Gomez, Lukasz Kaiser, and Illia Polosukhin. 2017. At- tention is All you Need. In Advances in Neural Information Processing Systems, I. Guyon, U. Von Luxburg, S. Bengio, H. Wallach, R. Fergus, S. ...

  36. [44]

    Da Yan, Wei Wang, and Xiaowen Chu. 2020. Demystifying Tensor Cores to Optimize Half-Precision Matrix Multiply. In2020 IEEE Interna- tional Parallel and Distributed Processing Symposium (IPDPS) . 634–643. https://doi.org/10.1109/IPDPS47924.2020.00071

  37. [45]

    Da Yan, Wei Wang, and Xiaowen Chu. 2020. Optimizing batched wino- grad convolution on GPUs. In Proceedings of the 25th ACM SIGPLAN Symposium on Principles and Practice of Parallel Programming (San Diego, California) (PPoPP ’20). Association for Computing Machinery, New York, N...

  38. [46]

    Rico Zhang, Biao an/cudad Sennrich. 2019. Root mean square layer normalization. Curran Associates Inc., Red Hook, NY, USA

  39. [47]

    Xiuxia Zhang, Guangming Tan, Shuangbai Xue, Jiajia Li, Keren Zhou, and Mingyu Chen. 2017. Understanding the GPU Microar- chitecture to Achieve Bare-Metal Performance Tuning. In Proceed- ings of the 22nd ACM SIGPLAN Symposium on Principles and Prac- tice of Parallel Programming...

  40. [81]

    https://proceedings.mlsys.org/paper_files/paper/2020/file/ 5b47430e24a5a1f9fe21f0e8eb814131-Paper.pdf

  41. [2018]

    arXiv:1804.06826 [cs.DC] https://arxiv.org/abs/1804.06826

    Dissecting the NVIDIA Volta GPU Architecture via Microbench- marking. arXiv:1804.06826 [cs.DC] https://arxiv.org/abs/1804.06826

  42. [2019]

    arXiv:1903.07486 [cs.DC]

    Dissecting the NVidia Turing T4 GPU via Microbenchmarking. arXiv:1903.07486 [cs.DC]

Pith tools

Reviewed August 10, 2026 · model on record in the stance chip above.