Enzyme

https://enzyme.mit.edu
Admin
Enzyme is an automatic differentiation framework that computes derivatives and gradients of...
Tokens:16,984
Snippets:193
Trust Score:10
Update:5 months ago
Show doc for...
Tokens:
Raw
### Compiling C++ Example with Enzyme

Source: https://enzyme.mit.edu/getting_started/CUDAGuide

This command shows how to compile the C++ example using clang with the Enzyme plugin. Ensure the plugin path is correct for your installation.

```bash
clang test2.cpp -fplugin=/path/to/ClangEnzyme-11.so -O2 

```

--------------------------------

### C++ Enzyme Setup Definitions

Source: https://enzyme.mit.edu/getting_started/Examples

These definitions are assumed to be prepended to each example snippet for conciseness.

```cpp
int enzyme_dup;
int enzyme_dupnoneed;
int enzyme_out;
int enzyme_const;

template < typename return_type, typename ... T >
return_type __enzyme_fwddiff(void*, T ... );

template < typename return_type, typename ... T >
return_type __enzyme_autodiff(void*, T ... );

```

--------------------------------

### Compiling CUDA Example with Enzyme

Source: https://enzyme.mit.edu/getting_started/CUDAGuide

This command shows how to compile the CUDA example using clang with the Enzyme plugin and specifies the GPU architecture. Ensure the CUDA library path is correct.

```bash
clang test3.cu -fplugin=/path/to/ClangEnzyme-11.so -O2 --cuda-gpu-arch=sm_70 -lcudart -L/usr/local/cuda-10.1/lib64

```

--------------------------------

### Install Enzyme using Homebrew

Source: https://enzyme.mit.edu/Installation

Use this command to install Enzyme via Homebrew. Homebrew provides a pre-built binary on macOS and builds from source on Linux.

```bash
brew install enzyme

```

--------------------------------

### Verify Enzyme Installation with Unit Tests

Source: https://enzyme.mit.edu/Installation

Run Enzyme's unit tests to ensure that your build of Enzyme and LLVM are interoperating correctly. This command should complete in under a minute.

```bash
ninja check-enzyme

```

--------------------------------

### Example Program Differentiating sumAndMul

Source: https://enzyme.mit.edu/getting_started/CallingConvention

Provides an example of differentiating the `sumAndMul` function, showcasing the use of duplicated, inactive, and output arguments. The `__enzyme_autodiff` call demonstrates how these argument types are passed.

```c
double sumAndMul(double* array, size_t size, double mul) {
  double sum = 0;
  for(int i=0; i<size; i++) {
    sum += array[i];
  }
  return sum * mul;
}

...    

double d_mul = __enzyme_autodiff<double>((void*)sumAndMul,
                     /*duplicated argument*/array, d_array,
                     /*inactive argument*/size,
                     /*output argument*/mul);

// link to complete example: https://fwd.gymni.ch/78Hsev


```

--------------------------------

### C++ Autodiff Example

Source: https://enzyme.mit.edu/explorer

Demonstrates automatic differentiation of a C++ function using Enzyme. This snippet shows how to define a function, its derivative, and use it in main.

```cpp
#include <stdio.h>
extern double __enzyme_autodiff(void *, double);
double square(double x) { return x * x; }
double dsquare(double x) {
// This returns the derivative of square or 2 * x
return __enzyme_autodiff((void *)square, x);
}
int main() {
for (double i = 1; i < 5; i++)
printf("square(%f)=%f, dsquare(%f)=%f\n", i, square(i), i, dsquare(i));
}
```

--------------------------------

### GPU Kernel and CPU Wrapper Functions

Source: https://enzyme.mit.edu/getting_started/CUDAGuide

Defines a GPU kernel, a CPU wrapper for it, and a main CPU function that repeatedly calls the wrapper. This setup is used to demonstrate heterogeneous AD challenges.

```c++
// GPU Kernel
__global__ 
void collide(float* src, float* dst) {
    size_t idx = threadIdx.x;
    if (idx < 100) {
        dsr[idx] += src[idx] * src[idx] - 3 * src[idx];
    }
}

// Wrapper CPU function which calls kernel
void kern(float* src, float* dst) {
    collide<<<1, 100>>>(src, dst);
}

// Main CPU code that calls wrapper function
void iter(int nTimeSteps, float* src, float* dst) {
    for (unsigned int i=0; i<nTimeSteps/2; i++) {
        kern(src, dst);
        kern(dst, src);
    }
}

```

--------------------------------

### CUDA Example with Enzyme Autodiff

Source: https://enzyme.mit.edu/getting_started/CUDAGuide

This C++ CUDA code demonstrates Enzyme's autodiff on the GPU. It requires specific CUDA versions and compiler flags. Note that __enzyme_autodiff should be invoked on __device__ code.

