Vecz’s design is based on the automatic whole function vectorization research by Ralf Karrenberg, titled “Automatic SIMD Vectorization of SSA-based Control Flow Graphs”, and a combination of other papers referenced at various places in this document. While the process followed by Vecz is not exactly the same, understanding the research would help to understand Vecz better.

Vecz relies on certain other classes and functions provided by compiler::utils. The BuiltinInfo interface is used to ascertain certain properties of OpenCL builtin functions such as whether a particular function has a vector equivalent, or whether it is a work item ID query. Specific implementations of the interface are provided for OpenCL.

Before running vecz, it is recommended to run the compiler::utils::OptimalBuiltinReplacementPass. This replaces certain builtin calls with LLVM instructions or intrinsics that perform the equivalent operation, which enables later optimizations to work with them (which applies both to LLVM’s own optimization passes that vecz runs, and to some of vecz’s own transform passes). Furthermore, it allows these builtins to be widened arbitrarily, without being limited to the widths available as OpenCL builtins.

If a target intends to use the compiler::utils::WorkItemLoopsPass after Vecz, it is important to ensure that, before vecz, all calls to barrier-like functions in the full nested kernel call-graph are given unique barrier IDs. Note that this effectively mandates full inlining of all functions containing barrier-like calls.

This is necessary because vectorization can considerably affect control flow, so the ordering of the barriers in the function may change. If the WorkItemLoopsPass needs to combine two different versions of the same kernel into a single scheduled kernel, it is vital that direct correspondence of the barrier calls is maintained.

Users can run the compiler::utils::PrepareBarriersPass, which satisfies these requirements.

The vecz::RunVeczPass does not delete the original scalar kernel after vectorization, nor does it transfer the scalar kernel name to the vectorized function.

Vecz provides an interface, vecz::TargetInfo, that allows the vectorizer to make target-dependent decisions. This is retrieved via an analysis pass: vecz::TargetInfoAnalysis. Targets may override the vecz::TargetInfo constructed by this Analysis. The interface has a default implementation, which may be overridden.

Targets can override:

• Builder functions for all of the various forms of memory operation that vecz can output: loads and stores in masked or unmasked forms; contiguous access; interleaved access; scatter/gather access. Targets may want to provide their own intrinsics for these operations, if they exist.

• Builder functions for special forms of vector shuffles on scalable vectors, namely “inner” and “outer” broadcasts (meaning the duplication of each vector element n times, and the duplication of the entire vector n times), as well as vectorized scalable extract and insert instructions (which vectorize to a pick or replacement of every nth element to/from a narrower vector). Since there are no LLVM instructions to efficiently perform these operations on scalable vectors, the default implementation involves writing to memory and reading it back, which is likely to be suboptimal.

• Builder function to construct the vector length argument of predicated vector instructions, on targets that support this feature.

• Functions to support the Interleaved Groups Combining Pass.

• A function to support the SIMD Width Analysis that returns the widest vectorization factor possible for a given set of values.

• A function to compute the preferred vector width of a given scalar type. The packetizer can use this to emit multiple vectors per scalar-kernel value, instead of a single wider vector. The default implementation is based on the bit width of a single vector register (from llvm::TargetTransformInfo).

For full details, see the documentation of the vecz/target_info.h header file.

• Clone the kernel function

• Analyze the kernel function to determine if it has a pattern not prone to being vectorized

• Run preparation passes

• Perform control-flow to data-flow conversion (to handle divergent CFGs)

• Scalarize the kernel if desired/needed

• Run pre-packetization “middle optimization” passes

• Determine the optimal SIMD width for the kernel

• Packetize the kernel

• Perform optimizations and cleanup

• Define internal builtins

Instructions in a kernel function may be uniform (i.e. they evaluate to the same value or have the same side-effects on all lanes) or varying. Varying instructions usually have a data dependency on either the global ID or the local ID of the work-item executing the kernel. As an example, in the following kernel the store to the global memory, as well as the calculation of the address for the store, are varying instructions depending on the global ID of the work-item. The multiplication of in by 2 on the other hand is uniform, since it is the same across all the work-items.

