build based on e8f16bf

dev/api/array/index.html

@@ -20,4 +20,4 @@
 3-element MtlVector{Int64, Metal.MTL.MTLResourceStorageModePrivate}:
  1
  2
- 3</code></pre></div><a class="docs-sourcelink" target="_blank" href="https://github.com/JuliaGPU/Metal.jl/blob/630278e9a9ef52b559590512918433989c354673/src/array.jl#L339-L380">source</a></section></article></article><nav class="docs-footer"><a class="docs-footer-prevpage" href="../kernel/">« Kernel programming</a><a class="docs-footer-nextpage" href="../mps/">Metal Performance Shaders »</a><div class="flexbox-break"></div><p class="footer-message">Powered by <a href="https://github.com/JuliaDocs/Documenter.jl">Documenter.jl</a> and the <a href="https://julialang.org/">Julia Programming Language</a>.</p></nav></div><div class="modal" id="documenter-settings"><div class="modal-background"></div><div class="modal-card"><header class="modal-card-head"><p class="modal-card-title">Settings</p><button class="delete"></button></header><section class="modal-card-body"><p><label class="label">Theme</label><div class="select"><select id="documenter-themepicker"><option value="documenter-light">documenter-light</option><option value="documenter-dark">documenter-dark</option></select></div></p><hr/><p>This document was generated with <a href="https://github.com/JuliaDocs/Documenter.jl">Documenter.jl</a> version 0.27.23 on <span class="colophon-date" title="Monday 16 October 2023 06:12">Monday 16 October 2023</span>. Using Julia version 1.8.5.</p></section><footer class="modal-card-foot"></footer></div></div></div></body></html>
+ 3</code></pre></div><a class="docs-sourcelink" target="_blank" href="https://github.com/JuliaGPU/Metal.jl/blob/e8f16bfc332889aef19b5b486b5a4fdc560a5e45/src/array.jl#L339-L380">source</a></section></article></article><nav class="docs-footer"><a class="docs-footer-prevpage" href="../kernel/">« Kernel programming</a><a class="docs-footer-nextpage" href="../mps/">Metal Performance Shaders »</a><div class="flexbox-break"></div><p class="footer-message">Powered by <a href="https://github.com/JuliaDocs/Documenter.jl">Documenter.jl</a> and the <a href="https://julialang.org/">Julia Programming Language</a>.</p></nav></div><div class="modal" id="documenter-settings"><div class="modal-background"></div><div class="modal-card"><header class="modal-card-head"><p class="modal-card-title">Settings</p><button class="delete"></button></header><section class="modal-card-body"><p><label class="label">Theme</label><div class="select"><select id="documenter-themepicker"><option value="documenter-light">documenter-light</option><option value="documenter-dark">documenter-dark</option></select></div></p><hr/><p>This document was generated with <a href="https://github.com/JuliaDocs/Documenter.jl">Documenter.jl</a> version 0.27.23 on <span class="colophon-date" title="Monday 30 October 2023 04:03">Monday 30 October 2023</span>. Using Julia version 1.8.5.</p></section><footer class="modal-card-foot"></footer></div></div></div></body></html>

dev/faq/contributing/index.html

