### Initialize HeteroCL and Define Compute Kernel

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_08_backend.html

Initializes HeteroCL, defines a placeholder for input data, and creates a compute kernel. This setup is used across different back-end examples.

```python
import heterocl as hcl
import numpy as np

hcl.init()
A = hcl.placeholder((10, 10), "A")


def kernel(A):
    return hcl.compute((8, 8), lambda y, x: A[y][x] + A[y + 2][x + 2], "B")

s = hcl.create_scheme(A, kernel)
s.downsize(kernel.B, hcl.UInt(4))
s = hcl.create_schedule_from_scheme(s)
s.partition(A)
s[kernel.B].pipeline(kernel.B.axis[1])

```

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

### HeteroCL IR Example

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_01_get_started.html

An example of the Intermediate Representation (IR) generated by HeteroCL, showing the lowered computation.

```llvm
module {
  func.func @top(%arg0: memref<1xi32>, %arg1: memref<10x10xi32>) -> memref<10x10xi32> attributes {itypes = "ss", otypes = "s"} {
    %0 = memref.alloc() {name = "B"} : memref<10x10xi32>
    affine.for %arg2 = 0 to 10 {
      affine.for %arg3 = 0 to 10 {
        %1 = affine.load %arg1[%arg2, %arg3] {from = "A"} : memref<10x10xi32>
        %2 = affine.load %arg0[0] {from = "a"} : memref<1xi32>
        %3 = arith.extsi %1 : i32 to i33
        %4 = arith.extsi %2 : i32 to i33
        %5 = arith.addi %3, %4 : i33
        %6 = arith.trunci %5 : i33 to i32
        affine.store %6, %0[%arg2, %arg3] {to = "B"} : memref<10x10xi32>
      } {loop_name = "y"}
    } {loop_name = "x", op_name = "B"}
    return %0 : memref<10x10xi32>
  }
}
```

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

### Initialize HeteroCL and Define Placeholder

Source: https://cornell-zhang.github.io/heterocl/_downloads/366e024301a866e5e7272eed4f1329df/tutorial_08_backend.ipynb

Initializes HeteroCL and defines a placeholder tensor 'A'. This is a common setup before defining compute kernels.

```python
import heterocl as hcl
import numpy as np

hcl.init()
A = hcl.placeholder((10, 10), "A")


def kernel(A):
    return hcl.compute((8, 8), lambda y, x: A[y][x] + A[y + 2][x + 2], "B")


s = hcl.create_scheme(A, kernel)
s.downsize(kernel.B, hcl.UInt(4))
s = hcl.create_schedule_from_scheme(s)
s.partition(A)
s[kernel.B].pipeline(kernel.B.axis[1])
```

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

### Verify HeteroCL Installation

Source: https://cornell-zhang.github.io/heterocl/setup/index.html

Run the provided pytest command to confirm that HeteroCL has been installed correctly.

```bash
$ python3 -m pytest test

```

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

### Install HeteroCL via Pip

Source: https://cornell-zhang.github.io/heterocl/setup/index.html

Install HeteroCL using pip after cloning the repository. Requires cmake (>=3.19) and python (>=3.7).

```bash
$ pip install .

```

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

### Intel HLS SYCL Kernel Example

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_08_backend.html

This C++ code demonstrates a SYCL kernel for Intel FPGAs, including device selection, buffer creation, and kernel submission with specific pragmas for optimization.

```cpp
//===------------------------------------------------------------*- C++ -*-//
//
// Automatically generated file for Intel High-level Synthesis (HLS).
//
//===----------------------------------------------------------------------===//
#include <CL/sycl.hpp>
#include <iostream>
#include <sycl/ext/intel/fpga_extensions.hpp>
#include <vector>

// dpc_common.hpp can be found in the dev-utilities include folder.
// e.g., $ONEAPI_ROOT/dev-utilities//include/dpc_common.hpp
#include "dpc_common.hpp"

using namespace sycl;

// Forward declare the kernel name in the global scope to reduce name mangling.
// This is an FPGA best practice that makes it easier to identify the kernel in
// the optimization reports.
class Top;


int main() {

  // Select either:
  //  - the FPGA emulator device (CPU emulation of the FPGA)
  //  - the FPGA device (a real FPGA)
#if defined(FPGA_EMULATOR)
  ext::intel::fpga_emulator_selector device_selector;
#else
  ext::intel::fpga_selector device_selector;
#endif

  try {

    // Create a queue bound to the chosen device.
    // If the device is unavailable, a SYCL runtime exception is thrown.
    queue q(device_selector, dpc_common::exception_handler);

    // Print out the device information.
    std::cout << "Running on device: "
              << q.get_device().get_info<info::device::name>() << "\n";

    {
      // Create buffers to share data between host and device.
      // The runtime will copy the necessary data to the FPGA device memory
      // when the kernel is launched.
      buffer<int32_t, 2> buf_v0(range(10, 10));
      buffer<ac_uint<4>, 2> buf_v1(range(8, 8));

      // Submit a command group to the device queue.
      q.submit([&](handler& h) {

        // The SYCL runtime uses the accessors to infer data dependencies.
        // A "read" accessor must wait for data to be copied to the device
        // before the kernel can start. A "write no_init" accessor does not.
        accessor v0(buf_v0, h, read_only);
        accessor v1(buf_v1, h);

        // The kernel uses single_task rather than parallel_for.
        // The task's for loop is executed in pipeline parallel on the FPGA,
        // exploiting the same parallelism as an equivalent parallel_for.
        //
        //    DPC++FPGA/Tutorials/Features/kernel_args_restrict
        h.single_task<Top>([=]() [[intel::kernel_args_restrict]] {
          for (int v2 = 0; v2 < 8; v2 += 1) {   // L472
            [[intel::initiation_interval(1)]]
            for (int v3 = 0; v3 < 8; v3 += 1) { // L472
              int32_t v4 = v0[v2][v3];  // L21
              int32_t v5 = v0[(v2 + 2)][(v3 + 2)];      // L21
              ac_int<33> v6 = v4;       // L472
              ac_int<33> v7 = v5;       // L472
              ac_int<33> v8 = v6 + v7;  // L21
              ac_int<4> v9 = v8;        // L472
              v1[v2][v3] = v9;  // L472
            }
          }

        });
      });
    }

    // The queue destructor is invoked when q passes out of scope.
    // q's destructor invokes q's exception handler on any device exceptions.
  }
  catch (sycl::exception const& e) {
    // Catches exceptions in the host code
    std::cerr << "Caught a SYCL host exception:\n" << e.what() << "\n";

    // Most likely the runtime couldn't find FPGA hardware!
    if (e.code().value() == CL_DEVICE_NOT_FOUND) {
      std::cerr << "If you are targeting an FPGA, please ensure that your "
                   "system has a correctly configured FPGA board.\n";
      std::cerr << "Run sys_check in the oneAPI root directory to verify.\n";
      std::cerr << "If you are targeting the FPGA emulator, compile with "
                   "-DFPGA_EMULATOR.\n";
    }
    std::terminate();
  }

  return 0;
}

```

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

