[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