kernel void fn(int in, global int *out) {
  size_t tid = get_global_id(0);
  in = in * 2;
  out[tid] = in;
}

Assuming to vectorize on the x dimension, after packetization, functions calls like get_global_id(0) or get_local_id(0) will return vectors of consecutive indices, which will allow us to packetize the store into a vector store.

On architectures that support both scalar and vector instructions we do not want to packetize uniform instructions, as each vector lane will perform the same operation on the same data and return the same result. Instead, we want to keep uniform instructions scalar (i.e. keep the original instructions untouched), and broadcast their value into a vector when necessary.

In order to differentiate between the varying and the uniform instructions, we have implemented an analysis pass called the Uniform Analysis. This analysis starts by finding “vector roots” in the function. By “roots” we mean instructions (usually function calls) that we know to be varying. For example, work-item ID related functions, or packetized arguments in non-kernel functions are some common vector roots. Each root and its users are recursively marked as varying. Marking a value happens before marking its users, so that use cycles (e.g. phi nodes) do not cause infinite recursion. The instructions remaining after this process are considered to be uniform.

In the previous example kernel, the vector root used is the call to the get_global_id(0) function. Starting from that point and then recursively going through its users and their users etc. We recursively mark the address calculation (getelementptr) for the output and the store to the output as varying too. The in value and its multiplication are not marked as varying since they are not using any varying values, they are being used by one. We should note here that under special cases, such as an alloca instruction that is stored into by a varying store, we might mark some instructions as varying just because they are used by a varying instruction but in the general case we do not.

The relevant classes and functions can be found in source/include/analysis/uniform_value_analysis.h and source/analysis/uniform_value_analysis.cpp.

Memory operations can access their data in several different patterns, as a function of the work item ID, categorized as:

• Uniform: data is accessed from the same address for all work items;

• Contiguous: data is accessed sequentially with no gaps;

• Strided: data is accessed linearly but spaced apart with a constant or uniform stride;

• Divergent: data is accessed with no discernible pattern.

The stride analysis traverses address computation expressions, to ascertain which kind of memory access is required, computing any constant strides encountered. Uniform but variable strides are not computed during the analysis, because doing so would require creating new instructions in the function, which is at odds with the idea of an analysis pass. However, it is usually sufficient to know that access is linear, without needing to know its actual value. When a transform pass wishes to make use of the actual value, it can call the manifest() function of the analysis, which will traverse its internal dataset and create any instructions required. Note that pointers to these instructions will survive until the analysis is invalidated.

This analysis uses the result of Uniform Analysis.

The relevant classes and functions can be found in source/include/analysis/stride_analysis.h and source/analysis/stride_analysis.cpp.

The packetizer needs to know which instructions require packetization in advance, for optimal functioning. An instruction that has been marked as varying by the Uniform Analysis may or may not require packetization, since some varying values will form an expression computing the address of a contiguous or strided memory operation. Therefore, this analysis starts at the function’s vector leaves, and works backwards through operands, recursively marking values for packetization. When a contiguous or strided memory operation is encountered, its address operand is not followed. This allows a more accurate estimation of packetization requirements prior to actual packetization, which is useful for the SIMD Width Analysis.

This analysis uses the result of Stride Analysis.

The relevant classes and functions can be found in source/include/analysis/packetization_analysis.h and source/analysis/packetization_analysis.cpp.

Another useful analysis we employ is the Control Flow Graph (CFG) Analysis. As the name suggests, it analyzes the control flow graph to store useful information about blocks and loops such as what loop does a block live in and what are the lcssa values of a loop (values that are live through a loop).

The analysis works by iterating over the basic blocks, create a tag for each block, and if the block belongs in a loop, then create a tag for that loop (if one does not yet exist) and mark the loop as owning the block.

If in the process of visiting the blocks we encounter a divergent branch, then we say that the CFG needs to be converted into a Data-Flow Graph (see the Control Flow Conversion section).

The relevant classes and functions can be found in source/include/analysis/control_flow_analysis.h and source/analysis/control_flow_analysis.cpp.

Formally known as Rewire Target Analysis.