### Initialize HeteroCL and Define Compute Example

Source: https://cornell-zhang.github.io/heterocl/_downloads/8321e56154d9f8dcb74584a3b20dd05d/tutorial_03_api.ipynb

Initializes HeteroCL, defines placeholder tensors A and B, and creates a compute function to add them elementwise, returning a new tensor C. The generated IR is then printed.

```python
hcl.init()

A = hcl.placeholder((10,), "A")
B = hcl.placeholder((10,), "B")


def compute_example(A, B):
    return hcl.compute(A.shape, lambda x: A[x] + B[x], "C")


s = hcl.create_schedule([A, B], compute_example)
print(hcl.lower(s))
```

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

### Initialize HeteroCL and Placeholder

Source: https://cornell-zhang.github.io/heterocl/_downloads/9a0f69545c1a21968941d255884310e7/tutorial_02_imperative.ipynb

Initializes HeteroCL and creates a placeholder tensor for subsequent operations. This is a common setup for HeteroCL programs.

```python
import heterocl as hcl

hcl.init()

A = hcl.placeholder((10,), "A")
```

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

### Initialize HeteroCL and Define Update Example

Source: https://cornell-zhang.github.io/heterocl/_downloads/8321e56154d9f8dcb74584a3b20dd05d/tutorial_03_api.ipynb

Initializes HeteroCL, defines placeholder tensors A, B, and C. It then uses `hcl.update` to add A and B elementwise into C in place. The generated IR is printed.

```python
hcl.init()
A = hcl.placeholder((10,), "A")
B = hcl.placeholder((10,), "B")
C = hcl.placeholder((10,), "C")


def update_example(A, B, C):
    hcl.update(C, lambda x: A[x] + B[x], "U")


s = hcl.create_schedule([A, B, C], update_example)
print(hcl.lower(s))
```

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

### Create a Schedule for Stencil Computation

Source: https://cornell-zhang.github.io/heterocl/_downloads/7f66f80a24c0571773a47ac3b94202bf/tutorial_09_stencil.ipynb

Initializes placeholder tensors for input and output images and creates a HeteroCL schedule for the jacobi stencil computation. This setup is for using the SODA DSL.

```python
dtype = hcl.Float()
input_image = hcl.placeholder((480, 640), name="input", dtype=dtype)
output_image = hcl.placeholder((480, 640), name="output", dtype=dtype)
soda_schedule = hcl.create_schedule([input_image, output_image], jacobi)
# soda_schedule[jacobi.output].stencil()
# print(hcl.build(soda_schedule, target='soda'))
```

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

### Export Final Environment Variables for HeteroCL

Source: https://cornell-zhang.github.io/heterocl/setup/index.html

Set the final PYTHONPATH and PATH variables in the HeteroCL root folder to complete the setup.

```bash
$ export PYTHONPATH=$(pwd):${PYTHONPATH}
$ export PATH=${LLVM_BUILD_DIR}/bin:${PATH}

```

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

### Install LLVM Python Requirements

Source: https://cornell-zhang.github.io/heterocl/setup/index.html

Install necessary Python packages for LLVM's Python bindings. Ensure you are in the llvm-project directory.

```bash
$ python3 -m pip install -r mlir/python/requirements.txt

```

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

### Clone HeteroCL Repository

Source: https://cornell-zhang.github.io/heterocl/setup/index.html

Clone the HeteroCL repository and its submodules. Ensure you have git installed.

```bash
$ git clone https://github.com/cornell-zhang/heterocl.git heterocl-mlir
$ cd heterocl-mlir
$ git submodule update --init --recursive

```

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

