{"status":"ok","message-type":"work","message-version":"1.0.0","message":{"indexed":{"date-parts":[[2025,8,21]],"date-time":"2025-08-21T11:40:12Z","timestamp":1755776412884,"version":"3.44.0"},"publisher-location":"New York, NY, USA","reference-count":49,"publisher":"ACM","license":[{"start":{"date-parts":[[2025,3,30]],"date-time":"2025-03-30T00:00:00Z","timestamp":1743292800000},"content-version":"vor","delay-in-days":0,"URL":"https:\/\/creativecommons.org\/licenses\/by\/4.0\/"}],"funder":[{"DOI":"10.13039\/501100006374","name":"NSF (National Science Foundation)","doi-asserted-by":"publisher","award":["2303735, 2238346"],"award-info":[{"award-number":["2303735, 2238346"]}],"id":[{"id":"10.13039\/501100006374","id-type":"DOI","asserted-by":"publisher"}]}],"content-domain":{"domain":["dl.acm.org"],"crossmark-restriction":true},"short-container-title":[],"published-print":{"date-parts":[[2025,3,30]]},"DOI":"10.1145\/3676641.3716281","type":"proceedings-article","created":{"date-parts":[[2025,3,27]],"date-time":"2025-03-27T16:47:32Z","timestamp":1743094052000},"page":"1382-1399","update-policy":"https:\/\/doi.org\/10.1145\/crossmark-policy","source":"Crossref","is-referenced-by-count":0,"title":["Virgo: Cluster-level Matrix Unit Integration in GPUs for Scalability and Energy Efficiency"],"prefix":"10.1145","author":[{"ORCID":"https:\/\/orcid.org\/0000-0002-0080-0548","authenticated-orcid":false,"given":"Hansung","family":"Kim","sequence":"first","affiliation":[{"name":"University of California, Berkeley, Berkeley, California, USA"}],"role":[{"role":"author","vocabulary":"crossref"}]},{"ORCID":"https:\/\/orcid.org\/0009-0001-3695-6499","authenticated-orcid":false,"given":"Ruohan Richard","family":"Yan","sequence":"additional","affiliation":[{"name":"University of California, Berkeley, Berkeley, California, USA"}],"role":[{"role":"author","vocabulary":"crossref"}]},{"ORCID":"https:\/\/orcid.org\/0000-0002-6317-6690","authenticated-orcid":false,"given":"Joshua","family":"You","sequence":"additional","affiliation":[{"name":"University of California, Berkeley, Berkeley, California, USA"}],"role":[{"role":"author","vocabulary":"crossref"}]},{"ORCID":"https:\/\/orcid.org\/0009-0008-9316-689X","authenticated-orcid":false,"given":"Tieliang Vamber","family":"Yang","sequence":"additional","affiliation":[{"name":"NVIDIA Corporation, Santa Clara, California, USA"}],"role":[{"role":"author","vocabulary":"crossref"}]},{"ORCID":"https:\/\/orcid.org\/0000-0003-1811-5407","authenticated-orcid":false,"given":"Yakun Sophia","family":"Shao","sequence":"additional","affiliation":[{"name":"University of California, Berkeley, Berkeley, California, USA"}],"role":[{"role":"author","vocabulary":"crossref"}]}],"member":"320","published-online":{"date-parts":[[2025,3,30]]},"reference":[{"key":"e_1_3_2_1_1_1","unstructured":"AMD. 2021. AMD CDNA 2 Architecture. https:\/\/www.amd.com\/content\/dam\/amd\/en\/documents\/instinct-business-docs\/white-papers\/amd-cdna2-white-paper.pdf."},{"key":"e_1_3_2_1_2_1","unstructured":"AMD. 2024a. AMD Matrix Instruction Calculator. https:\/\/github.com\/ROCm\/amd_matrix_instruction_calculator\/blob\/dae289b7f681776584749e3ae21f177b7f974d92\/README.md."},{"key":"e_1_3_2_1_3_1","unstructured":"AMD. 2024b. Composable Kernel. https:\/\/github.com\/ROCm\/composable_kernel."},{"key":"e_1_3_2_1_4_1","volume-title":"Krste Asanovi\u00e7, and Borivoje Nikoli'c.","author":"Amid Alon","year":"2020","unstructured":"Alon Amid, David Biancolin, Abraham Gonzalez, Daniel Grubb, Sagar Karandikar, Harrison Liew, Albert Magyar, Howard Mao, Albert Ou, Nathan Pemberton, Paul Rigge, Colin Schmidt, John Wright, Jerry Zhao, Yakun Sophia Shao, Krste Asanovi\u00e7, and Borivoje Nikoli'c. 2020. Chipyard: Integrated design, simulation, and implementation framework for custom socs. IEEE Micro, Vol. 40, 4 (2020)."},{"key":"e_1_3_2_1_5_1","unstructured":"Simran Arora Sabri Eyuboglu Michael Zhang Aman Timalsina Silas Alberti Dylan Zinsley James Zou Atri Rudra and Christopher R\u00e9. 2024. Simple linear attention language models balance the recall-throughput tradeoff. arxiv: 2402.18668 [cs.CL] https:\/\/arxiv.org\/abs\/2402.18668"},{"key":"e_1_3_2_1_6_1","volume-title":"The Rocket Chip Generator. EECS Department","author":"Asanovi\u00e7 Krste","year":"2016","unstructured":"Krste Asanovi\u00e7, Rimas Avizienis, Jonathan Bachrach, Scott Beamer, David Biancolin, Christopher Celio, Henry Cook, Daniel Dabbelt, John Hauser, Adam Izraelevitz, Sagar Karandikar, Ben Keller, Donggyu Kim, and John Koenig. 2016. The Rocket Chip Generator. EECS Department, University of California, Berkeley, Tech. Rep. UCB\/EECS-2016--17, Vol. 4 (2016)."},{"key":"e_1_3_2_1_7_1","unstructured":"Ganesh Bikshandi and Jay Shah. 2023. A Case Study in CUDA Kernel Fusion: Implementing FlashAttention-2 on NVIDIA Hopper Architecture using the CUTLASS Library. arxiv: 2312.11918 [cs.LG] https:\/\/arxiv.org\/abs\/2312.11918"},{"key":"e_1_3_2_1_8_1","volume-title":"Mintz","author":"Bridges Robert A.","year":"2016","unstructured":"Robert A. Bridges, Neena Imam, and Tiffany M. Mintz. 2016. Understanding GPU power: A survey of profiling, modeling, and simulation methods. ACM Computing Surveys (CSUR), Vol. 49, 3 (2016)."},{"key":"e_1_3_2_1_9_1","unstructured":"Tom Brown Benjamin Mann Nick Ryder Melanie Subbiah Jared Kaplan Prafulla Dhariwal Arvind Neelakantan Pranav Shyam Girish Sastry Amanda Askell et al. 2020. Language models are few-shot learners. Advances in neural information processing systems (NeurIPS) (2020)."},{"key":"e_1_3_2_1_10_1","doi-asserted-by":"publisher","DOI":"10.1109\/MM.2023.3256796nolinkurl10.1109\/MM.2023.3256796"},{"key":"e_1_3_2_1_11_1","volume-title":"Productive Design of Extensible On-Chip Memory Hierarchies. Ph.,D. Dissertation. EECS Department","author":"Cook Henry","year":"2016","unstructured":"Henry Cook. 2016. Productive Design of Extensible On-Chip Memory Hierarchies. Ph.,D. Dissertation. EECS Department, University of California, Berkeley. http:\/\/www2.eecs.berkeley.edu\/Pubs\/TechRpts\/2016\/EECS-2016--89.html"},{"key":"e_1_3_2_1_12_1","volume-title":"1st Workshop on Computer Architecture Research with RISC-V","volume":"23","author":"Cook Henry","year":"2017","unstructured":"Henry Cook, Wesley Terpstra, and Yunsup Lee. 2017. Diplomatic design patterns: A TileLink case study. In 1st Workshop on Computer Architecture Research with RISC-V, Vol. 23."},{"key":"e_1_3_2_1_13_1","unstructured":"NVIDIA Corporation. 2017. Nvidia Tesla V100 GPU architecture. https:\/\/images.nvidia.com\/content\/volta-architecture\/pdf\/volta-architecture-whitepaper.pdf."},{"key":"e_1_3_2_1_14_1","unstructured":"NVIDIA Corporation. 2022. NVIDIA A100 Tensor Core GPU Datasheet. https:\/\/www.nvidia.com\/content\/dam\/en-zz\/Solutions\/Data-Center\/a100\/pdf\/nvidia-a100-datasheet-nvidia-us-2188504-web.pdf."},{"key":"e_1_3_2_1_15_1","unstructured":"NVIDIA Corporation. 2024a. NVIDIA Ampere GPU Architecture Tuning Guide. https:\/\/docs.nvidia.com\/cuda\/pdf\/Ampere_Tuning_Guide.pdf."},{"key":"e_1_3_2_1_16_1","unstructured":"NVIDIA Corporation. 2024b. NVIDIA H100 Tensor Core GPU Datasheet. https:\/\/resources.nvidia.com\/en-us-tensor-core\/nvidia-tensor-core-gpu-datasheet."},{"key":"e_1_3_2_1_17_1","unstructured":"NVIDIA Corporation. 2024c. Parallel Thread Execution ISA Version 8.5. https:\/\/docs.nvidia.com\/cuda\/parallel-thread-execution\/index.html."},{"key":"e_1_3_2_1_18_1","volume-title":"Davidson and Sanjay Jinturkar","author":"Jack","year":"1994","unstructured":"Jack W. Davidson and Sanjay Jinturkar. 1994. Memory access coalescing: A technique for eliminating redundant memory accesses. Acm Sigplan Notices, Vol. 29, 6 (1994)."},{"key":"e_1_3_2_1_19_1","doi-asserted-by":"publisher","DOI":"10.1145\/3620665.3640367"},{"key":"e_1_3_2_1_20_1","doi-asserted-by":"publisher","DOI":"10.1145\/3620666.3651378"},{"key":"e_1_3_2_1_21_1","doi-asserted-by":"publisher","DOI":"10.1109\/DAC18074.2021.9586216"},{"key":"e_1_3_2_1_22_1","doi-asserted-by":"publisher","DOI":"10.1145\/3453953.3453972"},{"key":"e_1_3_2_1_23_1","unstructured":"John Hauser. 2019. Berkeley HardFloat. http:\/\/www.jhauser.us\/arithmetic\/HardFloat.html."},{"key":"e_1_3_2_1_24_1","doi-asserted-by":"publisher","DOI":"10.1109\/BigData47090.2019.9005632"},{"key":"e_1_3_2_1_25_1","doi-asserted-by":"publisher","DOI":"10.1145\/3613424.3623775"},{"key":"e_1_3_2_1_26_1","unstructured":"Intel. 2022. Introduction to the Xe-HPG Architecture. https:\/\/www.intel.com\/content\/www\/us\/en\/developer\/articles\/technical\/introduction-to-the-xe-hpg-architecture.html."},{"key":"e_1_3_2_1_27_1","doi-asserted-by":"publisher","DOI":"10.1109\/ISCA.2018.00014"},{"key":"e_1_3_2_1_28_1","volume-title":"Cutlass: Fast linear algebra in cuda c. NVIDIA Developer Blog","author":"Kerr Andrew","year":"2017","unstructured":"Andrew Kerr, Duane Merrill, Julien Demouth, and John Tran. 2017. Cutlass: Fast linear algebra in cuda c. NVIDIA Developer Blog (2017)."},{"volume-title":"Proceedings of the International Symposium on Microarchitecture (MICRO).","author":"Kim Hyeonjin","key":"e_1_3_2_1_29_1","unstructured":"Hyeonjin Kim, Sungwoo Ahn, Yunho Oh, Bogil Kim, Won Woo Ro, and William J. Song. 2020. Duplo: Lifting redundant memory accesses of deep neural networks for GPU tensor cores. In Proceedings of the International Symposium on Microarchitecture (MICRO)."},{"key":"e_1_3_2_1_30_1","doi-asserted-by":"publisher","DOI":"10.1109\/PACT58117.2023.00033"},{"key":"e_1_3_2_1_31_1","doi-asserted-by":"publisher","DOI":"10.1109\/HPCA53966.2022.00093"},{"key":"e_1_3_2_1_32_1","unstructured":"Weile Luo Ruibo Fan Zeyu Li Dayou Du Qiang Wang and Xiaowen Chu. 2024. Benchmarking and Dissecting the Nvidia Hopper GPU Architecture. https:\/\/arxiv.org\/abs\/2402.13499. arxiv: 2402.13499 [cs.AR]"},{"key":"e_1_3_2_1_33_1","doi-asserted-by":"publisher","DOI":"10.1109\/LCA.2023.3278652"},{"volume-title":"Proceedings of the International Symposium on Performance Analysis of Systems and Software (ISPASS).","author":"Raihan Md A.","key":"e_1_3_2_1_34_1","unstructured":"Md A. Raihan, Negar Goli, and Tor M. Aamodt. 2019. Modeling deep learning accelerator enabled gpus. In Proceedings of the International Symposium on Performance Analysis of Systems and Software (ISPASS)."},{"key":"e_1_3_2_1_35_1","doi-asserted-by":"crossref","unstructured":"Gabin Schieffer Daniel Medeiros Jennifer Faj Aniruddha Marathe and Ivy Peng. 2024. Characterizing the Performance Power Efficiency and Programmability of AMD Matrix Cores. Technical Report. Lawrence Livermore National Laboratory Livermore CA.","DOI":"10.1109\/ISPASS61541.2024.00022"},{"key":"e_1_3_2_1_36_1","doi-asserted-by":"publisher","DOI":"10.1109\/IJCNN55064.2022.9891914"},{"key":"e_1_3_2_1_37_1","unstructured":"Jay Shah Ganesh Bikshandi Ying Zhang Vijay Thakkar Pradeep Ramani and Tri Dao. 2024. FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision. arxiv: 2407.08608 [cs.LG] https:\/\/arxiv.org\/abs\/2407.08608"},{"key":"e_1_3_2_1_38_1","doi-asserted-by":"publisher","DOI":"10.1145\/1067649.801719"},{"key":"e_1_3_2_1_39_1","article-title":"Dissecting tensor cores via microbenchmarks: Latency, throughput and numeric behaviors","volume":"34","author":"Sun Wei","year":"2022","unstructured":"Wei Sun, Ang Li, Tong Geng, Sander Stuijk, and Henk Corporaal. 2022. Dissecting tensor cores via microbenchmarks: Latency, throughput and numeric behaviors. IEEE Transactions on Parallel and Distributed Systems, Vol. 34, 1 (2022).","journal-title":"IEEE Transactions on Parallel and Distributed Systems"},{"key":"e_1_3_2_1_40_1","doi-asserted-by":"publisher","DOI":"10.1145\/3613424.3614247"},{"key":"e_1_3_2_1_41_1","doi-asserted-by":"publisher","DOI":"10.1145\/2063384.2063431"},{"key":"e_1_3_2_1_42_1","doi-asserted-by":"publisher","DOI":"10.1145\/3466752.3480128"},{"key":"e_1_3_2_1_43_1","doi-asserted-by":"publisher","DOI":"10.1109\/ICPPW.2016.32"},{"key":"e_1_3_2_1_44_1","doi-asserted-by":"publisher","DOI":"10.1109\/ISCA52012.2021.00088"},{"key":"e_1_3_2_1_45_1","doi-asserted-by":"publisher","DOI":"10.1109\/IPDPS47924.2020.00071"},{"key":"e_1_3_2_1_46_1","unstructured":"Michael Zhang Kush Bhatia Hermann Kumbong and Christopher R\u00e9. 2024. The hedgehog & the porcupine: Expressive linear attentions with softmax mimicry. arxiv: 2402.04347 [cs.LG] https:\/\/arxiv.org\/abs\/2402.04347"},{"key":"e_1_3_2_1_47_1","doi-asserted-by":"publisher","DOI":"10.1145\/3018743.3018755"},{"key":"e_1_3_2_1_48_1","volume-title":"Proceedings of the International Symposium on Computer Architecture (ISCA).","author":"Zhang Yunan","year":"2022","unstructured":"Yunan Zhang, Po-An Tsai, and Hung-Wei Tseng. 2022. SIMD2: A generalized matrix instruction set for accelerating tensor computation beyond GEMM. In Proceedings of the International Symposium on Computer Architecture (ISCA)."},{"key":"e_1_3_2_1_49_1","doi-asserted-by":"publisher","DOI":"10.1145\/3620678.3624793"}],"event":{"name":"ASPLOS '25: 30th ACM International Conference on Architectural Support for Programming Languages and Operating Systems","sponsor":["SIGPLAN ACM Special Interest Group on Programming Languages","SIGOPS ACM Special Interest Group on Operating Systems","SIGARCH ACM Special Interest Group on Computer Architecture"],"location":"Rotterdam Netherlands","acronym":"ASPLOS '25"},"container-title":["Proceedings of the 30th ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Volume 2"],"original-title":[],"link":[{"URL":"https:\/\/dl.acm.org\/doi\/10.1145\/3676641.3716281","content-type":"unspecified","content-version":"vor","intended-application":"text-mining"},{"URL":"https:\/\/dl.acm.org\/doi\/pdf\/10.1145\/3676641.3716281","content-type":"unspecified","content-version":"vor","intended-application":"similarity-checking"}],"deposited":{"date-parts":[[2025,8,21]],"date-time":"2025-08-21T11:07:27Z","timestamp":1755774447000},"score":1,"resource":{"primary":{"URL":"https:\/\/dl.acm.org\/doi\/10.1145\/3676641.3716281"}},"subtitle":[],"short-title":[],"issued":{"date-parts":[[2025,3,30]]},"references-count":49,"alternative-id":["10.1145\/3676641.3716281","10.1145\/3676641"],"URL":"https:\/\/doi.org\/10.1145\/3676641.3716281","relation":{},"subject":[],"published":{"date-parts":[[2025,3,30]]},"assertion":[{"value":"2025-03-30","order":3,"name":"published","label":"Published","group":{"name":"publication_history","label":"Publication History"}}]}}