[llvm] bef5ed1 - [AMDGPU][Docs] Update Code Object V3 example to includes args section

Krzysztof Drewniak via llvm-commits llvm-commits at lists.llvm.org
Fri Jul 9 10:42:36 PDT 2021


Author: Krzysztof Drewniak
Date: 2021-07-09T17:42:29Z
New Revision: bef5ed1eeac35d5f59a16a07c7a465f35d30d136

URL: https://github.com/llvm/llvm-project/commit/bef5ed1eeac35d5f59a16a07c7a465f35d30d136
DIFF: https://github.com/llvm/llvm-project/commit/bef5ed1eeac35d5f59a16a07c7a465f35d30d136.diff

LOG: [AMDGPU][Docs] Update Code Object V3 example to includes args section

The documentation for the AMDGPU assembler's examples don't show the
.args section, which, if ommitted, will cause arguments to silently
not be passed into the kernel. This commit fixes this issue.

Reviewed By: #amdgpu, scott.linder

Differential Revision: https://reviews.llvm.org/D105222

Added: 
    

Modified: 
    llvm/docs/AMDGPUUsage.rst

Removed: 
    


################################################################################
diff  --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index 47963b198a33..242e44b84cdb 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -12187,9 +12187,23 @@ Here is an example of a minimal assembly source file, defining one HSA kernel:
        .sgpr_count: 2
        .vgpr_count: 3
        .max_flat_workgroup_size: 256
-   ...
+       .args:
+         - .size: 8
+           .offset: 0
+           .value_kind: global_buffer
+           .address_space: global
+           .actual_access: write_only
+   //...
    .end_amdgpu_metadata
 
+This kernel is equivalent to the following HIP program:
+
+.. code::
+   :number-lines:
+   __global__ void hello_world(float *p) {
+       *p = 3.14159f;
+   }
+
 If an assembly source file contains multiple kernels and/or functions, the
 :ref:`amdgpu-amdhsa-assembler-symbol-next_free_vgpr` and
 :ref:`amdgpu-amdhsa-assembler-symbol-next_free_sgpr` symbols may be reset using


        


More information about the llvm-commits mailing list