### Initialize HeteroCL and Define Mutate Example

Source: https://cornell-zhang.github.io/heterocl/_downloads/8321e56154d9f8dcb74584a3b20dd05d/tutorial_03_api.ipynb

Initializes HeteroCL, defines placeholder tensors A, B, and C. It uses `hcl.mutate` to perform elementwise addition of A and B into C. Note that direct assignment within the lambda is forbidden by Python syntax.

```python
hcl.init()
A = hcl.placeholder((10,), "A")
B = hcl.placeholder((10,), "B")
C = hcl.placeholder((10,), "C")


def mut_example(A, B, C):
    def loop_body(x):
        C[x] = A[x] + B[x]

    hcl.mutate((10,), lambda x: loop_body(x), "M")


s = hcl.create_schedule([A, B, C], mut_example)
print(hcl.lower(s))
```

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

### Sample Input, Filter, and Output for 2D Convolution

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_06_memory.html

Provides a concrete example of the input tensor 'A', the filter tensor 'F', and the resulting output tensor 'B' after performing a 2D convolution. This helps visualize the computation's effect.

```text
Input:
array([[0, 5, 9, 2, 7, 3],
       [1, 1, 2, 6, 9, 3],
       [4, 8, 3, 9, 8, 4],
       [0, 4, 6, 3, 2, 9],
       [5, 8, 8, 5, 5, 3],
       [9, 3, 8, 7, 5, 8]])
Filter:
array([[4, 6, 3],
       [0, 6, 2],
       [1, 7, 6]])
Output:
array([[145, 187, 237, 208],
       [134, 134, 180, 214],
       [218, 213, 185, 184],
       [184, 220, 175, 177]])
```

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

### HeteroCL Imperative DSL IR Example

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_02_imperative.html

This is the Intermediate Representation (IR) generated by HeteroCL for an imperative sorting algorithm. It showcases the use of affine maps, loops, memory operations, and conditional logic.

```mlir
#map = affine_map<(d0) -> (d0 + 1)>
module {
  func.func @top(%arg0: memref<10xi32>) attributes {itypes = "s", otypes = ""} {
    %c0 = arith.constant 0 : index
    affine.for %arg1 = 0 to 9 {
      %0 = affine.apply #map(%arg1) 
      %1 = memref.alloc() {name = "key"} : memref<1xi32>
      %2 = affine.load %arg0[%0] {from = "A"} : memref<10xi32>
      affine.store %2, %1[%c0] {to = "key"} : memref<1xi32>
      %3 = memref.alloc() {name = "j"} : memref<1xi32>
      %c1_i32 = arith.constant 1 : i32
      %4 = arith.index_cast %0 : index to i34
      %5 = arith.extsi %c1_i32 : i32 to i34
      %6 = arith.subi %4, %5 : i34
      %7 = arith.trunci %6 : i34 to i32
      affine.store %7, %3[%c0] {to = "j"} : memref<1xi32>
      scf.while : () -> () {
        %true = arith.constant {unsigned} true
        %14 = affine.load %3[0] {from = "j"} : memref<1xi32>
        %c0_i32 = arith.constant 0 : i32
        %15 = arith.cmpi sge, %14, %c0_i32 : i32
        %16 = arith.andi %true, %15 {unsigned} : i1
        %17 = affine.load %1[0] {from = "key"} : memref<1xi32>
        %18 = arith.index_cast %14 {unsigned} : i32 to index
        %19 = memref.load %arg0[%18] {from = "A"} : memref<10xi32>
        %20 = arith.cmpi slt, %17, %19 : i32
        %21 = arith.andi %16, %20 {unsigned} : i1
        scf.condition(%21)
      } do {
        %14 = affine.load %3[0] {from = "j"} : memref<1xi32>
        %15 = arith.index_cast %14 {unsigned} : i32 to index
        %16 = memref.load %arg0[%15] {from = "A"} : memref<10xi32>
        %17 = arith.extsi %14 : i32 to i33
        %18 = arith.extsi %c1_i32 : i32 to i33
        %19 = arith.addi %17, %18 : i33
        %20 = arith.index_cast %19 {unsigned} : i33 to index
        memref.store %16, %arg0[%20] {to = "A"} : memref<10xi32>
        %21 = affine.load %3[0] {from = "j"} : memref<1xi32>
        %22 = arith.extsi %21 : i32 to i33
        %23 = arith.subi %22, %18 : i33
        %24 = arith.trunci %23 : i33 to i32
        affine.store %24, %3[0] {to = "j"} : memref<1xi32>
        scf.yield
      }
      %8 = affine.load %1[0] {from = "key"} : memref<1xi32>
      %9 = affine.load %3[0] {from = "j"} : memref<1xi32>
      %10 = arith.extsi %9 : i32 to i33
      %11 = arith.extsi %c1_i32 : i32 to i33
      %12 = arith.addi %10, %11 : i33
      %13 = arith.index_cast %12 {unsigned} : i33 to index
      memref.store %8, %arg0[%13] {to = "A"} : memref<10xi32>
    } {loop_name = "loop_0", op_name = "i"}
    return
  }
}
```

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

### Initialize HeteroCL and Define Placeholder

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_02_imperative.html

