# Add lane support for SIMD/SIMT machines Section 2.5.1.3, pg 29ff Implicitly vectorized code executes multiple instances of a source code loop body or of a source code kernel function simultaneously in a single sequence of instructions operating on a vector of data elements (cf. SIMD: Single Instruction Multiple Data). This can be used to implement a SIMT (cf. Single Instruction Multiple Threads) execution model on SIMD hardware such as GPUs, e.g. for the OpenCL, SYCL, or HIP languages. It can also be used to implement SIMD loops on SIMD hardware such as GPUs or on traditional CPU hardware with vector extensions like Intel AVX, e.g. for OpenMP or for automatically vectorized code. In all these cases, the original source code has been written for a single thread in the case of SIMT or for a single vector element in the case of SIMD. The remainder of this text uses the term 'lane' to refer to a single SIMT thread or a single SIMD vector element. The number of lanes inside a vectorized code region is fixed. For SIMT, this typically covers an entire function. For SIMD, this typically covers an instance of a loop body. For example, the following OpenCL kernel function `vec_add` is logically invoked by a separate language thread for each index in a range of indices specified when the kernel is dispatched: __kernel void vec_add (__global char *dst, __global char *src) { int i = get_global_id(0); dst[i] += src[i]; } The following code fragment dispatches the `vec_add` kernel: size_t global_offset[1] = { 0 }; size_t global_size[1] = { NITEMS }; size_t group_size[1] = { NLANES }; clEnqueueNDRangeKernel(queue, vec_add_kernel, 1, global_offset, global_size, group_size, 0, NULL, NULL); Each invocation calls `get_global_id()` to determine the index it is assigned to work on, which returns a different value in each OpenCL thread. When compiling for GPUs, the compiler may map one OpenCL thread to one lane on the SIMD hardware. To do that, the compiler implicitly vectorizes the code to match the natural number of lanes on the target hardware. The same loop or the same function may be compiled with different vectorization factors resulting in different instances with a different number of lanes. In the SIMD case, we may end up with different instances of the same source loop using different vectorization factors resulting in a different number of lanes within the same function. For example, when the following C function is compiled with `gcc -O2 -fopenmp` on IA, the "for" statement is transformed into both a 16-wide loop instance and a 1-wide loop instance: void vec_add (char dst[], char src[], int len) { #pragma omp simd for (int i = 0; i < len; ++i) dst[i] += src[i]; } The vectorized loop instance packs 16 source elements into IA vector registers and processes all 16 source elements in one machine instruction. The code would load 16 adjacent source elements, each, add them, and store the resulting 16 source elements back. The trip count, which is the number of times the loop iterates during execution, is not known at compile-time. Therefore the compiler also generates a 1-wide instance of the loop that processes one source element at a time. Control flows from the vectorized loop instance to the scalar loop instance when the remaining trip count falls below a threshold determined by the compiler. When debugging the above kernel function or the above loop, the user would like to be able to inspect the variable `i` and the array elements `dst[i]` and `src[i]`. Since the code had been implicitly vectorized, multiple instances of the source code are executed in parallel. Debuggers may need to focus on a single lane at a time. To map such vectorized machine code back to scalar source code, debug information must describe the location of a given variable with respect to a given lane. DWARF-5 can describe the location of the vectorized object but it cannot describe the location of an instance of a variable inside that vectorized object. Further, type information would describe an individual instance, not the vectorized variable. DWARF-5 cannot describe the relationship between the two. To describe the location of a source variable inside an implicitly vectorized object, we propose a new operator: DW_OP_push_lane that allows the location of a variable to be described as function of the lane. The same scalar source code may be compiled using different vectorization factors resulting in a different number of lanes at different locations in the machine code. Debuggers may need to show the number of lanes for the current machine code location and only allow the user to focus on a lane within those boundaries. To describe this, we propose a new subprogram attribute: DW_AT_num_lanes that describes the number of lanes for implicitly vectorized code. Note that this only refers to implicitly vectorized code. For explicitly vectorized code, the source variables are themselves vectors and location descriptions refer to the vector object. Compilers may indicate this by describing the code region as `num_lanes == 1`. They would not use `DW_OP_push_lane` in DWARF expressions. --- Section 2.2, p.17. Add DW_AT_num_lanes | Number of implicitly vectorized lanes to table 2.2. Section 2.5.1.3, p.29. Add > 16\. `DW_OP_push_lane` > The `DW_OP_push_lane` operation pushes a lane index value of generic > type, which provides the context of the lane in which the > expression is being evaluated. See section 3.3.5 on page 79. > > *Producers that widen source code into vectorized machine code may > use this operation to describe the location of a source variable > as function of a single lane in the widened machine code.* > > *Consumers will supply the lane argument to obtain the location of > the instance of that source variable that corresponds to the > provided lane argument.* to the end of the list on p.33. Section 3.3.5, p.79. Add > *SIMD instructions process multiple data elements in one instruction. > The number of data elements that is processed with one instruction is > typically referred to as SIMD width. Each individual data element is > typically referred to as SIMD lane.* > > *When generating code for a SIMD architecture, compilers may need to > implicitly widen the source code to match the SIMD width of the > instruction set they are using. Source variables are widened into a > vector of variables, with one instance per SIMD lane.* > > A subroutine that is implicitly vectorized may have a `DW_AT_num_lanes` > attribute whose value describes the implicit vectorization factor and > the corresponding number of lanes in the generated code. The value of > this attribute is determined as described in Section 2.19 on page 55. > > To refer to individual lanes in such vectorized code, for example to > describe the location of widened source variables, producers may use > the `DW_OP_push_lane` operation to have consumers supply the current > focus lane for which to evaluate the expression. The pushed > lane index must be an unsigned integer value between zero (inclusive) > and the value of `DW_AT_num_lanes` (exclusive) at the current location. > See section 2.5.1.3 on page 29. > > If the attribute is omitted, its value is defined by the ABI. > > *If the source code had already been vectorized and is not further > widened by the compiler, the value should be one.* > > *This value does not only apply to vector instructions. If a loop or > function has been widened, the entire loop or function body shall be > annotated with the vectorization factor.* Section 7.5.4, p.207. Add DW_AT_num_lanes | TBD | constant, exprval, vallist to Table 7.5. Section 7.7, p.223. Add DW_OP_push_lane | TBD | 0 | to table 7.9. Section D.17 (new) Add > D.17 SIMD location descriptions > > The following example uses a hypothetical machine with 64b scalar > registers `r0`, `r1`, ..., and 256b vector registers `v0`, `v1`, ... that > supports SIMD instructions with different SIMD widths. Scalar > arguments are passed in scalar registers starting with `r0` for the > first argument. > > Consider the source code in figure D.73, which is implicitly widened > by a vectorization factor of 8 to match the 256b vector registers of > the target machine, resulting in the pseudo-code in figure D.74. > > Figure D.73: C OpenMP source code > > void vec_add (int dst[], int src[], int len) { > #pragma omp simd > for (int i = 0; i < len; ++i) > dst[i] += src[i]; > } > > Figure D.74: Pseudo Assembly code > > .l0: > move.64b r3, 0 ; i = 0 > .l1: ; implicitly 8-wide vectorized loop body > add.64b r4, r3, 8 ; inext = i + 8 > cmp.64b r4, r2 ; compare inext to len > jmp.ge .l2 ; jump to .l2 if inext >= len > load.256b v0, [r0+4*r3] ; v0[n] = dst[i+n] for n in [0..7] > .l1.1: > load.256b v1, [r1+4*r3] ; v1[n] = src[i+n] for n in [0..7] > .l1.2: ; add 8 elements > add.simd-8 v0, v0, v1 ; v0[n] = v0[n] + v1[n] for n in [0..7] > store.256b [r0+4*r3], v0 ; dst[i+n] = v0[n] for i in [0..7] > .l1.3: > mov.64b r3, r4 ; i = inext > jmp .l1 ; loop back for more > .l2: ; scalar loop body > add.64b r4, r3, 1 ; inext = i + 1 > cmp.64b r4, r2 ; compare inext to len > jmp.ge .l3 ; jump to .l3 if inext >= len > load.32b r5, [r0+4*r3] ; r5 = dst[i] > .l2.1: > load.32b r6, [r1+4*r3] ; r6 = src[i] > .l2.2: ; add a single element > add.32b r5, r5, r6 ; r5 = r5 + r6 > store.32b [r0+4*r3], r5 ; dst[i] = r5 > .l2.3: > mov.64b r3, r4 ; i = inext > jmp .l2 ; loop back for more > .l3: > return > .l4: > > The machine code contains two instances of the source loop: one > instance with SIMD width 8 at `.l1`, and one scalar instance at `.l2` to > handle any remaining elements. > > This function may be described in DWARF as shown in figure D.75. > > Figure D.75: Possible DWARF description of the function in D.73 compiled into pseudo machine code in D.74. > > DW_TAG_subprogram > DW_AT_name "vec_add" > DW_AT_num_lanes .vallist.0 > ... > DW_TAG_variable > DW_AT_name "i" > DW_AT_type int > DW_AT_location .loclist.1 > ... > > .vallist.0: > range [.l1, .l2) > DW_OP_lit8 > end-of-list > > .loclist.1: > range [.l0, .l1) > DW_OP_regx r3 > range [.l1, .l2) > DW_OP_bregx r3, 0 > DW_OP_push_lane > DW_OP_plus > DW_OP_stack_value > range [.l2, .l4) > DW_OP_regx r3 > end-of-list to the end of Appendix D on page 365. --- 2022-05-19: Revised. 2023-02-01: Revised. 2023-04-18: Revised based on feedback from 4/17/23 meeting. 2023-05-25: Revised based on feedback from 5/24/23 GPU meeting. 2023-06-12: Revised based on email feedback. 2023-09-11: Revised examples. 2024-04-29: Needs revision to use vallist class for `DW_AT_num_lanes` and remove `DW_OP_stack_value` from `loclist.0` in example. 2024-04-30: Revised based on feedback from 4/29/24 meeting. 2024-05-08: Revised example to fix typo and add comments. 2024-05-13: Accepted with new title. Changed from "SIMD Location Descriptions" to "Add lane support for SIMD/SIMT machines."