This analysis is used to find all the information regarding divergence in a CFG. It uses as a pre-requisite the Uniform Analysis to know which instructions are varying when we first evaluate the CFG. Those instructions are the basis to find divergent branches, which are branches whose operands are varying. We name blocks having such instructions div_causing. The analysis works by iterating over all the branches in the CFG until no more new div_causing blocks are found. When we find a div_causing block, we first compute the divergent path that this block creates. This divergent path contains all the blocks from the div_causing block to the post dominator of the latter. All the blocks that belong in a divergent path may need to have their instructions marked varying, in the case where they might be used outside the divergent path and thus need to be packetized. We then find all the join points of the div_causing block. Such blocks have disjoint paths from the div_causing block. These blocks are called blend blocks and will need to have their PHI nodes transformed into select instructions because their path will be linearized. These blocks also need to have their PHI nodes marked as varying. After we have processed the div_causing block, we must find if this block makes a loop divergent. A loop is divergent if it is possible that some work items leave the loop at an exit, while others keep iterating. In essence, this means that a divergent branch has no post-dominator in the loop. We also mark all the exits of the loop where some work items may leave as divergent because they will need to be linearized. Finally, we compute by_all blocks which are blocks that need not be predicated because no divergence is present when they are executed.

The relevant functions can be found in source/include/analysis/divergence_analysis.h and source/analysis/divergence_analysis.cpp.

This analysis is used to determine the set of live values at any point in the program. A value becomes “live” at its definition and remains live until all of its uses have been encountered. The result of this analysis provides an info object for every basic block in the program, that contains the “Live In” set (i.e. the values that are live at the start of the Basic Block, including all of its PHI nodes), and the “Live Out” set (i.e. the values that are still live at the end of the block). By iterating backwards over a Basic Block, starting from the Live Outs, one can determine the set of values that are live at any point in the program.

The implementation is based on Section 5.2 of the paper “Computing Liveness Sets for SSA-Form Programs.” by Florian Brandner, Benoit Boissinot, Alain Darte, Benoît Dupont de Dinechin, Fabrice Rastello.

This analysis is used by BOSCC, and by SIMD Width Analysis.

The relevant classes and functions can be found in source/include/analysis/liveness_analysis.h and source/analysis/liveness_analysis.cpp

This analysis is used to estimate the optimal vectorization width, depending on the contents of the kernel. The current strategy is a two-stage process: first we find the widest used varying value type, disregarding types that make up less than a small proportion of the program according to some tolerance threshold. The SIMD width computed is the number of elements of this type that will fit into a single vector register. Then we analyze the program using the Liveness Analysis to estimate the maximum SIMD width that will fit into vector registers, if it is wider than the previously computed result. This allows vectorization to produce values that will not necessarily fit into single vector registers, but will fit across multiple registers after legalization.

SIMD Width Analysis is performed only when vectorization is set to automatic (either by using the -cl-wfv=auto option or by passing bool Auto=true to Vectorizer::vectorize()). The analysis is performed after control flow conversion, so that any changes made by this or prior passes will be taken into account.

This analysis uses the result of Packetization Analysis.

The relevant classes and functions can be found in source/include/analysis/simd_width_analysis.h and source/analysis/simd_width_analysis.cpp

This is not an analysis in the traditional sense but more of a filter for kernels that we know we will not be able to vectorize. There are a number of cases that we cannot handle in Vecz, such as kernels containing specific atomic instructions, functions returning vector results, or specific OpenCL builtins. By checking for these conditions early in the vectorization process we can save on compile time and also avoid accidentally generating an incorrectly vectorized kernel.

The relevant functions can be found in source/include/vectorization_context.h and source/vectorization_context.cpp.

This is a utility class created to speed up CFG reachability queries required by BOSCC. It is not an analysis pass managed by LLVM, but must be created manually where required. The algorithm is based on an Open Proceedings paper entitled “Reachability Queries in Very Large Graphs: A Fast Refined Online Search Approach” by Renê R. Veloso, Loïc Cerf, Wagner Meira Jr, Mohammed J. Zaki. In addition to that approach, dominator and post-dominator trees are used to further accelerate the process.