Initializes HeteroCL and defines a placeholder array 'A' of size 10 for subsequent computations.

```python
import heterocl as hcl

hcl.init()

A = hcl.placeholder((10,), "A")

```

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

### Build for CPU Backend (LLVM)

Source: https://cornell-zhang.github.io/heterocl/_downloads/366e024301a866e5e7272eed4f1329df/tutorial_08_backend.ipynb

Builds the HeteroCL schedule for the CPU backend. The `target="llvm"` is the default and can be omitted. Customization primitives like `partition` and `pipeline` are ignored for CPU.

```python
f = hcl.build(s)  # equivalent to hcl.build(s, target="llvm")
```

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

### Build, Run, and Verify the Hardware Module

Source: https://cornell-zhang.github.io/heterocl/_downloads/7512888384f5e641cbb69eba842024c7/tutorial_07_module.ipynb

Builds the HeteroCL computation into an executable function, runs it with sample data, and verifies the output against NumPy's maximum function.

```python
f = hcl.build(s)

a = np.random.randint(100, size=(10,))
b = np.random.randint(100, size=(10,))
c = np.random.randint(100, size=(10,))
d = np.random.randint(100, size=(10,))
o = np.zeros(10)

hcl_A = hcl.asarray(a)
hcl_B = hcl.asarray(b)
hcl_C = hcl.asarray(c)
hcl_D = hcl.asarray(d)
hcl_O = hcl.asarray(o, dtype=hcl.Int())

f(hcl_A, hcl_B, hcl_C, hcl_D, hcl_O)

print("Input tensors:")
print(hcl_A)
print(hcl_B)
print(hcl_C)
print(hcl_D)
print("Output tensor:")
print(hcl_O)

# Test the correctness
m1 = np.maximum(a, b)
m2 = np.maximum(c, d)
m = np.maximum(m1, m2)
assert np.array_equal(hcl_O.asnumpy(), m)
```

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

### Initialize HeteroCL Environment

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_01_get_started.html

Initialize the HeteroCL environment. Optionally set the default data type (defaults to 32-bit integers).

```python
hcl.init()
```

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

### hcl.compute API Example

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_03_api.html

Use hcl.compute to define a new tensor elementwise. This API returns a new tensor.

```python
hcl.init()

A = hcl.placeholder((10,), "A")
B = hcl.placeholder((10,), "B")


def compute_example(A, B):
    return hcl.compute(A.shape, lambda x: A[x] + B[x], "C")


s = hcl.create_schedule([A, B], compute_example)
print(hcl.lower(s))
```

```ir
module {
  func.func @top(%arg0: memref<10xi32>, %arg1: memref<10xi32>) -> memref<10xi32> attributes {itypes = "ss", otypes = "s"} {
    %0 = memref.alloc() {name = "C"} : memref<10xi32>
    affine.for %arg2 = 0 to 10 {
      %1 = affine.load %arg0[%arg2] {from = "A"} : memref<10xi32>
      %2 = affine.load %arg1[%arg2] {from = "B"} : memref<10xi32>
      %3 = arith.extsi %1 : i32 to i33
      %4 = arith.extsi %2 : i32 to i33
      %5 = arith.addi %3, %4 : i33
      %6 = arith.trunci %5 : i33 to i32
      affine.store %6, %0[%arg2] {to = "C"} : memref<10xi32>
    } {loop_name = "x", op_name = "C"}
    return %0 : memref<10xi32>
  }
}
```

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

### Build and Test Quantized Application

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_05_dtype.html

Builds the executable with the applied quantization scheme and tests its output against the original floating-point output. It verifies the effect of quantization.

```python
sl = hcl.create_schedule_from_scheme(sm)
f = hcl.build(sl)

hcl_BQ = hcl.asarray(np.zeros(10), dtype=hcl.Fixed(10, 8))

f(hcl_A, hcl_BQ)

np_BQ = hcl_BQ.asnumpy()

print("Without quantization")
print(np_B)
print("Quantized to Fixed(10, 8)")
print(np_BQ)
```

```text
Without quantization
[ 0.28825521 -0.38279387  0.6224227   0.67729837  0.08186562 -0.25087348
  0.13945629 -0.75332499 -0.06783565 -0.50604028]
Quantized to Fixed(10, 8)
[ 0.28515625 -0.37890625  0.62109375  0.67578125  0.078125   -0.25
  0.13671875 -0.75       -0.06640625 -0.50390625]
```

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

### hcl.update API Example

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_03_api.html

Use hcl.update to modify an existing tensor elementwise. This API does not return a new tensor; its return value is None.

```python
hcl.init()
A = hcl.placeholder((10,), "A")
B = hcl.placeholder((10,), "B")
C = hcl.placeholder((10,), "C")


def update_example(A, B, C):
    hcl.update(C, lambda x: A[x] + B[x], "U")


s = hcl.create_schedule([A, B, C], update_example)
print(hcl.lower(s))
```

```ir
module {
  func.func @top(%arg0: memref<10xi32>, %arg1: memref<10xi32>, %arg2: memref<10xi32>) attributes {itypes = "sss", otypes = ""} {
    affine.for %arg3 = 0 to 10 {
      %0 = affine.load %arg0[%arg3] {from = "A"} : memref<10xi32>
      %1 = affine.load %arg1[%arg3] {from = "B"} : memref<10xi32>
      %2 = arith.extsi %0 : i32 to i33
      %3 = arith.extsi %1 : i32 to i33
      %4 = arith.addi %2, %3 : i33
      %5 = arith.trunci %4 : i33 to i32
      affine.store %5, %arg2[%arg3] {to = "C"} : memref<10xi32>
    } {loop_name = "x", op_name = "U"}
    return
  }
}
```

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