```cpp
#include <stdio.h>

void __device__ foo_impl(double* x_in, double *x_out) {
    x_out[0] = x_in[0] * x_in[0];    
}

typedef void (*f_ptr)(double*, double*);

extern void __device__ __enzyme_autodiff(f_ptr,
    int, double*, double*,
    int, double*, double*
);

void __global__ foo(double* x_in, double *x_out) {
    foo_impl(x_in, x_out);
}

int __device__ enzyme_dup;
int __device__ enzyme_out;
int __device__ enzyme_const;

void __global__ foo_grad(double* x, double *d_x, double *y, double *d_y) {

    __enzyme_autodiff(foo_impl, 
        enzyme_dup, x, d_x, 
        enzyme_dup, y, d_y);

}

int main() {

    double *x, *d_x, *y, *d_y; // device pointers

    cudaMalloc(&x, sizeof(*x));
    cudaMalloc(&d_x, sizeof(*d_x));
    cudaMalloc(&y, sizeof(*y));
    cudaMalloc(&d_y, sizeof(*d_y));

    double host_x = 1.4;
    double host_d_x = 0.0;
    double host_y;
    double host_d_y = 1.0;

    cudaMemcpy(x, &host_x, sizeof(*x), cudaMemcpyHostToDevice);
    cudaMemcpy(d_x, &host_d_x, sizeof(*d_x), cudaMemcpyHostToDevice);
    cudaMemcpy(y, &host_y, sizeof(*y), cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, &host_d_y, sizeof(*d_y), cudaMemcpyHostToDevice);

    // foo<<<1,1>>>(x, y); fwd-pass only
    foo_grad<<<1,1>>>(x, d_x, y, d_y); // fwd and bkwd pass

    cudaDeviceSynchronize(); // synchroniz

    cudaMemcpy(&host_x, x, sizeof(*x), cudaMemcpyDeviceToHost);
    cudaMemcpy(&host_d_x, d_x, sizeof(*d_x), cudaMemcpyDeviceToHost);
    cudaMemcpy(&host_y, y, sizeof(*y), cudaMemcpyDeviceToHost);
    cudaMemcpy(&host_d_y, d_y, sizeof(*d_y), cudaMemcpyDeviceToHost);

    printf("%f %f\n", host_x, host_y);
    printf("%f %f\n", host_d_x, host_d_y);

}


```

--------------------------------

### Example LLVM IR Output

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

Illustrates the structure of LLVM IR generated from C code, showing the 'square' function and the 'dsquare' function which calls Enzyme's autodiff intrinsic.

```llvm
; input.ll
...
define double @square(double %x) #0 {
entry:
  %mul = fmul double %x, %x
  ret double %mul
}

define double @dsquare(double %x) local_unnamed_addr #1 {
entry:
  %call = tail call double @__enzyme_autodiff(i8* bitcast (double (double)* @square to i8*), double %x) #4
  ret double %call
}
...

```

--------------------------------

### C++ Forward Mode Differentiation with Zero Derivatives

Source: https://enzyme.mit.edu/getting_started/Examples

Demonstrates how to selectively disable differentiation for specific inputs in forward mode by setting their derivative values to zero. Assumes the setup definitions are present.

```cpp
double dfdx = __enzyme_fwddiff<double>((void*)f, enzyme_dup, x, 1.0, enzyme_dup, y, 0.0);
double dfdy = __enzyme_fwddiff<double>((void*)f, enzyme_dup, x, 0.0, enzyme_dup, y, 1.0);
printf("dfdx = %f, dfdy = %f\n", dfdx, dfdy);
// prints dfdx = 2.000000, dfdy = 2.750000

```

--------------------------------

### Reference C++ Example for Enzyme Autodiff

Source: https://enzyme.mit.edu/getting_started/CUDAGuide

This C++ code demonstrates the basic usage of Enzyme's autodiff functionality on the CPU. It includes the necessary includes, function definition, and the main function to call __enzyme_autodiff.

```cpp
#include <stdio.h>

void foo(double* x_in, double *x_out) {
    x_out[0] = x_in[0] * x_in[0];
}


int enzyme_dup;
int enzyme_out;
int enzyme_const;

typedef void (*f_ptr)(double*, double*);

extern void __enzyme_autodiff(f_ptr, 
    int, double*, double*,
    int, double*, double*);

int main() {

    double x = 1.4;
    double d_x = 0.0;
    double y;
    double d_y = 1.0;

    __enzyme_autodiff(foo, 
        enzyme_dup, &x, &d_x, 
        enzyme_dup, &y, &d_y);

    printf("%f %f\n", x, y);
    printf("%f %f\n", d_x, d_y);

}


```

--------------------------------

### Run Enzyme Integration Tests

Source: https://enzyme.mit.edu/Installation

Execute Enzyme's C/C++ integration tests. Note that these tests require an existing installation of the Adept AD Engine and Eigen, and may take around 6 minutes to complete.

```bash
ninja check-enzyme-integration

```

--------------------------------

### C++ Forward Mode Differentiation with enzyme_const

Source: https://enzyme.mit.edu/getting_started/Examples

Shows how to use `enzyme_const` to mark arguments that should not be differentiated in forward mode. Assumes the setup definitions are present.

```cpp
double dfdx = __enzyme_fwddiff<double>((void*)f, enzyme_dup, x, 1.0, enzyme_const, y);
double dfdy = __enzyme_fwddiff<double>((void*)f, enzyme_const, x, enzyme_dup, y, 1.0);
printf("dfdx = %f, dfdy = %f\n", dfdx, dfdy);
// prints dfdx = 2.000000, dfdy = 2.750000

```

--------------------------------

### C Example: Differentiating a Square Function

Source: https://enzyme.mit.edu/

This C code demonstrates how to use Enzyme's `__enzyme_autodiff` function to compute the derivative of a simple square function. Ensure Enzyme is properly integrated into your build process.

```c
#include <stdio.h>

double square(double x) {
    return x * x;
}

double __enzyme_autodiff(void*, double);
int main() {
    double x = 3.14;
    // Evaluates to 2 * x = 6.28
    double grad_x = __enzyme_autodiff((void*)square, x);
    printf("square'(%f) = %f\n", x,  grad_x);
    return 0;
}


```

--------------------------------

### C++ Forward Mode Differentiation of f(x, y) = x*y + 1/y

Source: https://enzyme.mit.edu/getting_started/Examples

Applies Enzyme's forward mode differentiation to the f(x, y) function, differentiating with respect to both x and y. Assumes the setup definitions are present.