The reachability data structure must be constructed from a directed acyclic graph. Backedge traversal is a case not yet handled.

The relevant functions can be found in source/include/reachability.h and source/reachability.cpp.

We employ a set of preparation passes that includes both optimization passes, such as the Mem2Reg pass discussed later on, and passes necessary to generate IR that Vecz can handle:

• Dominator Tree Wrapper Pass (from LLVM)

• Loop Info Wrapper Pass (from LLVM)

• Switch Lowering Pass (from LLVM)

• Function Exit Nodes Unification Pass (from LLVM)

• Builtin Inlining Pass (from Vecz, described below)

• Promote Memory To Register Pass (from LLVM)

• Basic Mem2Reg Pass (from Vecz, described below)

• Instruction Combining Pass (from LLVM)

• Dead Code Elimination Pass (from LLVM)

• Pre-linearization Pass (from Vecz, described below)

• Instruction Combining Pass (again, to deal with Pre-linearization changes)

• CFG Simplification Pass (from LLVM)

• Unify Function Exit Nodes Pass (from LLVM)

• Loop Simplification Pass (from LLVM)

• Loop Rotate Pass (from LLVM)

• Simplify Infinite Loop Pass (from Vecz, described below)

• Induction Variable Simplification Pass (from LLVM)

• Early CSE Pass (from LLVM)

• LCSSA Pass (from LLVM - restores LCSSA form if broken by Early CSE)

The relevant classes and functions can be found in vectorizer.cpp, as well as builtin_inlining_pass.h, builtin_inlining_pass.cpp, basic_mem2reg_pass.h, and basic_mem2reg_pass.cpp for the Vecz passes.

This pass scans for PtrToInt casts that can be eliminated and converted into bitcasts or GEP instructions. It is able to eliminate a PtrToInt in the following cases:

• A PtrToInt followed by an IntToPtr, which is replaced by a pointer cast;

• A PtrToInt used by a PHI node, in which case the PHI node is replaced by one of the pointer type;

• A PtrToInt where the pointer type is i8*, followed by an integer add or subtract, in which case it is replaced by a GEP.

Removing intptr casts makes it possible for uniform pointer strides to be identified.

The pass can be found in source/transform/remove_intptr_pass.cpp

This pass looks for loads and stores of vector types that fit into a legal integer, where packetization would result in non-contiguous access, and replaces them with loads or stores of an integer scalar of the same size (where alignment requirements permit). This allows more efficient generation of scatter/gather or interleaved memory operations on these types.

The pass can be found in source/transform/squash_small_vectors_pass.cpp

This pass converts code that is already in a vector form into scalar code, or retains a partially scalarized code, so that the packetizer can produce vector IR at a vectorization width optimal for the target hardware.

The scalarization pass is divided into two stages: the analysis and the actual transformation stage.

In the analysis we mark the values that need scalarization, which includes vector leaves and non-vector instructions using vector operands. The non-vector instructions using vector operands are either ExtractElements with vector operands or BitCastInsts from vector to non-vector. Note that the analysis is not performed in an analysis pass, but a utility class that runs locally within the transform pass, since this information is not needed by any other pass.

If the vector leaf instruction or any of its arguments are of vector type and the primitive size is greater than the partial scalarization factor (called primitive size), the instruction is marked for needing scalarization. This marks the end of analysis.

Once we have the values that were marked for transformation by the analysis stage, the vector operands are first scalarized and then followed by the vector leaf instructions that need scalarization.

The utility classes can be found in source/include/transform/scalarizer.h and source/transform/scalarizer.cpp. The transform pass can be found in source/include/transform/scalarization_pass.h and source/transform/scalarization_pass.cpp.

The control flow conversion linearizes the divergent control flow executing both if and else condition blocks. In order to preserve safe access, the blocks are predicated with masks, to allow only legal access for calls and memory accesses that have side-effects.

Control flow conversion is the actual control-flow to data-flow conversion pass that uses information from the control flow analysis and divergence analysis. Conversion starts with generating masks, applying masks and generating selects, followed by linearizing the control flow and finally repair the SSA form that the linearization may have broken.