@@ -7,4 +7,4 @@
                         uint i [[thread_position_in_grid]])
 {
     atomic_store_explicit(&out[i], 0.0f, memory_order_relaxed);
-}</code></pre><p>To compile with Metal&#39;s tools and emit human-readable IR, run something roughly along the lines of: <code>xcrun metal -S -emit-llvm dummy_kernel.metal</code></p><p>This will create a <code>.ll</code> file that you can then parse for whatever information you need. Be sure to double-check the metadata at the bottom for any significant changes your functionality introduces.</p><p>Test with different types and configurations to see what changes are caused. Also ensure that when writing very simple kernels, whatever you&#39;re interested in doesn&#39;t get optimized away. Double-check that the kernel&#39;s IR makes sense for what you wrote.</p><h2 id="Metal-Performance-Shaders"><a class="docs-heading-anchor" href="#Metal-Performance-Shaders">Metal Performance Shaders</a><a id="Metal-Performance-Shaders-1"></a><a class="docs-heading-anchor-permalink" href="#Metal-Performance-Shaders" title="Permalink"></a></h2><p>Metal exposes a special interface to its library of optimized kernels. Rather than accepting the normal set of input GPU data structures, it requires special <code>MPS</code> datatypes that assume row-major memory layout. As this is not the Julia default, adapt accordingly. Adding MPS functionality should be mostly straightforward, so this can be an easy entry point to helping. To get started, you can have a look at the <a href="https://developer.apple.com/documentation/metalperformanceshaders?language=objc">Metal Performance Shaders Documentation</a> from Apple.</p><h2 id="Exposing-your-Interface"><a class="docs-heading-anchor" href="#Exposing-your-Interface">Exposing your Interface</a><a id="Exposing-your-Interface-1"></a><a class="docs-heading-anchor-permalink" href="#Exposing-your-Interface" title="Permalink"></a></h2><p>There are varying degrees of user-facing interfaces from Metal.jl. At the lowest level is <code>Metal.MTL.xxx</code>. This is for low-level functionality close to or at bare Objective-C, or things that a normal user wouldn&#39;t directly be using. <code>Metal.MPS.xxx</code> is for Metal Performance Shader specifics (like <code>MPSMatrix</code>). Next, is <code>Metal.xxx</code>. This is for higher-level, usually pure-Julian functionality (like <code>current_device()</code>). The only thing beyond this is exporting into the global namespace. That would be useful for uniquely-named functions/structures/macros with clear and common use-cases (<code>MtlArray</code> or <code>@metal</code>).</p><p>Additionally, you can override non-Metal.jl functions like <code>LinearAlgebra.mul!</code> seen <a href="https://github.com/JuliaGPU/Metal.jl/blob/main/lib/mps/linalg.jl#L34">here</a>. This is essentially (ab)using multiple dispatch to specialize for certain cases (usually for more performant execution).</p><p>If your function is only available from within GPU kernels (like thread indexing intrinsics). Be sure to properly annotate with <code>@device_function</code> to ensure that calling from the host doesn&#39;t kill your Julia process.</p><p>Generally, think about how frequently you expect your addition to be used, how complex its use-case is, and whether or not it clashes/reimplements/optimizes existing functionality from outside Metal.jl. Put it behind the corresponding interface.</p><h2 id="Creating-Tests"><a class="docs-heading-anchor" href="#Creating-Tests">Creating Tests</a><a id="Creating-Tests-1"></a><a class="docs-heading-anchor-permalink" href="#Creating-Tests" title="Permalink"></a></h2><p>As it&#39;s good practice, and JuliaGPU has great CI/CD workflows, your addition should have associated tests to ensure correctness and edge cases. Look to existing examples under the <code>test</code> folder for initial guidance, and be sure to create tests for all valid types. Any new Julia file in this folder will be ran as its own testset. If you feel your tests don&#39;t fit in any existing place, you&#39;ll probably want to create a new file with an appropriate name.</p><h2 id="Running-a-Subset-of-the-Existing-Tests"><a class="docs-heading-anchor" href="#Running-a-Subset-of-the-Existing-Tests">Running a Subset of the Existing Tests</a><a id="Running-a-Subset-of-the-Existing-Tests-1"></a><a class="docs-heading-anchor-permalink" href="#Running-a-Subset-of-the-Existing-Tests" title="Permalink"></a></h2><p>Sometimes you won&#39;t want to run the entire testsuite. You may just want to run the tests for your new functionality. To do that, you can either pass the name of the testset to the <code>test/runtests.jl</code> script: <code>julia --project=test test/runtests.jl metal</code> or you can isolate test files by running them alone after running the <code>test/setup.jl</code> script: <code>julia --project=test -L test/setup.jl test/metal.jl</code></p><h2 id="Thank-You-and-Good-Luck"><a class="docs-heading-anchor" href="#Thank-You-and-Good-Luck">Thank You and Good Luck</a><a id="Thank-You-and-Good-Luck-1"></a><a class="docs-heading-anchor-permalink" href="#Thank-You-and-Good-Luck" title="Permalink"></a></h2><p>Open-source projects like this only happen because people like you are willing to spend their free time helping out. Most anything you&#39;re able to do is helpful, but if you get stuck, seek guidance from Slack or Discourse. Don&#39;t feel like your contribution has to be perfect. If you put in effort and make progress, there will likely be some senior developer willing to polish your code before merging. Open-source software is a team effort...welcome to the team!</p></article><nav class="docs-footer"><a class="docs-footer-prevpage" href="../faq/">« Frequently Asked Questions</a><div class="flexbox-break"></div><p class="footer-message">Powered by <a href="https://github.com/JuliaDocs/Documenter.jl">Documenter.jl</a> and the <a href="https://julialang.org/">Julia Programming Language</a>.</p></nav></div><div class="modal" id="documenter-settings"><div class="modal-background"></div><div class="modal-card"><header class="modal-card-head"><p class="modal-card-title">Settings</p><button class="delete"></button></header><section class="modal-card-body"><p><label class="label">Theme</label><div class="select"><select id="documenter-themepicker"><option value="documenter-light">documenter-light</option><option value="documenter-dark">documenter-dark</option></select></div></p><hr/><p>This document was generated with <a href="https://github.com/JuliaDocs/Documenter.jl">Documenter.jl</a> version 0.27.23 on <span class="colophon-date" title="Monday 16 October 2023 06:12">Monday 16 October 2023</span>. Using Julia version 1.8.5.</p></section><footer class="modal-card-foot"></footer></div></div></div></body></html>
+}</code></pre><p>To compile with Metal&#39;s tools and emit human-readable IR, run something roughly along the lines of: <code>xcrun metal -S -emit-llvm dummy_kernel.metal</code></p><p>This will create a <code>.ll</code> file that you can then parse for whatever information you need. Be sure to double-check the metadata at the bottom for any significant changes your functionality introduces.</p><p>Test with different types and configurations to see what changes are caused. Also ensure that when writing very simple kernels, whatever you&#39;re interested in doesn&#39;t get optimized away. Double-check that the kernel&#39;s IR makes sense for what you wrote.</p><h2 id="Metal-Performance-Shaders"><a class="docs-heading-anchor" href="#Metal-Performance-Shaders">Metal Performance Shaders</a><a id="Metal-Performance-Shaders-1"></a><a class="docs-heading-anchor-permalink" href="#Metal-Performance-Shaders" title="Permalink"></a></h2><p>Metal exposes a special interface to its library of optimized kernels. Rather than accepting the normal set of input GPU data structures, it requires special <code>MPS</code> datatypes that assume row-major memory layout. As this is not the Julia default, adapt accordingly. Adding MPS functionality should be mostly straightforward, so this can be an easy entry point to helping. To get started, you can have a look at the <a href="https://developer.apple.com/documentation/metalperformanceshaders?language=objc">Metal Performance Shaders Documentation</a> from Apple.</p><h2 id="Exposing-your-Interface"><a class="docs-heading-anchor" href="#Exposing-your-Interface">Exposing your Interface</a><a id="Exposing-your-Interface-1"></a><a class="docs-heading-anchor-permalink" href="#Exposing-your-Interface" title="Permalink"></a></h2><p>There are varying degrees of user-facing interfaces from Metal.jl. At the lowest level is <code>Metal.MTL.xxx</code>. This is for low-level functionality close to or at bare Objective-C, or things that a normal user wouldn&#39;t directly be using. <code>Metal.MPS.xxx</code> is for Metal Performance Shader specifics (like <code>MPSMatrix</code>). Next, is <code>Metal.xxx</code>. This is for higher-level, usually pure-Julian functionality (like <code>current_device()</code>). The only thing beyond this is exporting into the global namespace. That would be useful for uniquely-named functions/structures/macros with clear and common use-cases (<code>MtlArray</code> or <code>@metal</code>).</p><p>Additionally, you can override non-Metal.jl functions like <code>LinearAlgebra.mul!</code> seen <a href="https://github.com/JuliaGPU/Metal.jl/blob/main/lib/mps/linalg.jl#L34">here</a>. This is essentially (ab)using multiple dispatch to specialize for certain cases (usually for more performant execution).</p><p>If your function is only available from within GPU kernels (like thread indexing intrinsics). Be sure to properly annotate with <code>@device_function</code> to ensure that calling from the host doesn&#39;t kill your Julia process.</p><p>Generally, think about how frequently you expect your addition to be used, how complex its use-case is, and whether or not it clashes/reimplements/optimizes existing functionality from outside Metal.jl. Put it behind the corresponding interface.</p><h2 id="Creating-Tests"><a class="docs-heading-anchor" href="#Creating-Tests">Creating Tests</a><a id="Creating-Tests-1"></a><a class="docs-heading-anchor-permalink" href="#Creating-Tests" title="Permalink"></a></h2><p>As it&#39;s good practice, and JuliaGPU has great CI/CD workflows, your addition should have associated tests to ensure correctness and edge cases. Look to existing examples under the <code>test</code> folder for initial guidance, and be sure to create tests for all valid types. Any new Julia file in this folder will be ran as its own testset. If you feel your tests don&#39;t fit in any existing place, you&#39;ll probably want to create a new file with an appropriate name.</p><h2 id="Running-a-Subset-of-the-Existing-Tests"><a class="docs-heading-anchor" href="#Running-a-Subset-of-the-Existing-Tests">Running a Subset of the Existing Tests</a><a id="Running-a-Subset-of-the-Existing-Tests-1"></a><a class="docs-heading-anchor-permalink" href="#Running-a-Subset-of-the-Existing-Tests" title="Permalink"></a></h2><p>Sometimes you won&#39;t want to run the entire testsuite. You may just want to run the tests for your new functionality. To do that, you can either pass the name of the testset to the <code>test/runtests.jl</code> script: <code>julia --project=test test/runtests.jl metal</code> or you can isolate test files by running them alone after running the <code>test/setup.jl</code> script: <code>julia --project=test -L test/setup.jl test/metal.jl</code></p><h2 id="Thank-You-and-Good-Luck"><a class="docs-heading-anchor" href="#Thank-You-and-Good-Luck">Thank You and Good Luck</a><a id="Thank-You-and-Good-Luck-1"></a><a class="docs-heading-anchor-permalink" href="#Thank-You-and-Good-Luck" title="Permalink"></a></h2><p>Open-source projects like this only happen because people like you are willing to spend their free time helping out. Most anything you&#39;re able to do is helpful, but if you get stuck, seek guidance from Slack or Discourse. Don&#39;t feel like your contribution has to be perfect. If you put in effort and make progress, there will likely be some senior developer willing to polish your code before merging. Open-source software is a team effort...welcome to the team!</p></article><nav class="docs-footer"><a class="docs-footer-prevpage" href="../faq/">« Frequently Asked Questions</a><div class="flexbox-break"></div><p class="footer-message">Powered by <a href="https://github.com/JuliaDocs/Documenter.jl">Documenter.jl</a> and the <a href="https://julialang.org/">Julia Programming Language</a>.</p></nav></div><div class="modal" id="documenter-settings"><div class="modal-background"></div><div class="modal-card"><header class="modal-card-head"><p class="modal-card-title">Settings</p><button class="delete"></button></header><section class="modal-card-body"><p><label class="label">Theme</label><div class="select"><select id="documenter-themepicker"><option value="documenter-light">documenter-light</option><option value="documenter-dark">documenter-dark</option></select></div></p><hr/><p>This document was generated with <a href="https://github.com/JuliaDocs/Documenter.jl">Documenter.jl</a> version 0.27.23 on <span class="colophon-date" title="Monday 30 October 2023 04:03">Monday 30 October 2023</span>. Using Julia version 1.8.5.</p></section><footer class="modal-card-foot"></footer></div></div></div></body></html>