```cpp
double x = 3.0;
double y = 2.0;
double dx = 1.0;
double dy = 1.0;
double df = __enzyme_fwddiff<double>((void*)f, enzyme_dup, x, dx, enzyme_dup, y, dy);
printf("f(x) = %f, df = %f", f(x, y), df)
// prints f(x) = 6.500000, df = 4.750000

```

--------------------------------

### C++ Forward Mode Differentiation of f(x) = x*x

Source: https://enzyme.mit.edu/getting_started/Examples

Applies Enzyme's forward mode differentiation to the f(x) = x*x function. Assumes the setup definitions are present.

```cpp
double x = 5.0;
double dx = 1.0;
double df_dx = __enzyme_fwddiff<double>((void*)f, enzyme_dup, x, dx);
printf("f(x) = %f, f'(x) = %f", f(x), df_dx);
// prints f(x) = 25.000000, f'(x) = 10.000000

```

--------------------------------

### Configure and Build Enzyme

Source: https://enzyme.mit.edu/Installation

After navigating to the Enzyme source directory and creating a build directory, configure Enzyme using CMake, pointing to your LLVM build. Then, build Enzyme using Ninja.

```bash
cd /path/to/Enzyme/enzyme
mkdir build && cd build
cmake -G Ninja .. -DLLVM_DIR=/path/to/llvm/lib/cmake/llvm
ninja

```

--------------------------------

### Download Enzyme Source Code

Source: https://enzyme.mit.edu/Installation

Clone the Enzyme repository from GitHub and navigate into the project directory. This is the first step for building from source.

```bash
git clone https://github.com/EnzymeAD/Enzyme
cd Enzyme

```

--------------------------------

### Compile LLVM IR to Executable

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

Compile the optimized LLVM IR (`output_opt.ll`) into a final executable binary using clang.

```bash
clang output_opt.ll -o a.exe
```

--------------------------------

### Initialize Shadow Arguments for Complex Data Structures

Source: https://enzyme.mit.edu/getting_started/CallingConvention

Illustrates initializing shadow arguments for complex data structures like linked lists by creating a corresponding shadow structure with initial values (e.g., 0.0 for doubles).

```c
struct List {
  double value;
  List* next;
};

double sumList(List* next);
List* mklist(double value, List* next);

List*   list = nullptr;
List* d_list = nullptr;

for(int i=0; i<5; ++i) {
    list = mklist(  i,   list);
  d_list = mklist(0.0, d_list);
}

__enzyme_autodiff(sumList, list, d_list);


```

--------------------------------

### Combine Optimization and Compilation

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

Combine the optimization and compilation steps into a single command for creating an executable.

```bash
clang output.ll -O3 -o a.exe
```

--------------------------------

### Enable Enzyme Preprocessing

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

Use the `enzyme-preopt=0` option to enable all preprocessing optimizations in the Enzyme pass.

```bash
$ opt input.ll -load-pass-plugin=./Enzyme/LLVMEnzyme-17.so -passes=enzyme -enzyme-preopt=0
```

--------------------------------

### Build LLVM from Source

Source: https://enzyme.mit.edu/Installation

This command configures and builds LLVM from source, which is a prerequisite for building Enzyme. It uses CMake with Ninja and specifies target architectures and enabled projects.

```bash
cd /path/to/llvm/source/
mkdir build && cd build
cmake -G Ninja ../llvm -DLLVM_TARGETS_TO_BUILD="host" -DLLVM_ENABLE_PROJECTS="clang;lld" -DLLVM_ENABLE_PLUGINS=ON -DCMAKE_BUILD_TYPE=Release -DLLVM_ENABLE_ASSERTIONS=ON
ninja

```

--------------------------------

### Run Enzyme LLVM Pass

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

Load Enzyme and run the differentiation transformation pass on LLVM IR. Ensure `opt` points to your LLVM build and the plugin path is correct.

```bash
opt input.ll --load-pass-plugin=/path/to/Enzyme/enzyme/build/Enzyme/LLVMEnzyme-<VERSION>.so -passes=enzyme -o output.ll -S
```

--------------------------------

### Run Enzyme Benchmarking Suite

Source: https://enzyme.mit.edu/Installation

Execute Enzyme's benchmarking suite, which includes tests from the ADBench suite. Be aware that these tests can take a significant amount of time, potentially up to an hour.

```bash
ninja bench-enzyme

```

--------------------------------

### Compile with Clang Enzyme Plugin

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

Automate the entire AD and compilation process using Enzyme's Clang plugin in a single command.

```bash
clang test.c -fplugin=/path/to/Enzyme/enzyme/build/Enzyme/ClangEnzyme-<VERSION>.so -o a.exe
```

--------------------------------

### Initialize Shadow Arguments for Arrays

Source: https://enzyme.mit.edu/getting_started/CallingConvention

Shows how to initialize shadow arguments, particularly for arrays, by zeroing them out before calling `__enzyme_autodiff`. Enzyme increments the shadow value by the resultant gradient.

```c
double   array[10] = { ... };
double d_array[10] = { 0.0 };

__enzyme_autodiff(sumSquare,
                  enzyme_dup, array, d_array);

printf("d(output)/darray[0] = %f\n", d_array[0]);


```

--------------------------------

### Forward Differentiation with Functor

Source: https://enzyme.mit.edu/getting_started/Examples

Demonstrates forward differentiation of a C++ functor using Enzyme. The functor must be passed via a wrapper function.

```cpp
struct MyObject {
    double operator()(double y) const { return x * y + 1.0 / y; }
    double x;
};

double wrapper(const MyObject & f, double y) {  return f(y); }

...
    
double dfdy = __enzyme_fwddiff<double>((void*)(wrapper<MyObject, double>), 
    enzyme_const, (void*)&obj,
    enzyme_dup, &y, &dy);
printf("dfdy = %f\n", dfdy);
// prints dfdy = 0.750000

```