### Build and Execute HeteroCL Program

Source: https://cornell-zhang.github.io/heterocl/_downloads/9a0f69545c1a21968941d255884310e7/tutorial_02_imperative.ipynb

Builds the executable from the schedule, creates a NumPy array, converts it to a HeteroCL array, executes the function, and converts the result back to a NumPy array for printing. This demonstrates the end-to-end workflow.

```python
f = hcl.build(s)

import numpy as np

hcl_A = hcl.asarray(np.random.randint(50, size=(10,)))

print("Before sorting:")
print(hcl_A)

f(hcl_A)

print("After sorting:")
np_A = hcl_A.asnumpy()
print(np_A)
```

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

### HeteroCL Data Types

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_05_dtype.html

Shows examples of various bit-accurate and floating-point data types supported by HeteroCL. If no argument is provided, the default bitwidth is 32.

```python
hcl.Int(15)  # 15-bit signed integer
hcl.UInt(24)  # 24-bit unsigned integer
hcl.Fixed(13, 5)  # 13-bit signed fixed point with 5 fractional bits
hcl.UFixed(44, 30)  # 44-bit unsigned fixed point with 30 fractional bits
hcl.Float(32)  # single-precision floating point
hcl.Float(64)  # double-precision floating point
```

```python
Float(64)
```

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

### Initialize HeteroCL and Define Computation

Source: https://cornell-zhang.github.io/heterocl/_downloads/cde269890b52fc0b2d1b779b005c3e15/tutorial_04_compute.ipynb

Initializes HeteroCL and defines a two-stage computation using placeholders and compute primitives. Stages must be named to be accessed by property.

```python
import heterocl as hcl
hcl.init()

A = hcl.placeholder((10, 100), "A")


def two_stage(A):
    B = hcl.compute(A.shape, lambda x, y: A[x, y] + 1, "B")
    C = hcl.compute(A.shape, lambda x, y: B[x, y] + 1, "C")
    return C


s = hcl.create_schedule([A], two_stage)
s_B = two_stage.B
```

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

### Reorder Axes with Split Primitives

Source: https://cornell-zhang.github.io/heterocl/_downloads/cde269890b52fc0b2d1b779b005c3e15/tutorial_04_compute.ipynb

Reorders axes after splitting a loop using `hcl.split`. This example demonstrates reordering the outer and inner loops created by a split operation.

```python
s = hcl.create_schedule([A], two_stage)
s[s_B].reorder(s_B.axis[1], x_out, x_in)

print(hcl.lower(s))
```

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

### Output Comparison: With and Without Window Buffer

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_06_memory.html

The printed output shows that the results are identical when computed with and without the window buffer, indicating that the reuse mechanism is functioning correctly and does not alter the final computation.

```text
Output without WB:
array([[145, 187, 237, 208],
       [134, 134, 180, 214],
       [218, 213, 185, 184],
       [184, 220, 175, 177]])
Output with WB:
array([[145, 187, 237, 208],
       [134, 134, 180, 214],
       [218, 213, 185, 184],
       [184, 220, 175, 177]])
```

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

### Create Schedule and Access Stage

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_04_compute.html

Create a HeteroCL schedule for the 'two_stage' computation and access the 'B' stage by its name.

```python
s = hcl.create_schedule([A], two_stage)
s_B = two_stage.B
```

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

### Create Schedule and Print AST/Lowered IR

Source: https://cornell-zhang.github.io/heterocl/_downloads/7512888384f5e641cbb69eba842024c7/tutorial_07_module.ipynb

Creates a HeteroCL schedule for the `maximum` function and prints its Abstract Syntax Tree (AST) and lowered Intermediate Representation (IR) for inspection.

```python
A = hcl.placeholder((10,), "A")
B = hcl.placeholder((10,), "B")
C = hcl.placeholder((10,), "C")
D = hcl.placeholder((10,), "D")

s = hcl.create_schedule([A, B, C, D], maximum)
print(s.ast)
print(hcl.lower(s))
```

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

### Create HeteroCL Schedule

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_01_get_started.html

Create a default schedule using `hcl.create_schedule`, providing the list of inputs and the algorithm function.

```python
s = hcl.create_schedule([a, A], simple_compute)
```

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

### hcl.mutate API Example

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_03_api.html

Use hcl.mutate to describe loops with vector code, useful for optimization when loop bodies lack common patterns or contain imperative DSL. Note that direct assignment statements within the lambda function are forbidden by Python syntax.

```python
hcl.init()
A = hcl.placeholder((10,), "A")
B = hcl.placeholder((10,), "B")
C = hcl.placeholder((10,), "C")


def mut_example(A, B, C):
    def loop_body(x):
        C[x] = A[x] + B[x]

    hcl.mutate((10,), lambda x: loop_body(x), "M")


s = hcl.create_schedule([A, B, C], mut_example)
print(hcl.lower(s))
```

