From d974601d75e8155c3575aa0866985ad5c69bbe36 Mon Sep 17 00:00:00 2001 From: scottstraughan <42965777+scottstraughan@users.noreply.github.com> Date: Fri, 12 Sep 2025 09:50:49 +0100 Subject: [PATCH 1/2] Added Ross blog. --- ...2-adventures-in-address-space-inference.md | 541 ++++++++++++++++++ .../MemLayout.png | Bin 0 -> 25803 bytes 2 files changed, 541 insertions(+) create mode 100644 _collections/_portal_posts/2025-09-12-adventures-in-address-space-inference.md create mode 100644 assets/images/portal/article-images/2025-09-12-adventures-in-address-space-inference/MemLayout.png diff --git a/_collections/_portal_posts/2025-09-12-adventures-in-address-space-inference.md b/_collections/_portal_posts/2025-09-12-adventures-in-address-space-inference.md new file mode 100644 index 0000000..de7cfe4 --- /dev/null +++ b/_collections/_portal_posts/2025-09-12-adventures-in-address-space-inference.md @@ -0,0 +1,541 @@ +--- +category: blogs +date: '2025-09-12T02:00:00.0' +hidden: false +layout: portal/portal-article-view +thumbnail: /assets/images/portal/article-images/2025-09-12-adventures-in-address-space-inference/MemLayout.jpg +title: 'Adventures in Address Space Inference' +user_id: 61238111182 +--- + +## Introduction + +SYCL is a language that allows standards compliant C++ code to be compiled for accelerator and GPU devices. As part of +this compilation, SYCL implementations use a type inference system to derive address space qualifications for pointers. +In this post I shall detail the journey Codeplay took to replace their SYCL implementation's inference engine and give +an overview of how the new system works. + +Of course, I do not wish to claim that this is all my own work; I was merely the person who happened to work on it the +longest. Other Codeplayers who worked on this are as followed: + +* Duncan Brawley +* Harald van Dijk +* Peter Žužek +* Victor Lomüller + +As well as many other people who I am not able to mention here. + +This post was entirely written by a human and contains authentic human mistakes. + +## Some Background +Codeplay's ComputeCpp and Intel's DPC++ compilers both implement the SYCL standard. This standard defines a "single +source" model where the same C++ source can be run on both the x86_64/ARM host device as well as any attached GPU or +accelerator chips. It also defines a number of helper functions and wrappers to make data and code transfers to the +device as easy as possible. + +### A Mental Model of the Memory Model +GPUs and other accelerator chips have complex memory configurations. Some memory on the device is only visible to a +single work item, whilst other memory is visible to everything but with higher latency. Since early versions of SYCL +inherited many things from OpenCL, it makes sense to look at the memory model used there. + +OpenCL code is run on a number of "work groups", which are then divided further down into "work items". Work items can +be thought of as similar to threads in host code that run in lockstep. With this architecture, we can see that there +are three different scopes that memory can exist in: +* **Private**: Available to only a single work item. +* **Local**: Available and shared between all work items in a single work group. +* **Global**: Shared between all work items regardless of work group. + +![A diagram of the OpenCL memory layout. There are two blocks labelled named "Work Group", each containing four circles (representing work items) linked to various memory regions. Each work item is linked to their own "Private" memory, a work-group-specific "Local Memory" and a single shared "Global Memory"]({{ '/assets/images/portal/article-images/2025-09-12-adventures-in-address-space-inference/MemLayout.jpg' | relative_url }}) + +Reads and writes to memory need to know what type of memory to access, so languages that compile to OpenCL extend +the type system with so called **address spaces** to contain this information. Instead of an `int *`, OpenCL C has +`__global int *`, `__local int *` and so on. This follows similar language rules to CV qualifiers like `const`. + +However, having to specify types with a language extensions turned out to be rather clunky and made it hard to share +code between the host and device. To solve this, in OpenCL 2.0 a new address space was added: `generic`. Pointers of any +address space can be implicitly cast to and from their equivalent `generic` pointer type, allowing developers to only +use `generic` pointers in their code and not worry about address spaces. This also enabled a feature called **USM** +(unified shared memory), where the host could call device allocation functions and get a vanilla C pointer that it can +pass to device code rather than having to fiddle around with OpenCL buffers. + +However, like many radical OpenCL 2.0 features, this became an optional feature in OpenCL 3.0. Since it also had reduced +performance on some devices, specifying or generating explicit address spaces where possible is preferred over using +the generic address space. + +### Introducing Inference +From the discussion above, it's easy to see how to get the best of both worlds: let the end user specify their code +without explicit address spaces and let the compiler do the heavy lifting and automagically work out the appropriate +address spaces. + +Among it's other SYCL features, ComputeCpp included a pass that did just this for device code. It would consume Clang +AST, infer the address space required for each pointer, and create new versions of each function with the correct pointer +address spaces. I didn't have a chance to take a deep dive into how this system worked exactly, since it was going to be +replaced anyway and had the menacing legacy code aura that threatens to consume all who touch it. + +For anyone interested, however, Victor Lomüller has done a talk about it, available at https://www.youtube.com/watch?v=YKX6EMEib4g . + +The main limitation with this system was that it required each pointer had to have precisely one address space +which was deduced as part of its initialization. For example: + +```c++ +int* a = some_func(); // the address space of a will be dependant on the return type of some_func + +int* b; // default address space is private +b = another_func(); // error if some_func doesn't return a pointer to private +``` + +This caused issues for modern SYCL code, which was moving towards USM pointers where the information required to infer +a pointer's address space can appear after its definition. Since Codeplay wanted to preserve support for devices that +didn't support generic pointers, setting the default address space as "generic" wasn't a viable solution. + +With pressure to support USM growing, Codeplay set up a crack team of software engineers to create a new address space +inference system. This new system would: + +* Run on LLVM IR, rather than the Clang AST due to counter-intuitive behaviour with template instantiation. +* Function as an LLVM pass or passes. +* Replace all pointers in types in the module with ones with an appropriate address space. +* Solve these address spaces wherever they appear, rather than assuming the first access is the correct address space. +* Detect when a pointer can point to multiple address spaces and convert that pointer to the generic address space if + possible. + +## Implementing the System +With our goals in place, we decided to spin up a new project named "Address Space Prototype" or "ASP-Infer" for short. +Of course, this name stuck even beyond the prototype stage, and we didn't have a chance to come up with a cool backronym +for it (maybe something snake related?). + +### Putting together a Prototype +The implementation was originally based on the paper "Complete and Easy Bidirectional Typechecking for Higher-Rank Polymorphism" by Jana Dunfield and Neelakantan R. Krishnaswami. However, we ended up diverging from it fairly quickly, +so our implementation and terminology is likely very different from theirs. + +To avoid having to provide a primer on LLVM IR in this article, I'll define a simple language that is close enough +to LLVM IR and give examples using that instead. We'll start with simple constructs, and then build up as we add +complexity. To start, let's define our language as having the following: +* Immutable values. +* Assignment expressions, where the right hand side is one of the following: + * Another value. + * An integer literal which may be implicitly converted to any pointer type. + * A function call to either `eq` (equality checking of pointers) or `ptr_write` (write to a pointer). +* C integer types, bools and address space qualified pointers to integers. + * Address spaces may be either concrete (one of `private`, `local` or `global`) or `unknown`. +* Implicit casts from integers to pointers. +* Labels, `goto` and `goto if`. +* Non-value returns. + +LLVM aficionados will realise that this is similar to the single assignment form of LLVM IR, especially once you +consider that LLVM instructions can generally be treated as function calls. + +Here is some nonsensical example code to use as an example: + +```c +local int *local_input = 0x1000; +unknown int *unknown_input = 0x3000; +global int *global_input = 0x2000; + +unknown int *target_dest = local_input; +unknown int *fancy_buffer = unknown_input; +unknown int *fancy_dest = global_input; + +bool check = eq(target_dest, fancy_buffer); +goto fancy_buffer_write if check else normal_buffer_write; + +fancy_buffer_write: + void = ptr_write(dest, 100); + return; + +normal_buffer_write: + void = ptr_write(fancy_dest, 200); + return; +``` + +Our job here is simple: Replace every `unknown` with one of `private`, `local`, `global` or `generic` and do it in such +a way that we don't introduce any type errors. We do this in two stages; the **solving** stage and the **duplication** +stage. Both of which are implemented as LLVM IR passes. + +To start off with, we first need to go to the analogy store and get a supply of bags for putting address spaces in. We +also gain the ability to create arbitrary or **unsolved** address spaces, which we will count up as `__a`, `__b`, `__c` +and so on. + +With these tools, we can loop through each statement in the module and do the following: +1. Look at the left hand side (the "returned" value). + 1. For each address space that is `unknown` in its type, create a new "unsolved" address space and then stamp that + unsolved address space onto the appropriate `unknown` value to keep track of the mapping. + 2. Put either the unsolved or concrete address space into a new bag and put that bag somewhere safe. +2. Then look at the function or assignment itself, and follow some appropriate rules: + 1. If it is a value assignment, put the address space of the right hand side into the bag that contains the address space + used for the left hand side. + 2. If it is the function `eq`, find the bags containing the address spaces of each of the two arguments. Then pour + the contents of one into the other. + 3. Any other functions do nothing special. + +Applying this on the previous example, we get the following: +```c +local int *local_input = 0x1000; +__a int *unknown_input = 0x3000; +global int *global_input = 0x2000; + +__b int *target_dest = local_input; +__c int *fancy_buffer = unknown_input; +__d int *fancy_dest = global_input; + +bool check = eq(target_dest, fancy_buffer); +goto fancy_buffer_write if check else normal_buffer_write; + +fancy_buffer_write: + void = ptr_write(dest, 100); + return; + +normal_buffer_write: + void = ptr_write(fancy_dest, 200); + return; +``` + +The three input values at the start create three bags; `{local}, {__a}, {global}`. The following three assignments add +new values to these bags: `{local, __b}, {__a, __c}, {global, __d}`. Then the equality function combines the first two +bags together resulting in `{local, __b, __a, __c}, {global, __d}`. The modified module and the bags are the result of +the **solving** stage. + +Of course, we can't emit the unsolved address spaces directly, so we need to identify which concrete address space to +use for each address space. This is done by the **duplication** pass, whose name will make sense later: +1. For each unsolved address space in the module, look up which bag it corresponds to and count how many concrete + address spaces it contains. + 1. If there are zero, there are no constraints on the address space so we leave it unsolved. + 2. If there is exactly one, then we convert it to that address space. + 3. If there are more than one, then we convert it to either the `generic` address space (if it is supported) or + raise an error. + +With our example, `__b`, `__a` and `__c` both share a bag with `local`, so get replaced with `local`. `__d` shares a +bag with `global` and so gets replaced with `global` itself. This results in the final value of our module being the +following: + +```c +local int *local_input = 0x1000; +local int *unknown_input = 0x3000; +global int *global_input = 0x2000; + +local int *target_dest = local_input; +local int *fancy_buffer = unknown_input; +global int *fancy_dest = global_input; + +bool check = eq(target_dest, fancy_buffer); +goto fancy_buffer_write if check else normal_buffer_write; + +fancy_buffer_write: + void = ptr_write(dest, 100); + return; + +normal_buffer_write: + void = ptr_write(fancy_dest, 200); + return; +``` + +After that, we go through the module again, and replace any pointers that are still unsolved with `private`. These are +pointers with no restrictions on their address spaces, so they can be any value. In our example we don't need to do this +as all pointers were solved. + +This whole process is able to fully resolve address spaces in code of any complexity that is written in our simple toy +language. This language is very restrictive though, so we need to expand both the language and our inference rules. + +### Figuring out Functions +In our rules for `eq` we define the function as having two input bags that are combined into one single bag. Let's steal +some syntax from C++ to represent this, and define a few more functions: + +```c +bool eq(A int *, A int *); +B int *memcpy(A int *, B int *, int); +A int *get_ptr_offset(A int *, int); +``` + +For each templated address space (`A`, `B`, etc.) we can replace it with any arbitrary address space, whether it be +concrete or unsolved. With these definitions for all functions, we can replace the hardcoded logic for `eq` in rule 2 +from the previous section with the following: + +2. Then look at the function or assignment itself, and follow some appropriate rules: + 1. If it is a value assignment, put the address space of the right hand side into the bag that contains the address space + used for the left hand side. + 2. If the right hand side is a function, look up the declaration of the function. + 1. For each parameter in the declaration that has a concrete address space, add that address space to the bag + of the appropriate value passed in. + 2. If the return value for the function declaration has a concrete address space, add that to the bag for the + assignment's left hand side. + 3. For each templated address space, combine the bags for the respective argument values and/or the expression's + left hand side. + +As an example, consider: + +```c +A int *do_something_strange(A int *, private int *); + +global int *first = 0x1000; +unknown int *second = 0x2000; + +unknown int *ret = do_something_strange(first, second); +``` + +When we hit `do_something_strange`, we have two bags: +* `{global}` linked to `first`. +* `{__a}` linked to `second`. + +We then create a third bag, `{__b}` for the return value of this function. Taking a look at the function signature, we +know that the second argument has the `private` address space, so we add it to `second`'s bag, resulting in +`{__b, private}`. We also notice that `A` is used both as the first argument and return value, so we combine the +`{global}` and `{__b}` bags together. + +Then the result of applying these solutions are as follows: + +```c +global int *do_something_strange(global int *, private int *); + +global int *first = 0x1000; +private int *second = 0x2000; + +global int *ret = do_something_strange(first, second); +``` + +There's some additional work to do with mangling the function name for the new types, but other than that this will be +able to solve any function call where a proper declaration can be generated. Since LLVM IR instructions also involve +input parameters and return values, this logic can apply to most IR instructions by pretending they are functions. + +### Discussing Duplication +We now have a system where we are given a templated function declaration and can use it to solve the address spaces of +surrounding code. But what about user specified functions? We can't assume the user will provide these templates +for all of their functions, and we want to allow the same function to have multiple solutions. + +We will expand our toy language to include the following: +* Functions can be defined with any number of arguments and a return type. +* A `return` expression that returns a given value from the function. + +Then, we will work on another example program: +```c +int *my_func(unknown int *a, unknown int *b, global int *iftrue, unknown int *iffalse) { + bool check = eq(a, b); + goto truebranch if check else falsebranch; + + truebranch: + return iftrue; + + falsebranch: + return iffalse; +} +``` + +Generating a templated declaration is relatively simple: We give the return value and arguments bags like any other +value and solve them in the same way. For `my_func`, this would result in this: + +```c +global int *my_func(__a int *a, __a int *b, global int *iftrue, global int *iffalse) { + // Constrains a and b to the same address space + bool check = eq(a, b); + goto truebranch if check else falsebranch; + + truebranch: + // Constrains the return branch to the same address space as iftrue + return iftrue; + + falsebranch: + // Constrains the return branch to the same address space as iffalse + return iffalse; +} +``` + +Looking at this function, we see that we can substitute `__a` with any address space we like and still have a valid +function. For stylistic reasons, we use upper case unsolved names, resulting in `my_func` having a declaration of: + +```c +global int *my_func(A int *, A int *, global int *, global int *); +``` + +By doing this, we can generate these declarations whenever we see a user specified function. These functions follow the +same rules as the built-in functions from the previous section and can be solved in the same way. Two major limitations +of this approach are that it does not support recursive or variadic functions, but neither of those are supported in +SYCL anyway. + +However, even though we have this definition, the function itself is still unsolved and uses the old unknown address +spaces. To continue the C++ analogy, we need to do template substitution with the function, making a new +"solved" function with the template arguments replaced. To do this we extend the duplication pass earlier to duplicate +functions (I told you the name would make sense later): + +1. For each unsolved address space in the expression, look up which bag it corresponds to and count how many concrete + address spaces it contains. + 1. If there are zero, there are no constraints on the address space so we leave it unsolved. + 2. If there is exactly one, then we convert it to that address space. + 3. If there are more than one, then we convert it to either the `generic` address space (if it is supported) or + raise an error. +2. If the expression is a user defined function, do the following: + 1. Create an exact copy of that function, and also all of the bags contained within it. + 2. Update the current expression to be a call to the newly cloned function. + 2. For each templated unsolved address space, look up the address space to instantiate it with from the call. + 3. Put that address space in the appropriate bag in the cloned function. + 4. Call this duplication pass on the newly cloned function. + +With that, we are now able to fully handle arbitrary user provided functions. For example, we can go from this: + +```c +int *my_func(unknown int *a, unknown int *b, global int *iftrue, unknown int *iffalse) { + bool check = eq(a, b); + goto truebranch if check else falsebranch; + truebranch: + return iftrue; + falsebranch: + return iffalse; +} + +void main() { + unknown int *iftrue = 0x8000; + unknown int *iffalse = 0x8008; + + private int *first_a = 0x1000; + unknown int *first_b = 0x2000; + my_func(first_a, first_b, iftrue, iffalse); + + local int *second_a = 0x1000; + unknown int *second_b = 0x2000; + my_func(second_a, second_b, iftrue, iffalse); +} +``` + +To this: +```c +global int *my_func__private(private int *a, private int *b, global int *iftrue, global int *iffalse) { + bool check = eq(a, b); + goto truebranch if check else falsebranch; + truebranch: + return iftrue; + falsebranch: + return iffalse; +} +global int *my_func__local(local int *a, local int *b, global int *iftrue, global int *iffalse) { + bool check = eq(a, b); + goto truebranch if check else falsebranch; + truebranch: + return iftrue; + falsebranch: + return iffalse; +} + +void main() { + global int *iftrue = 0x8000; + global int *iffalse = 0x8008; + + private int *first_a = 0x1000; + private int *first_b = 0x2000; + my_func__private(first_a, first_b, iftrue, iffalse); + + local int *second_a = 0x1000; + local int *second_b = 0x2000; + my_func__local(second_a, second_b, iftrue, iffalse); +} +``` + +### Time for Types +Until now we've only been looking at types that only have a single address. Real world code tends to be more complex, +with multiple levels of pointers, structs and arrays. Take this example, with unsolved address spaces: + +```c++ +struct Element { + __a int * __b *Active; + __c int *Stuff[5]; +} + +struct Container { + __d Element *Elem; + __e Container *Next; +} +``` + +Looking at the `Container` struct, we can generate a list of address spaces by walking through the type tree: + +``` + | Container =========================== | + | Elem ========================= | Next | + | *Elem ================== | | | + | Elem.Active | Elem.Stuff | | | +{ __b, __a, __c, __d, __e }; +``` + +This algorithm can also be reversed, allowing variants of the `Container` type to be rebuilt from a modified list. +However, one big limitation with this algorithm is that it cannot handle "recursive types" like `Container::Next` and +will treat the "inner" type as having the same assignments as the "outer" type. For our address space inference system, +this unfortunately means that all elements of a linked list or tree must share the same address space solutions. + +With the ability to produce and consume these lists, it is relatively easy to update the address space solver and +duplicator system to run on each list element-wise rather than only on a single address space. In theory, anyway. As +you can imagine, retrofitting this onto our existing code required quite a bit of troubleshooting and edge case +handling. + +With the expanded type support, we needed to handle some special operations relating to manipulating pointers. But those +were relatively simple: +* Dereferencing a pointer is popping the last element from the input's type list, and solving that against the output. +* Address-of-ing a value is pushing a new address space onto the input's type list and using that as the output. +* Offsetting into a pointer is parsing the input type and figuring out the appropriate member element to use as the + output. + +### Handing out Homework +There were a lot of other problems we needed to solve, however for the sake of brevity I'll do the time honoured +tradition of handwaving them away and saying "they are left as an exercise for the reader": + +* How do you handle global variables? +* How do you handle "PHI Nodes", which are what LLVM uses to represent loops? +* How do you handle casting two completely unrelated types being cast to each other? +* What if the user provides a function declaration with no body? +* Can we actually support variadic functions or recursion in some way? +* How do you handle C++ exceptions and `goto`? + +Some of these we were able to solve, others we decided were likely impossible. Enjoy. + +## The End of the Project +With the prototype version in place, we then started looking into upstreaming it to either LLVM or the DPC++ SYCL +implementation. Unfortunately, a number of issues arose as we started looking into doing so. + +### Removing Inference Mode +As a reminder, our inference system has two modes: +* **Inference mode**, where we allow a subset of SYCL where any unsolvable address spaces (either due to conflicts or + having unhandleable code structures) is a hard error. +* **Generic mode**, where those same unsolvable address spaces are lowered to generic pointers. + +OpenCL is somewhat of a relic in the modern world of CUDA and ROCm, both of which have native support for generic +pointers. Since generic pointers are commonplace, it is also supported in most modern OpenCL implementations as well. +This means that the desire for supporting devices without generic pointers has gone down over time, especially in SYCL, +which advertises itself as cross-vendor. + +In addition, the rules for inference mode are not specified by the SYCL spec and inconsistent between the original +inference system and the new one. The new inference passes run on LLVM IR, and the mapping between it and C++ are +non-trivial and hard to add to a specification document. + +Because of these two factors, the decision was made to remove the "inference" mode entirely and only support generic +mode. Inference passes will remove generic address spaces on a "best-effort" basis and leave existing generic address +spaces in the module. + +### Opaque Pointers +In the section discussing type inference earlier I talked about how we could use the type information in LLVM to work +with pointers and solve individual struct members. This was possible because LLVM stores full type information with +pointers as well as use explicit casts between different types. + +Of course, while we were working on this, upstream LLVM decided to remove all that information and treat all pointers +in IR as simply `void *`. They did this to make LLVM IR less verbose and more expressive, but it turned out we really +needed that inflexible verbosity. + +To solve this, we ended up rolling out a very simple type inference system on top of the address space inference. +This wasn't able to fully recreate the types of yore, but since we were no longer aiming for complete type removal, it +was good enough to solve most address spaces. + +### Upstream Inference +While we were working on our type inference system, LLVM rolled out their own version of it. It was initially only +available for Nvidia targets, but seemed trivial enough to extend to AMD as well. It was also enabled by default in +the DPC++ SYCL compiler implementation, so was known to function with SYCL code. + +It was a simpler and more naive pass than ours and didn't include our fancy duplication logic, so by pure numbers on +synthetic examples ours performed better. However, there wasn't that much of a performance difference in real world +code. + +Inside functions themselves, LLVM's inference pass was able to replace most address spaces and those that did remain +were not on hot enough paths to really matter. Our duplication system would have helped, except that most code compiled +for devices tends to be aggressively inlined to the point where most kernels are just a single giant function anyway. + +## Conclusion +With the above issues there was an increasingly growing "why are we doing this?" elephant in the room. Eventually, it +grew to the point where we decided to cut the project entirely and stop working on it. Which is a shame, as the project +resulted in an interesting and elegant solution. + +Hopefully though, the design can live on in the form of this blog post. Perhaps it will remain only a small curiosity +for people to skim through while drinking their coffee. Perhaps it will end up adapted to solve some other problem +somewhere. There's really no way to infer the future, after all. diff --git a/assets/images/portal/article-images/2025-09-12-adventures-in-address-space-inference/MemLayout.png b/assets/images/portal/article-images/2025-09-12-adventures-in-address-space-inference/MemLayout.png new file mode 100644 index 0000000000000000000000000000000000000000..29e840e100f6d891bfceb131f02436885f03e59a GIT binary patch literal 25803 zcmZ_0WmJ@H+x~s2poB<+BGNUafV6-#Gn5PsLzjTW5F)LF5=ssr5-QBl-HoKQ0)jL$ zw3GsZ^m`0m_x(K2`mgtc%e8QvvySuFxBa(mI}r~xRY-{Fi2nKK9}=jlBK)6!E@8kQ zF@me$Cq({c%m4ht^bb@~PWP$VW;$Ulb?2{r6U8ttsbIqQwI8r952)}Xtinz1Z4tA$ za2=u^xydakv+=fuJ<}dy%cLOdRD6ga%L=z*d(Yzd>?S2U2T>%E@~*+x?K8R+->tfp zLEU#F9(IFi`785vA2-+5innU!@_uaPlQid(CtRAa2q6*rLcA$YClX5@9Zv)RjeOC%S9ESE1~8Z{{mlTq zC5Mmr>WSYPj`)hJ&+IJwxj&Hd@4}#Y6w10&Z{A!*d_)re5Zm{*M<%kBvtnv+gmyH83~GL{nijP6PCuz)2WA)UEn0Mr^H57;tfVov(5O zAv9T4k^Wuy>eLghjmiclK^1rH62Dwa_MxVPBnFK>S@UGSZbdU?cxI$sPgCG*re_kQ z_KsQbTP(S3ZC_(#1nNF1%EAI6xmo9x`9tRX#ABr1%c){5JN2#itj{Nz-37+;@oD!y zzUAsE7h(H%0o6uTj)@N6iqqGIvR)qUzB}j57vtKSJ3ssOxO&3Q?)%5*%h6I>;r>s0 z;@OvX`UJ#+4*Sge>%G?dy<3h=J${cGq-IF_Da)8>;1A1`_{S(s9v6u3NMo1sKHHjg z!eC1$LQmEzlP$@}Y|s1x|Fko{y~gPOvMeAjcF(dnwDc6>n{tW4< zz4dROPd4mjg;|wfxjAMNLk3dCKJRT#eY@u{-0>xQw7?*hGvjGnv+LG$U9~Cqy2!(Q zpTtru&xe-CtWZuUd5L$5L0LSU+?h0h?bRZ<5e=rpM}ep7=Bc}bQlI3;XUp@|dtqPRYse?WK47c*(9alG?62mN|*5G&oxJ?CEZl z`AOU3@*a+5zq#WKpQS{Dey?w)_R6p6Y-#w6bw9y>%(b+nh&sPzm+W2c-MlJb8wf$%5gL{lYCuN6Ib7ywzQrK|T z{;U`o$3NvwKih~0YLN@zl&$+i2`PDRSz*!xkJjL$?E1wN%ZXY`y7EEPD!00{2R;QV z#ldg>;l8BKJ8oLt3D?bS5!R;16E|MhJ#(V2QL;n>zh z12c!&rwb&zsdFb?*KdpV$CD`v2c7wkH2QdqR=Kez)Ro~2(`-MMO+Jru?TS{84_$*W52amlCn2i3<${zw+E;7U`3?W0eR&yKfU-L{v$F^)12LgbXMT)keJ z#BzFm_&LbME7lHagdUkDb(rGDfTc8y)E$f=O7_8^*^~-$+(WLBijv;{`AlM`n{}Y& z{M1$0V_98Iv39DsZkf)bokZ*nF@NpHGg(`g=0m%7VS=#%#{?U-sM>AL2D{p04k8$LA^3NEKavA95p~J+ooEdCJ%-mTIf<8OnV<#*y^+q4t*4 z4TJBJHd;kGIhxhxdjbR4%h~r16edMmPG7rD=+N8Js=#3IN?-1SPqJr2#)7d|U8E2` zIce^ZRj66tb-o6zO0eWpY0aLtW`q;pko$R5R6W`2f3%ikx73?hnt(%{*vEJxzQlp5*22 z;E)@K=#FNpC9q*hIo192FVHXgu@}Fvd|wN8b-!A4@)4`kbL7w9mCI_7WK`zlm7(&7 zX~?cy)n@y>9bF#9%FYfkmQGq9 zEl93WrmulZ++?#Om-T6t5big&J+tF6-LUhF#K7ZX^crcmx;y;%V;B=Rh zTTX17#yq;8DNufvJ#U;D3GN@K#db{p`L(b^9lqa=^P9q2+XDzF|`aNze$$Ly_V9!6Agj$Y0-Of<~yHKmCW5`miUR$a}<2TMb8 z$@<$zY-iMF#`uSq432L_o#~E`{hQxU(yTeImwc|H@HS0p)ec>kheEbe(7z2TujV_) zftsklw^;bTs{wZqB0587$3v0tm$P?j2G4I76t$d*sc-Or*+wDjzXFihG@?m~WAj{PFHB7p`vt^hN=u}g`=$Eh=Q@1=RnuJG@2 zGsC%?JxHY2h8@mbOZw?}oa)q-KaGl(JEF;&4P!-1 zeFc}SwwexR{btQ3tDQ}5r#IfrtM<`zKIOC1)a?~cjAYe9qfYKZ-o!J8z)czO$3<~N>>nrh_qaPcD6#9_A})UL=xTJ^RG?H5*f57{}% z9aw@17Hpq7HyQY^NX?by38*ss31K|>SVxFf=FFF9wTMXCGC2bkcKmxaUNyT#kuF@e zB=V)Y95uSRW%#)VZ;p68q>vu>)jd=?xw<03YdNx~c^b683d~P?J-272lR6QSEb991 z_f8i%Wy{v?s_JtkLUgjy@1r*`%r<*@ETM<&n6h z6>uN^3(x!sI@_?`GjGBsK?6R-K_{m_eaCxx2&V|Fq6k^0z&={UPFZ6S5ob-rlB`-5 zRf*oZmy}lxUX{qyFx>a#)e1sSAVM|qiAnDqDB(Hh%VI5MkQsJ)m)wGq=6X@$;>8rJ zzWcww>e&7HIltV)`Ho6Y+d0a!>bC(n8M&s5lezekyc>5U--Hpj(%2Dw0n}oqUIL@;kBPc$D z>d~XBzS_f8*I_SoQw#1;(I`~BvQpf(cS@mBb0+!|VltWfZX&uAb$>(11>)ci9SF~9 zIirvN_p_loZ;!<)>;{A?>_1n>QT}O%2nSLOJmbPUS^Rv^avtC@Zuzp)H>lgToTy>7 z>W$iQ`DKn0BK2G$aMjrnqfuGd@Lt>d<-x2mXjI`stL;z3$Buy1mkhq$6!5v`fPwj! zj4svSkxD9})E>Gv{Z8VOwI8HcQ%9}H#JNTkv@}nyU0h7^R(#78f<>D~V20dBT{clB_cqL{)b+;SM*Au8eQupbz0(=#@ZBv~1o=x-FtoqiE zPm8dSIr*vJak!M6niVaXzi7Mg{2IeuR+Z>`d?wZR7P?*wn>IY*s~beXCxN`}pR96n zPTy`hZ)S%5taKRJSkmjo4_;qP_nsAe7qFi)5qP{A<=EbOjqx$7Gb2!aDI6H9kGd#% zo&DH8Ecuvcr}m$Z+|hJnG4b$dm(s^a2ivZE=<|c~Gjr)8aKL_JS&x;$qM9744W8dy zGgI`)4syGfcb?=PoRkEeIynhyWZa)HHp8eYpcT_=$!aDYc^X#mC4MDLbEivrivn3y z#h%R?bALF^y%(}PAU0Py@GQ5Kh?dvT;-GC%Vuk*;1vT+?Q?je8lTIZfhQ^zdRl?qL z0X46Q0_VeM=~TU=_^R*XRpVYa=K%gF6FNS2b_~`&2dj!i$@(G2uCQZDs$MokDJ6IA z$12>&%NoJy><3&01|`*@3XU(PLrCiUU@`Z1!r6<*icF?DPc|w(7wgcr*v~I}*W!oZ z-`}wRFjrOrbiVnmY45tEVHL&z8))L*_1MoEVly_@P2?K%ZEU_`=cIKH&>I2sXMP)H zpRV(V;Z#SgvQ^e=(Q^I=sz9mj3%NG_dcmGCEpVR~+8myvneJF5^kYR@+ILU&6)N+I z=)(0F_=%9&EN6T5G^k|qMf&u%S5&lxg*53QkkLnsOmQ-o6TuwHq(B6eXq5|Jq4w{`erRTMAFG>13y` znt_+Aumu4RzPVku!Z1={uVo@OrUzbPToaId6QcB5E6pO(*PH#MnF`)Z72RCeI#eB* z^_*`UBX@@)fGA!V{Uy_fJseJ4Vmqe?B@BO*HehtW`_x}7#~ zlPM#@E>m<2_iE+^kfy1FHD<3{O0 z`+O)oOo4Jk&r-?aH9qH(+fuscF!V{QxAe((c^LZgyStY1(s1neEBq2iq=)h{1dAa& z+m$1l6^ru4tOUGsaIDdEKMjrUafI870qgzVo-DJSr<0 zD=-k464P9+Q*fR;`&N0 zXipYat+3a2`e`CIS20ZG{+18}oL(vYVd5Uw?zKDqTs^iw34DYCjy!i)vtq`6wD=ut zJz?zZG>*G&_Cd(sG@4GNmw}5Q)OxwL2hALTI5K%G)t(U*xwtee19{`W$EtZZ@v96Y z6Ib4)b;c*R#5&`yg>gQB$r#_^jLsj1V?)P(ePVQNUt*e8u-axzM*Ck!7~YBfBKV13 z$1f&=b!}q-Po*YncDwwk;wbf}ZPUpdG}R{pDubN9T1a%{NLil_|HAwo6=w4c7%W&D zb%TL(ur{Rq^aE4U41?c>WmgR2wkIj{3;GMow6)cFn8Hy<=NEX>OqZOs~b1FLxJgHJEt zG>PxYuTM@26iD850x}BP%sv}F@0O?Ua8)#$ZExc^eHf`je>mnpe}rE6fY#0U<>njZ z`EyYE0O40DA%0bJ@6y&cvP&PPDUC{APT%ZCu7TDK1* z-Y8M#F|!(%nT-gU2rdY;Qa3{goZ70pcdKg5$A;n;hr@$jN+YzzP z#nYB2#hdVb?w3}T>+Ca3$WoV9DC>Mj4Ds2WF`65;*BS1Q^P1M{R$kL94lMijV3nEw ztNR-T78s0zFEQs83Pq*zD_TWlA@_$WVGXTHS;!ah!}DnG{L!LxE`nIA@5{aC-nu5L z)UBu7nOYCE;*0MiNLCYbsi#}8!9}_#6#8vsOWMV^yb8nS+R4^)2>m`H8zpxi5cxJv zMPL*=&95eR)1YN{ZhMk}mt()Zw4A95jTirEl5H>ct31)9u?8BvgZ3}DqlNsAisx_p zKZ|(IIbkx1r|+ygB1+g|w?-`Byzua`U;b-UJSez)rfVv3*xd06hB*(T*0^za@GH@2 zHwBMTJ%i4iQ3}VTV8|TpV9SIK*1isZ{e?7!sXU}F170^pw}O*O5z)ipS%|lAdZ@Bs zx@~!IRz54HFpUnz$s6gUr~i>{Er1RRshNSpai3^t1fO6Pgjh0@6TW52cA$d1+1ZA~ zMI`#ch0C= z6dTWO@B~K~a@SYQ^L3-5_B9mxu>?t=rjLLU5<4Vlh?eCwjc~!$AUuK%!{S#vt06oV z@q4mr;*MsN3hsJgpJGw3j4-G?An_W-hMT)?DZtfXvDzcq8qR|s?nR^|>#KnVqWCZ_H2hZ9(cGf7b0?HPX5FZ=6AR26{3L*Bfh)+KYoXeoT3$2W2zqr~) ze_MGZ8zKi&q!ES*4Xea)!9IZ&#N<|aQh^GLJmlfpE#bcW={M&38faalc;rpLPVN5h zZn-i%;uIz@qbroRtFlyLM2Ts5bMHJks(t#)YWIr*sf0RyO(v*SwG&_~b(Y+FfZI!j zACKSF?j@?J(NkV?{rI&ZyY;n|vuB$5)A$!W(QL+%{8{4pCE(ex^P=qr-cF$A9jIsz zOuw9M2b~`;ceBP6YFM|PzXKI18wfhlj5Jqf{^{OCo`J7{YmJ$-g_1${TRYE(Hv_BS)XfxV51k$619{aCN{u@oIG z>Cr;dhB0-Z&nh2JlrK)~1KIrc?DW_K-g3OvpaGPmb2=McDBjBFcdqqMHjQ`uxT}H- zpFuOv)v>Fibo85`Uwajl_Mhcb1qwuc$}m@_A#Cr8!OX|!SNH@y*M>3GPPIniAEm82 z!sDkMYY8W&IozKV_qiPIEG3a~YP_{^Q_`)l!B>wLtOSc7B&&}~g(3jomP%!4vX>~j z1C4j9%5k(8;k$+znuNYrEJG{mAV)P)ggT7{>pJitujwU*fvoLf__;chDovVl|8#F; zf8xMt3;B~5O+^qIIu(mh{YB^}EPk*UeB-u=qwWr)61V+|fVQ`)vdC;YY?nWFr8(InjRlm`IzOv{AZoT4>fZ$)9$Tp-ZwTfaQd+{!!qk zf(}cXXn6VAQnGo{$x4(!vs?bc$?j({vDv3tYk(`> z@#-iu&qhWH{_4=n`|?mg_ty|P6~HGE2vkBOgVNnJz< z*v*spOza4+@eo`l>gGy*e?$EHKWYf6jZ%fC6=u=t;0?XRc5W`CJllLG9+zto^;9Zd zT+wX8olWXLkb6o4&U&r$#_X8JvX!6hxF$k8?b+1pRZ>jv71H|_ratsVORV^tElrmCE%;KG2U3aOiGoEfGPY9>+ME7-y@J9p&%Gg)sCOmV!C?Ea>8DbO zWbL^R$yNs|l!LuHH{f``%7G?Q+c@p>9}dj>qEPsA1o#Hqqy0^%9)~`@nnKWH)b8TA zo=MK z=<|#Yb{o+f78{<=P~%pL1s5P{Q-tCB)qPcD6kFn8xGKhM-|egrt4`ZP^oA{|3MH6_i4qFKO7wTU;K}RF+CL46qp8l9^YPj-^2KoRdEQhxrAW zP8wVkGaqKR!dcr4N&QVst=cg+i3*8nD`f(OP35S*$wU;!fhxIG&HC7M3blV$Us*Ez zaZysgZi!k2^XT(=1^@2X^QkN+PXHd=&}b@Y0k~0{Z8^u#B^-57r;t{FnT+{Qo8~fl zQ+Rcnb=hWg&Az!T;D-A+SSq~eqMBJrFFC~!r012yw)cl*f_#DtT;@?Onl||O>H}Bm%)`a%>aCaU( z>%0HDikU=rYgooyJbrJ4T800VEP*1KrDo=diLgyK88o%etP3;Jd5_;&Y-nN2RZy}w zpbHa#D~I-xS)NMD$dsJ;U%yS2zc!b&bX}BHIc3mqHZdP)&?((}s&wyuQQ0>3hbuWZ z_n+_!tP!zArLqOF)qYdu^lj25 zE-;`YoZN1*(qVsxd+Bm3{^6k1_8aie{l-n5-8g1YT#F@A5Nf}RycLrxaGyY|SDw?O zNHLy+y)(`RMU))DDv*NBEiN8;1}-X1FwQ0e(?4D3;WYB@ns!@=>U@*mfv5^SQgI{2 z{`O_cp~g91r?4g+WgGnNh<7E8JCSSCbrsTbgo1Z`7UPytI+}v6&4<~e|0pVL<>C|5 zj2rr{qloBz=h8_GUD5+k8GNjs<(b$l1&R36J^BUSphe(z>5Nc(a@MDD)B^}J6+#D63&ARp3?;zz~9U&|Q! z=cuPx5$35S^Hp z?#O7>n@7Jemz`Y|15Mzljt0ff=o`I6($U_U_3=J{t;SAQ>&}kW`5^OB4EG+zPS~Y7 z_m^6>KW1PnpLa_B`6v9}aUaV(5M3)*KTOW#x_B<4td zW7v9KY7beyOkqnahOE$8N1a^ z68;@K^>pF2#;52zP4N_7!u*ELwNPaGR`c=JKKeren+7bB^&9zdb|aPx;jN9nk`b6j(RTiL@!%G5700aa#Az7x$zq0dXK z%tY#qQlx(z5-H-{!=C&G7|`Qv!A@g^W&gi;f+eu_l?oNaav=W8&sV~Oxuc5!Uf69l?rQD= z{$gNEkSBdJ7m-P0gI2MU3v#u$zj-^JO)<5%RyWrI;}m_xS-5j`NSC1NRk}lMjFwSfQaXQLIHgj5*iJB?MXj% zdZ;}kr+EDxy_i0m*ZDlra}AXOMCyjIuZ)2Eczvy4LI+1Fii-a{bdUQQ7e8&iG(3>m z2fPMs?jw;XM$cO@K!K^l3{~I!n0AI{sa}hrw4ZLBKO^%Fjl~(RuNY-#*BIb=lHVC9 zj|J$39*oztYIItufiO+nzF2-nZ6M=hbEDQO*WDrIE|R`@DgCtl^|7K<>9rMZXhJ-; z+pD9k&p;%4M^{Im{&Wo@d@?z|e*>+E7~{dcjIWy3sAh$n~zb?gg zKDV^d8bT`bsHhsq!O3sS=Te1ym8oUl?I%Oe+}^OQ8t5KpsuE(mG8ZXanQ$e*`cF;E zDPR7PWQ*2YKu0v&*|`hFjZ3l%zXL`RA@C=Hl5_G@5HFc63HsBOoajDv+ij?Yp3m7na@^zz5B>*o z)us3(NTn)FmwcmwEFjo7-g9Z*@j;1V(f2cI+Q%3QB^cr@yDmkY*`KFO_o;Sc)w(GR z>|LZTq=%<|cOL|R3GRz#CF&#aI1Exy`A!dqi8Q<_7i_{42Ohuhm^ z!QY(fjd(dJIXOI>sjOo>Bvw`TX2)bgwSWT-ZK%M^G?A5N2EXQW&q%h#k&EdlCe~mN z3WMS_7IuJtnt5_u76J{Z-l)dO!4>##Qno&qvPFBZ?9IrY{UXg-?j+=ZDxd+Uihtx0 zkOJ47>zRu`iwE@S4XQL&75YHF@$U0i;v`*7@$LY_{4Yf0UY&AE8!eSUbQxKfqdWMd z>$I5b?q4?|b7T8Q)mHvnqZlCUWkJH20PTQS#K@%+{f&k0O^kO-URPw0J|GUe0 zzI}|g3=aER>7bK(H8YI`UuUVSxo`ZyRMR{2V~4i)aKK_$J*x|IT)U|P2|KPpwiL+c zm2IC?b4*PBtJAOD@r=E7*Wy9$GBYcJC)qsE$4EsL2D`^iJoNg42+`nkY+2#g<-nZ5 zu7DQj_QG42F?XyCXeokLR~LTomI50OXn8B@B6=bGXx+L)RBztUP7@E2I)!|2EgKwy zgMbVgjEAVqIk*J-#bGU(k@u3{zp;c9T4bIkW(Zcezo(cP=DIuZMv%2U9{6HH<=62q z*7Tz{EX0Ss-{}uo63Rmjm`OBcqnmRb0X&#R>0=q5-RF$WV=rnQU@^04*ZxIQky&j=J%7 zwO;0$kQXR|jh;My)!=4kdJ((%?`OmxoQjezK#RFYOd|VTDj$@KU$=dW`>-?a6JHoBPk)ud;TB}u@YnHpvYrX3a8XESxQ8!9= z^^AcdMDoneQXe0)=G%Kc-y>X-1qOjklZM3?+121 zz2p-^)y0;&KiH@~$u5AqV^uNki@8D%S{s^!h-1q8xX{&KF9azVi`l^|N zUh{{@iFpP#dYySn>z)ba)w>b0(?P7UuE2=p%JsiA&Zf{Yd%Rv7eXsC7+K{rThc9$lNqeq-!KyTdd)4W}i4bJ925uEg41ola<&P(X)Wi+t z$R0jWuESIEe-Q43bAP`Dg_G1U`RQ5mH(;p^QnX+*H^|7zN4FccW&C_8ni##Oksv=8 zEr@1|13QAQK0TBw-hg93oGADiS2PjnV|_Q2CLBhf-k8)*N{euNZmC0Eb9;ypmI%b? z$NC-ddxF4W@4f$(%o>%&byxae+QSaWr=;*;#UL_m!IUS7s5biwQm`*3)>}~>`2Ph5 zY||Q`vRdVKmB#*YWY_bS=%jrUEL!o8m0+1{SC|U#?c`I&Vcze?|AVYhZ~asSDbxk6 z8Spu}k`QK!*zTEKo@G;<1mRK;cTaz&^z<3REbfAD3t3g1g;Ht~j8fVFM9X=bw;~Oj ziVCi-Ad}}uYpzb!l(6UU)l8SqX28D>vVJ37TIo-*CVlbntgvE`JBOQl%T(I}XWbNG zIs)lwW6I*YqdJ?@3sCAXc8Yr7XD>?HW0rV`RG8CuG-0vi{4@xQB`R7tx1A|-y?93l z`#0capGbRaMYMQ8XysW-bf?=|BU4?QC$f?6`JjZS9G;YQ&IIvmEZP=CKaQ|b0{y^O z?Z>%_2m7UOnYa*z^u=+iJ!m!`sOl{UMDKERq;?qSA)F#vtvk@6VI7x0YA@gI;eHc& z$zuAmMWXxT^n{R^25$k!m`;P+WmFQOAD#Vpf82Jdl|y(@L8v2(TOFkK-!cULvD^g~ zH64K(3p4@F7TtVig)<&Tv_Lae?*{5rX>9eZuTz6~q7H%=<6g~E?EM#w0Z*rrmnNl> z(VsA3Asp=wkAg&^y2t3)+0GMGe?ZsPH38tvGV%rAV|pk1(#3{_{lgd3h9@)$ay$dE z%nGn~f;U8TZ$rv*+RUhd+azp1Sox9y;4&_FZ&8T#{41|AtUT-{-|ex+DDNYRΝ$ zqdXO)ZW@EBTgPq0nEf^Y(9R(Iw+yO}Kyi(d%+l~0i+NoBC-M32r6H?mrB5^D2p?Ra z6hs88SM{{f9Z=;=yM=sz;E&>3G-qBpekL)8Va2`Z_Yc!7;_oTKK>3LfRTIb(hcN=c z9U2xYdopevLy}HNls){MC69TNZ38qBaI9)zmh$TjQvr8ha?1q}wZpL6Oo#Hv;}i%& zndhvWfA`+m*w>(t3uuIsz0+miOBLQ&Fq02_Z`>W^S-)9--OT7V%}C!xqDSd@M1t4( z$%?E7uu}$#8@3#ic#W(`#c-y-X!(n%i^ef|D6*vxaIhAUU-59v-=Ze9S*g={R~mW4 zagsB6yEz9(+oD+5r@cV(qslMw5~K+LfG0t_ATZzwcgDm5dG25s#lIn>6%`02B(1L$ z_zR6G$fL7B3|`m;edjSeh1!6RG@J)ImncK6F7&*eVn^F2e-vzel=V$Z_{W@ke+)P}l?KX9u!Q?7vrs z+Km!muz7B^<3@LepW)a3UVgx|9$*)$NK5Sf=;A!o8-|}yN^hOf zWMxg;BPKa7>Evw^W?ZP_wa8akA+3M`xn$o9o(QZgAxXf9ShDOClsh=3z~*D>Y3^>6 z-0IOBE6ka?s+}hd%Y29TI zEdC|JpXrOZss*>4%nH4B8UXSp1Goq^otVqxhN}4&ww_-&zkQ{DYQJsJQT6#Px+YSfKKWb_|f-@smn3bUU2)u9;1f;k`&dB3D zDo`jO(fQP12eBfT0&NMab&;I1Vxi3OMnyVJ6r8w$H#lrI?0-^3I-vxY(H33Kc=KeB zP>(K?kq5aGMyAP<(bfeT_V51wxNTH>XzG?GQEEqbbW443`rV>uoAJkhWG6@fFhN8m zn4eQKze{zg@43z>GtK{;mZCs{YbTxR%_Cd_MQ=g1 z^b1gkG?p|Ncg|`LBfQ<;ZCEMWUMe-!x)co~3qohG8{AH%OII@B){;G~450C*`@EcC zjA-_+Bbn%s*A$&%`fjfv>tlbmtzT%UtpUk13Qvr z30hZmrSjV8)zkQBwIMWHfSjTP1nhTXJa8l%DFOX&qya>nh#fN>qma3FnTj{;Zr;7S zmZR6{#lQ5wlWn=d9M?<_w+iC;(dPC>M*MCo!*z5hJ|#g%hn0-OFxyJW=|?MW8wfQE zVm1tutOZ+Sv9er&bK`+W{@JjG(M5`d$Rf_|{qu7bS#XAO3zMuG_{g=(mY?yWMwVDV z+`MZ^P9f1DQf7J407TDFs|@*q^kYX7g-u7Y8>#`1*kLZG$WN zzs2Pgkqee0^UV#bnNjL3fv1IiI1{2QuD;CB_xBImla%yTuLA!sthPd!$-tStTdd^>MyIs-D-m%twN;Yr6thnHr8iu80tek`r zHo-Sc2_t+jS}#VM&WJ2|K3)3JQ5Xob;i#Z66#E}0NPNKrrGX~C&WsB(rT)HPf~v0! zlaUwzciiH{TiF~`##Et5wp+ikeY&iFr9kRBh!c>E z-ld_~x_yHbrGvN{^FCW#-aDq9eWa+Q1@{*fgHfE z4OP$>4#bT%X|{b$cPRH(wSqiakYGy?$(9~KzeaAKeAYFSEdTog2h#VfyzzDMB)TD0 zdhi>^IK3`VganC!jb-PM7&q|m8-XS0`9HpI>*A)5{Zv1_<1UH~|AI zPkzYbl!mHkVI->rq%%4zmUx+e{W*?Sh)Z`RL*Jn9sJ}TeRHpx=jPuR4P-R-$F7tR3 z3o){0x8n)!SW7?=;rtULb;qn=)F zfRi8~9IOco9kh%iwR+!^t+$2&mvhNlVv?Tz3g^*32TtM79=d!WQPjOD>6sfKVt%HU z#4E^F9z{X30#e?S36C5oHM@a|>qE^#NygC}&@%^hun zWA-?gor0fgD<_oDM)6kktz;cE8Wogf=G!h4-AvX%Qt06Gnp=XNhrZ!{+0<*FS7~>a<)pc-M=eX@k`J2LPHozZdlx9 z$=qvH&!wW|`S5V(eYK=`+D^>-c=;p>&h_v>GQ5%jXAaJ7@hVmE+_|&;8ck`{YY?TN zX-M>9yKz2-El8(c&F(K*sRurb#$5Bf`O{c(IZdtmzdF874;W=otw)uK5+Ur=``F@LvE877*XI zl{U82D(k;d)b#Aipi^|AbtgF_gFtMT_R>kyF_>t13+x`lmX)Y|#PUbR{d^`CetVBg z^*4LyK&Ach&e%($)3~1=wk}MgVqoTcL_eNW9g!IRHzRgC8a&j}Si}3O5o8>HgIQ#{ z7bolov$0y@xvv_$lEpVHIUY2hL6OHqmh1-`_NL-pfmC@zox6Q%KZM*k2jJ*2o;&}2 zYxX#gNVQ5s?BY{y7@ZY91mml*HZ|L?t57qjXMBecOqwI+4vWkCC66tVhU z!QLTH!Tl&YgoKPpkZ8h{WrZ5=H?TjZ^KSork4lp$hUCtY_L9N(;6y)OT!Gfem?-6C50Z2Sv(@GFASo(LI6(1IFu=PcsRf)(W9R&`)78WoXZ zH}}V35!l@BMw_tVdumtjcK(f8L8Y#Ti_;7?!vQu%J2Hh;lOmbfXAMAbO>HD4Dg<3A z2F*ws#*H7K1LiKB&Emi4PP=H<3A_;=I_HJDm4&_lRSMa9N(1(<3qi zIQC;Z$OjJu;}biYf6lliM_fltne4`lV#-7-kcQsp^#7&2##=1PZt9AQnq^g%L}s}; z<5E*9k~M6*_wv=bRr9po$%yijvV{Zz=xl%iSpS-%Mx@Zh%>~&Cp*CzgiaxOV90$9R_K8JKy~c2Pli`ZnMjbtFf< z@c<`wJp^5UR{f+dz0&96f8iQ$ev|KSqR|^Ro*#CO(8b?T@ZJJlbk*J#xa#fFlqG}-= zV8iQ+v)zG0vnO5i@gLT5GG~L_+sS`(Z^j}g8P7h{GG1WDSe{0!J(-t>XfGt9nr}T^`(!!|NSPIggRi;8|M^cA^lPrR-TCPO zX?hNCB}n$$*}Qe>ezaaxQw^Q>SA`_Oe+D9aHcl-bJ+v?64v)uiwHJpI+Ab!^EDy z^z%w+Pd2wpO>R?Jy#~aeH3nGd|3bxlnXeBYl~!El=uu3HhkGsLmU2+5%nS4WMHs`l zS!7@*(dOTxI7tWdz;Q_*2fgd*{*whbTzThO6al7AwSt$rg_>0%Pswi=7^+WW={4!ZwL=Vo@65zTz9|0#Z}`RoO-1&>IUS;Dpa0G|rP5^V*r$8x3!3J? zwFhzFo@Y4UJ>Cbd{`a1g5s-r18GpBBTgCv$^E~bj!zy4|Fz=leXM4O{1V+`^d!I<+ED%gf zg{Q(b(0e4&`y`_=AT<{N<6k9FVXx^JzoI)DoyDRGp;rf@Rq60k{_ZDYA(pH3DHZ{$ zN^$YmC_c)|zXZ-EUz8HZ9ai4n|C#B7Li3uRP^%2zGBB0^!*oY#ozfT((*4GvI|at0 zRN2b>Ym=s6a!ib~1Ck|H)u6TV9 zUIiAm&+%rBk}E)3Z7)6@eA2YVs@{7=ZKR6Soi5d*!a+Qt?QbRKk6f8q_vLn28K`nH zf7j#^R%!Fa=m#NPu4@(c9_9u7XqN);w}yko_asw99s9LN^u-xi$h&sE%YdLU&vv*UZ|RuE%Vun;YK!VUN9;UrDB-?;c;9#G-R4Rq&+vwp&edGQ z;XTsq7v4TimEk7AI=Ozhp^o#DJyGHk^?iDTp5 z$lu>S89scec6^cmZ*K~xRHvAMfd6_tXqy280fSo?=K(H*I3j!;_!Snbe;z2K*e}MYX7D^kLdS(KiUR?~mPPtnhgL0Dq>M zIkT+;Ogn;HzYGQ#bjBqEwON}WL0(ZEm?AE_bwM<*jy3I1@>J5sk{jt6Jpz~B=Bi1f z&xkE2c+pf+=?@EYie@EM=YLasTn>`l$L_-=F*gO<)Kf(V%*-$drdp;spOrhi6ZY>? zHf!dZxvGTiaHP+xi5pT^fD5ixd?)PGuH7o%L9V2u(2>ty)s84;8!%!+P;4po5Ml4F z>B+qnIn+-k_|xs#CXzKN?BM+=GROzDVf$FejBcV>N?CKCCF&9I8xm4CX{}fkenN86 ztY-apH3|a5Bd4`Np1!nfFT7UHLY4RK<&F5_G|yib!+sSG3q}@oTmx2R0Oj4hY;u+H z7XJcsy0FYL6qlsmheDzGl$}HK6jfq|%7EJ+VLL7y>%ilwxyh$qmMr*`#4K^KNc=;x zWTnzC$8O$btL5X%*AhDA7sXG0j27_CyG><;PXmO0SeI z))8gzUA1Hm(N%V2!!fwAK=L0Bsm|4cIGpDbTbfyG-IXx{GJbXuYarlW{o6v$@^mSl zqT?y}H6C9%0z2v*#|l&Gw(8eWD7rUXxDk^ggZEff1{ppcM(@9|FTWV+eFD>DW>((l zTzq<=wlGg@1Sjm<;TLzHdh|&+FM)vr^Ov#Wg_m$sf-pLd8|tOku0SkX7bz&~yuc76 zRf=Tu!%ywgH|my{*;UEZ1m2-=CX-s)?(#szgy%b-a^DLz8h*+4(P63;;-+#NykAg17okZ*wI^Pa zsQB)^IXSu}&5zzYS}&QZ4g3GvI`4QY|G$raMI@1|BH2fdm5kCbvYpIhiy|d6Ba*#m zwu6i!&ar1^%E(SsP9X_7Mtn2kYbE1;U+VY0@5g=L|9YHje6H*Id_M2-dOZQg;T4q` z^(yB|s1JQNl@Z5jKF@bhqp|0z@VOd!$7N%JQptJQ@#S)7o>f;Y216U22c4AkAvCn- zhs;gX=d`ucCWh$huAhjkFL@W40Hu}Qvt5sniUu*iTm0aqQ^(M$@oL8Q^Gg)}(fO&( zRsMfEzZ2o9*S%izTen|S049Jv-~q0Cmw?7^>3kQ5cZvJI1-*H0W;4GmpOwf){k|qP zPM%d-`{CORDa;z$0|AAdjqFs9_$I$);Gp6k5o;hGd#>c?>{d?8n1f%c$x*hgdOZkt zeNxxt5ha+<^b(BqsfdLo8@MoDT@rtOR6zyg9hp0mie?x4emPS`bI5#NOr{8Wb9b54 zv7W79<_e0*YxWXD$*Qxg>;Fgl_o7ew1{fvslpnI3i{$c9u84Tg?5)e5Qoq#nYu+rN zTbl3HF5|l+tm5s16XG2Eo48yzmVIEz$`I)~)>piKEgDpcyj>%X8xykaw~tNvviD8P zGJA6!`lTArP$tPK`XX#AgyujhH=c>z?<@U*Ut$kgL>P)&Ynn+>GR%m89b^rU5o_24 zg!X9ka5)LcK^qC@e{%CS<(WKyou}0xjOJZ9RMtZ1(q(-cJYn6izoDrizOHLll6-R{ zCOJ&%bzqRZv~{c&{wHDIqPVZ@L+&>hYi%%k=S^(xBgSvcL=7pDnghBrtHqD3E3OUl z;P_)6$vz1IZ@yZdnvgxRsOIup$+trFxK-+k!Yo&PS9^lnx=T>*i0+cdoRhw}%Qn=~ z_8&xZm8Q#TjBFKB6N1jJ`3Z(i$g$hebzgp6B4`-R@pMN@EPi$NkGQ08DC6B^S6tlX z&g+<8D_PUK2mk2xCS*T?g3y6iP${yyS$C~{&HJ}}-yCoSE`0qnTizN2F}Y5?b$`;G zkP_}scOb_>8dEXiP&_!M8XGKgYM9BP9tR4#sV;Gn^f^RM;GGlw{_-@-BBjjk=41_b zq;yK`G+!kL_#r*Af;OH*bOI@|DO@7)8f!skWlWqpAC=#776dBmrRR8`Is3MUr@92L4zIhOZw zw0S^If-m29$W=dY0#5`6+0Lz6bW<-AEP|rdBHt!hLDd5_7A~a_!3zpJ+0Q(slTZWw zZE;-S1&d>G`2BkVt*N_s#}r(y96LYt7W0}0Vh0hhJJ;=(qPC*Z0*@9wlAmdAprP(p3_`+t6y=A;?4w(Gx+TP{uMbbfsm+y|l-Wp5pq z>SmqC^TWryB4}zH6s{^kyqf~yep=PPx6NpWbZ}=K0;8Q3D?f!(ryQmcOa1B-fx|g_ zyD(ik6qB^)#OatYc!h~#m8t&u%HAo1qey8k^aXBTcK1vXWOo+L@ievALdZox8yfmD zbN`Wu-z1w7lGZS0|B313L5eD8e2GR-u$;UW6fyEq2qiJ@c-hf-%GJoHQvCWzle4H|(N z-l%+avdqg_Utng-t1;OQ-Er~x8)MHuDvtx^O`RgJdzi7{3(_V=awkfp^MHn+BAXBs}a%E*Wu@$ zaPwkq867Wb;E^-^{lVktKxXHQg^Ujmb|*64vO&Tac^?VS{2W=tEs70X5Z{C8Ot1N7XYELR`iRE-Ro6ZlLS+92u4njd{dh9e z4pxzqzAn|c!cM*P>v?R<_!}QgAGvqAlf4-#bMj3U0Ev- z$MIPiqJLMMAf`V9g7{jRQ@r~_&_4!ZIUN)A$kf~1m$L1n;ua1yG_G&?0wb^@zre`W~9FEW=))`?Gu7egSLJ1M5wN>MziO6u94Wv)E z1>8Xnj=54w_*E3)-YZo`p->#2)sE1RL{c=>a+$I^w^gcm@wm@RDXo1Q7N6^$bI1h^ zym~KUfe1Uwc^8a$0ZPKC>pj%+ImuME!kco~>N@+Ox1CA}y&E_<&#jF%Tb+Cn1wR%TTOk0?L^^ldY1^?#vA~jRb#JZPL3&r3cG>fp?eMePuLW8M-Z%UE ze<(4LxCauHjLw*oBNs17AE17lB_}72QhObv=5laXWMk5+XfsC)}Z0J`Up3e!9dTrs1$*A zo<2y+dr4RpuvbRcio?bkZ$8;p-sZ!c_3uV22~xM7MN)u8gM@A|BWIA&O-(&5*Z+~& zp@XG7Zo}ZuW%K1x7yG?vd>SeltPyg9^k0|y00yVL#x%)7S8R;IB!}m&H5#&)YKKqm z-f2s}(EHYD7!K5=0LZT`iYd z_idlGU)qkol7OCZRY)*11=+z{0JND;zG4=dQ>f@re{i^XCHad^^|QQ}^qi+TeSQSW z(6ZJ!_f#4b8=gy#P-;pu^4rb^`4+Mte7l|}?#wiL23a{EP?l&dBzIp`6=de!I*#S%yugL2!UNZzpAd9n}G6iI^E!)$Gx zf-E19nlmM076s8=mX!VACq*WYQTJ2?##j>{E4wvII+nJ48tG#r#yjuX5DJg^$)61C zYuzRtPZO6XyGC{{7c%h^NcUuZ@)K4_YMaa{xn>;+Zh1e^MkdmNU+y55b{VKn=0R=w z>F1=dnI!-`riobwO^|FJsWs~prp>~i#4E+q93w^!QC)cYa9f7Z@^}O0-SEo2ncm96 zF|}q0*}B{dIM6Ed`ajFRoW_vANpb!A~c$GWgmP>{>Qn_OQer zesT80$`#$hd!i2)?!VH#iT_x7?Yx8wN$KQ?8U|s4uolyTUNvw%nandP7oD;=Jz!rL ztkEE{Isdtg|E}kUr;6e4chz&b6ji?LaqE>PVu^Yx7vSCf+t3)(N)e#n#~q=nSMCtC1J&_kbrxq1mU%(vWc zyo#gb6dqSw`@9oEhMLIX>D#RR`;^9N5LZ8Y>fQ=H1KZDT$OwG|O~3avQNoWqbl1N{ zuU9ltnjv93zOoTno*yTqe4tKJpv8ZRUy&(CbnBRj%LDw(2!9Px5zK)YO%mLH$~JAt z5ouD$5kjWLTNVb$Qv=KQv}Dni){FhyY%5Tc_=thqG+L|T)h4rT(8HQtBS}^Lxl_Rt zh09v6qx7R`_B%<{VEX^in2#3Zrm?U{Rm| zZ3W*qd9XPnGj6{XoCq`S1hcKUJQ}S$r|MtMrWQ=eq-uNOprc~C)K?hW`N7=8M+eK; zQ3Rg+xCZ(teY2Nw>Hx)y4H^@OC7m)V9nGRd>OHGqcKobQXUWr8DXK|XV6k%kld0Yc zn+`ilj`DMmn?2WAY$89oJ{GeABeL{JWlJZ1$6Phchf7ef0^3PgoVQ@rHCwVH6bj*5y*b{GR~jWdbo+UWj~Yf& zzD!D7kT`e&-Afa`KZR@Djp5ue3HkUb|MxYWQs@_woYadtVl@ox_c-Pc34W5`_RqL8Nb3WM;o^P}}dCb@d zTU396KugyjfV{YiF+~RCNPWf&f`-1ljJs?_Lg+{M1Wp{T$lcDmY8?2CvBm2>S;%X1 z<2CYnZK(cw+`?W{*{0=#^%b2M98Qo?)vJXN!I~F0WxId#U=4t9btIThqxv3bcdDp6s%=>!a}L=M zf#IrdlR=QZH`niiBH;5Qn#S{C&dJsky5H6%A5&NV`$_y0q8_$~R9!qK3)Msy{lna6kC>NWn_(keDm9j+aW%BYG{#=*?;$A~_TSWP7)w9Svx~k*FrJ(+L z=oZn7n7uy@h+?Ln!}ySiB#eiOLNnt3J?xmzl%3&gVZ%Sbc2CR-nW_7g0pn zQwS6Lib|H>R-rAR=Z?eo8X*u01{5{lR6F?`G9)rL&HKH`+z28zp3cJKQ}m>?*W6ak zF&NuU-HVcirQJop7Z-nE0$W2j=Ey)OlN&bcxi+@vZ!WuLK~?go+#XgTM_X}Z-eKQ6 zK2uLh|KMZKqx}zf>?O`dc^7}q(e>I7oEyT6rK8wbcYIBMj{)$YUyzWmdI2WlDC2Sf zdd4cI;avW0zx=XO!>zi1e=Mf$3D{qGe8Gc9nRtj1J&V8F^gs>mHFYzube7U4Qo}C$ zvor4Xm5*8-MaXA-(A+}TQO7E9psLaJ1h*d+Tg5-fnwLD(AHc*(bj7|?+Lk5;(hUTN zN|2~|_)&BCl$0b>ed1QT%wSMaAcYjy<99^Ys9Xvvy&I<b zBlWf2v(XQ7=_1Fu(=Oc0``PE@XLYV%dCs>6DjV*@KL^SO^I!}Ga7%P-Ey+U=A+!;A zYm?bI26Bo5OYG&F1FK|PKR2(~vwKBM2%O0rSoP{6IqKjf!cF-R)ga zex-xiSMt*nXjF{Re|LenjJYn2-t40{2q6bO!)jNT#>X{)50eSx9CyDu=*4D0LA!k$ zz~A)Esx(EEMJD7J%)sR>=1(3hI}myr-~0?@t3G5I_X9j>Mw$7Gbl?(7oyA@RyZ|C; zL`)JliVa%tyBu?d&Ekv@8n(5!{G@11r%7PjjRDQFL=jX*x%Lv;U`@X~^DU#;iHA18 z3lF162(&o6yw|+5;Iyzum4ZcZc>of01ez}0H75I+WiSY zuAQRf+Fi|;V?DsMx3lDbeqp5UJP?5!5%#iU1u0CJfV7Wm?E#@<4ch&xrAc#yu5D6b zaYY2J`AF=Qsp-paAaIrfQcCPP@^Kp1Fk-d!t4DAVEqt|-}& z+2Qw_M#zn@sP3gL)Q-GG`tS}=bVn>6=Mi88N{D$tK3yssKV|Dy)KM*pFLPac9J+;2 z&uLxa=<5?Y&e5~%faQI_6eB!?q(9~#6F`waWU9ye6ut4+Ulezd{}CX(6;A{ZvWPBs z$QyrvRrZ4Xci}qoh3m$jT)NBwpB(nc4E)_?nB=|m|9EG){Jm;UQYsU zzHTY=qI*$IYs4i0c8fuux(rd{Yb5b&3Y#cMHp5zkxOhrBh5$_Z4n#WI zA(xIa?0^qR7~O9~k<%8`dLRh*N#a2Lp+Tm*vTg^^6SB9C(y)T^#S`+tUZi0d$@&JP zEERtr79WCdNVRVbTfysJw@Yc)y^J%-fs5iX??==84uxVN!B^Zb+H(kh7Sf>aD#1A8 z?qH2n<11$xR};;iYOyKQO7e7>F%Ci&8r z3Tb7(S32WZNjM zDBx0{<55$mK5%=SnkPu|g87O0G?1>Z9y}06VOtkvm(3s<&jU6~ciWQ>hXU*klS$)v zZNFSHwX{YNI#*a;1h2qSoh*j=tTGUgOIguV|6z4>$#2NiH?v(&msURrzFk2}Ty;zl znJVo|sAAmDie&Vimn(!Lm60RGv48Ir;Sz}Vi(ubF!4_3WF+G?2XAk>oM4p1PyIbzx zQR=#Q39eS8MegOZ@}1jAS%374wn=zorULyWZLBmYqP%|YZHFxLzo7u%Q?kXME^|J? z>E1_fxMT%pkTM*esWUkCdmtBr^UE{a2h47^hfP_$rIhMYZRROEQAk!_12r}e_i2OW zsrO~Y@WUrk*;L7F`U0XanzyIVtP{85{8o9r^WP>iXw#oaJUP?8vjsC}!Ban~uuQj4 var3#se*%h!Y+O8Y8*>p4INNi!_s$P(SL?EBGs8bA`Ac0{OQ}r3BIv&WOSUF; literal 0 HcmV?d00001 From 6d68047098fc8cf558e0203524d58528e989a5a0 Mon Sep 17 00:00:00 2001 From: scottstraughan <42965777+scottstraughan@users.noreply.github.com> Date: Fri, 12 Sep 2025 09:50:59 +0100 Subject: [PATCH 2/2] Modifications to maxime blog. --- ...lashattention-performance-on-intel-gpu.md} | 151 +++++++++--------- 1 file changed, 79 insertions(+), 72 deletions(-) rename _collections/_portal_posts/{2025-09-02-gpu-tensor-core-and-data-feeding.md => 2025-09-02-improving-triton-flashattention-performance-on-intel-gpu.md} (80%) diff --git a/_collections/_portal_posts/2025-09-02-gpu-tensor-core-and-data-feeding.md b/_collections/_portal_posts/2025-09-02-improving-triton-flashattention-performance-on-intel-gpu.md similarity index 80% rename from _collections/_portal_posts/2025-09-02-gpu-tensor-core-and-data-feeding.md rename to _collections/_portal_posts/2025-09-02-improving-triton-flashattention-performance-on-intel-gpu.md index b5b1636..d56005c 100644 --- a/_collections/_portal_posts/2025-09-02-gpu-tensor-core-and-data-feeding.md +++ b/_collections/_portal_posts/2025-09-02-improving-triton-flashattention-performance-on-intel-gpu.md @@ -4,42 +4,44 @@ date: '2025-09-02T02:00:00.0' hidden: false layout: portal/portal-article-view thumbnail: /assets/images/portal/article-images/2025-09-02-intel-gpu/ComputeUnit.jpg -title: 'GPU Tensor Core and Data Feeding' +title: 'Improving Triton FlashAttention Performance on Intel GPU' user_id: 72195828122 --- +## GPU tensor core and data feeding + Tensor Cores are specialized processing units within a GPU. Introduced in 2017 by Nvidia in their Volta GPU architecture, this specialized hardware is used to accelerate matrix operations, such as those used in AI, deep learning, and high-performance computing tasks. Tensor Cores have also improved computation throughput without degrading accuracy (or only slightly) by combining lower-precision data types for operands with higher-precision accumulation, which is known as mixed-precision. - -Mixed-precision computation has unlocked AI computation performance, as AI models are able to provide results like +Mixed-precision computation has unlocked AI computation performance, as AI models are able to provide results similar to higher-precision calculations while reducing memory usage. But this last point is out of the scope for this blog post. In short, Tensor Cores have completely revolutionised the calculation of AI models by enabling high-performance, high-throughput matrix multiplication: -each generation of GPUs increased the number of floating-point operations (FLOPs) the hardware can compute per second by -order of magnitude, as shown in the table below. +each generation of GPUs increased the number of floating-point operations (FLOPs) the hardware is able to compute per +second by order of magnitude, as shown in the table below. -| GPU | GP100 (Pascal) | GV100 (Volta) | A100 (Ampere) | H200 (Hopper) | B200 (Blackwell) | -|--------------|----------------|---------------|--------------------------------------|---------------|------------------| -| Tensor Cores | N/A | 640 | 432 | 528 | TBD | -| FP16 Compute | 21.2 TFLOPs | 32.8 TFLOPs | 624 TFLOPs | 1979 TFLOPs | 10,000 TFLOPs | -| FP32 Compute | 10.6 TFLOPs | 16.4 TFLOPs | 156 TFLOPs
(19.5 TFLOPs standard) | 67 TFLOPs | 90 TFLOPs | -| FP64 Compute | 5.30 TFLOPs | 8.2 TFLOPs | 19.5 TFLOPs
(9.7 TFLOPs standard) | 34 TFLOPs | 45 TFLOPs | +| GPU | GP100 (Pascal) | V100 (Volta) | A100 (Ampere) | H100 (Hopper) | HGX B200 (Blackwell) | +|---------------|----------------|--------------|--------------------------------------|---------------|----------------------| +| Tensor Cores | N/A | 640 | 432 | 528 | TBD | +| Float 16 bits | 21.2 TFLOPs | 130 TFLOPs | 312 TFLOPs | 989.4 TFLOPs | 2,200 TFLOPs | +| Float 32 bits | 10.6 TFLOPs | 16.4 TFLOPs | 156 TFLOPs
(19.5 TFLOPs standard) | 494.7 TFLOPs | 1,100 TFLOPs | +| Float 64 bits | 5.30 TFLOPs | 8.2 TFLOPs | 19.5 TFLOPs
(9.7 TFLOPs standard) | 66.9 TFLOPs | 37 TFLOPs | -source : https://wccftech.com/nvidia-hopper-h100-gpu-more-powerful-latest-specifications-up-to-67-tflops-fp32-compute/ +sources: [GP100](https://images.nvidia.com/content/pdf/tesla/whitepaper/pascal-architecture-whitepaper.pdf), [V100](https://images.nvidia.com/content/technologies/volta/pdf/volta-v100-datasheet-update-us-1165301-r5.pdf), [A100](https://images.nvidia.com/aem-dam/en-zz/Solutions/data-center/nvidia-ampere-architecture-whitepaper.pdf), [H100](https://resources.nvidia.com/en-us-hopper-architecture/nvidia-h100-tensor-c) +and [B200](https://resources.nvidia.com/en-us-blackwell-architecture) But, now that the hardware is capable of computing a large amount of data per second, the question is: "How can we feed -these engines with enough data to really take advantage of their computing capability?". +these engines with enough data to really take advantage of their computing capability?" Indeed, if we are not able to continuously supply a large amount of data to these Tensor Cores, as powerful the compute unit may be, it will have to wait for data to be brought to the computing engine (from memory), and computing performance will be limited by this. -Another question is how a programmer who is not a hardware expert can take advantage of these “very” advanced -hardware features. +Another question is how a programmer who is not a hardware expert can take advantage of these very advanced hardware +features. One possible answer to this question might be to use Triton. Triton is "a Python-based programming environment for productively writing custom DNN compute kernels capable of running at maximum throughput on modern GPU hardware", as described on @@ -56,22 +58,24 @@ Let's first step back from the concrete implementation of a kernel and look at G If we consider simple matrix multiplication : $C += A \times B$, we need to provide Tensor Cores with two input operands A and B and store the output C. -In classical GPU architecture, the operands and outputs are provided to Tensor Cores (also called Vector Engines in -Intel GPUs) using registers. +In classical GPU architecture, the operands and outputs are provided to Tensor Cores (also called Matrix Engines or XMX +in Intel GPUs) using registers. Registers are a kind of small and fast memory bank (called Register File) located just beside the compute engine, as this can be seen on the following diagrams showing selected parts of an Intel GPU architecture. -![Xe2 GPU Vector engine Illustration]({{ '/assets/images/portal/article-images/2025-09-02-intel-gpu/ComputeUnit.jpg' | relative_url }})
+![Xe2 GPU Vector engine Illustration]({{ '/assets/images/portal/article-images/2025-09-02-intel-gpu/ComputeUnit.jpg' | +relative_url }})
*Illustration of an Intel Xe2 GPU Vector engine architecture (simplified)* -![XeCore GPU Illustration]({{ '/assets/images/portal/article-images/2025-09-02-intel-gpu/XeCore.jpg' | relative_url }})
+![XeCore GPU Illustration]({{ '/assets/images/portal/article-images/2025-09-02-intel-gpu/XeCore.jpg' | +relative_url }})
*Illustration of an Intel XeCore architecture (simplified)* -Basically, the tensor core reads operands A and B from the *register file* and then writes the accumulated output C -back to the *file register*. +Basically, the tensor core reads operands A and B from a the *Register File* and then writes the accumulated output C +back to the *Register File*. However, as we have seen in [Introduction](#gpu-tensor-core-and-data-feeding), Tensor Cores have improved significantly, -making it possible to compute more data per second. But, this implies that we need to feed Tensor Cores with an +making it possible to compute more and more data per second. But, this implies that we need to feed Tensor Cores with an increasing amount of data to take advantage of their computing power. This raises two issues: @@ -82,7 +86,7 @@ To address the first point, recent GPUs incorporate dedicated engines to load/st Memory. Therefore, these *Direct Memory Access (DMA)* engines (called *Tensor Memory Accelerator TMA* in Nvidia architectures) enable to hide the latency of accessing distant memory. As for increasing throughput, the common idea is to achieve that by sharing as much data as possible between Streaming -Processor (SM). +Processors (SM). Indeed, recent Nvidia architectures (Hopper and later) comes with a *Thread Block Cluster* feature that allows experienced users to reduce the amount of data fetched from distant memory. In this post, we won't go into more detail about these features and how to take advantage of them, but we can recommend @@ -94,15 +98,15 @@ and [Thread Block Cluster](https://research.colfax-intl.com/cutlass-tutorial-gem As mentioned above, *register files* have a limited size and must contain all the *live* variables we need to execute the user kernel. So, if the user kernel requires more *live* variables than the number of available registers, the compiler has to resort -to register spilling. This technique consists in freeing register by pushing its contents into memory and then loading +to register spilling. This technique consists in freeing Register by pushing its contents into memory and then loading it back before the contents can be used. -As one can imagine, these extra data movements back and forth to memory can severely impact computational performance. +As can one imagine, this extra data movements back and forth to memory can severely impact computational performance. Although compilers do their best to avoid register to be spilled by optimizing the generated code to reuse available -registers as much as possible, sometimes, if the amount of *live* variables is too large, the compiler cannot do its +registers as much as possible, sometimes, if the amount of *live* variables are too large, the compiler cannot do its magic and registers have to be spilled. However, some improvements and techniques can be considered before relying on low-level compilers for "last mile" -optisations. +optimisations. To this end, the 4th and subsequent generations of Nvidia Tensor Cores can load data (i.e., A and B operands) directly from the Shared Local Memory (SMEM). @@ -150,57 +154,58 @@ So, in practice, the compiler will decompose the `tl.load` into a two-step proce 1. From Global Memory to L1 Cache (or SMEM) 2. From L1 Cache (or SMEM) to Registers -In Intel architectures (PVC and BMG), a hardware-managed cache, i.e., the "L1 Data cache", is used to bring data closer -to the compute unit then data evictions are managed by the hardware in case of conflicts, quite similarly to what is -done for CPUs. +In Intel architectures (Ponte Vecchio GPU Max (PVC) and Battlemage Arc B580 GPU), a hardware-managed cache, i.e., the " +L1 Data cache", is used to to bring data closer to the compute unit then data evictions are managed by the hardware in +case of conflicts, quite similarly to what is done for CPUs. The first step of our loading process is therefore achieved by a `TritonIntelGPU::PrefetchOp` which prefetches the data from Global Memory to the L1 Cache, then the second step is carried out by the `Triton::LoadOp` which loads data into Registers, hopefully from the L1 cache if the data is still available in cache (cache hit). The diagram below illustrates this process: -![Intel Backend Memory Semantic]({{ '/assets/images/portal/article-images/2025-09-02-intel-gpu/IntelMemory.jpg' | relative_url }})
+![Intel Backend Memory Semantic]({{ '/assets/images/portal/article-images/2025-09-02-intel-gpu/IntelMemory.jpg' | +relative_url }})
*Intel Backend Memory Semantic (synchronous)* -Nvidia has chosen to leverage the Share Local Memory (SMEM) instead of the cache. SMEM is indeed a scratch pad memory +Nvidia has chosen to leverage the Shared Local Memory (SMEM) instead of the cache. SMEM is indeed a scratch pad memory explicitly managed by the software. Hence, to accommodate the Nvidia backend, we find in the *Triton GPU dialect* some operations to manage the SMEM, such as `TritonGPU::LocalAllocOp` `TritonGPU::LocalDeallocOp` to allocate and deallocate a memory buffer in SMEM, but also `TritonGPU::LocalLoadOp` and `TritonGPU::LocalStoreOp` to handle the data transfers between SMEM and Registers. -Consequently, the triton process for loading and storing data (synchronously) in the Nvidia architecture is as follow: +Consequently, the Triton process for loading and storing data (synchronously) in the Nvidia architecture is as follows: -![Nvidia Backend Memory Semantic]({{ '/assets/images/portal/article-images/2025-09-02-intel-gpu/NvidiaMemory.jpg' | relative_url }})
+![Nvidia Backend Memory Semantic]({{ '/assets/images/portal/article-images/2025-09-02-intel-gpu/NvidiaMemory.jpg' | +relative_url }})
*Nvidia Backend Memory Semantic (synchronous)* - --- **NOTE** -It worth noting here, that Nvidia Tensor Core version 4 and later, are special operations that do not require there -operand to be in Registers but operands can (or have to be) in SMEM for the *mma* operation to take place. Consequently, -the compiler does not need to explictly load the tensor from SMEM to register, but the *mma* operation mamange its -operands itself. +It's worth noting here that Nvidia Tensor Core version 4 and later are special operations that do not require their +operand to be in Registers, but operands can (or have to be) in SMEM for the *mma* operation to take place. +Consequently, the compiler does not need to explicitly load the tensor from SMEM to register, but the *mma* operation +manages its operands itself. --- ### Variable liveness and Register reservation -We say that a variable is *live* at a given point of a program if the variable contains a value that may be use in the +We say that a variable is *live* at a given point of a program if the variable contains a value that may be used in the future. -In the following example, the variable A is *live* from line 1 to line 7, where the last used of the variable A is -found. +In the following example, the variable A is *live* from line 1 to line 7, where the last use of the variable A is found. As for the variable B, its liveness only spans from line 4 to line 5. When register assignment is performed during compilation, the compiler attempts to keep A in registers for all its -livespan. +lifespan. So, in our example, if A needs $NumReg_A$ registers to be stored, this means that $NumReg_A$ registers will be reserved for A across the loop, and thus the compiler needs to fit the variables used between line 1 and 7 in $N - NumReg_A$ registers, with $N$ being the total number of registers available. -![variable liveness simple example]({{ '/assets/images/portal/article-images/2025-09-02-intel-gpu/liveness_example_annotated.jpg' | relative_url }})
+![variable liveness simple example]({{ ' +/assets/images/portal/article-images/2025-09-02-intel-gpu/liveness_example_annotated.jpg' | relative_url }})
*Variable liveness simple example* It is therefore easy to understand that in such a kernel, if the variable A is large and the kernel processing between -lines 2 and 7 is also register consuming, the compiler may have hard time to allocate registers while avoiding register -spills. +lines 2 and 7 is also consuming register, the compiler may have a hard time to allocate registers while avoiding +register spills. This is exactly what happens in the widespread case of [FlashAttention version 2](https://arxiv.org/abs/2307.08691). @@ -252,16 +257,16 @@ def FlashAttention2_forward(Q, K, V): In the second version of the implementation of the FlashAttention model, the loop order has been reversed to promote data locality. As long as there is enough local memory (or registers) to contain all the needed data, this algorithm works fine and -provide significant performance improvement compared to FlashAttention v1 (in the paper, the authors mention 2x faster -for the Cutlass implementation and 1.3-1.5√ó faster in Triton on an Nvidia Amper GPU A100). +provides significant performance improvements compared to FlashAttention v1 (in the paper, the authors mention 2x faster +for the Cutlass implementation and 1.3-1.5× faster in Triton on an Nvidia Ampere GPU A100). Deployed on a GPU target, line 4-10 constitutes the computing kernel that is dispatched to a Thread Block/Work-Group ( i.e. a SM/XeCore). But as you can see, variable Q is loaded before the loop (line 4) and remains *live* across the loop. The long lifespan of variable Q is even more problematic in the causal variation of the FlashAttention implementation. The causal variation is defined in the paper as : -> One common use case of attention is in auto-regressive language modelling, where we -> need to apply a causal mask to the attention matrix S (i.e., any entry Sùëñ ùëó with ùëó > ùëñ is set to ‚àí‚àû). +> One common use case of attention is in auto-regressive language modeling, where we +> need to apply a causal mask to the attention matrix S (i.e., any entry S𝑖 𝑗 with 𝑗 > 𝑖 is set to −∞). The Triton implementation of FlashAttention v2 with causal mask is as follow: @@ -362,9 +367,9 @@ another for calculating the data to which the mask does need to be applied). Our point is that variable `q` that loads a chunk of the Q matrix is *live* for the instruction `tl.load(Q)` until the second loop where variable `q` is read for the last time. When the target GPU architecture is able to load its operands directly from SMEM or TMEM (as discussed in the previous -section), this is less of problem because these memories are larger than the register files. -But when the target GPU does not have this capability, such as Intel PVC and BMG GPUs, the variable has to reside in -registers. +section), this is less of a problem because these memories are larger than the register files. +But when the target GPU does not have this capability, such as Intel Ponte Veccio and Battlemage GPUs, the variable has +to reside in registers. Thus, registers should be dedicated to saving this variable all along the kernel execution. Consequently, if variable `q` is large, many registers will be reserved for saving this variable and the register allocator will be forced to spill registers. @@ -373,7 +378,7 @@ allocator will be forced to spill registers. The register allocator's difficulty in assigning registers without spilling comes from the fact that some variables are *live* for a long period of time, which reduces the number of registers available for other variables. -Even worse, the high pressure on registers can prevent the compiler for applying other optimisations such as loop +Even worse, the high pressure on registers can prevent the compiler from applying other optimisations such as loop unrolling (which requires additional registers). As a consequence, to reduce the liveness of variables, when possible, relaxes the constraints on register allocations @@ -385,9 +390,10 @@ In the [XPU backend](https://github.com/intel/intel-xpu-backend-for-triton/tree/ added an [optimization pass](https://github.com/intel/intel-xpu-backend-for-triton/blob/main/third_party/intel/lib/TritonIntelGPUTransforms/ReduceVariableLiveness.cpp) which aims to reduce variable liveness where possible. -To this ends, the pass attempts to bring load operations closer to the actual uses of the loaded data. +To this end, the pass attempts to bring load operations closer to the actual uses of the loaded data. -![Reduce Variable Liveness pass diagram]({{ '/assets/images/portal/article-images/2025-09-02-intel-gpu/liveness-pass-diagram.jpg' | relative_url }})
+![Reduce Variable Liveness pass diagram]({{ ' +/assets/images/portal/article-images/2025-09-02-intel-gpu/liveness-pass-diagram.jpg' | relative_url }})
*Reduce Variable Liveness pass diagram* The diagram above shows how the compiler pass works to reduce the liveness of `DotOp` operands. @@ -397,11 +403,11 @@ This analysis computes the liveness of each variable in the source code. This wi variables at a specific point in the source code. In the second step, the pass looks for `DotOp` (i.e., the matrix multiplication operation) in a `For` loop. Currently, -the pass only considers `DotOp` in `For` loops as achor because it is a resource-consuming operation that is critical to -the performance of AI kernels. -But the pass can be extend to other cases in the future. +the pass only considers `DotOp` in `For` loops as anchor because it is a resource-consuming operation that is critical +to the performance of AI kernels. +But the pass can be extended to other cases in the future. -The third steps is to retrieve the `loadOp` that loads the `DotOp` operands. +The third step is to retrieve the `loadOp` that loads the `DotOp` operands. In brief, the pass rolls back the def-chain of the `DotOp` operands and returns the `loadOp` when it is found. Next, the pass checks if the `loadOp` is eligible to be moved. To be a candidate, a few conditions must be met: @@ -412,7 +418,7 @@ Next, the pass checks if the `loadOp` is eligible to be moved. To be a candidate - Empirically, we observe that loading a large amount of data is not the only criteria to determine if moving the `loadOp` is needed. The Triton language defines kernel at block/work-group level, but this group is then handled by multiple warps/sub-group (and threads/work-items). This sub-division into warps has an impact on the way data is - loaded and on the register assignment policy. As a result, the amount of loaded data must be large, but also, the + loaded and on the register assignment policy. As a result, the amount of loaded data must be large, but, also, the shape of the data on the dimension that is not split between warps must be large enough for the proposed optimisation to be relevant. - The `loadOp` must be outside the loop body, i.e., the operand must be a *live-in* variable of the loop. @@ -422,28 +428,30 @@ Otherwise, if these conditions are met for at least one of the operands, the pas care of moving the `loadOp`. If the `loadOp` has only one user (i.e., the `DotOp`), the load operation is sinked into the loop and a prefetch operation (`prefetchOp`) is inserted where the `loadOp` was initially located, as shown on the diagram. -The prefetch operation fetches the data from global memory into the cache. As a result, when the actual load take place, -the data is loaded from the cache and not from the global memory. +The prefetch operation fetches the data from global memory into the cache. As a result, when the actual load takes +place, the data is loaded from the cache and not from the global memory. The case where the `loadOp` has more than one user, is a little more complex as the `loadOp` cannot simply be sunk into -the loop, but the pass must ensure that a load operation take place before accessing the data. +the loop, but the pass must ensure that a load operation takes place before accessing the data. As loading data from the cache is not expensive, we chose to add another `loadOp` for subsequent uses. Hence, the liveness of the tensor is still reduced and the low-level compiler (*igc* for intel target) is able to perform its optimisations with less constraints on registers. #### Performance improvement -We have evaluated the performance of Triton FlashAttention v2 on Intel PVC GPU. +We have evaluated the performance of Triton FlashAttention v2 on Intel GPU Max PVC. The following plots show the normalised performance of the FlashAttention kernel with the *reduce-liveness-pass* enabled for different input configurations. -![Normalized performance PVC1100]({{ '/assets/images/portal/article-images/2025-09-02-intel-gpu/pvc1100_new.png' | relative_url }})
+![Normalized performance PVC1100](./illustrations/pvc1100_new.png)
*FlashAttention v2 Normalized performance PVC1100* -![Normalized performance PVC1550]({{ '/assets/images/portal/article-images/2025-09-02-intel-gpu/pvc1550_new.png' | relative_url }})
+![Normalized performance PVC1550](./illustrations/pvc1550_new.png)
*FlashAttention v2 Normalized performance PVC1550* +The testbed used for these evaluations and a disclaimer can be found [at the bottom](#disclaimer) of this blog post. + We can see that the pass has improved the performance for several configurations on all the targets evaluated by more -than 5%, and up to 10% for a some of them. +than 5%, and up to 10% for some of them. As expected the inputs configurations impacted by this optimisation are those with : - causal mask applied to the computed data @@ -463,16 +471,15 @@ We also exposed that low-level compilers may lack some knowledge about the sourc generating optimised code. In our use case, the *igc* compiler tried to assign registers based on the lifespan of the variables without knowing that the lifespan of some variables could be reduced to avoid register spilling. -Consequently, we take advantage of the progressive lowering of MLIR-based compiler to add an optimisation pass to the +Consequently, we take advantage of the progressive lowering of MLIR-based compilers to add an optimisation pass to the intel XPU backend of the Triton compiler. This pass aims to reduce the liveness of variables under certain conditions. -As a result, the performance of FlashAttention on PVC GPUs has been improved by up to 10% for certain input -configurations. +As a result, the performance of FlashAttention could be improved by up to 10% fo certain input configurations. A limitation of this pass is that we assume the data to be loaded is available in the L1 cache, so the load operations are cheap and can be easily moved around in the code. However, this might not be the case if cache conflicts occurred and the data was evicted from the cache before being -loaded into registers. This is likely to happen for GPU with small cache. If this happen, the load operation becomes +loaded into registers. This is likely to happen for GPUs with small cache. If this happens, the load operation becomes expensive and sinking a `loadOp` inside the loop body is far from a good idea. A future extension of the pass could therefore consider first loading the data into the SMEM, which is explicitly managed by the software, and then loading the data from the SMEM into registers instead of relying on the cache.