{"id":"https://openalex.org/W7224006175","doi":"https://doi.org/10.1109/iiswc71020.2026.00039","title":"When and Why to Use TMA: A Workload-Grounded Characterization of Tensor Memory Acceleration on NVIDIA Hopper GPUs","display_name":"When and Why to Use TMA: A Workload-Grounded Characterization of Tensor Memory Acceleration on NVIDIA Hopper GPUs","publication_year":2026,"publication_date":"2026-09-01","ids":{"openalex":"https://openalex.org/W7224006175","doi":"https://doi.org/10.1109/iiswc71020.2026.00039"},"language":"en","primary_location":{"id":"doi:10.1109/iiswc71020.2026.00039","doi":"https://doi.org/10.1109/iiswc71020.2026.00039","is_oa":false,"landing_page_url":"https://doi.org/10.1109/iiswc71020.2026.00039","pdf_url":null,"source":{"id":"https://openalex.org/S4306418766","display_name":"IEEE International Symposium on Workload Characterization","issn_l":null,"issn":null,"is_oa":false,"is_in_doaj":false,"is_core":false,"listed_in":[],"host_organization":null,"host_organization_name":null,"host_organization_lineage":[],"host_organization_lineage_names":[],"type":"conference"},"license":null,"license_id":null,"version":"publishedVersion","is_accepted":true,"is_published":true,"raw_source_name":"2026 IEEE International Symposium on Workload Characterization (IISWC)","raw_type":"proceedings-article"},"volume":null,"type":"conference-paper","indexed_in":["arxiv","crossref","datacite"],"open_access":{"is_oa":true,"oa_status":"green","oa_url":"https://arxiv.org/pdf/2610.34907","any_repository_has_fulltext":true},"authorships":[{"author_position":"first","author":{"id":"https://openalex.org/A5080345475","display_name":"Christin David Bose","orcid":null,"observed_orcids":[]},"institutions":[{"id":"https://openalex.org/I219193219","display_name":"Purdue University West Lafayette","ror":"https://ror.org/02dqehb95","country_code":"US","type":"education","lineage":["https://openalex.org/I219193219"]}],"countries":["US"],"is_corresponding":false,"raw_author_name":"Christin David Bose","raw_affiliation_strings":["Purdue University"],"raw_orcid":null,"affiliations":[{"raw_affiliation_string":"Purdue University","institution_ids":["https://openalex.org/I219193219"]}]},{"author_position":"middle","author":{"id":"https://openalex.org/A5157987024","display_name":"Cesar Avalos","orcid":"https://orcid.org/0009-0005-9989-0369","observed_orcids":["https://orcid.org/0009-0005-9989-0369"]},"institutions":[{"id":"https://openalex.org/I219193219","display_name":"Purdue University West Lafayette","ror":"https://ror.org/02dqehb95","country_code":"US","type":"education","lineage":["https://openalex.org/I219193219"]}],"countries":["US"],"is_corresponding":false,"raw_author_name":"Cesar Avalos Baddouh","raw_affiliation_strings":["Purdue University"],"raw_orcid":null,"affiliations":[{"raw_affiliation_string":"Purdue University","institution_ids":["https://openalex.org/I219193219"]}]},{"author_position":"middle","author":{"id":"https://openalex.org/A5047571426","display_name":"Junrui Pan","orcid":null,"observed_orcids":[]},"institutions":[{"id":"https://openalex.org/I219193219","display_name":"Purdue University West Lafayette","ror":"https://ror.org/02dqehb95","country_code":"US","type":"education","lineage":["https://openalex.org/I219193219"]}],"countries":["US"],"is_corresponding":false,"raw_author_name":"Junrui Pan","raw_affiliation_strings":["Purdue University"],"raw_orcid":null,"affiliations":[{"raw_affiliation_string":"Purdue University","institution_ids":["https://openalex.org/I219193219"]}]},{"author_position":"last","author":{"id":"https://openalex.org/A5159930392","display_name":"Timothy Rogers","orcid":null,"observed_orcids":[]},"institutions":[{"id":"https://openalex.org/I219193219","display_name":"Purdue University West Lafayette","ror":"https://ror.org/02dqehb95","country_code":"US","type":"education","lineage":["https://openalex.org/I219193219"]}],"countries":["US"],"is_corresponding":false,"raw_author_name":"Timothy G. Rogers","raw_affiliation_strings":["Purdue University"],"raw_orcid":null,"affiliations":[{"raw_affiliation_string":"Purdue University","institution_ids":["https://openalex.org/I219193219"]}]}],"institutions":[],"countries_distinct_count":1,"institutions_distinct_count":1,"corresponding_author_ids":[],"corresponding_institution_ids":["https://openalex.org/I219193219"],"apc_list":null,"apc_paid":null,"fwci":0.0,"has_fulltext":true,"cited_by_count":0,"citation_normalized_percentile":{"value":0.86484193,"is_in_top_1_percent":false,"is_in_top_10_percent":false},"cited_by_percentile_year":null,"biblio":{"volume":null,"issue":null,"first_page":"351","last_page":"362"},"is_retracted":false,"is_paratext":false,"is_xpac":false,"primary_topic":{"id":"https://openalex.org/T10054","display_name":"Parallel Computing and Optimization Techniques","score":1.0,"subfield":{"id":"https://openalex.org/subfields/1708","display_name":"Hardware and Architecture"},"field":{"id":"https://openalex.org/fields/17","display_name":"Computer Science"},"domain":{"id":"https://openalex.org/domains/3","display_name":"Physical Sciences"}},"topics":[{"id":"https://openalex.org/T10054","display_name":"Parallel Computing and Optimization Techniques","score":1.0,"subfield":{"id":"https://openalex.org/subfields/1708","display_name":"Hardware and Architecture"},"field":{"id":"https://openalex.org/fields/17","display_name":"Computer Science"},"domain":{"id":"https://openalex.org/domains/3","display_name":"Physical Sciences"}},{"id":"https://openalex.org/T10481","display_name":"Computer Graphics and Visualization Techniques","score":0.0,"subfield":{"id":"https://openalex.org/subfields/1704","display_name":"Computer Graphics and Computer-Aided Design"},"field":{"id":"https://openalex.org/fields/17","display_name":"Computer Science"},"domain":{"id":"https://openalex.org/domains/3","display_name":"Physical Sciences"}},{"id":"https://openalex.org/T10904","display_name":"Embedded Systems Design Techniques","score":0.0,"subfield":{"id":"https://openalex.org/subfields/1708","display_name":"Hardware and Architecture"},"field":{"id":"https://openalex.org/fields/17","display_name":"Computer Science"},"domain":{"id":"https://openalex.org/domains/3","display_name":"Physical Sciences"}}],"keywords":[{"id":"https://openalex.org/keywords/distributed-shared-memory","display_name":"distributed shared memory","score":0.6809999942779541},{"id":"https://openalex.org/keywords/workload-analysis","display_name":"workload characterization","score":0.5339999794960022}],"concepts":[{"id":"https://openalex.org/C41008148","wikidata":"https://www.wikidata.org/wiki/Q21198","display_name":"Computer science","level":0,"score":0.7786999940872192},{"id":"https://openalex.org/C32295351","wikidata":"https://www.wikidata.org/wiki/Q899288","display_name":"Multicast","level":2,"score":0.6484000086784363},{"id":"https://openalex.org/C55526617","wikidata":"https://www.wikidata.org/wiki/Q719375","display_name":"Operand","level":2,"score":0.5698000192642212},{"id":"https://openalex.org/C173608175","wikidata":"https://www.wikidata.org/wiki/Q232661","display_name":"Parallel computing","level":1,"score":0.4659000039100647},{"id":"https://openalex.org/C74193536","wikidata":"https://www.wikidata.org/wiki/Q574844","display_name":"Kernel (algebra)","level":2,"score":0.4636000096797943},{"id":"https://openalex.org/C188045654","wikidata":"https://www.wikidata.org/wiki/Q17148339","display_name":"Memory bandwidth","level":2,"score":0.4113999903202057},{"id":"https://openalex.org/C2776257435","wikidata":"https://www.wikidata.org/wiki/Q1576430","display_name":"Bandwidth (computing)","level":2,"score":0.39910000562667847},{"id":"https://openalex.org/C2780841128","wikidata":"https://www.wikidata.org/wiki/Q5073781","display_name":"Characterization (materials science)","level":2,"score":0.358599990606308},{"id":"https://openalex.org/C157764524","wikidata":"https://www.wikidata.org/wiki/Q1383412","display_name":"Throughput","level":3,"score":0.3521000146865845},{"id":"https://openalex.org/C2777338717","wikidata":"https://www.wikidata.org/wiki/Q1762621","display_name":"Vendor","level":2,"score":0.34599998593330383},{"id":"https://openalex.org/C96324660","wikidata":"https://www.wikidata.org/wiki/Q205446","display_name":"Dataflow","level":2,"score":0.3456999957561493},{"id":"https://openalex.org/C80444323","wikidata":"https://www.wikidata.org/wiki/Q2878974","display_name":"Theoretical computer science","level":1,"score":0.34209999442100525},{"id":"https://openalex.org/C56666940","wikidata":"https://www.wikidata.org/wiki/Q788790","display_name":"Documentation","level":2,"score":0.3416000008583069},{"id":"https://openalex.org/C155281189","wikidata":"https://www.wikidata.org/wiki/Q3518150","display_name":"Tensor (intrinsic definition)","level":2,"score":0.33980000019073486},{"id":"https://openalex.org/C52146309","wikidata":"https://www.wikidata.org/wiki/Q7431116","display_name":"Schema (genetic algorithms)","level":2,"score":0.336899995803833},{"id":"https://openalex.org/C206952183","wikidata":"https://www.wikidata.org/wiki/Q1193100","display_name":"Preemption","level":2,"score":0.3337000012397766},{"id":"https://openalex.org/C48044578","wikidata":"https://www.wikidata.org/wiki/Q727490","display_name":"Scalability","level":2,"score":0.3077000081539154},{"id":"https://openalex.org/C2776299755","wikidata":"https://www.wikidata.org/wiki/Q432449","display_name":"Carry (investment)","level":2,"score":0.30410000681877136},{"id":"https://openalex.org/C113843644","wikidata":"https://www.wikidata.org/wiki/Q901882","display_name":"Interface (matter)","level":4,"score":0.29829999804496765},{"id":"https://openalex.org/C3018587665","wikidata":"https://www.wikidata.org/wiki/Q7268696","display_name":"Qualitative analysis","level":3,"score":0.2800999879837036},{"id":"https://openalex.org/C2781357197","wikidata":"https://www.wikidata.org/wiki/Q5757597","display_name":"High memory","level":2,"score":0.2702000141143799},{"id":"https://openalex.org/C11413529","wikidata":"https://www.wikidata.org/wiki/Q8366","display_name":"Algorithm","level":1,"score":0.26910001039505005},{"id":"https://openalex.org/C97686452","wikidata":"https://www.wikidata.org/wiki/Q7604153","display_name":"Static analysis","level":2,"score":0.25690001249313354},{"id":"https://openalex.org/C120314980","wikidata":"https://www.wikidata.org/wiki/Q180634","display_name":"Distributed computing","level":1,"score":0.25279998779296875}],"mesh":[],"locations_count":3,"locations":[{"id":"doi:10.1109/iiswc71020.2026.00039","doi":"https://doi.org/10.1109/iiswc71020.2026.00039","is_oa":false,"landing_page_url":"https://doi.org/10.1109/iiswc71020.2026.00039","pdf_url":null,"source":{"id":"https://openalex.org/S4306418766","display_name":"IEEE International Symposium on Workload Characterization","issn_l":null,"issn":null,"is_oa":false,"is_in_doaj":false,"is_core":false,"listed_in":[],"host_organization":null,"host_organization_name":null,"host_organization_lineage":[],"host_organization_lineage_names":[],"type":"conference"},"license":null,"license_id":null,"version":"publishedVersion","is_accepted":true,"is_published":true,"raw_source_name":"2026 IEEE International Symposium on Workload Characterization (IISWC)","raw_type":"proceedings-article"},{"id":"pmh:oai:arXiv.org:2610.34907","doi":null,"is_oa":true,"landing_page_url":"https://arxiv.org/abs/2610.34907","pdf_url":"https://arxiv.org/pdf/2610.34907","source":{"id":"https://openalex.org/S4306400194","display_name":"arXiv (Cornell University)","issn_l":"2331-8422","issn":["2331-8422"],"is_oa":true,"is_in_doaj":false,"is_core":false,"listed_in":[],"host_organization":"https://openalex.org/I205783295","host_organization_name":"Cornell University","host_organization_lineage":["https://openalex.org/I205783295"],"host_organization_lineage_names":[],"type":"repository"},"license":null,"license_id":null,"version":"submittedVersion","is_accepted":false,"is_published":false,"raw_source_name":null,"raw_type":"text"},{"id":"doi:10.48550/arxiv.2610.34907","doi":"https://doi.org/10.48550/arxiv.2610.34907","is_oa":true,"landing_page_url":"https://doi.org/10.48550/arxiv.2610.34907","pdf_url":null,"source":{"id":"https://openalex.org/S4306400194","display_name":"arXiv (Cornell University)","issn_l":"2331-8422","issn":["2331-8422"],"is_oa":true,"is_in_doaj":false,"is_core":false,"listed_in":[],"host_organization":"https://openalex.org/I205783295","host_organization_name":"Cornell University","host_organization_lineage":["https://openalex.org/I205783295"],"host_organization_lineage_names":[],"type":"repository"},"license":"cc-by-sa","license_id":"https://openalex.org/licenses/cc-by-sa","version":null,"is_accepted":false,"is_published":false,"raw_source_name":null,"raw_type":"Preprint"}],"best_oa_location":{"id":"pmh:oai:arXiv.org:2610.34907","doi":null,"is_oa":true,"landing_page_url":"https://arxiv.org/abs/2610.34907","pdf_url":"https://arxiv.org/pdf/2610.34907","source":{"id":"https://openalex.org/S4306400194","display_name":"arXiv (Cornell University)","issn_l":"2331-8422","issn":["2331-8422"],"is_oa":true,"is_in_doaj":false,"is_core":false,"listed_in":[],"host_organization":"https://openalex.org/I205783295","host_organization_name":"Cornell University","host_organization_lineage":["https://openalex.org/I205783295"],"host_organization_lineage_names":[],"type":"repository"},"license":null,"license_id":null,"version":"submittedVersion","is_accepted":false,"is_published":false,"raw_source_name":null,"raw_type":"text"},"sustainable_development_goals":[],"study_designs":[],"awards":[{"id":"https://openalex.org/G1664597935","display_name":"CAREER: Accessible Accelerators: Leveraging Productive Software on Efficient Hardware","funder_award_id":"1943379","funder_id":"https://openalex.org/F4320306076","funder_display_name":"National Science Foundation"},{"id":"https://openalex.org/G3010069967","display_name":"Collaborative Research: CIRC: ENS: Enabling Detailed, Open-Source Accelerator Modeling","funder_award_id":"2450447","funder_id":"https://openalex.org/F4320306076","funder_display_name":"National Science Foundation"}],"funders":[{"id":"https://openalex.org/F4320306076","display_name":"National Science Foundation","ror":"https://ror.org/021nxhr62"}],"has_content":{"grobid_xml":true,"pdf":true},"content_urls":{"pdf":"https://content.openalex.org/works/W7224006175.pdf","grobid_xml":"https://content.openalex.org/works/W7224006175.grobid-xml"},"referenced_works_count":0,"referenced_works":[],"related_works":["https://openalex.org/W4309672181","https://openalex.org/W4393366869","https://openalex.org/W7220408924","https://openalex.org/W7202249660","https://openalex.org/W7220963610","https://openalex.org/W7054396888","https://openalex.org/W4390317766","https://openalex.org/W4213440247","https://openalex.org/W7054661169","https://openalex.org/W4410321788"],"abstract_inverted_index":{"The":[0],"Tensor":[1],"Memory":[2],"Accelerator":[3],"(TMA),":[4],"introduced":[5],"on":[6,104,139,160,177,188,204],"NVIDIA":[7],"Hopper,":[8,189],"is":[9,19],"used":[10],"pervasively":[11],"in":[12,25,66,80],"production":[13],"GPU":[14],"kernels":[15],"yet":[16],"its":[17],"usage":[18],"conspicuously":[20],"selective:":[21],"expert":[22],"kernel":[23,38],"developers":[24],"CUTLASS,":[26],"FlashAttention-3,":[27],"and":[28,59,87,112,130,152,168,182,190,221],"FlashInfer":[29],"apply":[30],"TMA":[31,64,144,179],"to":[32,42,195],"some":[33],"tensor":[34],"operands":[35],"within":[36],"a":[37,77,81,100,200],"but":[39,68],"fall":[40],"back":[41],"cp.async,":[43],"direct":[44],"global":[45],"loads,":[46],"or":[47,91,186],"distributed":[48],"shared":[49],"memory":[50,78],"(DSM)":[51],"for":[52,135],"others.":[53],"Vendor":[54],"documentation":[55],"describes":[56],"the":[57,71,108,120,136,154,193,208],"API,":[58],"recent":[60],"microbenchmarking":[61],"work":[62],"characterizes":[63],"throughput":[65],"isolation,":[67],"neither":[69],"answers":[70],"question":[72],"practitioners":[73],"actually":[74],"face:":[75],"given":[76],"operand":[79],"kernel,":[82],"should":[83],"it":[84],"use":[85],"TMA,":[86,109],"if":[88],"so,":[89],"with":[90,99,219],"without":[92],"cluster":[93],"multicast?":[94],"We":[95,106,141],"address":[96],"this":[97],"gap":[98],"workload-grounded":[101],"characterization":[102],"study":[103,203],"H100.":[105],"catalog":[107],"multicast,":[110],"DSM,":[111],"cp.async":[113],"choices":[114],"made":[115],"inside":[116],"FlashAttention-3":[117,155],"forward,":[118],"isolate":[119,153],"discriminating":[121],"features":[122],"through":[123],"targeted":[124],"microbenchmarks":[125],"(transfer-size":[126],"crossover,":[127],"multicast":[128,145,184],"scaling),":[129],"produce":[131],"rules":[132,211],"of":[133],"thumb":[134],"load-mechanism":[137],"choice":[138],"Hopper.":[140],"further":[142],"compare":[143],"against":[146],"DSM":[147],"as":[148],"competing":[149],"cluster-broadcast":[150],"primitives,":[151],"paged-key-value":[156],"(KV)":[157],"TMA-vs-cp.async":[158],"dispatch":[159,210],"H100":[161],"across":[162,213],"page":[163],"size,":[164],"datatype,":[165],"model":[166],"shape,":[167],"execution":[169],"phase.":[170],"Our":[171],"results":[172],"yield":[173],"concrete,":[174],"predictive":[175],"guidance":[176],"when":[178,183],"pays":[180],"off":[181],"helps":[185],"hurts":[187],"we":[191],"extend":[192],"analysis":[194],"server":[196],"Blackwell":[197],"(SM100)":[198],"via":[199],"CUTLASS":[201],"GEMM":[202],"B200":[205],"that":[206],"shows":[207],"qualitative":[209],"carry":[212],"generations":[214],"while":[215],"absolute":[216],"crossovers":[217],"shift":[218],"tensor-core":[220],"bandwidth":[222],"scaling.":[223]},"counts_by_year":[],"updated_date":"2026-11-05T12:25:35.769452","created_date":"2026-10-28T00:00:00"}