```ir
module {
  func.func @top(%arg0: memref<10xi32>, %arg1: memref<10xi32>, %arg2: memref<10xi32>) attributes {itypes = "sss", otypes = ""} {
    affine.for %arg3 = 0 to 10 {
      %0 = affine.load %arg0[%arg3] {from = "A"} : memref<10xi32>
      %1 = affine.load %arg1[%arg3] {from = "B"} : memref<10xi32>
      %2 = arith.extsi %0 : i32 to i33
      %3 = arith.extsi %1 : i32 to i33
      %4 = arith.addi %2, %3 : i33
      %5 = arith.trunci %4 : i33 to i32
      affine.store %5, %arg2[%arg3] {to = "C"} : memref<10xi32>
    } {loop_name = "x", op_name = "M"}
    return
  }
}
```

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

### Get Bit/Slice from Number in HeteroCL

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_02_imperative.html

Demonstrates how to extract the least significant bit (LSB) and a slice of bits from an array using HeteroCL's compute API. Slicing follows Python's convention: lower bound inclusive, upper bound exclusive.

```python
hcl.init()
A = hcl.placeholder((10,), "A")


def kernel(A):
    # get the LSB of A
    B = hcl.compute(A.shape, lambda x: A[x][0], "B")
    # get the lower 4-bit of A
    C = hcl.compute(A.shape, lambda x: A[x][0:4], "C")
    return B, C

```

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

### Build for CPU Back End

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_08_backend.html

Builds the HeteroCL schedule for the CPU back end. This is the default behavior if no target is specified. Note that `partition` and `pipeline` optimizations have no effect on the CPU back end; `parallel` can be used instead.

```python
f = hcl.build(s)  # equivalent to hcl.build(s, target="llvm")

```

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

### Get LSB and Lower 4-bit Slice from Array

Source: https://cornell-zhang.github.io/heterocl/_downloads/9a0f69545c1a21968941d255884310e7/tutorial_02_imperative.ipynb

Extracts the least significant bit (LSB) and a 4-bit slice from each element of an input array A. Slicing follows Python's convention: lower bound inclusive, upper bound exclusive. Requires HeteroCL initialization and placeholder definition.

```python
hcl.init()
A = hcl.placeholder((10,), "A")


def kernel(A):
    # get the LSB of A
    B = hcl.compute(A.shape, lambda x: A[x][0], "B")
    # get the lower 4-bit of A
    C = hcl.compute(A.shape, lambda x: A[x][0:4], "C")
    return B, C
```

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

### Build and Run HeteroCL Computation

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_07_module.html

This snippet demonstrates how to build the compiled function from the schedule and then execute it with NumPy arrays. It includes converting NumPy arrays to HeteroCL arrays, running the function, and printing the input and output tensors.

```python
f = hcl.build(s)

a = np.random.randint(100, size=(10,))
b = np.random.randint(100, size=(10,))
c = np.random.randint(100, size=(10,))
d = np.random.randint(100, size=(10,))
o = np.zeros(10)

hcl_A = hcl.asarray(a)
hcl_B = hcl.asarray(b)
hcl_C = hcl.asarray(c)
hcl_D = hcl.asarray(d)
hcl_O = hcl.asarray(o, dtype=hcl.Int())

f(hcl_A, hcl_B, hcl_C, hcl_D, hcl_O)

print("Input tensors:")
print(hcl_A)
print(hcl_B)
print(hcl_C)
print(hcl_D)
print("Output tensor:")
print(hcl_O)

```

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

### Configure and Build LLVM with Python Bindings

Source: https://cornell-zhang.github.io/heterocl/setup/index.html

Configure LLVM using CMake with specific options for MLIR and Python bindings, then build it. Requires a Unix Makefiles generator.

```bash
$ mkdir build && cd build
$ cmake -G "Unix Makefiles" ../llvm \
     -DLLVM_ENABLE_PROJECTS=mlir \
     -DLLVM_BUILD_EXAMPLES=ON \
     -DLLVM_TARGETS_TO_BUILD="host" \
     -DCMAKE_BUILD_TYPE=Release \
     -DLLVM_ENABLE_ASSERTIONS=ON \
     -DLLVM_INSTALL_UTILS=ON \
     -DMLIR_ENABLE_BINDINGS_PYTHON=ON \
     -DPython3_EXECUTABLE=`which python3`
$ make -j8

```

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

### Clone and Build LLVM for Developers

Source: https://cornell-zhang.github.io/heterocl/setup/index.html

Clone the LLVM repository and checkout the required version (v15.0.0). This is for developers who need to build LLVM manually.

```bash
$ git clone https://github.com/llvm/llvm-project.git
$ cd llvm-project
$ git checkout tags/llvmorg-15.0.0

```

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

### Partition Tensor in HeteroCL

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_06_memory.html

Demonstrates how to partition a tensor using `s.partition(A)`. This primitive modifies the property of a tensor and is applied directly to the schedule.

```python
hcl.init()

A = hcl.placeholder((10, 10), "A")


def kernel(A):
    return hcl.compute(A.shape, lambda x, y: A[x][y] + 1, "B")

s = hcl.create_schedule(A, kernel)
s.partition(A)
print(hcl.lower(s))
```

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

### Further Optimization: Partitioning and Pipelining in HeteroCL

Source: https://cornell-zhang.github.io/heterocl/_downloads/cd2a6c3261bac33ba8fc1663396dae4d/tutorial_06_memory.ipynb

