<figure><img src="image_1.png" /></figure>

HAL
open science

# Advances in Semantic Patching for HPC-oriented Refactorings with Coccinelle

Michele Martone, Julia Lawall

<figure><img src="image_2.png" /></figure>

## To cite this version:

Michele Martone, Julia Lawall. Advances in Semantic Patching for HPC-oriented Refactorings with Coccinelle. 30th International Workshop on High-Level Parallel Programming Models and Supportive Environments, Jun 2025, Milan, Italy. hal-05337107

HAL Id: hal-05337107

https://inria.hal.science/hal-05337107v1

Submitted on 29 Oct 2025

HAL is a multi-disciplinary open access
archive for the deposit and dissemination of sci-
entific research documents, whether they are pub-
lished or not. The documents may come from
teaching and research institutions in France or
abroad, or from public or private research centers.

L'archive ouverte pluridisciplinaire HAL, est
destinée au dépôt et à la diffusion de documents
scientifiques de niveau recherche, publiés ou non,
émanant des établissements d'enseignement et de
recherche français ou étrangers, des laboratoires
publics ou privés.

<figure><img src="image_3.png" /></figure>

Distributed under a Creative Commons Attribution 4.0 International License
# Advances in Semantic Patching for HPC-oriented Refactorings with Coccinelle

Michele MARTONE* and Julia LAWALL†
[0000-0003-3239-8554] [0000-0002-1684-1264]

* Leibniz Supercomputing Centre, Garching bei Muenchen, Germany Email: michele.martone@lrz.de

† Inria, Paris, France Email: julia.lawall@inria.fr

**Abstract**—Currently, the most energy-efficient hardware platforms for large-scale floating point-intensive calculations (also known as High Performance Computing, or HPC) are graphical processing units (GPUs). However, porting existing scientific codes to GPUs can be far from trivial. This article summarizes our recent advances in enabling machine-assisted, HPC-oriented refactorings with reference to existing APIs and programming idioms available in C and C++. The tool we are extending and using for the purpose is called Coccinelle. An important workflow we aim to support is that of writing and maintaining tersely written application code, while deferring circumstantial, ad-hoc, performance-related changes to specific, separate rules called “semantic patches”. GPUs currently offer very limited debugging facilities. The approach we are developing aims at preserving intelligibility, longevity, and relatedly, debuggability of existing code on CPUs, while at the same time enabling HPC-oriented code evolutions such as introducing support for GPUs, in a scriptable and possibly parametric manner. This article sketches a number of self-contained use cases, including further HPC-oriented cases which are independent from GPUs.

**Index Terms**—semantic patching, transpiler, source-to-source translation, HPC, refactoring, OPENMP, C/C++, SMPL, AoS, SoA, GPU, *diff*, *patch*, PYTHON, OCAML, HIP, CUDA, STL.

## I. INTRODUCTION

In today's computing environment, the best hardware on which to execute demanding HPC applications, in terms of both execution time and energy costs, is often the GPU, rather than the CPU. Nevertheless, much HPC software was written in the CPU era. Continuing to use these valuable applications effectively thus requires transforming their source code to use GPU libraries. Such transformations must be done systematically. The large size and complexity of the codebases involved suggest that it is desirable to do as much as possible of these transformations automatically.

In the Linux kernel developer community, the tool Coccinelle<fnref n="1" /> is regularly used to perform large scale transformations to the Linux kernel code base. Coccinelle targets transformation of C code and has been designed around the notion of a semantic patch. A semantic patch is a change specification written in terms of lines to add and remove, like a traditional patch, but is applied taking into account the type system and the control flow of the target programming language. Beyond the Linux kernel, Coccinelle is also used by developers of some other systems software such as SYSTEMD<fnref n="2" /> and ZEPHYR

OS<fnref n="3" />. Recently [1], Coccinelle has also been applied to the cosmological software GADGET, with the goal of the targeted transformation of specific data from the array of structures (AoS) to the structure of arrays (SoA) representation.

This article is motivated by a new training course [2] on Coccinelle that has been offered at the Leibniz Supercomputing Centre (LRZ), to facilitate the use of Coccinelle by HPC programmers. This training contains new introductory didactic material relevant to HPC, such as operations on floating-point numeric arrays and loop constructs with parallel APIs, and includes an overview of possible applications. The present contribution does not have the pretense of exhaustively explaining Coccinelle's transformation specification language Semantic Patch Language (SMPL), but will comment on a selection of relatively simple and intuitive scripts enacting transformations of high interest to HPC practitioners. Whatever the transformation we show, the main target application context is that of automation of large-scale changes, that is, changes that would otherwise be prohibitively costly to perform directly by hand, or too error-prone to be desirably implemented via an ad-hoc tool written in a general-purpose programming language.

The present article is structured as follows. Sec. II introduces Coccinelle's context and past use cases, the concept of refactoring and what other tools or programming techniques target similar goals. Then, Sec. III proceeds with a bird's-eye view of eleven use cases that have been enabled by recent Coccinelle features. Sec. IV then offers a high level discussion about the approach. Sec. V concludes.

## II. LITERATURE OVERVIEW

We now contextualize Coccinelle and the present work within related literature and techniques.

A Semantic patch, written in the Semantic Patch Language (SMPL), as introduced by Coccinelle [3], refers to a code-change specification that has a form similar to a POSIX *unified code patch*, indicating removed and added lines of code. Instead of matching text only, as for a POSIX unified code patch, a semantic patch is matched against a program's abstract syntax tree (AST) and control-flow graph (CFG). Furthermore, change specifications are abstracted by *metavariables* making it possible for subterms to match arbitrary subterms such as types, expressions, and statements. The use

