2024-10-15 23:12:17 +08:00
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "https://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
< html xmlns = "http://www.w3.org/1999/xhtml" lang = "en-US" >
< head >
< meta http-equiv = "Content-Type" content = "text/xhtml;charset=UTF-8" / >
< meta http-equiv = "X-UA-Compatible" content = "IE=11" / >
2025-01-10 05:56:20 +08:00
< meta name = "generator" content = "Doxygen 1.13.1" / >
2024-10-15 23:12:17 +08:00
< meta name = "viewport" content = "width=device-width, initial-scale=1" / >
< title > MLX: mlx/backend/metal/jit/steel_gemm.h File Reference< / title >
< link href = "tabs.css" rel = "stylesheet" type = "text/css" / >
< script type = "text/javascript" src = "jquery.js" > < / script >
< script type = "text/javascript" src = "dynsections.js" > < / script >
< script type = "text/javascript" src = "clipboard.js" > < / script >
< link href = "navtree.css" rel = "stylesheet" type = "text/css" / >
2025-01-10 05:56:20 +08:00
< script type = "text/javascript" src = "navtreedata.js" > < / script >
< script type = "text/javascript" src = "navtree.js" > < / script >
2024-10-15 23:12:17 +08:00
< script type = "text/javascript" src = "resize.js" > < / script >
< script type = "text/javascript" src = "cookie.js" > < / script >
< link href = "search/search.css" rel = "stylesheet" type = "text/css" / >
< script type = "text/javascript" src = "search/searchdata.js" > < / script >
< script type = "text/javascript" src = "search/search.js" > < / script >
2025-01-10 05:56:20 +08:00
< script type = "text/javascript" >
/* @license magnet:?xt=urn:btih:d3d9a9a6595521f9666a5e94cc830dab83b65699& dn=expat.txt MIT */
$(function() { init_search(); });
/* @license-end */
< / script >
2024-10-15 23:12:17 +08:00
< link href = "doxygen.css" rel = "stylesheet" type = "text/css" / >
< / head >
< body >
< div id = "top" > <!-- do not remove this div, it is closed by doxygen! -->
< div id = "titlearea" >
< table cellspacing = "0" cellpadding = "0" >
< tbody >
< tr id = "projectrow" >
< td id = "projectalign" >
< div id = "projectname" > MLX
< / div >
< / td >
2025-01-10 05:56:20 +08:00
< td > < div id = "MSearchBox" class = "MSearchBoxInactive" >
< span class = "left" >
< span id = "MSearchSelect" onmouseover = "return searchBox.OnSearchSelectShow()" onmouseout = "return searchBox.OnSearchSelectHide()" >   < / span >
< input type = "text" id = "MSearchField" value = "" placeholder = "Search" accesskey = "S"
onfocus="searchBox.OnSearchFieldFocus(true)"
onblur="searchBox.OnSearchFieldFocus(false)"
onkeyup="searchBox.OnSearchFieldChange(event)"/>
< / span > < span class = "right" >
< a id = "MSearchClose" href = "javascript:searchBox.CloseResultsWindow()" > < img id = "MSearchCloseImg" border = "0" src = "search/close.svg" alt = "" / > < / a >
< / span >
< / div >
< / td >
2024-10-15 23:12:17 +08:00
< / tr >
< / tbody >
< / table >
< / div >
<!-- end header part -->
2025-01-10 05:56:20 +08:00
<!-- Generated by Doxygen 1.13.1 -->
2024-10-15 23:12:17 +08:00
< script type = "text/javascript" >
/* @license magnet:?xt=urn:btih:d3d9a9a6595521f9666a5e94cc830dab83b65699& dn=expat.txt MIT */
var searchBox = new SearchBox("searchBox", "search/",'.html');
/* @license-end */
< / script >
< script type = "text/javascript" >
/* @license magnet:?xt=urn:btih:d3d9a9a6595521f9666a5e94cc830dab83b65699& dn=expat.txt MIT */
$(function() { codefold.init(0); });
/* @license-end */
< / script >
2025-01-10 05:56:20 +08:00
< / div > <!-- top -->
< div id = "side-nav" class = "ui-resizable side-nav-resizable" >
< div id = "nav-tree" >
< div id = "nav-tree-contents" >
< div id = "nav-sync" class = "sync" > < / div >
< / div >
< / div >
< div id = "splitbar" style = "-moz-user-select:none;"
class="ui-resizable-handle">
< / div >
< / div >
2024-10-15 23:12:17 +08:00
< script type = "text/javascript" >
/* @license magnet:?xt=urn:btih:d3d9a9a6595521f9666a5e94cc830dab83b65699& dn=expat.txt MIT */
2025-01-10 05:56:20 +08:00
$(function(){initNavTree('steel__gemm_8h.html',''); initResizable(true); });
2024-10-15 23:12:17 +08:00
/* @license-end */
< / script >
2025-01-10 05:56:20 +08:00
< div id = "doc-content" >
2024-10-15 23:12:17 +08:00
<!-- window showing the filter options -->
< div id = "MSearchSelectWindow"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
onkeydown="return searchBox.OnSearchSelectKey(event)">
< / div >
<!-- iframe showing the search results (closed by default) -->
< div id = "MSearchResultsWindow" >
< div id = "MSearchResults" >
< div class = "SRPage" >
< div id = "SRIndex" >
< div id = "SRResults" > < / div >
< div class = "SRStatus" id = "Loading" > Loading...< / div >
< div class = "SRStatus" id = "Searching" > Searching...< / div >
< div class = "SRStatus" id = "NoMatches" > No Matches< / div >
< / div >
< / div >
< / div >
< / div >
< div class = "header" >
< div class = "summary" >
< a href = "#var-members" > Variables< / a > < / div >
< div class = "headertitle" > < div class = "title" > steel_gemm.h File Reference< / div > < / div >
< / div > <!-- header -->
< div class = "contents" >
< p > < a href = "steel__gemm_8h_source.html" > Go to the source code of this file.< / a > < / p >
< table class = "memberdecls" >
< tr class = "heading" > < td colspan = "2" > < h2 class = "groupheader" > < a id = "var-members" name = "var-members" > < / a >
Variables< / h2 > < / td > < / tr >
< tr class = "memitem:a4c6009fd5357b730805f2fd4ba6e093e" id = "r_a4c6009fd5357b730805f2fd4ba6e093e" > < td class = "memItemLeft" align = "right" valign = "top" > constexpr std::string_view  < / td > < td class = "memItemRight" valign = "bottom" > < a class = "el" href = "#a4c6009fd5357b730805f2fd4ba6e093e" > steel_gemm_fused_kernels< / a > < / td > < / tr >
< tr class = "separator:a4c6009fd5357b730805f2fd4ba6e093e" > < td class = "memSeparator" colspan = "2" >   < / td > < / tr >
< tr class = "memitem:a62a358fd3ec5365081920d07aceb581c" id = "r_a62a358fd3ec5365081920d07aceb581c" > < td class = "memItemLeft" align = "right" valign = "top" > constexpr std::string_view  < / td > < td class = "memItemRight" valign = "bottom" > < a class = "el" href = "#a62a358fd3ec5365081920d07aceb581c" > steel_gemm_masked_kernels< / a > < / td > < / tr >
< tr class = "separator:a62a358fd3ec5365081920d07aceb581c" > < td class = "memSeparator" colspan = "2" >   < / td > < / tr >
< tr class = "memitem:a92108ab01d826e38bca83d8569b947d9" id = "r_a92108ab01d826e38bca83d8569b947d9" > < td class = "memItemLeft" align = "right" valign = "top" > constexpr std::string_view  < / td > < td class = "memItemRight" valign = "bottom" > < a class = "el" href = "#a92108ab01d826e38bca83d8569b947d9" > steel_gemm_splitk_kernels< / a > < / td > < / tr >
< tr class = "separator:a92108ab01d826e38bca83d8569b947d9" > < td class = "memSeparator" colspan = "2" >   < / td > < / tr >
< tr class = "memitem:a144a64b8d94f0371fb144e2cc308fcf9" id = "r_a144a64b8d94f0371fb144e2cc308fcf9" > < td class = "memItemLeft" align = "right" valign = "top" > constexpr std::string_view  < / td > < td class = "memItemRight" valign = "bottom" > < a class = "el" href = "#a144a64b8d94f0371fb144e2cc308fcf9" > steel_gemm_splitk_accum_kernels< / a > < / td > < / tr >
< tr class = "separator:a144a64b8d94f0371fb144e2cc308fcf9" > < td class = "memSeparator" colspan = "2" >   < / td > < / tr >
< tr class = "memitem:a40a86e1381c241aba8511e51a981a4bf" id = "r_a40a86e1381c241aba8511e51a981a4bf" > < td class = "memItemLeft" align = "right" valign = "top" > constexpr std::string_view  < / td > < td class = "memItemRight" valign = "bottom" > < a class = "el" href = "#a40a86e1381c241aba8511e51a981a4bf" > steel_gemm_splitk_accum_axbpy_kernels< / a > < / td > < / tr >
< tr class = "separator:a40a86e1381c241aba8511e51a981a4bf" > < td class = "memSeparator" colspan = "2" >   < / td > < / tr >
< / table >
< h2 class = "groupheader" > Variable Documentation< / h2 >
< a id = "a4c6009fd5357b730805f2fd4ba6e093e" name = "a4c6009fd5357b730805f2fd4ba6e093e" > < / a >
< h2 class = "memtitle" > < span class = "permalink" > < a href = "#a4c6009fd5357b730805f2fd4ba6e093e" > ◆   < / a > < / span > steel_gemm_fused_kernels< / h2 >
< div class = "memitem" >
< div class = "memproto" >
< table class = "mlabels" >
< tr >
< td class = "mlabels-left" >
< table class = "memname" >
< tr >
< td class = "memname" > std::string_view steel_gemm_fused_kernels< / td >
< / tr >
< / table >
< / td >
< td class = "mlabels-right" >
2025-01-10 05:56:20 +08:00
< span class = "mlabels" > < span class = "mlabel constexpr" > constexpr< / span > < / span > < / td >
2024-10-15 23:12:17 +08:00
< / tr >
< / table >
< / div > < div class = "memdoc" >
< b > Initial value:< / b > < div class = "fragment" > < div class = "line" > = R< span class = "stringliteral" > " (< / span > < / div >
< div class = "line" > < span class = "stringliteral" > template [[host_name(" {name}" )]]< / span > < / div >
< div class = "line" > < span class = "stringliteral" > [[kernel]] void gemm< {itype}, {bm}, {bn}, {bk}, {wm}, {wn}, {trans_a}, {trans_b}, float> (< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const device {itype} *A [[buffer(0)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const device {itype} *B [[buffer(1)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const device {itype} *C [[buffer(2), function_constant(use_out_source)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > device {itype} *D [[buffer(3)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant GEMMParams* params [[buffer(4)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant GEMMAddMMParams* addmm_params [[buffer(5), function_constant(use_out_source)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant int* batch_shape [[buffer(6)]],< / span > < / div >
2025-01-10 05:56:20 +08:00
< div class = "line" > < span class = "stringliteral" > const constant int64_t* batch_strides [[buffer(7)]],< / span > < / div >
2024-10-15 23:12:17 +08:00
< div class = "line" > < span class = "stringliteral" > const constant uint32_t* lhs_indices [[buffer(10), function_constant(do_gather)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant uint32_t* rhs_indices [[buffer(11), function_constant(do_gather)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant uint32_t* C_indices [[buffer(12), function_constant(gather_bias)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant int* operand_shape [[buffer(13), function_constant(do_gather)]],< / span > < / div >
2025-01-10 05:56:20 +08:00
< div class = "line" > < span class = "stringliteral" > const constant int64_t* operand_strides [[buffer(14), function_constant(do_gather)]],< / span > < / div >
2024-10-15 23:12:17 +08:00
< div class = "line" > < span class = "stringliteral" > const constant packed_int3& operand_batch_ndim [[buffer(15), function_constant(do_gather)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > uint simd_lane_id [[thread_index_in_simdgroup]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > uint simd_group_id [[simdgroup_index_in_threadgroup]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > uint3 tid [[threadgroup_position_in_grid]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > uint3 lid [[thread_position_in_threadgroup]]);< / span > < / div >
< div class = "line" > < span class = "stringliteral" > )" < / span > < / div >
< / div > <!-- fragment -->
< / div >
< / div >
< a id = "a62a358fd3ec5365081920d07aceb581c" name = "a62a358fd3ec5365081920d07aceb581c" > < / a >
< h2 class = "memtitle" > < span class = "permalink" > < a href = "#a62a358fd3ec5365081920d07aceb581c" > ◆   < / a > < / span > steel_gemm_masked_kernels< / h2 >
< div class = "memitem" >
< div class = "memproto" >
< table class = "mlabels" >
< tr >
< td class = "mlabels-left" >
< table class = "memname" >
< tr >
< td class = "memname" > std::string_view steel_gemm_masked_kernels< / td >
< / tr >
< / table >
< / td >
< td class = "mlabels-right" >
2025-01-10 05:56:20 +08:00
< span class = "mlabels" > < span class = "mlabel constexpr" > constexpr< / span > < / span > < / td >
2024-10-15 23:12:17 +08:00
< / tr >
< / table >
< / div > < div class = "memdoc" >
< b > Initial value:< / b > < div class = "fragment" > < div class = "line" > = R< span class = "stringliteral" > " (< / span > < / div >
< div class = "line" > < span class = "stringliteral" > template [[host_name(" {name}" )]] [[kernel]] void< / span > < / div >
< div class = "line" > < span class = "stringliteral" > block_masked_gemm< < / span > < / div >
< div class = "line" > < span class = "stringliteral" > {itype},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {outmasktype},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {opmasktype},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {bm},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {bn},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {bk},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {wm},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {wn},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {trans_a},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {trans_b},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {mn_aligned},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {k_aligned}> (< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const device {itype}* A [[buffer(0)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const device {itype}* B [[buffer(1)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > device {itype}* D [[buffer(3)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant GEMMParams* params [[buffer(4)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant int* batch_shape [[buffer(6)]],< / span > < / div >
2025-01-10 05:56:20 +08:00
< div class = "line" > < span class = "stringliteral" > const constant int64_t* batch_strides [[buffer(7)]],< / span > < / div >
2024-10-15 23:12:17 +08:00
< div class = "line" > < span class = "stringliteral" > const device {outmasktype}* out_mask [[buffer(10)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const device {opmasktype}* lhs_mask [[buffer(11)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const device {opmasktype}* rhs_mask [[buffer(12)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant int* mask_strides [[buffer(13)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > uint simd_lane_id [[thread_index_in_simdgroup]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > uint simd_group_id [[simdgroup_index_in_threadgroup]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > uint3 tid [[threadgroup_position_in_grid]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > uint3 lid [[thread_position_in_threadgroup]]);< / span > < / div >
< div class = "line" > < span class = "stringliteral" > )" < / span > < / div >
< / div > <!-- fragment -->
< / div >
< / div >
< a id = "a40a86e1381c241aba8511e51a981a4bf" name = "a40a86e1381c241aba8511e51a981a4bf" > < / a >
< h2 class = "memtitle" > < span class = "permalink" > < a href = "#a40a86e1381c241aba8511e51a981a4bf" > ◆   < / a > < / span > steel_gemm_splitk_accum_axbpy_kernels< / h2 >
< div class = "memitem" >
< div class = "memproto" >
< table class = "mlabels" >
< tr >
< td class = "mlabels-left" >
< table class = "memname" >
< tr >
< td class = "memname" > std::string_view steel_gemm_splitk_accum_axbpy_kernels< / td >
< / tr >
< / table >
< / td >
< td class = "mlabels-right" >
2025-01-10 05:56:20 +08:00
< span class = "mlabels" > < span class = "mlabel constexpr" > constexpr< / span > < / span > < / td >
2024-10-15 23:12:17 +08:00
< / tr >
< / table >
< / div > < div class = "memdoc" >
< b > Initial value:< / b > < div class = "fragment" > < div class = "line" > = R< span class = "stringliteral" > " (< / span > < / div >
< div class = "line" > < span class = "stringliteral" > template [[host_name(" {name}" )]] [[kernel]] void< / span > < / div >
< div class = "line" > < span class = "stringliteral" > gemm_splitk_accum_axpby< {atype}, {otype}> (< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const device {atype}* C_split [[buffer(0)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > device {otype}* D [[buffer(1)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant int& k_partitions [[buffer(2)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant int& partition_stride [[buffer(3)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant int& ldd [[buffer(4)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const device {otype}* C [[buffer(5)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant int& ldc [[buffer(6)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant int& fdc [[buffer(7)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant float& alpha [[buffer(8)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant float& beta [[buffer(9)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > uint2 gid [[thread_position_in_grid]]);< / span > < / div >
< div class = "line" > < span class = "stringliteral" > )" < / span > < / div >
< / div > <!-- fragment -->
< / div >
< / div >
< a id = "a144a64b8d94f0371fb144e2cc308fcf9" name = "a144a64b8d94f0371fb144e2cc308fcf9" > < / a >
< h2 class = "memtitle" > < span class = "permalink" > < a href = "#a144a64b8d94f0371fb144e2cc308fcf9" > ◆   < / a > < / span > steel_gemm_splitk_accum_kernels< / h2 >
< div class = "memitem" >
< div class = "memproto" >
< table class = "mlabels" >
< tr >
< td class = "mlabels-left" >
< table class = "memname" >
< tr >
< td class = "memname" > std::string_view steel_gemm_splitk_accum_kernels< / td >
< / tr >
< / table >
< / td >
< td class = "mlabels-right" >
2025-01-10 05:56:20 +08:00
< span class = "mlabels" > < span class = "mlabel constexpr" > constexpr< / span > < / span > < / td >
2024-10-15 23:12:17 +08:00
< / tr >
< / table >
< / div > < div class = "memdoc" >
< b > Initial value:< / b > < div class = "fragment" > < div class = "line" > = R< span class = "stringliteral" > " (< / span > < / div >
< div class = "line" > < span class = "stringliteral" > template [[host_name(" {name}" )]] [[kernel]] void< / span > < / div >
< div class = "line" > < span class = "stringliteral" > gemm_splitk_accum< {atype}, {otype}> (< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const device {atype}* C_split [[buffer(0)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > device {otype}* D [[buffer(1)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant int& k_partitions [[buffer(2)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant int& partition_stride [[buffer(3)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant int& ldd [[buffer(4)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > uint2 gid [[thread_position_in_grid]]);< / span > < / div >
< div class = "line" > < span class = "stringliteral" > )" < / span > < / div >
< / div > <!-- fragment -->
< / div >
< / div >
< a id = "a92108ab01d826e38bca83d8569b947d9" name = "a92108ab01d826e38bca83d8569b947d9" > < / a >
< h2 class = "memtitle" > < span class = "permalink" > < a href = "#a92108ab01d826e38bca83d8569b947d9" > ◆   < / a > < / span > steel_gemm_splitk_kernels< / h2 >
< div class = "memitem" >
< div class = "memproto" >
< table class = "mlabels" >
< tr >
< td class = "mlabels-left" >
< table class = "memname" >
< tr >
< td class = "memname" > std::string_view steel_gemm_splitk_kernels< / td >
< / tr >
< / table >
< / td >
< td class = "mlabels-right" >
2025-01-10 05:56:20 +08:00
< span class = "mlabels" > < span class = "mlabel constexpr" > constexpr< / span > < / span > < / td >
2024-10-15 23:12:17 +08:00
< / tr >
< / table >
< / div > < div class = "memdoc" >
< b > Initial value:< / b > < div class = "fragment" > < div class = "line" > = R< span class = "stringliteral" > " (< / span > < / div >
< div class = "line" > < span class = "stringliteral" > template [[host_name(" {name}" )]] [[kernel]] void< / span > < / div >
< div class = "line" > < span class = "stringliteral" > gemm_splitk< < / span > < / div >
< div class = "line" > < span class = "stringliteral" > {itype},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {otype},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {bm},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {bn},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {bk},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {wm},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {wn},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {trans_a},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {trans_b},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {mn_aligned},< / span > < / div >
< div class = "line" > < span class = "stringliteral" > {k_aligned}> (< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const device {itype}* A [[buffer(0)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const device {itype}* B [[buffer(1)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > device {otype}* C [[buffer(2)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > const constant GEMMSpiltKParams* params [[buffer(3)]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > uint simd_lane_id [[thread_index_in_simdgroup]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > uint simd_group_id [[simdgroup_index_in_threadgroup]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > uint3 tid [[threadgroup_position_in_grid]],< / span > < / div >
< div class = "line" > < span class = "stringliteral" > uint3 lid [[thread_position_in_threadgroup]]);< / span > < / div >
< div class = "line" > < span class = "stringliteral" > )" < / span > < / div >
< / div > <!-- fragment -->
< / div >
< / div >
< / div > <!-- contents -->
< / div > <!-- doc - content -->
2025-01-10 05:56:20 +08:00
<!-- start footer part -->
< div id = "nav-path" class = "navpath" > <!-- id is needed for treeview function! -->
< ul >
< li class = "navelem" > < a class = "el" href = "dir_938ab0ecf10b8b860ff766c820f665fd.html" > mlx< / a > < / li > < li class = "navelem" > < a class = "el" href = "dir_1d446c9bd3c99228254c9484e0bc5c06.html" > backend< / a > < / li > < li class = "navelem" > < a class = "el" href = "dir_d0c977ea65824390717cdb7efc36c157.html" > metal< / a > < / li > < li class = "navelem" > < a class = "el" href = "dir_4336740ec0075891704443b417fef6cb.html" > jit< / a > < / li > < li class = "navelem" > < a class = "el" href = "steel__gemm_8h.html" > steel_gemm.h< / a > < / li >
< li class = "footer" > Generated by < a href = "https://www.doxygen.org/index.html" > < img class = "footer" src = "doxygen.svg" width = "104" height = "31" alt = "doxygen" / > < / a > 1.13.1 < / li >
< / ul >
< / div >
2024-10-15 23:12:17 +08:00
< / body >
< / html >