This code snippet shows how to further optimize the design by partitioning the linebuffer (LB) and window buffer (WB) in specific dimensions, partitioning the filter (F), and pipelining the kernel's computation axis. This aims to maximize memory bandwidth and achieve an initiation interval (II) of 1.

```python
s_final = hcl.create_schedule([A, F], kernel)
LB = s_final.reuse_at(A, s_final[kernel.B], kernel.B.axis[0], "LB")
WB = s_final.reuse_at(LB, s_final[kernel.B], kernel.B.axis[1], "WB")
s_final.partition(LB, dim=1)
s_final.partition(WB)
s_final.partition(F)
s_final[kernel.B].pipeline(kernel.B.axis[1])
print(hcl.lower(s_final))
```

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

### Build HeteroCL Executable

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_01_get_started.html

Build the executable program using `hcl.build(s)`. The default target is `llvm` for CPU execution.

```python
f = hcl.build(s)
```

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

### Quantizing to Fixed(10, 8) in HeteroCL

Source: https://cornell-zhang.github.io/heterocl/_downloads/1378a995fdcf33d9885392f73005f7d6/tutorial_05_dtype.ipynb

Demonstrates creating a schedule, building it, and applying quantization to a NumPy array using HeteroCL's `Fixed` data type. This is useful for reducing precision and memory usage in hardware.

```python
sl = hcl.create_schedule_from_scheme(sm)
f = hcl.build(sl)

hcl_BQ = hcl.asarray(np.zeros(10), dtype=hcl.Fixed(10, 8))

f(hcl_A, hcl_BQ)

np_BQ = hcl_BQ.asnumpy()

print("Without quantization")
print(np_B)
print("Quantized to Fixed(10, 8)")
print(np_BQ)
```

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

### Apply hcl.for_ to Imperative DSL

Source: https://cornell-zhang.github.io/heterocl/_downloads/cde269890b52fc0b2d1b779b005c3e15/tutorial_04_compute.ipynb

Shows how to use hcl.for_ within an imperative function and apply HeteroCL's scheduling and optimization techniques. Requires HeteroCL initialization and placeholder definition.

```python
hcl.init()

A = hcl.placeholder((10,))


def custom_imperative(A):
    with hcl.for_(0, 10, tag="S", name="i") as i:
        A[i] = i - 10


s = hcl.create_schedule([A], custom_imperative)
s_S = custom_imperative.S
i_out, i_in = s[s_S].split(s_S.i, 2)

print(hcl.lower(s))
```

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

### Run HeteroCL Executable

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_01_get_started.html

Execute a prepared HeteroCL function with input and output tensors.

```python
f(hcl_a, hcl_A, hcl_B)
```

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

### Execute and Verify Downsized Module

Source: https://cornell-zhang.github.io/heterocl/_downloads/7512888384f5e641cbb69eba842024c7/tutorial_07_module.ipynb

Execute a module with downsized data types and verify the correctness of the results, considering the modulo arithmetic introduced by downsizing.

```python
hcl_A = hcl.asarray(a, hcl.UInt(4))
hcl_B = hcl.asarray(b, hcl.UInt(4))
hcl_C = hcl.asarray(c, hcl.UInt(4))
hcl_D = hcl.asarray(d, hcl.UInt(4))
hcl_O = hcl.asarray(o)

f(hcl_A, hcl_B, hcl_C, hcl_D, hcl_O)

print("Downsized output tensor:")
print(hcl_O)
```

```python
# Test the correctness
m1 = np.maximum(a % 16, b % 16)
m2 = np.maximum(c % 16, d % 16)
m = np.maximum(m1 % 16, m2 % 16)
assert np.array_equal(hcl_O.asnumpy(), m)
```

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

### Define and Build a Custom Hardware Accelerator in HeteroCL

Source: https://cornell-zhang.github.io/heterocl/_downloads/8321e56154d9f8dcb74584a3b20dd05d/tutorial_03_api.ipynb

This snippet shows how to define a custom hardware accelerator using HeteroCL. It includes defining placeholders, a loop body for finding the two largest elements, creating a schedule, and building the optimized function. Finally, it demonstrates how to run the function with NumPy arrays and assert the correctness of the output.

```python
hcl.init()
A = hcl.placeholder((10,), "A")
M = hcl.placeholder((2,), "M")

def find_max_two(A, M):
    def loop_body(x):
        with hcl.if_(A[x] > M[0]):
            with hcl.if_(A[x] > M[1]):
                M[0] = M[1]
                M[1] = A[x]
            with hcl.else_():
                M[0] = A[x]

    hcl.mutate(A.shape, lambda x: loop_body(x))


s = hcl.create_schedule([A, M], find_max_two)
f = hcl.build(s)

import numpy as np

hcl_A = hcl.asarray(np.random.randint(50, size=(10,)))
hcl_M = hcl.asarray(np.array([-1, -1]))

f(hcl_A, hcl_M)

np_A = hcl_A.asnumpy()
np_M = hcl_M.asnumpy()

print(np_A)
print(np_M)

assert np.array_equal(np_M, np.sort(np_A)[-2:])
```

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

### Define and Build a Hardware Module with Return Statement

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_07_module.html

Defines a hardware module 'maximum' using hcl.def_ and builds it using hcl.build. This module computes the element-wise maximum of four input tensors.