<footnote n="1">¹https://github.com/coccinelle/coccinelle</footnote>

<footnote n="2">²https://github.com/systemd/systemd/tree/main/coccinelle</footnote>

<footnote n="3">³https://github.com/zephyrproject-rtos/zephyr/tree/main/scripts/coccinelle</footnote>
of VARIABLES makes a semantic patch generic, allowing a
single change specification to be applied across a code base.
Matching against the AST, rather than the source code text as
done for a POSIX patch, ensures that the language syntax is
respected, including matching complete variable names, rather
than substrings, and respecting the relative precedence of infix
binary operators. Finally, respecting the CFG makes it possible
for the matching of a semantic patch to follow the control flow
of execution, for example around a loop, thus respecting the
execution semantics.

Coccinelle was created to provide the LINUX kernel devel-
opers with a means to succinctly describe transformation rules
and apply them. Exchange of POSIX patches is a common
working practice within the LINUX kernel community, and, by
analogy, Coccinelle has become the de-facto tool to express
complex refactorings within the LINUX kernel community, by
means of files containing semantic patches.

Since the 1990s, the term "Refactoring" [4] has been used
to denote semantics-preserving stepwise code manipulations
that introduce or enforce certain desirable qualities (e.g.,
absence of *code smells*) that would otherwise stem from
best practices of design and implementation. Existing IDEs
(Integrated Development Environments) often support a fixed
number of specific *refactorings* that are expected to be of
general interest. Many of these are quite simple. However,
as they are implemented using a general-purpose language,
such as C++, they can be arbitrarily complex. Nevertheless,
as the set of refactorings offered by an IDE is fixed, it cannot
be easily adapted to the specific properties of a software
project. Furthermore, some desirable transformations may not
precisely preserve the semantics, putting them outside of the
scope of refactorings. For example, a change introducing
parallelism may affect the result due to changing the order of
arithmetic operations; such changes can be tolerable but are
not always desirable. To counter such situations, Coccinelle
allows fine-grained control of the set of code fragments that
are to be transformed.

The C *preprocessor* (Sec. 6.10 in [5]) enables several tech-
niques for selective compilation, code selection, and avoiding
redundancy. That works by introducing macros to abstract
code from semantics (e.g., masking loops within macro def-
initions), and can involve a significant cost in clarity and an
increased risk of bugs, especially when nesting several levels
of macros. For these reasons, despite supporting some changes
in code semantics equivalent to the *HPC-oriented refactorings*
we are interested in, preprocessor tricks can quickly reach their
limits, and lead to problems. Extensive usage of preprocessor
macros is regarded as dangerous and is therefore discouraged.

The LLVM project offers a templated C++-based API that
allows direct and low-level matching and manipulation of the
AST [6], and output in the original language (C or C++). This
allows matching and modifying potentially vast swathes of
code by following an imperative programming model. Such
a framework can be used to implement refactorings. This
solution requires matching constructs at the exact *grammar*
*production* level, which in many cases may be unnecessarily

picky, in that it requires writing *refactoring programs* that are
more verbose (and that require more effort) than is desirable.
Consequently, it is mostly authors of IDEs or underlying tools
such as *clang-tidy* who use this API to prepare catalogues
of refactorings.

The ROSE compiler infrastructure [7][8] has a broad scope
of applications that includes HPC-oriented source-to-source
translation, and has been created for the needs of the United
States of America's Department of Energy. Differently than
SMPL for Coccinelle, and similarly to LLVM, ROSE offers a
C++ interface to express refactorings.

Most recently, machine learning-backed *code assistant chatbot* services started to be offered by a few major corporations (e.g., GITHUB COPILOT by Microsoft, CHATGPT or CODEGPT by OPENAI, etc.). Such chatbots are operated by the user by means of natural language in written form. They can generate code or manipulate existing code in a way requested by the user. Their inner mechanisms (source code and training datasets) are often somewhat opaque, just as *explainability* of machine learning techniques remains a challenge. These approaches are substantially different from the one proposed here, in that they offer full automation and little user control. Our approach does not involve any form of machine learning.

Several HPC libraries offer abstractions by creating a sep-
aration layer between an actual data structure (and where it
resides) and the syntax used to access it. This is one of the
features of e.g., KOKKOS [9], where the user can make a
compile-time choice between *Structure of Arrays (SoA)*-like
and *Array of Structures (AoS)*-like accesses with a unified
syntax. Nevertheless, it is likely that most numeric expressions
will require a rewrite in order to use KOKKOS array access
notation.

As an alternative to large-scale expression rewriting, modern C++ allows introduction of a number of classes equipped with *inline accessor member functions*, and custom *operator overloads* in a way that AoS access syntax can be fully preserved across the code, and SoA accesses occur under the hood. Provided compiler optimization is active (so as to cancel the underlying overhead), such techniques ease achieving vectorizable arrays access, and restrict code changes mostly to data structure definitions. Compiling the sources without any optimization and with debug flags enabled will mostly likely retain object code at each inline function call, in places where optimization flags would have reduced it to a mere memory location access. These consequences of the inherently more complicated logic underneath can slow down the debuggable code to a degree where reaching the point of a crash can be problematic.

