{"id":"https://openalex.org/W7167707177","doi":"https://doi.org/10.1145/3811257.3811269","title":"Porting ThunderKittens from CUDA to SYCL for Intel GPU: Process, Challenges, and Lessons Learned","display_name":"Porting ThunderKittens from CUDA to SYCL for Intel GPU: Process, Challenges, and Lessons Learned","publication_year":2026,"publication_date":"2026-05-06","ids":{"openalex":"https://openalex.org/W7167707177","doi":"https://doi.org/10.1145/3811257.3811269"},"language":null,"primary_location":{"id":"doi:10.1145/3811257.3811269","is_oa":true,"landing_page_url":"https://doi.org/10.1145/3811257.3811269","pdf_url":null,"source":null,"license":"cc-by","license_id":"https://openalex.org/licenses/cc-by","version":"publishedVersion","is_accepted":true,"is_published":true,"raw_source_name":"Proceedings of the International Workshop on OpenCL and SYCL","raw_type":"proceedings-article"},"type":"conference-paper","indexed_in":["crossref"],"open_access":{"is_oa":true,"oa_status":"gold","oa_url":"https://doi.org/10.1145/3811257.3811269","any_repository_has_fulltext":null},"authorships":[{"author_position":"first","author":{"id":"https://openalex.org/A5138450603","display_name":"Yehong Jiang","orcid":null},"institutions":[{"id":"https://openalex.org/I1343180700","display_name":"Intel (United States)","ror":"https://ror.org/01ek73717","country_code":"US","type":"company","lineage":["https://openalex.org/I1343180700"]},{"id":"https://openalex.org/I63966007","display_name":"Massachusetts Institute of Technology","ror":"https://ror.org/042nb2s44","country_code":"US","type":"education","lineage":["https://openalex.org/I63966007"]}],"countries":["US"],"is_corresponding":false,"raw_author_name":"Yehong Jiang","raw_affiliation_strings":["Intel Corporation, Santa Clara, CA, USA and Department of Electrical Engineering and Computer Science, Massachusetts Institute of Technology, CAMBRIDGE, Massachusetts, USA"],"raw_orcid":"https://orcid.org/0009-0006-1340-8170","affiliations":[{"raw_affiliation_string":"Intel Corporation, Santa Clara, CA, USA and Department of Electrical Engineering and Computer Science, Massachusetts Institute of Technology, CAMBRIDGE, Massachusetts, USA","institution_ids":["https://openalex.org/I1343180700","https://openalex.org/I63966007"]}]},{"author_position":"middle","author":{"id":"https://openalex.org/A5101825806","display_name":"Sheng Chen","orcid":"https://orcid.org/0009-0000-3770-2004"},"institutions":[{"id":"https://openalex.org/I1343180700","display_name":"Intel (United States)","ror":"https://ror.org/01ek73717","country_code":"US","type":"company","lineage":["https://openalex.org/I1343180700"]}],"countries":["US"],"is_corresponding":false,"raw_author_name":"Shen Chen","raw_affiliation_strings":["Intel Corporation, Shanghai, Massachusetts, China"],"raw_orcid":"https://orcid.org/0009-0000-3770-2004","affiliations":[{"raw_affiliation_string":"Intel Corporation, Shanghai, Massachusetts, China","institution_ids":["https://openalex.org/I1343180700"]}]},{"author_position":"middle","author":{"id":"https://openalex.org/A5100815128","display_name":"Fangwen Fu","orcid":null},"institutions":[{"id":"https://openalex.org/I1343180700","display_name":"Intel (United States)","ror":"https://ror.org/01ek73717","country_code":"US","type":"company","lineage":["https://openalex.org/I1343180700"]}],"countries":["US"],"is_corresponding":false,"raw_author_name":"Fangwen Fu","raw_affiliation_strings":["Intel Corporation, Santa Clara, CA, USA"],"raw_orcid":"https://orcid.org/0000-0003-3548-6066","affiliations":[{"raw_affiliation_string":"Intel Corporation, Santa Clara, CA, USA","institution_ids":["https://openalex.org/I1343180700"]}]},{"author_position":"middle","author":{"id":"https://openalex.org/A5140229978","display_name":"Vincent Lu","orcid":"https://orcid.org/0009-0007-8746-8448"},"institutions":[{"id":"https://openalex.org/I1343180700","display_name":"Intel (United States)","ror":"https://ror.org/01ek73717","country_code":"US","type":"company","lineage":["https://openalex.org/I1343180700"]}],"countries":["US"],"is_corresponding":false,"raw_author_name":"Vincent Lu","raw_affiliation_strings":["Intel Corporation, Santa Clara, CA, USA"],"raw_orcid":"https://orcid.org/0009-0007-8746-8448","affiliations":[{"raw_affiliation_string":"Intel Corporation, Santa Clara, CA, USA","institution_ids":["https://openalex.org/I1343180700"]}]},{"author_position":"middle","author":{"id":"https://openalex.org/A5042651234","display_name":"Yen-Kuang Chen","orcid":"https://orcid.org/0000-0003-4546-9497"},"institutions":[{"id":"https://openalex.org/I1343180700","display_name":"Intel (United States)","ror":"https://ror.org/01ek73717","country_code":"US","type":"company","lineage":["https://openalex.org/I1343180700"]}],"countries":["US"],"is_corresponding":false,"raw_author_name":"Yen-Kuang Chen","raw_affiliation_strings":["Intel Corporation, Santa Clara, CA, USA"],"raw_orcid":"https://orcid.org/0000-0003-4546-9497","affiliations":[{"raw_affiliation_string":"Intel Corporation, Santa Clara, CA, USA","institution_ids":["https://openalex.org/I1343180700"]}]},{"author_position":"middle","author":{"id":"https://openalex.org/A5140248455","display_name":"Hong Wang","orcid":"https://orcid.org/0009-0008-8204-6454"},"institutions":[{"id":"https://openalex.org/I1343180700","display_name":"Intel (United States)","ror":"https://ror.org/01ek73717","country_code":"US","type":"company","lineage":["https://openalex.org/I1343180700"]}],"countries":["US"],"is_corresponding":false,"raw_author_name":"Hong Wang","raw_affiliation_strings":["Intel Corporation, Santa Clara, CA, USA"],"raw_orcid":"https://orcid.org/0009-0008-8204-6454","affiliations":[{"raw_affiliation_string":"Intel Corporation, Santa Clara, CA, USA","institution_ids":["https://openalex.org/I1343180700"]}]},{"author_position":"last","author":{"id":"https://openalex.org/A5020276955","display_name":"Xinmin Tian","orcid":"https://orcid.org/0000-0001-6228-924X"},"institutions":[{"id":"https://openalex.org/I1343180700","display_name":"Intel (United States)","ror":"https://ror.org/01ek73717","country_code":"US","type":"company","lineage":["https://openalex.org/I1343180700"]}],"countries":["US"],"is_corresponding":false,"raw_author_name":"Xinmin Tian","raw_affiliation_strings":["Intel Corporation, Santa Clara, CA, USA"],"raw_orcid":"https://orcid.org/0000-0001-6228-924X","affiliations":[{"raw_affiliation_string":"Intel Corporation, Santa Clara, CA, USA","institution_ids":["https://openalex.org/I1343180700"]}]}],"institutions":[],"countries_distinct_count":1,"institutions_distinct_count":2,"corresponding_author_ids":[],"corresponding_institution_ids":[],"apc_list":null,"apc_paid":null,"fwci":null,"has_fulltext":false,"cited_by_count":0,"citation_normalized_percentile":null,"cited_by_percentile_year":null,"biblio":{"volume":null,"issue":null,"first_page":"1","last_page":"5"},"is_retracted":false,"is_paratext":false,"is_xpac":false,"primary_topic":{"id":"https://openalex.org/T10054","display_name":"Parallel Computing and Optimization Techniques","score":0.9355999827384949,"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":0.9355999827384949,"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/T12326","display_name":"Network Packet Processing and Optimization","score":0.010900000110268593,"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/T10904","display_name":"Embedded Systems Design Techniques","score":0.010700000450015068,"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/porting","display_name":"Porting","score":0.9186000227928162},{"id":"https://openalex.org/keywords/software-portability","display_name":"Software portability","score":0.7505999803543091},{"id":"https://openalex.org/keywords/cuda","display_name":"CUDA","score":0.569100022315979},{"id":"https://openalex.org/keywords/task","display_name":"Task (project management)","score":0.4293000102043152},{"id":"https://openalex.org/keywords/coprocessor","display_name":"Coprocessor","score":0.42820000648498535},{"id":"https://openalex.org/keywords/implementation","display_name":"Implementation","score":0.36649999022483826},{"id":"https://openalex.org/keywords/kernel","display_name":"Kernel (algebra)","score":0.3513999879360199},{"id":"https://openalex.org/keywords/latency","display_name":"Latency (audio)","score":0.30730000138282776}],"concepts":[{"id":"https://openalex.org/C106251023","wikidata":"https://www.wikidata.org/wiki/Q851989","display_name":"Porting","level":3,"score":0.9186000227928162},{"id":"https://openalex.org/C41008148","wikidata":"https://www.wikidata.org/wiki/Q21198","display_name":"Computer science","level":0,"score":0.8672000169754028},{"id":"https://openalex.org/C63000827","wikidata":"https://www.wikidata.org/wiki/Q3080428","display_name":"Software portability","level":2,"score":0.7505999803543091},{"id":"https://openalex.org/C2778119891","wikidata":"https://www.wikidata.org/wiki/Q477690","display_name":"CUDA","level":2,"score":0.569100022315979},{"id":"https://openalex.org/C173608175","wikidata":"https://www.wikidata.org/wiki/Q232661","display_name":"Parallel computing","level":1,"score":0.49070000648498535},{"id":"https://openalex.org/C111919701","wikidata":"https://www.wikidata.org/wiki/Q9135","display_name":"Operating system","level":1,"score":0.43209999799728394},{"id":"https://openalex.org/C2780451532","wikidata":"https://www.wikidata.org/wiki/Q759676","display_name":"Task (project management)","level":2,"score":0.4293000102043152},{"id":"https://openalex.org/C86111242","wikidata":"https://www.wikidata.org/wiki/Q859595","display_name":"Coprocessor","level":2,"score":0.42820000648498535},{"id":"https://openalex.org/C118524514","wikidata":"https://www.wikidata.org/wiki/Q173212","display_name":"Computer architecture","level":1,"score":0.412200003862381},{"id":"https://openalex.org/C26713055","wikidata":"https://www.wikidata.org/wiki/Q245962","display_name":"Implementation","level":2,"score":0.36649999022483826},{"id":"https://openalex.org/C199360897","wikidata":"https://www.wikidata.org/wiki/Q9143","display_name":"Programming language","level":1,"score":0.3540000021457672},{"id":"https://openalex.org/C74193536","wikidata":"https://www.wikidata.org/wiki/Q574844","display_name":"Kernel (algebra)","level":2,"score":0.3513999879360199},{"id":"https://openalex.org/C149635348","wikidata":"https://www.wikidata.org/wiki/Q193040","display_name":"Embedded system","level":1,"score":0.32820001244544983},{"id":"https://openalex.org/C82876162","wikidata":"https://www.wikidata.org/wiki/Q17096504","display_name":"Latency (audio)","level":2,"score":0.30730000138282776},{"id":"https://openalex.org/C170723468","wikidata":"https://www.wikidata.org/wiki/Q182933","display_name":"x86","level":3,"score":0.3061000108718872},{"id":"https://openalex.org/C32802771","wikidata":"https://www.wikidata.org/wiki/Q2443617","display_name":"Port (circuit theory)","level":2,"score":0.3034000098705292},{"id":"https://openalex.org/C50630238","wikidata":"https://www.wikidata.org/wiki/Q971505","display_name":"General-purpose computing on graphics processing units","level":3,"score":0.3012000024318695},{"id":"https://openalex.org/C202491316","wikidata":"https://www.wikidata.org/wiki/Q272683","display_name":"Instruction set","level":2,"score":0.3003000020980835},{"id":"https://openalex.org/C75608658","wikidata":"https://www.wikidata.org/wiki/Q44395","display_name":"Pascal (unit)","level":2,"score":0.2969000041484833},{"id":"https://openalex.org/C169590947","wikidata":"https://www.wikidata.org/wiki/Q47506","display_name":"Compiler","level":2,"score":0.2939999997615814},{"id":"https://openalex.org/C154556556","wikidata":"https://www.wikidata.org/wiki/Q192969","display_name":"Computer multitasking","level":2,"score":0.28999999165534973},{"id":"https://openalex.org/C43364308","wikidata":"https://www.wikidata.org/wiki/Q8799","display_name":"Byte","level":2,"score":0.28189998865127563},{"id":"https://openalex.org/C34165917","wikidata":"https://www.wikidata.org/wiki/Q188267","display_name":"Programming paradigm","level":2,"score":0.2648000121116638},{"id":"https://openalex.org/C2777735758","wikidata":"https://www.wikidata.org/wiki/Q817765","display_name":"Path (computing)","level":2,"score":0.2581999897956848},{"id":"https://openalex.org/C2780615836","wikidata":"https://www.wikidata.org/wiki/Q2471869","display_name":"USable","level":2,"score":0.2581999897956848},{"id":"https://openalex.org/C166178840","wikidata":"https://www.wikidata.org/wiki/Q1571735","display_name":"Single-chip Cloud Computer","level":5,"score":0.2540999948978424}],"mesh":[],"locations_count":1,"locations":[{"id":"doi:10.1145/3811257.3811269","is_oa":true,"landing_page_url":"https://doi.org/10.1145/3811257.3811269","pdf_url":null,"source":null,"license":"cc-by","license_id":"https://openalex.org/licenses/cc-by","version":"publishedVersion","is_accepted":true,"is_published":true,"raw_source_name":"Proceedings of the International Workshop on OpenCL and SYCL","raw_type":"proceedings-article"}],"best_oa_location":{"id":"doi:10.1145/3811257.3811269","is_oa":true,"landing_page_url":"https://doi.org/10.1145/3811257.3811269","pdf_url":null,"source":null,"license":"cc-by","license_id":"https://openalex.org/licenses/cc-by","version":"publishedVersion","is_accepted":true,"is_published":true,"raw_source_name":"Proceedings of the International Workshop on OpenCL and SYCL","raw_type":"proceedings-article"},"sustainable_development_goals":[],"awards":[],"funders":[],"has_content":{"pdf":false,"grobid_xml":false},"content_urls":null,"referenced_works_count":1,"referenced_works":["https://openalex.org/W4416162128"],"related_works":[],"abstract_inverted_index":{"This":[0,51],"paper":[1],"presents":[2],"our":[3,111],"experience":[4],"porting":[5,90],"ThunderKittens,":[6],"a":[7,23],"CUDA-based":[8],"GPU":[9,141],"programming":[10],"framework":[11],"for":[12,18,48],"deep":[13],"learning":[14],"kernels,":[15],"to":[16,54,71,138],"SYCL":[17,112],"Intel":[19,49,72],"GPUs.":[20,50],"We":[21],"follow":[22],"three-stage":[24],"process:":[25],"automated":[26],"conversion,":[27],"manual":[28,89],"adaptation,":[29],"and":[30,42,58,74,87,123],"architecture-specific":[31],"optimization.":[32],"The":[33],"main":[34],"task":[35],"is":[36,91,136],"translating":[37],"ThunderKittens\u2019":[38],"C++":[39],"template":[40],"library":[41],"replacing":[43,61],"PTX-based":[44],"implementations":[45],"with":[46,66],"alternatives":[47],"required":[52,137],"changes":[53],"both":[55],"register":[56,67],"layout":[57],"memory":[59],"movement,":[60],"CUDA\u2019s":[62],"shared-memory-centric":[63],"execution":[64],"path":[65],"layouts":[68],"better":[69],"suited":[70],"GPUs":[73],"L1-prefetch-based":[75],"latency":[76],"hiding.":[77],"Our":[78],"results":[79,133],"show":[80,134],"where":[81,88],"current":[82],"CUDA-to-SYCL":[83],"translation":[84],"tools":[85],"help":[86],"still":[92],"required.":[93],"Although":[94],"full":[95],"performance":[96,122],"portability":[97],"remains":[98],"difficult":[99],"across":[100,109,143],"architectures,":[101],"the":[102,127],"same":[103],"high-level":[104],"abstractions":[105,142],"can":[106],"be":[107],"preserved":[108],"platforms:":[110],"port":[113,139],"comes":[114],"within":[115],"5%":[116],"of":[117],"Intel\u2019s":[118],"hand-optimized":[119],"oneDNN":[120],"GEMM":[121],"exceeds":[124],"it":[125],"on":[126],"fused":[128],"kernels":[129],"we":[130],"evaluated.":[131],"These":[132],"what":[135],"framework-level":[140],"architectures.":[144]},"counts_by_year":[],"updated_date":"2026-07-14T23:27:15.235271","created_date":"2026-07-09T00:00:00"}