--------------------------------

### Classify Arguments with LLVM Metadata

Source: https://enzyme.mit.edu/getting_started/CallingConvention

Illustrates how to explicitly classify arguments for differentiation using LLVM metadata. Use `"enzyme_dup"`, `"enzyme_const"`, and `"enzyme_out"` for duplicated, inactive, and output arguments respectively.

```llvm
%d_mul = tail call double @__enzyme_autodiff(double (double*, i64, double)* @sumAndMul, metadata !"enzyme_dup", double* %array, double* %d_array, metadata !"enzyme_const", i64 %size, metadata !"enzyme_out", double %mul)


```

--------------------------------

### Simplified Optimized LLVM IR

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

The LLVM IR after applying -O2 optimization, showing a simplified gradient function.

```llvm
; output_opt.ll
define double @dsquare(double %x) local_unnamed_addr #0 {
entry:
  %factor.i = fmul fast double %x, 2.000000e+00
  ret double %factor.i
}

```

--------------------------------

### Compile Optimized LLVM IR

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

Compile the LLVM IR after optimization passes to create an executable.

```bash
opt output.ll -O2 -o output_opt.ll -S
```

--------------------------------

### Print Functions Before and After Preprocessing

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

The `-enzyme-print` option displays functions being differentiated at various stages: before preprocessing, after preprocessing, and after synthesis by Enzyme. This is primarily for debugging the AD process.

```bash
$ opt input.ll -load-pass-plugin=./Enzyme/LLVMEnzyme-17.so.so -passes=enzyme -enzyme-print

```

```llvm
prefn:

; Function Attrs: norecurse nounwind readnone uwtable
define double @square(double %x) #0 {
entry:
  %mul = fmul double %x, %x
  ret double %mul
}

```

--------------------------------

### Generate LLVM IR from C/C++ Code

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

Generates LLVM IR from a C source file using clang. Use optimization flags like -O2 for better performance, and disable vectorization/unrolling before AD.

```bash
clang test.c -S -emit-llvm -o input.ll -O2 -fno-vectorize -fno-slp-vectorize -fno-unroll-loops

```

--------------------------------

### Classify Arguments with Global Variables

Source: https://enzyme.mit.edu/getting_started/CallingConvention

Demonstrates classifying arguments for differentiation by using global variables like `enzyme_dup`, `enzyme_out`, and `enzyme_const`. This method provides a simpler syntax for frontend developers.

```c
int enzyme_dup;
int enzyme_out;
int enzyme_const;

int main() {
  double d_mul = __enzyme_autodiff<double>((void*)sumAndMul,
                       enzyme_dup  , array, d_array,
                       enzyme_const, size,
                       enzyme_out  , mul);
  return 0;
}


```

--------------------------------

### Enable Compressed Bool Cache

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

Use the `enzyme-smallbool=1` option to enable Enzyme's cache to store 8 boolean values per byte, reducing memory usage.

```bash
$ opt input.ll -load-pass-plugin=./Enzyme/LLVMEnzyme-17.so -passes=enzyme -enzyme-smallbool=1
```

--------------------------------

### Reverse Mode Differentiation with Input Lambda and Output Arguments

Source: https://enzyme.mit.edu/getting_started/Examples

Shows reverse-mode differentiation where a lambda is provided as input to `__enzyme_autodiff`, and both function output and derivatives are captured. `enzyme_dupnoneed` can be used for arguments where only the derivative is required.

```cpp
// lambda is an input to __enzyme_autodiff 
double lambda = 2.0;
#if 1
double2 mu = __enzyme_autodiff<double2>((void*)f, enzyme_out, x, 
                                          enzyme_out, y, 
                                          enzyme_dup, &z, &lambda); 
printf("z = %f, mu.x = %f, mu.y = %f\n", z, mu.x, mu.y);
// prints z = 6.500000, mu.x = 4.000000, mu.y = 5.500000
#else
double2 mu = __enzyme_autodiff<double2>((void*)f, enzyme_out, x, 
                                          enzyme_out, y, 
                                          enzyme_dupnoneed, &z, &lambda); 
printf("z = %f, mu.x = %f, mu.y = %f\n", z, mu.x, mu.y);
// prints z = 0.000000, mu.x = 4.000000, mu.y = 5.500000
#endif
```

--------------------------------

### Invoke Enzyme Autodiff with Function Pointers

Source: https://enzyme.mit.edu/getting_started/CallingConvention

Demonstrates how to use `__enzyme_autodiff` with function pointers for differentiating functions with float and double types. Ensure the correct `__enzyme_autodiff` variant is used for the function's type.

```c
#include <stdio.h>
template<typename T>
T square(T x) { return x * x; }

float __enzyme_autodiffFloat(float (*)(float), float);
double __enzyme_autodiffDouble(double (*)(double), double);

int main() {
  printf("float  d/dx %f\n", __enzyme_autodiffFloat(square<float>, 1.0f));
  printf("double d/dx %f\n", __enzyme_autodiffDouble(square<double>, 1.0));
  return 0;
}


```

--------------------------------

### Forward Mode Differentiation of Function Templates

Source: https://enzyme.mit.edu/getting_started/Examples

Demonstrates how to differentiate function templates with Enzyme. Explicitly specify the template arguments when passing the function to `__enzyme_fwddiff`.

```cpp
template < typename T >
void f(T x, T y, T & output) { output = x * y + 1.0 / y; }

Then the first argument looks like
__enzyme_fwddiff<void>((void*)f<double>, enzyme_dup, x, dx, 
                                         enzyme_dup, y, dy, 
                                         enzyme_dupnoneed, &z, &dz);
```

--------------------------------

### Enable Loose Type Analysis

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