In our past work [1], we described a use case where we
performed an AoS → SoA change explicitly on the entire
GADGET cosmological code [10]. That intervention was rec-
ommended by colleagues (see their pilot study in [11]) in order
to achieve a higher rate of compiler auto-vectorization [12,
Ch. 8]. In addition to the data structure definitions (a mere
few hundred lines of code, which we could have changed
by hand), the rules we created patched many tens of array- accessing expressions within each of thousands of loops. We left the original code to be straight C, and performed our changes by means of a collection of Coccinelle rules. If using no optimization flags and/or using debug flags there is no penalty, because there is no syntax abstraction obscuring the semantics.

In this example, Coccinelle makes it possible to provide the domain scientists with transformation rules that allow them to continue developing the original AoS code, which can be a bit clearer to understand than SoA code. A refactored SoA copy of the code can be obtained by invoking Coccinelle on it. In such a workflow, one may want to use version control for the semantic patches only, which are much terser and less redundant than the transformed code. The approach also allows easily fine-grained control of where the transformation is applied: specified quantities can be kept in AoS form if this is desired for modularization or organizational reasons. This style of working could allow developing multiple concurrent experimental derivations out of a codebase, without the burden of having to maintain several branches.

### III. ENABLED HPC REFACTORINGS

Coccinelle has been recently enhanced with features that allow expressing HPC-oriented refactorings that were previously not expressible. Here is an overview.

**Interfacing with an instrumentation API:** Several tools for instrumentation (e.g., LIKWID,<fnref n="4" /> SCORE-P,<fnref n="5" /> CALIPER.<fnref n="6" /> etc.) offer an API (also referred as a *marker API*) to collect code performance metrics at runtime. Introduction and removal of instrumentation syntax is one of the simplest possible tasks that we present. To achieve that, one would specify an SMPL pattern describing the code that should be measured by the marker API, for instance by enclosing it with calls to start or stop measurement collection. As an example, the following semantic patch introduces calls to the LIKWID macros.

```c
1 @@ @@
2 #include <omp.h>
3 + #include <likwid-marker.h>
4
5 @@ @@
6 #pragma omp ...
7 {
8  + LIKWIDze Marker StArT (______func______);
9  ...
10 + LIKWIDze Marker StArT (______func______);
11 }
```

This semantic patch uses two rules. The first one introduces the LIKWID header just after the `#include <omp.h>` line. The second rule locates (any) `#pragma omp` line followed by a block within braces, and inserts instrumentation to start/stop performance collection, passing the compiler-provided `__func__` string as a parameter specifying the current function name as a computation phase label. In practice,

one could refine the pattern to be more selective in choosing such code locations. For instance, restricting to specific numeric kernels, or around the specific areas of the code that one is concentrating on at a given time. The parametrical nature of the match specification allows flexibility and experimentation in deciding which areas of the code to instrument, perhaps only transitorily.

**OPENMP's declare variant:** Since version 5, the specification document of OPENMP [13] introduces a notation for the definition and automatic selection of hardware-optimized functions, referred to as *variants*. A *variant* is only compiled if its specific *instruction set architecture* (ISA) is selected at compile time. If not, only the source code corresponding to the base version is compiled. This notation supports variants via one `#pragma omp declare variant ...` line per variant just before the base function definition. Each such a clause specifies which function is to be treated as a *variant*. We illustrate here an approach to introducing variants on a large scale with the help of Coccinelle.

Let us assume that we have a source file with an arbitrary number of functions that we wish to diversify in this way. For each candidate function (in the example here, we select those whose name matches the regular expression `"kernel"`), two clones are created (their names being derived from the base function name) and declared as variants by means of `#pragma omp declare variant` clauses just above the base function. The function clones can be created easily by matching function definitions with the SMPL code `T f (PL) { SL }` (line 13). Such an arrangement matches a previously defined `type T`, a function name (`identifier`) `f`, and its `parameter list PL`. These components are then used to insert two function clones (lines 9 and 10 of the listing, prefixed by `+`). The clones' names are produced via `fresh identifier` (lines 6 and 7), extending (via the `#` operator) the original function name (as matched by `f`). The names of the clones of each original function are constructed by prefixing respectively `"avx512_"` and `"avx10_"` to `f`, and the `match` clause of each clone indicates to the compiler under which target architecture it shall be used.

```c
1 @@
2 type T;
3 identifier f = "kernel";
4 parameter list PL;
5 statement list SL;
6 fresh identifier f512 = "avx512_"##f;
7 fresh identifier f10 = "avx10_"##f;
8 @@
9 + T f512 (PL) { SL }
10 + T f10 (PL) { SL }
11 +#pragma omp declare variant(v512_f) match(
12   device={isa{"core-avx512"}})
13 +#pragma omp declare variant(v10_f) match(
14   device={isa{"core-avx10"}})
15 T f (PL) { SL }
```

At this stage, the two function clones corresponding to `f512` and `f10` are identical to the base function identified

<footnote n="4">⁴https://github.com/RRZE-HPC/likwid</footnote>

<footnote n="5">⁵https://score-p.org/</footnote>

<footnote n="6">⁶https://github.com/LLNL/Caliper</footnote>
by f. To obtain the hardware-specific variants we seek, we would still have to write a few extra rules that enact specific transformations on them. Such rules could reside in this semantic patch and *inherit* the f512 and f10 metavariables, or be completely independent and reside in a different semantic patch file, only to be applied in a separate step, exploiting the pragma and the specific function name (which suggests it is a clone) to use an architecture-specific rule.

In either case, modifying the clones would involve rules introducing changes that are optimal for the respective architectures (here we alluded to either AVX-10 or AVX-512). Such rules would modify statements; e.g., introduce calls to *intrinsics* or arrange for specific properties or types in variable definitions. Such rules must necessarily be very program-specific, and their discussion does not belong here.

