v1.2.4: Ternary Normal Form

27 December 2024 During the hackathon with Nvidia, we profiled the code and discovered that the representation of propagators was problematic. Let's take an example with the constraint x * x + y < 5 which is internally represented by the following AST:

Because the constraints can have any arity and be arbitrarily deep, we have no choice but to represent this tree with pointers and dynamic memory allocation. In Turbo, we used variant and shared_ptr to represent propagators. The propagators are executed using a switch statement of the following form:

switch(term.index()) {
  case IVar:
  case IConstant:
  case INeg:
  case IAdd:
  case ISub:
  case IMul:
  // ...

In the profiling report, it is shown that switch(term.index()) generates a lot of uncoalesced accesses. Previously, we have reduced divergent branches by sorting propagators with the same shape (see sorting propagators). But the main issue was the uncoalesced accesses to the propagators.

To address this, we propose decomposing the problem into constraints that can be represented without pointers (fixed depth) and that are small (to enable coalesced memory load). For this purpose, we introduce a new problem decomposition call the ternary normal form (TNF) where all propagators are of the form x = y <op> z with x,y,z variables. Furthermore, we take a minimal set of operators that includes the arithmetic operators +,-,*,/,min,max and two comparison operators <=, =. All the other constraints are rewritten into those operators. The constraint x * x + y < z in TNF is:

t1 = x * x
t2 = t1 * y
ONE = (t2 <= z)

Where t1,t2,ONE are new variables where ONE is initialized to 1. Note that the symbol = is relational, hence using ONE = (t2 <= z) forces the constraint on the right to be true. This is called a "reified constraint" in the jargon of constraint programming, but in TNF it has the same representation as any other constraint. TNF allows us to represent compactly a propagator using:

struct bytecode_type {
  Sig op;
  AVar x;
  AVar y;
  AVar z;
};

AVar is an integer representing the index of the variable in the store and Sig is also an integer representing the operator.

Problems in Ternary Normal Form

One obvious drawback of this approach is the increased number of temporary variables and constraints. While the increase in the number of constraints is relatively low and stable across problems (between 1x and 4x), the increase in the number of variables ranges between 1x and 139x in the case of the Wordpress problem. This variability warrants further investigation in the future to explore ways of reducing the number of generated variables.

Benchmarks

The following table compares hybrid v1.2.3 and hybrid v1.2.4, both with 128 blocks.

Let's analyze the speedup of this new version v1.2.4. The number of fixpoint iterations per second has increased by 254% which is a very large gain. However, the increase in nodes per second is only 29%, and it increases in only 11 out of 16 problems. There are two reasons behind this discrepancy:

1. The convergence of the fixpoint loop is two times slower than in v1.2.3, likely due to improper sorting of the TNF propagators. In all 16 problems, the new version requires more iterations to converge compared to v1.2.3.
2. The increase in variables results in fewer problem being stored in shared memory: only 3 out of 16 in comparison to 11 out of 16 before.

The first point is somewhat mitigated by the number of fixpoint iterations per second which almost double, demonstrating the effectiveness of TNF for GPU architecture. Perhaps more surprisingly, the memory footprint of the propagators is much lower than in v1.2.3, with an average reduction of 80%. Although we have more propagators, each one requires only 16 bytes, as opposed to v1.2.3, where propagator representation involved additional overhead from the shared pointer and variant data structures.

As for the full GPU version, we notice a similar behavior (we compare against full GPU v1.2.0 because v1.2.{1-3} were only modifying the hybrid version):

Profiling

The hackathon reminded me that we should always prioritize the next improvement based on what the profiler reveals. While this is a well-known principle, it's easy to forget.

During the hackathon, we mainly used the tool ncu to profile our kernels. It creates a report that can be visualized with ncu-ui and provides various metrics, such as the overall utilization of the SMs, the number of uncoalesced accesses and their locations in the code, the L1/L2 cache hit rates, warp state statistics, and more. It was in this report we identified the issues with the previous AST-based representation of the propagators. Among the remaining issues, the report tells us that warps are spending most of their time waiting on __syncthreads() barriers or global memory loads.

Beyond those low-level GPU optimizations, we can also examine the overall behavior of the constraint solver to identify which parts of the algorithm need further optimization. To achieve this, we equipped Turbo with time counters to track where most of the time is spent in the code. Let's first focus on the full GPU architecture (-arch gpu).

Clearly, the preprocessing time and the time spent eliminating the entailed propagators (from v1.2.3) are negligible. The orange section represents the propagation time, and the green section represents the search time. Prior to optimizing the propagation, the search accounted for about 10-20% of the total time. However, it has now become a clear candidate for further optimization. In the full GPU architecture, the search is handled by a single GPU thread which is why it is quite slow.

Next, we look at the hybrid architecture (-or 128 -arch hybrid).

Firstly, it is clear that searching on the CPU and memory transfers are not (yet) the bottleneck of the hybrid algorithm. The time spent in the fixpoint loop is somewhat similar to the full-GPU architecture. The main difference is the significant amount of time spent waiting for the CPU, as represented by the red section. One possible explanation is that we have too many CPU threads, while the CPU has few cores (only 10). To confirm this hypothesis, we should test this version on the GH200, which is powered by a 72-cores ARM CPU.

Based on these charts, there are at least two opportunities to improve efficiency:

1. Continue optimizing the fixpoint loop (this would benefit both architectures, and potentially provide even greater improvements for the hybrid architecture on a CPU with many cores).
2. Parallelize the search (this would benefit only the full-GPU architecture).

From a research perspective it is more interesting to focus on the fixpoint loop parallelization. Furthermore, improving the full-GPU version---which does not compile on modern GPUs---is less exciting.