Use the `enzyme-loose-types=1` option to allow Enzyme to make educated guesses about unprovable value types, useful for bootstrapping but not for production.

```bash
$ opt input.ll -load-pass-plugin=./Enzyme/LLVMEnzyme-17.so.so -passes=enzyme -enzyme-loose-types=1
```

--------------------------------

### BibTeX Citation for Enzyme (GPU+Optimizations)

Source: https://enzyme.mit.edu/

This BibTeX entry is for citing the Enzyme paper focusing on GPU and optimizations. Use this when discussing Enzyme's performance on GPUs.

```bibtex
@inproceedings{10.1145/3458817.3476165,
author = {Moses, William S. and Churavy, Valentin and Paehler, Ludger and H"{u}ckelheim, Jan and Narayanan, Sri Hari Krishna and Schanen, Michel and Doerfert, Johannes},
title = {Reverse-Mode Automatic Differentiation and Optimization of GPU Kernels via Enzyme},
year = {2021},
isbn = {9781450384421},
publisher = {Association for Computing Machinery},
address = {New York, NY, USA},
url = {https://doi.org/10.1145/3458817.3476165},
doi = {10.1145/3458817.3476165},
booktitle = {Proceedings of the International Conference for High Performance Computing, Networking, Storage and Analysis},
articleno = {61},
numpages = {16},
keywords = {CUDA, LLVM, ROCm, HPC, AD, GPU, automatic differentiation},
location = {St. Louis, Missouri},
series = {SC '21}
}

```

--------------------------------

### Device Function and Modified Kernel/Wrapper

Source: https://enzyme.mit.edu/getting_started/CUDAGuide

Refactors the GPU kernel logic into a device function 'collide_body' to enable differentiation. The original kernel and wrapper are also updated.

```c++
// We move the body of collide into a separate device function collide_body to allow us
// to pass collide_body to various differentiation methods. This is necessary as differentiation
// can only be done on device, not global kernel functions.
__device__
void collide_body(float* src, float* dst) {
    size_t idx = threadIdx.x;
    if (idx < 100) {
        dst[idx] += src[idx] * src[idx] - 3 * src[idx];
    }
}

// GPU Kernel
__global__
void collide(float* src, float* dst) {
    collide_body(src, dst);
}

// Wrapper CPU function which calls kernel
__attribute__((noinline))
void kern(float* src, float* dst) {
    collide<<<1, 100>>>(src, dst);
}

// Main CPU code that calls wrapper function
void iter(int nTimeSteps, float* src, float* dst) {
    for (unsigned int i=0; i<nTimeSteps/2; i++) {
        kern(src, dst);
        kern(dst, src);
    }
}


```

--------------------------------

### Enzyme Error Message for Heterogeneous AD

Source: https://enzyme.mit.edu/getting_started/CUDAGuide

Illustrates the typical error message encountered when Enzyme attempts to differentiate a function containing a CUDA call directly, highlighting the need for custom derivative handling.

```text
declare dso_local i32 @__cudaPushCallConfiguration(i64, i32, i64, i32, i64, i8*) local_unnamed_addr #2

clang-13: /home/wmoses/git/Enzyme/enzyme/Enzyme/EnzymeLogic.cpp:1459: const AugmentedReturn& EnzymeLogic::CreateAugmentedPrimal(llvm::Function*, DIFFE_TYPE, const std::vector<DIFFE_TYPE>&, llvm::TargetLibraryInfo&, TypeAnalysis&, bool, const FnTypeInfo&, std::map<llvm::Argument*, bool>, bool, bool, bool, bool): Assertion `0 && "attempting to differentiate function without definition"' failed.
PLEASE submit a bug report to https://bugs.llvm.org/ and include the crash backtrace, preprocessed source, and associated run script.


```

--------------------------------

### Disable Enzyme Preprocessing

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

Use the `enzyme-preopt=1` option to disable preprocessing optimizations in the Enzyme pass, except for the minimum necessary.

```bash
$ opt input.ll -load-pass-plugin=./Enzyme/LLVMEnzyme-17.so -passes=enzyme -enzyme-preopt=1
```

--------------------------------

### Print Activity Analysis Results

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

Use `-enzyme-print-activity` to view the results of activity analysis as they are derived. The output is specific to the analysis pass and intended for developers.

```bash
$ opt input.ll -load-pass-plugin=./Enzyme/LLVMEnzyme-17.so.so -passes=enzyme -enzyme-print-activity

```

```text
in new function diffesquare nonconstant arg double %0
 VALUE nonconst from arg nonconst double %x
checking if is constant[3]   %mul = fmul double %x, %x
 < UPSEARCH3>  %mul = fmul double %x, %x
 VALUE nonconst from arg nonconst double %x
nonconstant(3)  up-inst   %mul = fmul double %x, %x op double %x
 </UPSEARCH3>  %mul = fmul double %x, %x
couldnt decide nonconstants(3):  %mul = fmul double %x, %x
 Value nonconstant (couldn't disprove)[3]  %mul = fmul double %x, %x

```

--------------------------------

### Optimized LLVM IR after Enzyme

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

The LLVM IR generated by Enzyme, after applying optimization passes like -O2, shows inlined and simplified gradient functions.

```llvm
; output.ll
define internal { double } @diffesquare(double %x, double %differeturn) #0 {
entry:
  %"mul'de" = alloca double
  store double 0.000000e+00, double* %"mul'de"
  %"x'de" = alloca double
  store double 0.000000e+00, double* %"x'de"
  br label %invertentry

invertentry:
  store double %differeturn, double* %"mul'de"
  %0 = load double, double* %"mul'de"
  %m0diffex = fmul fast double %0, %x
  %m1diffex = fmul fast double %0, %x
  store double 0.000000e+00, double* %"mul'de"
  %1 = load double, double* %"x'de"
  %2 = fadd fast double %1, %m0diffex
  store double %2, double* %"x'de"
  %3 = load double, double* %"x'de"
  %4 = fadd fast double %3, %m1diffex
  store double %4, double* %"x'de"
  %5 = load double, double* %"x'de"
  %6 = insertvalue { double } undef, double %5, 0
  ret { double } %6
}