**Function cloning and introduction of attributes for function multiversioning:** Independent of OpenMP, the GCC and CLANG compilers offer a C/C++ extension to create a form of function variant, called *function multiversioning*, by means of specific *attributes*. The first one (*target_clone*ses) allows implicit creation of processor-specific copies of a function [14], as well as automatic selection of these at runtime. The second one (*target*) requires multiple function definitions, each of which can also contain several custom compilation options. These extensions ease the production of binary libraries with object code optimized for different processors within the same base architecture family. Here in this example, the base function is being prefixed by `__attribute__((target("default")))` and variations of this function are marked as such, e.g., `__attribute__((target("avx512")))`. To automate creation of such function clones with Coccinelle we proceed as we have shown for `#pragma omp declare variant` (previous listing). In a second step (following listing) we match on the attribute expression containing `"avx512"`; the rules which would follow this match (not shown here) would have to to introduce specific specializations (e.g. usage of intrinsics).

1. `@@`
2. `identifier f;`
3. `type T;`
4. `@@`
5. `__attribute__((target(...,"avx512",...)))`
6. `T f(...)`
7. `{`
8. `+ // add and modify avx512-specific code only`
9. `...`
10. `}`

If matching is needed on compiler-specific attributes that are not prefixed by `__``, Coccinelle would require them to be specified as such via *attribute name*.

**Bloat and clone removal:** Given a project with a long development history, we may be interested in removing obsolete code. Taking the last two examples as a reference, we may want to delete *function specializations* based on their occurrence in the code. Suppose we have been using the `__attribute__`-based technique sketched in the last

paragraph, and we wish to remove code specializations corresponding to the AVX512 ISA extension and its predecessor AVX2. Then a semantic patch for `cleanup` could proceed as follows.

1. `@c@`
2. `type T;`
3. `function f;`
4. `parameter list PL;`
5. `@@`
6. `__attribute__((target(`
7. `{`
8. `"avx512"`
9. `|`
10. `"avx2"`
11. `)}`
12. `)))`
13. `T f(PL) { ... }`
14. `@d@`
15. `type c.T;`
16. `function c.f;`
17. `parameter list c.PL;`
18. `@@`
19. `__attribute__((target("default")))`
20. `T f(PL) { ... }`

Two logical steps are needed, corresponding to two rules; here, `c` and `d`. Rule `c` removes all functions that are specific to one of the two mentioned instruction sets, identifiable by being prefixed by a `__attribute__((target.... Rule `d` merely removes the `__attribute__((target( "default")))` specification, which precedes the unspecialized function definition, which is not deleted (in practice one would want to remove this attribute only once no specialization is there anymore). Rule `c` consists of a *disjunction* matching the different attribute values with branches enclosed by parentheses and separated by `|` (all in the first column, to avoid confusion with the C/C++ bit-or syntax). Rule `d` reuses entities inherited from rule `c`, via the *metavariables* `c.T`, `c.f`, and `c.PL`. Many bloat removal interventions are possible. Of the most pervasive and systematic ones, we imagine the location and removal of code associated with specific attributes or compiler-specific pragmas.

**Removal of explicit loop unrolling:** We might obtain a codebase that we are unfamiliar with, containing thousands of lines of code generated by a script, with loops unrolled, but no access to the original generator. Removal of many such explicit loops (or conceptually similar transformations) can be tricky, so here we show how we may approach this with Coccinelle. Let us assume that we want to target loops unrolled four times. More precisely, loops containing the same variation of a statement repeated four times, and each time with a different indexing of a specific subexpression, without local declarations, and most importantly, with the statement *unknown*. The first semantic patch we show here consists of a rule named `p0` that locates four statements `(A, B, C, D)` in a loop (lines 14–22) with a specific loop header (lines 8–13) and checks that, given a loop variable *matched* by *metavariable* `i`, these *contain* respectively `i+0`, `i+1`, ... (lines 14, 16, 18, 20). Such a rule can match many simply unrolled loops, and with
this arrangement of statements, removing B, C, and D removes 37 } any unrolling.

1 @p0@
2 type T;
3 identifier i,1;
4 constant k={4};
5 statement A,B,C,D;
6 @@
7 + #pragma omp unroll partial (4)
8 for (T i=0; i
9         +k-1
10        < 1 ;
11 -         i+=k
12 +         ++i
13     )
14 {
15   \( A \& i+0 \) \(
16 - B \& i+1
17   \) \(
18 - C \& i+2
19   \) \(
20 - D \& i+3
21   \)
22 }

In `p0` we introduce the `#pragma omp unroll` partial clause, available since `OPENMP 5.1`, which requests that the compiler enact loop unrolling, avoiding the need to modify the source code.

1 @p1@
2 type T;
3 identifier i,1;
4 constant k={4};
5 statement A,B,C,D;
6 @@
7 for (T i=0; i+k-1 < l; i+=k)
8 {
9   \( A \& i+0 \) \( B \&
10 - i+1
11 + i+0
12   \) \( C \&
13 - i+2
14 + i+0
15   \) \( D \&
16 - i+3
17 + i+0
18   \)
19 }
20
21 @r1@
22 type T;
23 identifier i,1;
24 constant k={4};
25 statement p1.A;
26 @@
27 + #pragma omp unroll partial (4)
28 for (T i=0; i
29         +k-1
30        < 1 ;
31 -         i+=k
32 +         ++i
33     )
34 {
35   A
36 - A A A

In absence of further constraints, any sequence of four statements respectively accessing `i+0`, `i+1`... may be matched by `p0`, even if the sequence does not represent an unrolled loop, i.e., the statements are not identical modulo `i+0`, `i+1`, etc. In certain codes this matching ambiguity may constitute a problem; for these cases rules `p1` and `r1` improve over `p0`.

Just like rule `p0`, `p1` may also match more than we are interested into, because there is no guarantee that the statements are equal after substitution of the `i+0`... subexpressions. But the construction of `p1` is different: after locating four statements within a loop, on lines (10–11, 13–14, 16–17), it replaces the occurrences of `i+0`, `i+1`... by `i+0` with the goal of making the four statements identical. The second rule `r1` will only match if all the substitutions of `p1` are enacted and if they lead to a repetition of the first statement `A`. That can only occur if the statements differed by the indexing expression only. In that case, `p1` will delete the last three statements `A`, and only keep the first one (the one on `i+0`).

If `p1` patches and transforms the code, but the resulting statements are not identical, then `r1` will not match and the loop will retain the four modified statements in the loop body, resulting in code that is incorrect. To address this issue, we could introduce a third rule that undoes the transformations of `p1` when `r1` is not applied. Once again, we showed a transformation that requires some knowledge of the code, namely on its conventions (whether variable declarations are in the unrolled loop or outside, etc.). The code may also need some preparatory modifications, enforcing certain conventions, before large-scale changes can be performed.

**Advanced expression modification (e.g., `mdspan`):** Typically, a data layout change in a number-crunching codebase involves updating a large number of expressions accessing numeric arrays. The following rule `tomultiindex` replaces triple nested square bracket expressions by C++23's triple index expressions.

1 # spatch --c++=23
2 @tomultiindex@
3 symbol a;
4 expression x,y,z;
5 @@
6 - a[x][y][z]
7 + a[x, y, z]

This rule focuses on an array named `a`, but does not introduce any requirement on the indexing expressions `x`, `y`, and `z`, which can be arbitrarily complex. Modifications of this sort are the most pervasive – they usually aim at the entire code – and probably are also of most interest to the readers of this article. Of course, in order to be robust in production, the name of the redeclared arrays (here a hardcoded `a`) should probably follow a match in a global declaration, or perhaps a function-level parameter definition.

**Translation of very similar APIs:** Some APIs may be so similar to each other, that their mutual translation would
mostly consist of token-to-token correspondence among two
enumerable sets. This is the case of NVIDIA's CUDA [15] and
AMD's HIP [16], for which specific translation tools exists
(e.g., hipify-perl, based on PERL, and hipify-clang,
based on LLVM). This sort of translation is also easy by
means of Coccinelle.

1 @initialize:python@ @@
2 C2HF = { "curandUniform_double":
3 "rocrandUniform_double" }
4
5 @cfe@
6 identifier fn;
7 expression list el;
8 position p;
9 @@
10 fn@p(el)
11
12 @script:python cf2hf@
13 fn << cfe.fn;
14 nf;
15 @@
16 coccinelle.nf =
17 coci.make_ident(C2HF[fn]);
18
19 @hfe@
20 identifier cfe.fn;
21 identifier cf2hf.nf;
22 position cfe.p;
23 @@
24 - fn@p
25 + nf
26 (...)

A Coccinelle toy semantic patch for API translation, as
illustrated above, could use a PYTHON rule<fnref n="7" /> to declare a
dictionary (defined in the first, anonymous, special rule starting
on line 1) to store the mappings of functions (this is exactly
how hipify-perl does it, albeit without using an AST).

1 @initialize:python@ @@
2 C2HT = { "__half": "rocblashalf" }
3 @cte@
4 type c_t;
5 identifier i;
6 @@
7 c_t i;
8
9 @script:python ct2hf@
10 c_t << cte.c_t;
11 h_t;
12 @@
13 coccinelle.h_t = \
14 coci.make_type(C2HT[c_t]);
15
16 @hte@
17 type ct2hf.h_t;
18 type cte.c_t;
19 identifier cte.i;
20 @@
21 - c_t i;
22 + h_t i;

To translate the types as well, we proceed similarly, using
a rule `cte` to identify functions; a rule `ct2hf` to use the
dictionary; and a rule `hte` to enact the change.

For simplicity, here we translate only one type and one
function. Rule `cfe` identifies CUDA function calls; rule
`cf2hf` uses the dictionary with the function translations to
HIP; finally, rule `hfe` enacts the change.

A complete version of these function- and type-level semantic patches would need to have the entire list of functions and types involved in the two APIs. That would be the same as the approach of hipify-perl, but with the advantage of enactment of the changes at the AST level.

Among the further requirements of CUDA to HIP translation (which Coccinelle fullfills) is the ability to recognize the *triple chevron* syntax (<<<>>>) used by CUDA to launch GPU kernels, and replace it with corresponding HIP library function calls, as HIP does not use the *triple chevron*.

1 #sparc --c++
2 @@
3 identifier k;
4 expression b,t,x,y;
5 expression list el;
6 @@
7 - k<<<b,t,x,y>>> (el)
8 + hipLaunchKernelGGL(k,b,t,x,y,el)

The rule above translates the invocation of a user-specified
CUDA kernel (k) into HIP notation, which passes it to
hipLaunchKernelGGL instead, along with all the existing
arguments. Notice that the rule is quite general, as it applies
at the expression level and not at a statement level (no per-
line replacement). Translation to other hip-prefixed kernel
functions is to be addressed with trivial separate rules.

**Translation of directive-based APIs:** For very simple programs, translation of e.g., OPENACC [17] into OPENMP may proceed on a line-by-line basis (a #pragma directive occupies at least one line at a time) and without depending on the context.

1 @moa@
2 pragmainfo pi;
3 @@
4 #pragma acc pi
5
6 @script:python o2o@
7 pi << moa.pi;
8 po;
9 @@
10 // Here we could have a small parser and
 translator using pi, but for simplicity we
 are just returning a hardcoded clause
11 coccinelle-po =
12 coci.make_pragmainfo
13 ("kernels copy(a)");
14
15 @@
16 pragmainfo moa.pi;
17 pragmainfo o2o-po;
18 @@
19 - #pragma acc pi
20 + #pragma omp po

<footnote n="7">⁷Coccinelle also provides an interface to OCAML.</footnote>
Moreover, the majority of projects adhere to using a specific subset of such directive-based APIs. So an approach of translation on a line basis may apply to a large number of cases. A skeleton semantic patch for such an approach is shown in the listing above. There, the middle rule (`o2o`) invokes `PYTHON` code. Please notice that in this example no variable initialization by means of `initialize:python` is needed. Such a `PYTHON` rule could invoke a line-oriented parser-based translator implemented in place or in a separate `PYTHON` module. In extracting the pragma lines, Coccinelle is liberal in accepting whitespaces, and it does not break on line continuations, as an ad-hoc line-oriented script may do, so the parser would receive correct input. In addition to the individual `#pragma` lines, more accurate translations would also require analyzing the context: that is also loops and local declarations. We are investigating such a translation-based approach for `OPENACC-to-OPENMP` in a student project. More complex uses of `PYTHON` would require stateful rule scripts, marked as such via the `initialize` and `finalize` keywords.

The logic of translation could be not much different than that followed by the script provided by INTEL<fnref n="8" /> to assist with similar tasks, although the script provided by INTEL lacks a proper parser or AST representation.

**Ease introduction of modern C++ STL constructs:**
C++ is a huge language and it allows vastly different styles of programming. For consistency reasons, most projects restrict the programmer to use a subset of its features. Recently, the C++ committee has also been involved in suggesting programming guidelines [18]. These may help avoid so-called *code smells*, keeping code malleable and less error-prone. The following example replaces a so-called *raw loop* with an STL-backed one (`std::find`). In this loop, apart from setting a flag (`result = true`), we could only perhaps expect some diagnostics (line 12), which in this example we delete.

```cpp
1 #spatch --c++=17
2 @r1@
3 type T;
4 constant k;
5 identifier elem result, arrid;
6 @@
7 - bool result = false;
8 ... 
9 - for (T &elem : arrid )
10 -   if (\( elem == k \| k == elem \) )
11 -   {
12 -     ...
13 -     result = true;
14 -     break;
15 -   }
16 + const bool result =
17 +   (find( begin (arrid), end (arrid), k ) !=
18 +     end (arrid));
19
20 @ah depends on r1@
21 @@
22 #include <iostream>
23 #include <algorithm>
```

```cpp
24 #include <functional>
```

While not being particularly HPC-specific, this example shows that if there is sufficient regularity in our code, we may match specific recurring code portions and replace them by function calls (`find`). That is exactly what HPC-oriented C++ APIs usually require. The semantic patch we have shown also illustrates how intra-function declarations and constructs can be matched and manipulated. Notice how the `depends on` clause on the second rule `ah` makes it apply conditionally on the first rule `r1`.

**Introduction of APIs enclosing lambdas:** Certain APIs (e.g.: `KOKKOS` [9], `RAJA` [19], ISO C++ standard parallelism [20], and `SYCL` [21]) require enclosing the numerical kernels within C++ `lambda functions`, thus introducing one abstraction layer masking the exact parallel execution model on the hardware, and allowing selection via e.g. definitions or template variables just before the lambda is instantiated.

```cpp
1 #spatch --c++
2 @r0@ @@
3 + #include <Kokkos_Core.hpp>
4 #include <cmath>
5
6 @r1@
7 statement fb, fc;
8 expression n;
9 identifier c = {i, j};
10 position p;
11 @@
12 (
13 fc@p
14 &
15 for (...; c<n; ...) fb
16 )
17
18 @script: python r2@
19 fb << r1.fb;
20 lb;
21 rp;
22 @@
23 coccinelle.lb =
24 "KOKKOS_LAMBDA(const int i)" + fb;
25 coccinelle.rp =
26 "RangePolicy<HostExecutionSpace>(0,n)";
27
28 @r3@
29 statement r1.fc;
30 position r1.p;
31 identifier r2.lb;
32 identifier r2.rp;
33 @@
34 (
35 fc@p
36 &
37 (
38 - for (...; ...; ...) { ... result += ...; }
39 + parallel_reduce rp, lb);
40 |
41 - for (...; ...; ...) { ... }
42 + parallel_for rp, lb);
43 )
44 )
```

<footnote n="8">⁸https://github.com/intel/intel-application-migration-tool-for-openacc-to-openmp/</footnote>
As of the most recent release of Coccinelle (version 1.3), lambda manipulations are not yet fully supported. We nevertheless illustrate a loophole in Coccinelle that can be used to accomplish certain transformations until the missing syntax becomes supported.

The above semantic patch illustrates this technique being applied to a specific exercise<fnref n="9" /> from the KOKKOS tutorial. Rule `r0` inserts the KOKKOS header next to a header we may already include. Rule `r1` locates `for` loops to be transformed into KOKKOS loops; after inspecting the source code associated to that exercise we established that the relevant loops are exactly those with index variables `i` and `j` (other loops exist, but we wish to keep them as they are).

In identifying the loops in this program we are content to use any `expression n` as loop bound. We also declare two `statement` metavariables; one (`fc`) for the entire loop construct, and the other (`fb`) for the loop body. Rule `r1` enacts no change: it only matches entities for reuse in the rule following it. Rule `r2`, which uses PYTHON scripting, makes an anonymous lambda function out of the entire block `fb` representing the body of the `for(...)` loop, and saves that in a PYTHON string. Rule `r3` matches different possible loop bodies to recognize the different KOKKOS parallelization constructs they belong to (normal parallelism vs. reduction). Rule `r3` exploits the aforementioned intentional loophole in Coccinelle, which allows treating any string as an `identifier` metavariable, and passes the string in calls to KOKKOS' `parallel_for` and `parallel_reduce`, in a context that would otherwise accept identifiers.

This *dirty hack* of passing statements through SMPL identifiers allows a limited manipulation of lambdas with Coccinelle. We plan to implement proper lambda manipulation support soon.

**Workarounds for occasional compiler bugs:** A few years ago<fnref n="10" />, failure of multiple checks was reported in the test suite of the LIBRSB library, which led to the suspicion of a bug there [22]. The failing involved double precision complex conjugate computations. A brief investigation showed that the bug was in the *vectorizer* of the GCC compiler,<fnref n="11" /> version 11.2 (and a few other ones after). Considering that the library was being packaged on several LINUX distributions, the bug needed a portable fix that would possibly only be activated when using the affected compiler versions. The decision of the LIBRSB author was to modify the build system (Makefiles & co.) to trigger a semantic patch-based transformation, conditional on the compiler version. This semantic patch applies only to the functions that have been identified as being affected by the compiler bug, and adds specific `#pragma` lines lowering the optimization level, so to get correct numerical results. A regular expression can be used to identify the affected

functions because the code uses a specific naming convention for the functions doing double precision complex conjugation.

```
@pragma_inject@
identifier i = "rsb_BCSR_spmv_sasa_double_complex_[CH]_t [NTC]_r1_cl_uu_s[HS]_dE_uG";
type T;
@@
+ #pragma GCC push_options
+ #pragma GCC optimize "-O3", "-fno-tree-loop-vectorize"
T i(...)
{
    ...
}
+ #pragma GCC pop_options
```

The semantic patch impacts a dozen functions among a few hundred; its attractiveness derives mostly from its inobtrusiveness and for being triggerable on demand. A similar approach could be used by other projects to implement *transitory changes* conditionally via *transitory semantic patches*.

## IV. DISCUSSION

**Acceptance obstacles:** Code owners or maintainers may resist large-scale modifications of the codebase, especially if such restructurings originate from a different team. The habit of writing comprehensive test suites in the scientific software communities is slowly but steadily progressing; this can surely facilitate reviewing a large “refactoring contribution” on the basis of the tests’ results, notwithstanding that even the best test suite cannot be considered complete. Apart from simple mistrust, other factors hindering acceptance can be aesthetic (with some relation to intelligibility) or doubts about the maintainability of the new code. We believe that the spread and adoption of sound *research software engineering* practices can reduce acceptance impediments while raising confidence in the code behaviour and its malleability.

**Replayable refactorings:** Coccinelle can enable developing restructurings over time by making it possible to keep an existing code separate from the change specifications. That is, as an existing code base can evolve, and semantic patches can be developed in parallel. This gives the ability to apply a refactoring on demand, and not having to maintain overly similar branches in parallel. Once the *refactoring specification* is thoroughly validated (apart from tests, that may include inspection and study), the authors may be willing to accept the resulting patched version. Note that the large-scale modifications may also impact the results of tests. For instance, parameters order may change, or involved variables’ types. In that case we would want to update the unit test source files accordingly as well.

**The future of C/C++:** Recent trends in programming languages and organizational security discourage the use of C/C++, in favour of intrinsically safer languages. Given the sheer size of the installed user base, any migration to a new safer programming language will take a long time. During that time, tooling for programming responsibly, to mitigate existing

<footnote n="9">⁹Exercises/01/Begin/exercise_1_begin.cpp from https://github.com/kokkos/kokkos-tutorials, revision fd4852c.</footnote>

<footnote n="10">¹⁰Private communication by the Debian-side maintainer to the first author, who is also author of LIBRSB.</footnote>

<footnote n="11">¹¹https://gcc.gnu.org/bugzilla/show_bug.cgi?id=103995</footnote>
risks or weaknesses, and perhaps migrating to other languages, will continue to be needed.

Even if exhibiting a lower urgency than in other categories of software (e.g., medical, avionics), HPC software deserves correctness and reliability. If such code is ever ported to or developed in RUST, the resulting code may still be able to benefit from the use of Coccinelle, as a RUST port of Coccinelle is under way.¹²

And while Coccinelle is not a tool for translating between languages, it is possible to conceive an iterative migration process consisting of partial rewrites of program modules or classes, and subsequent replacement of old constructs with new ones referring to new functions and classes in the different language. The viability of such an approach is clearly highly project-dependent.

Whatever the programming model or next best performance portability library of interest may be, our work aims at aiding otherwise unfeasible or work-intensive, large-scale code changes in a possibly reversible (or replayable) manner.

## V. CONCLUSIONS

We have succinctly described eleven use cases illustrating the construction of the Coccinelle semantic patch language which can be of interest to HPC practitioners.

We like to regard such cases as individual “components” in typically larger, more involved HPC refactorings. These refactorings are primarily aimed at separating the expertise of machine-specific experts from that of the domain-specific expert.

The techniques we have shown are more effective and productive when the codebase is *tidy* and regular. In itself, such properties are desirable also when an expert is refactoring directly with a code editor/IDE and basic text-editing tools.

A challenge in working with C++ is the large size and the continual, vigorous evolution of the language. Development of support for C++ code in Coccinelle is ongoing.

## VI. ACKNOWLEDGEMENTS

This work was partially funded by the German Federal Ministry of Education and Research (BMBF) and the Bavarian Ministry of State for Education, Culture, Science and Art (StMBW) within the project SIVEGCS. Further support by the French Ministry of Europe and Foreign Affairs (MEAE) and the StMBW through BAYFRANCE'23 is gratefully acknowledged.

## REFERENCES

[1] M. Martone and J. Lawall, "Refactoring for performance with semantic patching: Case study with recipes," in *High Performance Computing - ISC High Performance Digital 2021 International Workshops, Frankfurt am Main, Germany, June 24 - July 2, 2021, Revised Selected Papers*, ser. Lecture Notes in Computer Science, H. Jagode, H. Anzt, H. Ltaief, and P. Luszczek, Eds., vol. 12761. Springer, 2021, pp. 226-232. [Online]. Available: https://doi.org/10.1007/978-3-030-90539-2_15

[2] M. Martone, "Introduction to Semantic Patching of C and C++ Programs with Coccinelle," https://doi.org/10.5281/zenodo.14728519.

[12] https://rust-for-linux.com/coccinelle-for-rust

[3] J. Lawall and G. Muller, "Coccinelle: 10 years of automated evolution in the Linux kernel," in *2018 USENIX Annual Technical Conference, USENIX ATC*, 2018, pp. 601-614. [Online]. Available: https://www.usenix.org/conference/atc18/presentation/lawall

[4] M. Fowler, K. Beck, J. Brant, W. Opdyke, and D. Roberts, "Refactoring: Improving the Design of Existing Code." Addison-Wesley Professional, 1999, ISBN: 0-201-48567-2.

[5] "The C23 standard draft," https://www.open-std.org/jtc1/sc22/wg14/ www/docs/n3096.pdf.

[6] "CLANG's refactoring engine," https://clang.llvm.org/docs/ RefactoringEngine.html.

[7] D. Quinlan and C. Liao, "The ROSE source-to-source compiler infrastructure," in *Cetus users and compiler infrastructure workshop, in conjunction with PACT*, vol. 2011, 2011, p. 1. [Online]. Available: https://engineering.purdue.edu/Cetus/cetusworkshop/papers/4-1.pdf

[8] "The ROSE compiler webpage," https://rosecompiler.org/.

[9] H. C. Edwards, C. R. Trott, and D. Sunderland, "Kokkos: Enabling manycore performance portability through polymorphic memory access patterns," *J. Parallel Distributed Comput.*, vol. 74, no. 12, pp. 3202-3216, 2014, https://doi.org/10.1016/j.jpdc.2014.07.003. [Online]. Available: https://doi.org/10.1016/j.jpdc.2014.07.003

[10] V. Springel, "The cosmological simulation code GADGET-2," *MNRAS*, vol. 364, pp. 1105-1134, 2005.

[11] F. Baruffa, L. Iapichino, N. J. Hammer, and V. Karakasis, "Performance optimisation of smoothed particle hydrodynamics algorithms for multi/many-core architectures," *CoRR*, vol. abs/1612.06090, 2016. [Online]. Available: http://arxiv.org/abs/1612.06090

[12] S. J. Pennycook, C. J. Hughes, and M. Smelyanskiy, "Chapter 8 - optimizing gather/scatter patterns," in *High Performance Parallelism Pearls*, J. Reinders and J. Jeffers, Eds. Boston: Morgan Kaufmann, 2015, pp. 143-157. [Online]. Available: http://www.sciencedirect.com/ science/article/pii/B978012802118700008X

[13] "OpenMP 6.0 Specification," https://www.openmp.org/wp-content/ uploads/OpenMP-API-Specification-6-0.pdf.

[14] "GCC: Function multiversioning (attributes target and target_clone s)," https://gcc.gnu.org/onlinedocs/gcc/Common- Function-Attributes.html.

[15] "Nvidia CUDA C++ Programming Guide," https://docs.nvidia.com/ cuda/pdf/CUDA_C_Programming_Guide.pdf.

[16] "AMD HIP (Heterogeneous-Compute Interface for Portability)," https: //rocmdocs.amd.com/projects/HIP/.

[17] "OpenACC 3.2 Specification," https://www.openacc.org/sites/default/ files/inline-images/Specification/OpenACC-3.2-final.pdf.

[18] "C++ core guidelines," https://isocpp.github.io/CppCoreGuidelines/ CppCoreGuidelines.

[19] D. Beckingsale, T. R. W. Scogland, J. Burmark, R. Hornung, H. Jones, W. Killian, A. J. Kunen, O. Pearce, P. Robinson, and B. S. Ryujin, "RAJA: portable performance for large-scale scientific applications," in *2019 IEEE/ACM International Workshop on Performance, Portability and Productivity in HPC*, P3HPC@SC 2019, Denver, CO, USA, November 22, 2019. IEEE, 2019, pp. 71-81. [Online]. Available: https://doi.org/10.1109/P3HPC49587.2019.00012

[20] "The C++23 standard draft," https://www.open-std.org/JTC1/SC22/ WG21/docs/papers/2023/n4950.pdf.

[21] "SYCL: C++ Programming for Heterogeneous Parallel Computing," https://www.khronos.org/sycl/.

[22] "LIBRSB: Sparse BLAS based on the Recursive Sparse Blocks Format," https://librsb.sf.net/.