The mask for every basic block is generated starting with the entry mask, which is followed by a branch mask for cases where the entry mask is a phi node of its predecessors. Then special masks in the loop are generated to handle run time control flow divergence, namely; the loop active mask, combined exit mask. Next, the loop live values need to reflect the right values for early exited lanes. The masks are then applied to prevent side-effects for the inactive instances. In case of a call to memory operation, it is replaced with a corresponding masked internal builtin call Defining Internal Builtins. The phi nodes are then transformed into selects, to enable control-flow to data-flow conversion.

The CFG is then linearized, where necessary. It is actually partially linearized to retain uniform branches that we know need not be linearized. We apply the partial linearization by identifying every divergent blocks thanks to the divergence analysis to know which blocks should be linearized, and which blocks may remain untouched. A divergent block is called a div causing block. To linearize the CFG, we keep a deferral list that represents all the blocks that lost their predecessor’s edge because of divergence. When we reach a block, if the block is a div causing block, then it can only have one successor, otherwise the block can keep the same amount of successors it has. To know which block should be the successor of another block, we choose between the current successors, and the deferral list available for that block. The choice is then made based on the Dominance-Compact Block Indexing, which assigns each block a unique index. Dominance compactness means that for any block, all other blocks dominated by that block follow on in a contiguous sequence. This is constructed by a depth-first traversal of the dominator tree, visiting children in CFG reverse post-order. (In actual fact, Loop-Compactness takes precedence over Dominance-Compactness; the latter usually implies the former, but certain loops with multiple exits can break this, so special care has to be taken.) Using that index to choose the successor guarantees that if an edge A to B existed in the original graph, an edge A to X to B will exist in the linearized graph, thus conserving dominance.

Once the CFG is linearized, we may have introduced new edges that were not there previously, which may have broken the SSA form. Therefore, we must repair the SSA form by introducing blend instructions (in the form of phi nodes) at the new converging points.

The partial linearization implementation was inspired from the paper Automatic SIMD Vectorization of SSA-based Control Flow Graphs by Ralf Karrenberg and Partial Control-Flow Linearization by Simon Moll & Sebastian Hack.

The pass can be found in source/include/transform/control_flow_conversion_pass.h and source/transform/control_flow_conversion_pass.cpp.

if (some_condition) {
  return;
}

During packetization, instructions that define varying values or produce varying side-effects are turned into instructions that define the same value or produce the same effects for each SIMD instance. For example, an add instruction that takes two i32 operands is turned into another add instruction that takes two <N x i32> operands (where N is the SIMD width). This is done recursively in three steps; first, we packetize any branch condition that requires packetization, and then, starting at the “vector leaves” of the function, the rest of the instructions. Vector leaves are instructions that allow varying values to “escape” from the function. Some examples of leaves include:

• Store instructions, when the value to store and/or the pointer is varying

• Call instructions, when varying operands are present or when the call has no use

• Return instructions

After those two steps, then we proceed to packetize any remaining phi nodes, explained in more details in a following subsection.

During the packetization process, we might run into cases where we cannot packetize a varying instruction but instead we have to instantiate it. By instantiation we mean repeating an instruction N times, one for each SIMD lane. A common example would be calls to the printf function, as in the following kernel:

kernel void fn(global int *in, global int *out) {
  size_t tid = get_global_id(0);
  int load_in = in[tid];
  int result = load_in * tid;
  printf("in[%d] = %d\n", tid, result);
  out[tid] = result;
}

In this kernel, the call to the printf function will be repeated N times, with its arguments adjusted to use the correct global ID for each lane. On the other hand, the load and store instructions, as well as the multiplication, will be packetized into vector instructions of width N. More details on instantiation can be found in the relevant section.

As we have already mentioned, the packetization (or instantiation) starts from the vector leaves and recursively continues into their operands. As a matter of fact, the operands are packetized before the instruction itself; in order to generate the correct instruction, we first need to have the correct operands. This process stops when we reach either:

1. An operand that is uniform, such as constants or kernel arguments.

2. A vector root such as get_global_id(0).

3. A pointer that we can handle in its scalar form.