define double @dsquare(double %x) local_unnamed_addr #1 {
entry:
  %0 = call { double } @diffesquare(double %x, double 1.000000e+00)
  %1 = extractvalue { double } %0, 0
  ret double %1
}

```

--------------------------------

### Forward Differentiation of Lambda (Member Function Wrapper)

Source: https://enzyme.mit.edu/getting_started/Examples

Uses a member function wrapper to differentiate a lambda function with Enzyme. This approach is necessary when the lambda captures variables or for explicit type control.

```cpp
double dfdy = __enzyme_fwddiff<double>((void*)(wrapper<decltype(f), double, double>), 
        enzyme_const, (void*)&f,
        enzyme_const, &x,
        enzyme_dup, &y, &dy);
printf("dfdy = %f\n", dfdy);
// dfdy = 0.750000

```

--------------------------------

### Wrapper Function for Member Function Differentiation

Source: https://enzyme.mit.edu/getting_started/Examples

Creates a free function that calls a member function, enabling differentiation with Enzyme.

```cpp
double wrapper(MyObject obj, double y) {
    return obj.f(y);
}

```

--------------------------------

### Forward Mode Differentiation with `enzyme_dupnoneed`

Source: https://enzyme.mit.edu/getting_started/Examples

Illustrates forward-mode differentiation using `__enzyme_fwddiff`. The `enzyme_dupnoneed` descriptor can be used when only the derivative is needed, allowing the compiler to optimize away unnecessary function value calculations.

```cpp
// these will be overwritten by __enzyme_fwddiff
double z = 0;
double dz = 0;
#if 1
__enzyme_fwddiff<void>((void*)f, enzyme_dup, x, dx, 
                                 enzyme_dup, y, dy, 
                                 enzyme_dup, &z, &dz);
printf("f(x,y) = %f, df = %f\n", z, dz);
// prints f(x,y) = 6.500000, df = 7.500000
#else
__enzyme_fwddiff<void>((void*)f, enzyme_dup, x, dx, 
                                 enzyme_dup, y, dy, 
                                 enzyme_dupnoneed, &z, &dz);
printf("f(x,y) = %f, df = %f\n", z, dz);
// prints f(x,y) = 0.000000, df = 7.500000
#endif
```

--------------------------------

### Force Inlining with Enzyme

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

Use the `enzyme-inline=1` option to forcibly inline all subfunction calls during the Enzyme pass.

```bash
$ opt input.ll -load-pass-plugin=./Enzyme/LLVMEnzyme-17.so -passes=enzyme -enzyme-inline=1
```

--------------------------------

### Enzyme Autodiff Variadic Function Signatures

Source: https://enzyme.mit.edu/getting_started/CallingConvention

Shows variadic and templated function signatures for `__enzyme_autodiff`, enabling flexible integration with different argument types. These signatures allow Enzyme to register multiple function signatures.

```c
void __enzyme_autodiff(...);

template<typename RT, typename... Args>
RT __enzyme_autodiff(void*, Args...);


```

--------------------------------

### Assume Inactivity of Undefined Functions

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

Use the `-enzyme-emptyfn-inactive=1` option to assume that calls to undefined functions are inactive. This can improve performance and avoid compile-time errors but should be used cautiously.

```bash
$ opt input.ll -load-pass-plugin=./Enzyme/LLVMEnzyme-17.so.so -passes=enzyme -enzyme-emptyfn-inactive=1

```

--------------------------------

### Wrapper Function for Complicated Types

Source: https://enzyme.mit.edu/getting_started/CallingConvention

Create a wrapper function to explicitly pass complicated types as duplicated arguments. The wrapper takes a pointer and calls the original function with the dereferenced value.

```c
class MyClass;

MyClass compute(MyClass&);

void wrapper(MyClass* in, MyClass* out) {
  *out = compute(*in);
}

MyClass d_compute(MyClass& in) {
  MyClass d_in(0.0);
  MyClass out;
  MyClass d_out(1.0);
  __enzyme_autodiff<void>((void*)wrapper, &in, d_in, out, d_out);
  return d_in;
}

```

--------------------------------

### Custom Derivative Registration Functions

Source: https://enzyme.mit.edu/getting_started/CUDAGuide

Declares Enzyme's custom forward and reverse pass functions, '__enzyme_augmentfwd' and '__enzyme_reverse', which are essential for handling derivatives across devices.

```c++
template <typename... Args>
void __enzyme_autodiff(Args...);

void grad_iter(int nTimeSteps, float* src, float* dsrc, float* dst, float* ddst) {
    __enzyme_autodiff(iter, nTimeSteps, src, dsrc, dst, ddst);
}

// A function similar to __enzyme_autodiff, except it only calls the augmented forward pass, returning
// a tape structure to hold any values that may be overwritten and needed for the reverse.
template <typename... Args>
__device__ void* __enzyme_augmentfwd(Args...);

// A function similar to __enzyme_autodiff, except it only calls the revese pass, taking in the tape
// as its last argument.
template <typename... Args>
__device__ void __enzyme_reverse(Args...);


```

--------------------------------

### Specify LLVM LIT path for CMake

Source: https://enzyme.mit.edu/Installation

If CMake fails to find `llvm-lit`, the LLVM unit tester, this flag can be used to explicitly provide the path to `lit.py`, `lit`, or `llvm-lit`.

```bash
cmake .. -DLLVM_EXTERNAL_LIT=/path/to/lit/lit.py

