zhangyiss/mlx
1
0
Fork 0
mirror of https://github.com/ml-explore/mlx.git synced 2025-09-18 18:28:12 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

rebase

Browse Source
This commit is contained in:
CircleCI Docs
2024-11-22 20:24:16 +00:00
parent 6433a22733
commit 8ba8a39e43
905 changed files with 30035 additions and 16934 deletions
Download Patch File Download Diff File Expand all files Collapse all files

18
docs/build/html/kernels_2gemv__masked_8h_source.html vendored
View File

@@ -777,12 +777,12 @@ $(function(){ initResizable(false); });
<div class="line"><a id="l00663" name="l00663"></a><span class="lineno"> 663</span> </div>
<div class="line"><a id="l00664" name="l00664"></a><span class="lineno"> 664</span> <span class="comment">// Update batch offsets</span></div>
<div class="line"><a id="l00665" name="l00665"></a><span class="lineno"> 665</span> <span class="keywordflow">if</span> (kDoNCBatch) {</div>
<div class="line"><a id="l00666" name="l00666"></a><span class="lineno"> 666</span> in_vec += <a class="code hl_function" href="backend_2metal_2kernels_2utils_8h.html#a8fd0c8fc6058e650fc99bca8b6acd7d1">elem_to_loc</a>(tid.z, batch_shape, vector_batch_stride, batch_ndim);</div>
<div class="line"><a id="l00667" name="l00667"></a><span class="lineno"> 667</span> mat += <a class="code hl_function" href="backend_2metal_2kernels_2utils_8h.html#a8fd0c8fc6058e650fc99bca8b6acd7d1">elem_to_loc</a>(tid.z, batch_shape, matrix_batch_stride, batch_ndim);</div>
<div class="line"><a id="l00666" name="l00666"></a><span class="lineno"> 666</span> in_vec += <a class="code hl_function" href="backend_2metal_2kernels_2utils_8h.html#a22eaa505dbc7dd2a63a895f2e16712f5">elem_to_loc</a>(tid.z, batch_shape, vector_batch_stride, batch_ndim);</div>
<div class="line"><a id="l00667" name="l00667"></a><span class="lineno"> 667</span> mat += <a class="code hl_function" href="backend_2metal_2kernels_2utils_8h.html#a22eaa505dbc7dd2a63a895f2e16712f5">elem_to_loc</a>(tid.z, batch_shape, matrix_batch_stride, batch_ndim);</div>
<div class="line"><a id="l00668" name="l00668"></a><span class="lineno"> 668</span> </div>
<div class="line"><a id="l00669" name="l00669"></a><span class="lineno"> 669</span> <span class="keywordflow">if</span> (has_output_mask) {</div>
<div class="line"><a id="l00670" name="l00670"></a><span class="lineno"> 670</span> out_mask +=</div>
<div class="line"><a id="l00671" name="l00671"></a><span class="lineno"> 671</span> <a class="code hl_function" href="backend_2metal_2kernels_2utils_8h.html#a8fd0c8fc6058e650fc99bca8b6acd7d1">elem_to_loc</a>(tid.z, batch_shape, mask_batch_strides, batch_ndim);</div>
<div class="line"><a id="l00671" name="l00671"></a><span class="lineno"> 671</span> <a class="code hl_function" href="backend_2metal_2kernels_2utils_8h.html#a22eaa505dbc7dd2a63a895f2e16712f5">elem_to_loc</a>(tid.z, batch_shape, mask_batch_strides, batch_ndim);</div>
<div class="line"><a id="l00672" name="l00672"></a><span class="lineno"> 672</span> mask_batch_strides += batch_ndim;</div>
<div class="line"><a id="l00673" name="l00673"></a><span class="lineno"> 673</span> }</div>
<div class="line"><a id="l00674" name="l00674"></a><span class="lineno"> 674</span> </div>
@@ -876,12 +876,12 @@ $(function(){ initResizable(false); });
<div class="line"><a id="l00763" name="l00763"></a><span class="lineno"> 763</span> </div>
<div class="line"><a id="l00764" name="l00764"></a><span class="lineno"> 764</span> <span class="comment">// Update batch offsets</span></div>
<div class="line"><a id="l00765" name="l00765"></a><span class="lineno"> 765</span> <span class="keywordflow">if</span> (kDoNCBatch) {</div>
<div class="line"><a id="l00766" name="l00766"></a><span class="lineno"> 766</span> in_vec += <a class="code hl_function" href="backend_2metal_2kernels_2utils_8h.html#a8fd0c8fc6058e650fc99bca8b6acd7d1">elem_to_loc</a>(tid.z, batch_shape, vector_batch_stride, batch_ndim);</div>
<div class="line"><a id="l00767" name="l00767"></a><span class="lineno"> 767</span> mat += <a class="code hl_function" href="backend_2metal_2kernels_2utils_8h.html#a8fd0c8fc6058e650fc99bca8b6acd7d1">elem_to_loc</a>(tid.z, batch_shape, matrix_batch_stride, batch_ndim);</div>
<div class="line"><a id="l00766" name="l00766"></a><span class="lineno"> 766</span> in_vec += <a class="code hl_function" href="backend_2metal_2kernels_2utils_8h.html#a22eaa505dbc7dd2a63a895f2e16712f5">elem_to_loc</a>(tid.z, batch_shape, vector_batch_stride, batch_ndim);</div>
<div class="line"><a id="l00767" name="l00767"></a><span class="lineno"> 767</span> mat += <a class="code hl_function" href="backend_2metal_2kernels_2utils_8h.html#a22eaa505dbc7dd2a63a895f2e16712f5">elem_to_loc</a>(tid.z, batch_shape, matrix_batch_stride, batch_ndim);</div>
<div class="line"><a id="l00768" name="l00768"></a><span class="lineno"> 768</span> </div>
<div class="line"><a id="l00769" name="l00769"></a><span class="lineno"> 769</span> <span class="keywordflow">if</span> (has_output_mask) {</div>
<div class="line"><a id="l00770" name="l00770"></a><span class="lineno"> 770</span> out_mask +=</div>
<div class="line"><a id="l00771" name="l00771"></a><span class="lineno"> 771</span> <a class="code hl_function" href="backend_2metal_2kernels_2utils_8h.html#a8fd0c8fc6058e650fc99bca8b6acd7d1">elem_to_loc</a>(tid.z, batch_shape, mask_batch_strides, batch_ndim);</div>
<div class="line"><a id="l00771" name="l00771"></a><span class="lineno"> 771</span> <a class="code hl_function" href="backend_2metal_2kernels_2utils_8h.html#a22eaa505dbc7dd2a63a895f2e16712f5">elem_to_loc</a>(tid.z, batch_shape, mask_batch_strides, batch_ndim);</div>
<div class="line"><a id="l00772" name="l00772"></a><span class="lineno"> 772</span> mask_batch_strides += batch_ndim;</div>
<div class="line"><a id="l00773" name="l00773"></a><span class="lineno"> 773</span> }</div>
<div class="line"><a id="l00774" name="l00774"></a><span class="lineno"> 774</span> </div>
@@ -933,13 +933,13 @@ $(function(){ initResizable(false); });
</div>
<div class="ttc" id="abackend_2metal_2kernels_2steel_2utils_8h_html"><div class="ttname"><a href="backend_2metal_2kernels_2steel_2utils_8h.html">utils.h</a></div></div>
<div class="ttc" id="abackend_2metal_2kernels_2steel_2utils_8h_html_aaf4974425147d6f26d031691e321637f"><div class="ttname"><a href="backend_2metal_2kernels_2steel_2utils_8h.html#aaf4974425147d6f26d031691e321637f">elem_to_loc_broadcast</a></div><div class="ttdeci">METAL_FUNC ulong2 elem_to_loc_broadcast(uint elem, constant const int *shape, constant const size_t *a_strides, constant const size_t *b_strides, int ndim)</div><div class="ttdef"><b>Definition</b> utils.h:7</div></div>
<div class="ttc" id="abackend_2metal_2kernels_2utils_8h_html_a8fd0c8fc6058e650fc99bca8b6acd7d1"><div class="ttname"><a href="backend_2metal_2kernels_2utils_8h.html#a8fd0c8fc6058e650fc99bca8b6acd7d1">elem_to_loc</a></div><div class="ttdeci">METAL_FUNC stride_t elem_to_loc(uint elem, constant const int *shape, constant const stride_t *strides, int ndim)</div><div class="ttdef"><b>Definition</b> utils.h:87</div></div>
<div class="ttc" id="abackend_2metal_2kernels_2utils_8h_html_a22eaa505dbc7dd2a63a895f2e16712f5"><div class="ttname"><a href="backend_2metal_2kernels_2utils_8h.html#a22eaa505dbc7dd2a63a895f2e16712f5">elem_to_loc</a></div><div class="ttdeci">METAL_FUNC IdxT elem_to_loc(uint elem, constant const int *shape, constant const StrideT *strides, int ndim)</div><div class="ttdef"><b>Definition</b> utils.h:93</div></div>
<div class="ttc" id="akernels_2gemv__masked_8h_html_a0386011c52d03e60885a31e6fbd903dd"><div class="ttname"><a href="kernels_2gemv__masked_8h.html#a0386011c52d03e60885a31e6fbd903dd">MLX_MTL_CONST</a></div><div class="ttdeci">#define MLX_MTL_CONST</div><div class="ttdef"><b>Definition</b> gemv_masked.h:7</div></div>
<div class="ttc" id="akernels_2gemv__masked_8h_html_a069b682d7d21827461544817d722bfd3"><div class="ttname"><a href="kernels_2gemv__masked_8h.html#a069b682d7d21827461544817d722bfd3">MLX_MTL_PRAGMA_UNROLL</a></div><div class="ttdeci">#define MLX_MTL_PRAGMA_UNROLL</div><div class="ttdef"><b>Definition</b> gemv_masked.h:8</div></div>
<div class="ttc" id="akernels_2gemv__masked_8h_html_a0c8d353fc453e448b2d0ed9a19431b63"><div class="ttname"><a href="kernels_2gemv__masked_8h.html#a0c8d353fc453e448b2d0ed9a19431b63">gemv_t_masked</a></div><div class="ttdeci">void gemv_t_masked(const device T *mat, const device T *in_vec, device T *out_vec, const constant int &amp;in_vec_size, const constant int &amp;out_vec_size, const constant int &amp;marix_ld, const constant int &amp;batch_ndim, const constant int *batch_shape, const constant size_t *vector_batch_stride, const constant size_t *matrix_batch_stride, const device out_mask_t *out_mask, const device op_mask_t *mat_mask, const device op_mask_t *vec_mask, const constant int *mask_strides, const constant size_t *mask_batch_strides, uint3 tid, uint3 lid, uint simd_gid, uint simd_lid)</div><div class="ttdoc">Vector matrix multiplication.</div><div class="ttdef"><b>Definition</b> gemv_masked.h:736</div></div>
<div class="ttc" id="akernels_2gemv__masked_8h_html_ab3070d14cdecb1dd7dc220a551da6b7b"><div class="ttname"><a href="kernels_2gemv__masked_8h.html#ab3070d14cdecb1dd7dc220a551da6b7b">gemv_masked</a></div><div class="ttdeci">void gemv_masked(const device T *mat, const device T *in_vec, device T *out_vec, const constant int &amp;in_vec_size, const constant int &amp;out_vec_size, const constant int &amp;marix_ld, const constant int &amp;batch_ndim, const constant int *batch_shape, const constant size_t *vector_batch_stride, const constant size_t *matrix_batch_stride, const device out_mask_t *out_mask, const device op_mask_t *mat_mask, const device op_mask_t *vec_mask, const constant int *mask_strides, const constant size_t *mask_batch_strides, uint3 tid, uint3 lid, uint simd_gid, uint simd_lid)</div><div class="ttdoc">Matrix vector multiplication.</div><div class="ttdef"><b>Definition</b> gemv_masked.h:636</div></div>
<div class="ttc" id="anamespacemetal_html"><div class="ttname"><a href="namespacemetal.html">metal</a></div><div class="ttdef"><b>Definition</b> bf16.h:265</div></div>
<div class="ttc" id="anamespacemetal_html_af6e2dd7ae087aba6abac4f0350b7611c"><div class="ttname"><a href="namespacemetal.html#af6e2dd7ae087aba6abac4f0350b7611c">metal::simd_shuffle_down</a></div><div class="ttdeci">METAL_FUNC bfloat16_t simd_shuffle_down(bfloat16_t data, ushort delta)</div><div class="ttdef"><b>Definition</b> bf16_math.h:391</div></div>
<div class="ttc" id="anamespacemetal_html"><div class="ttname"><a href="namespacemetal.html">metal</a></div><div class="ttdef"><b>Definition</b> bf16_math.h:226</div></div>
<div class="ttc" id="anamespacemetal_html_af6e2dd7ae087aba6abac4f0350b7611c"><div class="ttname"><a href="namespacemetal.html#af6e2dd7ae087aba6abac4f0350b7611c">metal::simd_shuffle_down</a></div><div class="ttdeci">METAL_FUNC bfloat16_t simd_shuffle_down(bfloat16_t data, ushort delta)</div><div class="ttdef"><b>Definition</b> bf16_math.h:377</div></div>
<div class="ttc" id="astruct___no_mask_html"><div class="ttname"><a href="struct___no_mask.html">_NoMask</a></div><div class="ttdef"><b>Definition</b> gemv_masked.h:10</div></div>
<div class="ttc" id="astruct___no_mask_html_a0c4a4557d5c97ceafe3a2c4e521cdf7e"><div class="ttname"><a href="struct___no_mask.html#a0c4a4557d5c97ceafe3a2c4e521cdf7e">_NoMask::x</a></div><div class="ttdeci">char x</div><div class="ttdef"><b>Definition</b> gemv_masked.h:11</div></div>
<div class="ttc" id="astruct_g_e_m_v_kernel_html"><div class="ttname"><a href="struct_g_e_m_v_kernel.html">GEMVKernel</a></div><div class="ttdef"><b>Definition</b> gemv_masked.h:48</div></div>