```python
import heterocl as hcl
import numpy as np

a = np.array([ 4, 22, 59, 10, 80, 83, 77, 53, 39, 49])
b = np.array([70,  5, 93, 12, 51, 64, 20, 21, 87, 62])
c = np.array([18, 13, 44, 52, 12, 30, 28, 63, 16, 43])
d = np.array([23, 98, 81, 48,  8, 34, 32, 76, 26, 20])

hcl.init()
A = hcl.placeholder((10,), name='A')
B = hcl.placeholder((10,), name='B')
C = hcl.placeholder((10,), name='C')
D = hcl.placeholder((10,), name='D')
O = hcl.placeholder((10,), name='O')

@hcl.def_([A.shape, B.shape, C.shape, D.shape, O.shape])
def maximum(A, B, C, D, O):
    m1 = hcl.placeholder(A.shape, dtype=A.dtype)
    m2 = hcl.placeholder(C.shape, dtype=C.dtype)
    with hcl.for_(0, A.shape[0]) as i:
        m1[i] = hcl.maximum(A[i], B[i])
        m2[i] = hcl.maximum(C[i], D[i])
        O[i] = hcl.maximum(m1[i], m2[i])

s = hcl.create_schedule([A, B, C, D, O], maximum)
f = hcl.build(s)

hcl_A = hcl.asarray(a)
hcl_B = hcl.asarray(b)
hcl_C = hcl.asarray(c)
hcl_D = hcl.asarray(d)
hcl_O = hcl.asarray(o)

f(hcl_A, hcl_B, hcl_C, hcl_D, hcl_O)

print("Output tensor:")
print(hcl_O)

# Test the correctness
m1 = np.maximum(a, b)
m2 = np.maximum(c, d)
m = np.maximum(m1, m2)
assert np.array_equal(hcl_O.asnumpy(), m)
```

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

### Import HeteroCL and NumPy

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_09_stencil.html

Import necessary libraries for HeteroCL and numerical operations.

```python
import numpy as np

import heterocl as hcl
```

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

### Set Environment Variables for HeteroCL

Source: https://cornell-zhang.github.io/heterocl/setup/index.html

Add these lines to your .bashrc to ensure HeteroCL and LLVM are correctly configured in your environment.

```bash
$ export LLVM_BUILD_DIR=$(pwd)/hcl-dialect/externals/llvm-project/build
$ export PATH=${LLVM_BUILD_DIR}/bin:${PATH}

```

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

### Build Application Without Quantization

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_05_dtype.html

Builds an executable for a computation without applying any quantization scheme. This serves as a baseline for comparison.

```python
A = hcl.placeholder((10,))

def quantization(A):
    return hcl.compute(A.shape, lambda x: hcl.tanh(A[x]), "B")

s = hcl.create_schedule([A], quantization)
f = hcl.build(s)

import numpy as np

hcl_A = hcl.asarray(np.random.rand(10) * 2 - 1)
hcl_B = hcl.asarray(np.zeros(10))

f(hcl_A, hcl_B)

np_A = hcl_A.asnumpy()
np_B = hcl_B.asnumpy()

print("Before tanh")
print(np_A)
print("After tanh")
print(np_B)
```

```text
Before tanh
[ 0.29666233 -0.4033291   0.7289502   0.8241058   0.08204924 -0.25634474
  0.14037104 -0.98059881 -0.06793999 -0.5573926 ]
After tanh
[ 0.28825521 -0.38279387  0.6224227   0.67729837  0.08186562 -0.25087348
  0.13945629 -0.75332499 -0.06783565 -0.50604028]
```

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

### Prepare Inputs/Outputs with NumPy and asarray

Source: https://cornell-zhang.github.io/heterocl/gallery/tutorial_01_get_started.html

Prepare inputs and outputs for the executable using `hcl.asarray`. This function transforms NumPy arrays into HeteroCL containers. Random integers are generated for the input tensor `A`.

```python
import numpy as np

hcl_a = 10
np_A = np.random.randint(100, size=A.shape)
hcl_A = hcl.asarray(np_A)
hcl_B = hcl.asarray(np.zeros(A.shape))
```

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

### 2D Image Blur with Reuse Buffers in HeteroCL

Source: https://cornell-zhang.github.io/heterocl/_downloads/cd2a6c3261bac33ba8fc1663396dae4d/tutorial_06_memory.ipynb

This code demonstrates how to infer reuse buffers for explicit reduction operations in HeteroCL, specifically for a 2D image blur. It initializes HeteroCL, defines a placeholder for input data, and creates a compute kernel for the blur operation. The schedule is then created, and reuse buffers are applied to the input data along the y and x axes before printing the lower-level representation.

```python
hcl.init()
A = hcl.placeholder((10, 10), "A")

def kernel_blur(A):
    return hcl.compute(
        (8, 8), lambda y, x: A[y, x] + A[y + 1, x + 1] + A[y + 2, x + 2], "B"
    )

s_blur = hcl.create_schedule(A, kernel_blur)
B = kernel_blur.B
RB_y = s_blur.reuse_at(A, s_blur[B], B.axis[0], "RB_y")
RB_x = s_blur.reuse_at(RB_y, s_blur[B], B.axis[1], "RB_x")
print(hcl.lower(s_blur))
```