```

--------------------------------

### Reverse Mode Differentiation with Struct Output

Source: https://enzyme.mit.edu/getting_started/Examples

Demonstrates reverse-mode differentiation using `__enzyme_autodiff` where the output is a struct. Use C++17 structured binding to access struct members. Avoid using `std::tuple` for output due to potential memory layout inconsistencies.

```cpp
struct double2{ double x, y; };

...

double x = 3.0;
double y = 2.0;
auto [mu_x, mu_y] = __enzyme_autodiff<double2>((void*)f, enzyme_out, x, enzyme_out, y); 
printf("mu_x = %f, mu_y = %f\n", mu_x, mu_y);
// prints mu_x = 2.000000, mu_y = 2.750000
```

--------------------------------

### Assume Inactivity of Unmarked Globals

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

Use the `-enzyme-globals-default-inactive=1` option to assume that global variables without a defined shadow global are inactive. Exercise caution with this option as it may lead to incorrect derivative computations.

```bash
$ opt input.ll -load-pass-plugin=./Enzyme/LLVMEnzyme-17.so.so -passes=enzyme -enzyme-globals-default-inactive=1

```

--------------------------------

### Forward Differentiation of Lambda (C++20 Wrapper)

Source: https://enzyme.mit.edu/getting_started/Examples

Employs a C++20 template wrapper to differentiate a lambda function with Enzyme. This method utilizes non-type template parameters for passing the lambda.

```cpp
// C++20 or later
template < auto obj, typename ... arg_types >
auto wrapper(arg_types && ... args) {
    return obj(args ... );
}
   
...
   
double dfdy = __enzyme_fwddiff<double>((void*)(wrapper<f, double, double>), 
    enzyme_const, &x,
    enzyme_dup, &y, &dy);
printf("dfdy = %f\n", dfdy);
// dfdy = 0.750000