In the first case, we create a packetized version of the operand by simply broadcasting its value into a vector, so that each element of the vector contains the same value. Since the value is uniform, we do not need to proceed and packetize its operands. The second case is handled specially by using the scalar value to create a vector of sequentially increasing values.

The third case is also special, because it depends on the access pattern of the pointer. If the pointer is varying and we can determine a stride to the access pattern then we do not need to packetize the pointer. Instead, we can use the same pointer value as the base for a vector memory instruction. More details on this can be found in the Packetizing Memory Operations subsection.

Given all of these, we can now see how the example kernel given above will be vectorized, with a vector width of 4. Of course, Vecz is not a source-to-source transformation pass but the following kernel captures the equivalent IR changes that will be performed in an easier to read format:

kernel void fn(global int *in, global int *out) {
  size_t tid = get_global_id(0);
  size_t4 tid4 = {tid, tid, tid, tid} + {0, 1, 2, 3};
  int4 load_in4 = in[tid];
  int4 result4 = load_in4 * tid4;
  printf("in[%d] = %d\n", tid4.s0, result4.s0);
  printf("in[%d] = %d\n", tid4.s1, result4.s1);
  printf("in[%d] = %d\n", tid4.s2, result4.s2);
  printf("in[%d] = %d\n", tid4.s3, result4.s3);
  out[tid] = result4;
}

Notice how the address for the vector load and store are still calculated using the scalar ID variable (tid), since the kernel accesses the elements of the array consecutively and thus we can use a vector load and a vector store with the same base address.

Regardless of how a varying instruction has been handled, after we have packetized or instantiated it, we mark it for deletion. After we have gone through all the vector leaves, we proceed to delete all the instructions that we marked for deletion, as long as they have no remaining users.

The packetization pass can be found in source/include/transform/packetization_pass.h and source/transform/packetization_pass.cpp.

This is the general approach taken to packetize instructions but some cases need to be handled specially. We will now explain in more depth some special packetization cases.

... = in[get_global_id(0)]

%call = call spir_func i64 @_Z13get_global_idj(i32 0)
%arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %in, i64 %call
%0 = load i32, i32 addrspace(1)* %arrayidx, align 4

%call = call spir_func i64 @_Z13get_global_idj(i32 0) #2
%arrayidx0 = getelementptr inbounds i32, i32 addrspace(1)* %in1, i64 %call, i64 0
%arrayidx1 = getelementptr inbounds i32, i32 addrspace(1)* %in1, i64 %call, i64 1
%arrayidx2 = getelementptr inbounds i32, i32 addrspace(1)* %in1, i64 %call, i64 2
%arrayidx3 = getelementptr inbounds i32, i32 addrspace(1)* %in1, i64 %call, i64 3
%0 = load i32, i32 addrspace(1)* %arrayidx0, align 4
%1 = load i32, i32 addrspace(1)* %arrayidx1, align 4
%2 = load i32, i32 addrspace(1)* %arrayidx2, align 4
%3 = load i32, i32 addrspace(1)* %arrayidx3, align 4

define spir_func i32 @_Z7isequalff(float %x, float %y) {
entry:
  %cmp.i = fcmp oeq float %x, %y
  %conv.i = zext i1 %cmp.i to i32
  ret i32 %conv.i
}

declare spir_func <4 x i32> @__vecz_v4__Z7isequalff(<4 x float> %x, <4 x float> %y)

define spir_func <4 x i32> @__vecz_v4__Z7isequalff(<4 x float> %x, <4 x float> %y) {
entry:
  %placeholder_x = extractelement <4 x float> %x, i32 0
  %placeholder_y = extractelement <4 x float> %y, i32 0
  %cmp.i = fcmp oeq float %placeholder_x, %placeholder_y
  %conv.i = zext i1 %cmp.i to i32
  ret i32 %conv.i
}

define spir_func <4 x i32> @__vecz_v4__Z7isequalff(<4 x float> %x, <4 x float> %y) {
entry:
  %cmp.i1 = fcmp oeq <4 x float> %x, %y
  %conv.i2 = zext <4 x i1> %cmp.i1 to <4 x i32>
  ret <4 x i32> %conv.i2
}