```

--------------------------------

### Limit Inlining with Enzyme

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

Use `enzyme-inline=1` with `enzyme-inline-count` to limit the number of subfunction calls inlined by the Enzyme pass.

```bash
$ opt input.ll -load-pass-plugin=./Enzyme/LLVMEnzyme-17.so -passes=enzyme -enzyme-inline=1 -enzyme-inline-count=100
```

--------------------------------

### BibTeX Citation for Enzyme (Parallel Programs)

Source: https://enzyme.mit.edu/

This BibTeX entry is for citing the Enzyme paper on the automatic differentiation of arbitrary parallel programs. Cite this when discussing Enzyme's capabilities with parallel paradigms.

```bibtex
@inproceedings{10.5555/3571885.3571964,
author = {Moses, William S. and Narayanan, Sri Hari Krishna and Paehler, Ludger and Churavy, Valentin and Schanen, Michel and H"{u}ckelheim, Jan and Doerfert, Johannes and Hovland, Paul},
title = {Scalable Automatic Differentiation of Multiple Parallel Paradigms through Compiler Augmentation},
year = {2022},
isbn = {9784665454445},
publisher = {IEEE Press},
booktitle = {Proceedings of the International Conference on High Performance Computing, Networking, Storage and Analysis},
articleno = {60},
numpages = {18},
keywords = {automatic differentiation, tasks, OpenMP, compiler, Julia, parallel, Enzyme, C++, RAJA, hybrid parallelization, MPI, distributed, LLVM},
location = {Dallas, Texas},
series = {SC '22}
}

```

--------------------------------

### Return-By-Reference Function Definition

Source: https://enzyme.mit.edu/getting_started/Examples

Defines a function that uses an out-parameter for its return value. This is an alternative to return-by-value.

```cpp
void f(double x, double y, double & output) { output = x * y + 1.0 / y; }
```

--------------------------------

### GPU Kernel for Forward Pass Augmentation

Source: https://enzyme.mit.edu/getting_started/CUDAGuide

A wrapper GPU kernel that utilizes '__enzyme_augmentfwd' to compute the augmented forward pass of 'collide_body'. It stores the resulting tape per thread.

```c++
// A wrapper GPU kernel for calling the forward pass of collide. The wrapper code stores
// the tape generated by Enzyme into a unique location per thread
__global__ void aug_collide(float* src, float* dsrc, float* dst, float* ddst, void** tape)
{
    size_t idx = threadIdx.x;
    tape[idx] = __enzyme_augmentfwd((void*)collide_body, src, dsrc, dst, ddst);
}


```

--------------------------------

### Forward Differentiation of Lambda (Direct Pass)

Source: https://enzyme.mit.edu/getting_started/Examples

Performs forward differentiation on a lambda function by passing it directly to Enzyme. This method works for lambdas that do not capture variables.

```cpp
double dfdx = __enzyme_fwddiff<double>((void*)+f, 
  enzyme_dup, x, dx,
  enzyme_const, y);
printf("dfdx = %f\n", dfdx);
// dfdx = 6.200000

```

--------------------------------

### Custom Gradient Metadata Declaration

Source: https://enzyme.mit.edu/getting_started/CallingConvention

Declare custom gradients by attaching metadata to the original function. This metadata links to the augmented forward pass and the gradient reverse pass functions.

```llvm
declare !enzyme_augment !{{ {}, double } (double)* @augment_add2} !enzyme_gradient !{{ double } (double, double, {})* @gradient_add2} double @add2(double %x)

```

--------------------------------

### Lambda Function Compiler Expansion

Source: https://enzyme.mit.edu/getting_started/Examples

Illustrates the C++ compiler's internal representation of a lambda function as a functor object. This is what Enzyme interacts with.

```cpp
// what the compiler "sees" when we type 
// auto f = [](double x, double y) { return x * y + 1.0 / y; };

class __lambda_f {
  public: 
    inline /*constexpr */ double operator()(double x, double y) const {
      return (x * y) + (1.0 / y);
    }
    
    using retType_f = double (*)(double, double);
    inline constexpr operator retType_f () const noexcept {
      return __invoke;
    };
    
    private: 
    static inline /*constexpr */ double __invoke(double x, double y) {
      return __lambda_f{}.operator()(x, y);
    }
};

__lambda_f f = __lambda_f{};

```

--------------------------------

### Duplicated Argument with No Value Needed

Source: https://enzyme.mit.edu/getting_started/CallingConvention

Use `enzyme_dupnoneed` when only the gradient of an argument is required, not its forward pass value. This enables additional optimizations.

```c
void neuralNet(double* loss, double* W, double* b, double* input);

void main() {
  ...
  double loss;
  double d_loss = 1.0;
  __enzyme_autodiff<void>((void*)neuralNet,
                    enzyme_dupnoneed, &loss, &d_loss,
                    enzyme_dup,       W, d_W,
                    enzyme_dup,       b, d_b,
                    enzyme_const,     input);
  // This value is undefined behavior if using diffe_dupnoneed, otherwise
  // it is the same as it would be from calling neuralNet normally.
  printf("loss=%f\n", loss);

  printf("d_b[0]=%f\n", d_b[0]);
}

```

--------------------------------

### Print Type Analysis Results

Source: https://enzyme.mit.edu/getting_started/UsingEnzyme

The `-enzyme-print-type` option shows the results of type analysis during derivation. This output is specific to the analysis pass and is intended for developers.

```bash
$ opt input.ll -load-pass-plugin=./Enzyme/LLVMEnzyme-17.so.so -passes=enzyme -enzyme-print-type

```

```text
analyzing function square
 + knowndata: double %x : {[-1]:Float@double} - {}
 + retdata: {}
updating analysis of val: double %x current: {} new {[-1]:Float@double}
updating analysis of val: double %x current: {[-1]:Float@double} new {[-1]:Float@double} from double %x
updating analysis of val:   %mul = fmul double %x, %x current: {} new {}
updating analysis of val: double %x current: {[-1]:Float@double} new {[-1]:Float@double} from   %mul = fmul double %x, %x
updating analysis of val: double %x current: {[-1]:Float@double} new {[-1]:Float@double} from   %mul = fmul double %x, %x
updating analysis of val:   %mul = fmul double %x, %x current: {} new {[-1]:Float@double} from   %mul = fmul double %x, %x

```

--------------------------------

### Capturing Lambda Compiler Expansion

Source: https://enzyme.mit.edu/getting_started/Examples

Shows the C++ compiler's internal representation of a lambda function that captures variables. The captured variable becomes a member of the functor object.

```cpp
// what the compiler "sees" when we type 
// double x = 2.0;
// auto f = [x](double y) { return x * y + 1.0 / y; };

class __lambda_f  {
 public: 
  inline /*constexpr */ double operator()(double y) const {
    return (x * y) + (1.0 / y);
  }
    
  __lambda_f(double & _x) : x{_x}{}
    
 private: 
  double x;
};

double x = 2.0;
__lambda_f f = __lambda_f{x};

```

--------------------------------

### Global Variable Shadow Specification

Source: https://enzyme.mit.edu/getting_started/CallingConvention

Specify the shadow of active global variables in LLVM by attaching metadata. This metadata indicates the shadow global.

```llvm
@global = external local_unnamed_addr global double, align 8, !enzyme_shadow !{double* @dglobal}
@dglobal = external local_unnamed_addr global double, align

```

--------------------------------

### Register Custom Gradient Kernels

Source: https://enzyme.mit.edu/getting_started/CUDAGuide

Register custom forward and reverse pass functions for a kernel. This allows Enzyme to use user-defined implementations for differentiation.

```c
void* __enzyme_register_gradient_kern[3] = { (void*)kern, (void*)aug_kern, (void*)rev_kern };
```

--------------------------------

### GPU Kernel for Reverse Pass Execution

Source: https://enzyme.mit.edu/getting_started/CUDAGuide

A wrapper GPU kernel that uses '__enzyme_reverse' to execute the reverse pass of 'collide_body', retrieving the tape generated during the forward pass for each thread.

```c++
// A wrapper GPU kernel for calling the reverse pass of collide. The wrapper code retrieves
// the corresponding tape per thread being executed.
__global__ void rev_collide( float* src, float* dsrc, float* dst, float* ddst, void** tape)
{
    size_t idx = threadIdx.x;
    __enzyme_reverse((void*)collide_body, src, dsrc, dst, ddst, tape[idx]);
}


```

--------------------------------

### CPU Kern Augmented Forward Pass

Source: https://enzyme.mit.edu/getting_started/CUDAGuide

CPU-side function for the augmented forward pass. It allocates tape memory and calls the augmented collide GPU kernel.

```c
void* aug_kern(float* src, float* dsrc, float* dst, float* ddst) {
    float* tape;
    cudaMalloc(&tape, sizeof(float) * /*total number of threads*/100);
    aug_collide<<<1, 100>>>(src, dsrc, dst, ddst, tape);
    return (void*)tape;
}
```

--------------------------------

### Custom Gradient Augmented Forward Pass

Source: https://enzyme.mit.edu/getting_started/CallingConvention

Define an augmented forward pass function that saves necessary state for the reverse pass. The return type includes the tape and original return value.

```llvm
define internal { {}, double } @augment_add2(double %x) {
entry:
  %add = fadd fast double %x, 2.000000e+00
  %struct1 = insertvalue { {}, double } undef, double %add, 1
  ret { {}, double } %struct1
}

```
Tokens:
Raw