After the vectorization process is completed, we run some additional passes to further optimize and cleanup the code:

• Inline Post Vectorization Pass (from Vecz)

• CFG Simplification Pass (from LLVM)

• Global Value Numbering Pass (from LLVM)

• Dead Code Elimination Pass (from LLVM)

• Interleaved Group Combining pass (from Vecz)

• Instruction Combining pass (from LLVM)

• Masked Memory Operations Simplification pass (from Vecz)

• Internal Builtin Definition pass (from Vecz)

The passes can be found in the files source/include/transform/passes.h and source/transform/passes.cpp.

-------------------------------------------------
|A1|B1|C1|D1|A2|B2|C2|D2|A3|B3|C3|D3|A4|B4|C4|D4|
-------------------------------------------------

-------------------------------------------------
|a1|a2|a3|a4|b1|b2|b3|b4|c1|c2|c3|c4|d1|d2|d3|d4|
-------------------------------------------------

This section covers various necessary utilities that the main vectorizer passes are using.

kernel void fn(global int *in) {
  size_t tid = get_global_id(0);
  global int* ptr = (in + tid) * 2;
  ...
}

LLVM has support for collecting counters (called statistics in LLVM) from the passes. Vecz is among the passes that can produce such information. This can be done in two ways.

First, the official way is to use opt with the -stats option. This will print the statistics from all the passes that have any.

The second way is to pass the -cl-llvm-stats option to the oneAPI Construction Kit. This will do pretty much the same work that the -stats option does, but it can be used in cases where it is not possible to use -stats.

Vecz utilizes the Remarks system available in LLVM, mostly to warn about vectorization failures. The remarks can be enabled by passing the -pass-remarks=vecz and -pass-remarks-missed=vecz command line option to opt, or the -v flag to oclc.

The command line tool veczc is a standalone compiler that is used to vectorize LLVM bitcode binary files. Its main use is in our vecz LIT-based testing (see modules/compiler/vecz/test).

It has the following arguments:

• -o file output bitcode file

• -w width the width to vectorize the code to

• -d dimension the dimension index to vectorize the code on

• -k name the function names to select for vectorization. It can appear multiple times, in one of several forms. In the standard form, simply passing the names of kernels will ensure those kernels are vectorized by the globally selected parameters e.g. veczc -k foo -k bar ... selects both the foo and bar kernels for vectorization. The more complex form allows specifying the vectorization parameters per-kernel in multiplicate: e.g. veczc -k foo:4,8,16 ... will generate multiple vectorized versions of foo at vectorization factors of 4, 8, and 16. All other paremeters will be inherited from the global configuration. The complete syntax for the kernel specification switch (k) value is as follows:

  <kernel_spec> ::= <kernel_name> ':' <spec>
  <kernel_spec> ::= <kernel_name>
  <spec> ::= <vf><dimension>(opt)<width>(opt)<scalable_spec>(opt)<auto>(opt)
  <spec> ::= <spec> ',' <spec> // multiple specs are comma-separated
  <number> ::= [0-9]+ // a decimal integer
  <kernel_name> ::= [a-zA-Z_][a-zA-Z_0-9]+ // As in the simple form - the name of the kernel to vectorize
  <dim> ::= '.' [123] // Vectorize only the given dimension
  <vf> ::= <number> // vectorize by the given factor
  <vf> ::= 'a' // automatic vectorization factor
  <simd_width> ::= '@' <number> // Assume local size (SIMD width) is the given number
  <scalable_spec> ::= 's' // Turn on scalable vector support
  
  

  n.b. (There should be no whitespace as this interface is designed for easy nonquoted use in common shells)

It supports bitcode files with the following target triples:

• spir-unknown-unknown 32-bit SPIR binaries

• spir64-unknown-unknown 64-bit SPIR binaries

Because veczc doesn’t load all of the builtins prior to vectorization, declarations of scalar or vector versions of any builtins used in the input file must be present, otherwise scalarization or packetization will not be able to materialize the scalarized/vectorized builtin calls and veczc will fail with an error message.