SimpleNDArray Part 1 - A Simple CUDA-Accelerated Array Library

SimpleNDArray is a zero-dependency, CUDA-accelerated array library for Python.

Caution

SimpleNDArray is in active early development. There will be bugs, API changes, and performance improvements as I implement more features.

Purpose

Previously, in my undergrad, I created the SimpleTensor library with a friend, which fully implemented automatic differentiation. However, it used NumPy/CuPy as the underlying array libraries to perform the operations. I wanted to do it completely from scratch[1].

I created this library to ensure I fully understand CUDA and deep learning performance optimizations and libraries from the ground up. This includes the strided dense array construct, Python C bindings, dynamic dispatch (a simplified version of PyTorch’s dispatch table), CUDA, and fast CUDA kernels. To make sure that I actually learn, AI usage will be limited to only the unit tests and other boring tasks (like fighting pyrefly).

I am not trying to beat PyTorch at performance or features (impossible), and I am not trying to beat tinygrad at simplicity. I simply want my own testbed to write and experiment with MLOps end-to-end. It may not be as polished as production libraries, but its purpose is to help me learn from the ground up.

Also, I personally hate how every framework nowadays takes several seconds/minutes to warm up on every run because of JIT compilation/autotuning. This is why I wanted a simple framework where you simply define the kernels you want, compile them once and cache them, and just go. Even the 5 seconds it takes to import torch has gotten infuriating. There’s also the 5GB+ of packages that need to be installed to use any of these libraries, which is far from ideal.

Roadmap

• Add important array operations. Currently, element-wise is completed, and reduction is a work in progress. I will add matmuls, convolutions, flash attention, and maybe other attention functions (linear, sparse, etc.)

• Clean up codebase (unify Buffer/BufferCuda, simpler dispatch, etc.)

• Add benchmark sweep vs. PyTorch Cuda on Ampere (maybe Hopper and/or Blackwell down the line)

• Add automatic differentiation (see SimpleTensor) using a Tensor class which wraps the Array class

• Add distributed support (MPI collective ops via NCCL/custom implementations).

• I will likely not be exploring JIT compilation/kernel generation/autotuning too much, but I am not sure yet. Making my own triton + torch.compile from scratch would likely be a much larger project than this one.

Quick Overview

        ---
title: SimpleNDArray Architecture
---
graph TB
  User(["User API"])

  subgraph ArrayLayer["Array Layer"]
    direction LR
    FC["Array.relu() .\_\_add\_\_() .arange() .reshape()"]
    A["Array
data: Buffer | BufferCuda
shape, stride, offset"]
  end

  subgraph Storage["Storage"]
    B["Buffer (CPU) - array.array"]
    BC["BufferCuda (GPU) - cudaMalloc"]
  end

  PK["Python Modules and Kernels
@module.compile_fn(...) 
def my_kernel(...)"]
  AST["AST Parse + Type Substitution"]
  RT["Runtime: .c/.cu generation
gcc/nvcc -> .so"]
  CM["Compiled Modules"]
  D["(DType, Device, Kernel) Dispatch"]
  K["Kernel Call
arange_int(out, offset, stride, n)"]

  User --> FC
  User --> PK
  FC --> A
  A --> Storage
  Storage --> D
  A --> D

  PK --> AST --> RT --> CM
  D --> CM
  CM --> K
    

SimpleNDArray provides a PyTorch-like API, but is built on fully custom CUDA kernels and a custom runtime. You can register your custom kernels entirely in Python or just use the ones provided. Operations on these Arrays will be automatically dispatched to the appropriate kernel based on the array’s device and dtype. SimpleNDArray will provide all of the standard functions like element-wise (+, -, *, /, log, exp, relu, etc.), reductions (sum, mean, max, min, etc.), products (matmul, conv, etc.), flash attention, and more. All without any external dependencies.

Related Projects

1. NumPy: This is the pioneer of the strided dense multidimensional array for Python. It features dynamic dispatch (with Array API/DLPack support), strided arrays, and bindings to OpenBLAS/LAPACK for near-speed-of-light operations.

2. CuPy: Similar to NumPy, but extends its scope to fast GPU implementations of matrix operations, a drop-in NumPy/scipy API, and distributed arrays.

3. PyTorch: Currently the de facto deep learning library. PyTorch includes features like different array layouts (dense, sparse, jagged), accelerator backends (cpu, cuda, rocm, xpu, tpu, etc.), automatic differentiation, deep learning features (modules, operations, optimizers, data loaders, etc.), and JIT compilation.

4. tinygrad: Also a simple, Python-based zero-dependency deep learning library with features similar to PyTorch/JAX, but with fully lazy tensors and JIT compilation all the way down.

5. JAX: Another deep learning/scientific computing library similar to PyTorch, but enforces functional purity and other restrictions to ensure JIT compilation can be done effectively through MLIR. Also has cool features like higher-order gradients, JVP, vmap, pmap, etc.

Getting Started: The Transpiler and Runtime

Since we are not writing a fully Python-only library, we will need a way to turn our Python code into performant C++/CUDA code. For this, I have made a Python->C/CUDA transpiler using the Python typing system.

Here is an example:

from array import array

from simplendarray import PythonModule

mod = PythonModule()


@mod.compile_fn(pybind=True)
def vector_sum(x: list[int], n: int, y: list[int]):
    y[0] = 0
    for i in range(n):
        y[0] += x[i]


mod.compile()

arr = array("i", range(6))
out = array("i", [0])
mod.vector_sum(arr.buffer_info()[0], 6, out.buffer_info()[0])
print(out)
print(sum(range(6)))

Output:

array('i', [15])
15

This Python script did the following things:

• Transpiled a simple vector_sum function into C at runtime

• Generated Python bindings automatically

• Compiled the Python C extension into a .so file

• Imported this extension at runtime

• Ran the vector sum on a Python array (the built-in Python array module, if you are not aware) and compared it to the ground truth

Let’s take a look at the generated C code:

Show C code

#define PY_SSIZE_T_CLEAN
#include <Python.h>

/* Forward declarations */
void vector_sum(int *, int, int *);

/* Transpiled C functions */
void vector_sum(int *x, int n, int *y) {
  (y[0]) = 0;
  for (int i = 0; i < n; i++) {
    (y[0]) += (x[i]);
  }
}

/* Python wrapper functions */

static PyObject *vector_sum_wrapper(PyObject *self, PyObject *args) {
  unsigned long long x;
  int n;
  unsigned long long y;
  if (!PyArg_ParseTuple(args, "KiK", &x, &n, &y)) {
    return NULL;
  }
  vector_sum((int *)x, n, (int *)y);
  Py_RETURN_NONE;
}

/* Method table */
static PyMethodDef _ndarray_rt_module_methods[] = {
    {"vector_sum", vector_sum_wrapper, METH_VARARGS,
     "Transpiled function vector_sum"},
    {NULL, NULL, 0, NULL}};

static struct PyModuleDef _ndarray_rt_module_module = {
    PyModuleDef_HEAD_INIT, "_ndarray_rt_module", NULL, -1,
    _ndarray_rt_module_methods};

PyMODINIT_FUNC PyInit__ndarray_rt_module(void) {
  return PyModule_Create(&_ndarray_rt_module_module);
}

This gives us a good starting point for writing CUDA kernels in Python without external Python DSLs like triton, CuTeDSL, helion, etc.

This transpiler uses the Python ast module to:

• Gather type hints for every variable and translate them into a C type

• Turn Python constructs (loops, conditionals, operations, etc.) into their C equivalent

• Automatically generate Python bindings to these C functions (set pybind=True)

Note that I chose the Python list type to translate directly to a pointer type because of the same indexing semantics. The arr.buffer_info()[0] gives the memory address of the start of the array as a Python int, which is converted to an unsigned long long (u64) in the Python binding, and is finally passed into the vector_sum binding as an int *.

However, this is not enough to write performant CUDA kernels. For example, how do we declare shared memory or statically sized arrays? What about function attributes like static, __global__, __device__, etc.?

The Python type system does not support this. For function attributes, the closest thing we have is async def, and we cannot add our own keywords. So for function attributes, we define them in mod.compile_fn(..., c_attrs=[...]). For variable attributes and arrays, we can use typing.Annotated. The transpiler can find these special annotations and replace them with the correct C syntax.

Finally, I needed generic support to prevent rewriting the same kernel for every single dtype. I used Python generics for this purpose. Python generics act as a string replacement as opposed to true generics. Here is a full example of all of these features being used in my CUDA element-wise kernels.

Show Python Code

from typing import Annotated, Callable

from simplendarray.dtypes import all_dtypes, all_float_dtypes, cname, ctype, i64, u64
from simplendarray.transpiler.runtime import DType, SpecItem

from ._element_wise_cuda_stubs import _ElementWiseModuleClass
from .helpers import __syncthreads, blockDim, blockIdx, printf, threadIdx

void = None

element_wise_module_cuda = _ElementWiseModuleClass(
    includes=["#include <math.h>", "#include <cuda_runtime.h>"],
    stub_path=__file__,
    stub_var="element_wise_module",
)


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_prim_add_{cname(dt)}") for dt in all_dtypes], c_attrs=["__device__"]
)
def _add[T: DType](x: T, y: T) -> T:
    return x + y


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_prim_sub_{cname(dt)}") for dt in all_dtypes], c_attrs=["__device__"]
)
def _sub[T: DType](x: T, y: T) -> T:
    return x - y


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_prim_mul_{cname(dt)}") for dt in all_dtypes], c_attrs=["__device__"]
)
def _mul[T: DType](x: T, y: T) -> T:
    return x * y


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_prim_div_{cname(dt)}") for dt in all_dtypes], c_attrs=["__device__"]
)
def _div[T: DType](x: T, y: T) -> T:
    return x / y


def atan2[T: DType](x: T, y: T) -> T: ...


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_prim_atan2_{cname(dt)}") for dt in all_float_dtypes], c_attrs=["__device__"]
)
def _atan2[T: DType](y: T, x: T) -> T:
    return atan2(y, x)


numerical_binary_kernel_specs: list[SpecItem] = []
for dt in all_dtypes:
    for op in ["add", "sub", "mul", "div", "atan2"]:
        if op == "atan2" and ctype(dt) not in ("float", "double"):
            continue
        numerical_binary_kernel_specs.append(
            SpecItem({"T": ctype(dt), "Op": f"_prim_{op}_{cname(dt)}"}, f"_{op}_{cname(dt)}")
        )


@element_wise_module_cuda.compile_fn(numerical_binary_kernel_specs, c_attrs=["__global__"])
def element_wise_binary_kernel[T: DType, Op: Callable](
    a: list[T],
    a_off: i64,
    a_stride: i64,
    b: list[T],
    b_off: i64,
    b_stride: i64,
    c: list[T],
    c_off: i64,
    c_stride: i64,
    n: i64,
) -> void:
    i: u64 = threadIdx.x + blockIdx.x * blockDim.x
    if i < n:
        c[c_off + i * c_stride] = Op(a[a_off + i * a_stride], b[b_off + i * b_stride])  # pyrefly: ignore [not-callable]


numerical_binary_specs: list[SpecItem] = []
for dt in all_dtypes:
    for op in ["add", "sub", "mul", "div", "atan2"]:
        if op == "atan2" and ctype(dt) not in ("float", "double"):
            continue
        numerical_binary_specs.append(
            SpecItem(
                {"T": ctype(dt), "Op": f"_{op}_{cname(dt)}"},
                f"element_wise_binary_{cname(dt)}__{op}",
            )
        )


@element_wise_module_cuda.compile_fn(numerical_binary_specs, pybind=True)
def element_wise_binary[T: DType, Op: Callable](
    a: list[T],
    a_off: i64,
    a_stride: i64,
    b: list[T],
    b_off: i64,
    b_stride: i64,
    c: list[T],
    c_off: i64,
    c_stride: i64,
    n: i64,
) -> None:
    NUM_THREADS: int = 1024
    NUM_BLOCKS: int = (n + NUM_THREADS - 1) // NUM_THREADS
    Op[[[NUM_BLOCKS, NUM_THREADS]]](a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n)  # pyrefly: ignore[unsupported-operation]


def _math_id[T](x: T) -> T: ...


exp = _math_id
exp2 = _math_id
log = _math_id
log2 = _math_id
log10 = _math_id
sqrt = _math_id
sin = _math_id
cos = _math_id
tan = _math_id
asin = _math_id
acos = _math_id
atan = _math_id
sinh = _math_id
cosh = _math_id
tanh = _math_id


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_exp_{cname(dt)}") for dt in all_float_dtypes], c_attrs=["__device__"]
)
def _exp[T: DType](x: T) -> T:
    return exp(x)


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_exp2_{cname(dt)}") for dt in all_float_dtypes],
    c_attrs=["__device__"],
)
def _exp2[T: DType](x: T) -> T:
    return exp2(x)


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_log_{cname(dt)}") for dt in all_float_dtypes], c_attrs=["__device__"]
)
def _log[T: DType](x: T) -> T:
    return log(x)


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_log2_{cname(dt)}") for dt in all_float_dtypes],
    c_attrs=["__device__"],
)
def _log2[T: DType](x: T) -> T:
    return log2(x)


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_log10_{cname(dt)}") for dt in all_float_dtypes],
    c_attrs=["__device__"],
)
def _log10[T: DType](x: T) -> T:
    return log10(x)


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_relu_{cname(dt)}") for dt in all_float_dtypes],
    c_attrs=["__device__"],
)
def _relu[T: DType](x: T) -> T:
    return x if x > 0 else 0  # pyrefly: ignore [bad-return,unsupported-operation]


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_square_{cname(dt)}") for dt in all_float_dtypes],
    c_attrs=["__device__"],
)
def _square[T: DType](x: T) -> T:
    return x * x


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_sqrt_{cname(dt)}") for dt in all_float_dtypes],
    c_attrs=["__device__"],
)
def _sqrt[T: DType](x: T) -> T:
    return sqrt(x)


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_sin_{cname(dt)}") for dt in all_float_dtypes], c_attrs=["__device__"]
)
def _sin[T: DType](x: T) -> T:
    return sin(x)


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_cos_{cname(dt)}") for dt in all_float_dtypes], c_attrs=["__device__"]
)
def _cos[T: DType](x: T) -> T:
    return cos(x)


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_tan_{cname(dt)}") for dt in all_float_dtypes], c_attrs=["__device__"]
)
def _tan[T: DType](x: T) -> T:
    return tan(x)


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_asin_{cname(dt)}") for dt in all_float_dtypes],
    c_attrs=["__device__"],
)
def _asin[T: DType](x: T) -> T:
    return asin(x)


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_acos_{cname(dt)}") for dt in all_float_dtypes],
    c_attrs=["__device__"],
)
def _acos[T: DType](x: T) -> T:
    return acos(x)


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_atan_{cname(dt)}") for dt in all_float_dtypes],
    c_attrs=["__device__"],
)
def _atan[T: DType](x: T) -> T:
    return atan(x)


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_sinh_{cname(dt)}") for dt in all_float_dtypes],
    c_attrs=["__device__"],
)
def _sinh[T: DType](x: T) -> T:
    return sinh(x)


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_cosh_{cname(dt)}") for dt in all_float_dtypes],
    c_attrs=["__device__"],
)
def _cosh[T: DType](x: T) -> T:
    return cosh(x)


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"_primitive_tanh_{cname(dt)}") for dt in all_float_dtypes],
    c_attrs=["__device__"],
)
def _tanh[T: DType](x: T) -> T:
    return tanh(x)


unary_ops = [
    "exp",
    "exp2",
    "log",
    "log2",
    "log10",
    "relu",
    "square",
    "sqrt",
    "sin",
    "cos",
    "tan",
    "asin",
    "acos",
    "atan",
    "sinh",
    "cosh",
    "tanh",
]

numerical_unary_kernel_specs: list[SpecItem] = []
for dt in all_float_dtypes:
    for op in unary_ops:
        numerical_unary_kernel_specs.append(
            SpecItem({"T": ctype(dt), "Op": f"_primitive_{op}_{cname(dt)}"}, f"_{op}_{cname(dt)}")
        )


@element_wise_module_cuda.compile_fn(numerical_unary_kernel_specs, c_attrs=["__global__"])
def element_wise_unary_kernel[T: DType, Op: Callable](
    a: list[T],
    a_off: i64,
    a_stride: i64,
    c: list[T],
    c_off: i64,
    c_stride: i64,
    n: i64,
) -> void:
    i: u64 = threadIdx.x + blockIdx.x * blockDim.x
    if i < n:
        c[c_off + i * c_stride] = Op(a[a_off + i * a_stride])  # pyrefly: ignore [not-callable]


numerical_unary_specs: list[SpecItem] = []
for dt in all_float_dtypes:
    for op in unary_ops:
        numerical_unary_specs.append(
            SpecItem({"T": ctype(dt), "Op": f"_{op}_{cname(dt)}"}, f"element_wise_unary_{cname(dt)}__{op}")
        )


@element_wise_module_cuda.compile_fn(numerical_unary_specs, pybind=True)
def element_wise_unary[T: DType, Op: Callable](
    a: list[T], a_off: i64, a_stride: i64, c: list[T], c_off: i64, c_stride: i64, n: i64
) -> None:
    NUM_THREADS: int = 1024
    NUM_BLOCKS: int = (n + NUM_THREADS - 1) // NUM_THREADS
    Op[[[NUM_BLOCKS, NUM_THREADS]]](a, a_off, a_stride, c, c_off, c_stride, n)  # pyrefly: ignore[unsupported-operation]


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"arange_kernel_{cname(dt)}") for dt in all_dtypes], c_attrs=["__global__"]
)
def arange_kernel[T: DType](out: list[T], out_offset: i64, out_stride: i64, numel: i64) -> void:
    i: u64 = threadIdx.x + blockIdx.x * blockDim.x
    if i < numel:
        out[out_offset + i * out_stride] = i  # pyrefly: ignore [unsupported-operation]


@element_wise_module_cuda.compile_fn(
    [SpecItem({"T": ctype(dt), "Kernel": f"arange_kernel_{cname(dt)}"}, f"arange_{cname(dt)}") for dt in all_dtypes],
    pybind=True,
)
def arange[T: DType, Kernel: Callable](out: list[T], out_offset: i64, out_stride: i64, numel: i64) -> None:
    NUM_THREADS: int = 1024
    NUM_BLOCKS: int = (numel + NUM_THREADS - 1) // NUM_THREADS
    Kernel[[[NUM_BLOCKS, NUM_THREADS]]](out, out_offset, out_stride, numel)  # pyrefly: ignore[unsupported-operation]


@element_wise_module_cuda.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"reshape_copy_kernel_{cname(dt)}") for dt in all_dtypes], c_attrs=["__global__"]
)
def reshape_copy_kernel[T: DType](
    inp: list[T],
    inp_strides: list[i64],
    inp_shape: list[i64],
    inp_offset: i64,
    inp_ndim: i64,
    out: list[T],
    out_strides: list[i64],
    out_shape: list[i64],
    out_offset: i64,
    out_ndim: i64,
    numel: i64,
) -> void:
    inp_idx: i64 = inp_offset
    out_idx: i64 = out_offset
    contiguous_inp_shape: Annotated[list[i64], "__shared__", 16] = []
    shared_inp_shape: Annotated[list[i64], "__shared__", 16] = []
    shared_inp_strides: Annotated[list[i64], "__shared__", 16] = []
    contiguous_out_shape: Annotated[list[i64], "__shared__", 16] = []
    shared_out_shape: Annotated[list[i64], "__shared__", 16] = []
    shared_out_strides: Annotated[list[i64], "__shared__", 16] = []

    i: u64 = threadIdx.x + blockIdx.x * blockDim.x
    if threadIdx.x < inp_ndim:
        shared_inp_shape[threadIdx.x] = inp_shape[threadIdx.x]
        shared_inp_strides[threadIdx.x] = inp_strides[threadIdx.x]
    if threadIdx.x < out_ndim:
        shared_out_shape[threadIdx.x] = out_shape[threadIdx.x]
        shared_out_strides[threadIdx.x] = out_strides[threadIdx.x]
    __syncthreads()
    if threadIdx.x == 0:
        contiguous_inp_shape[inp_ndim - 1] = 1
        for i_in in range(inp_ndim - 2, -1, -1):
            contiguous_inp_shape[i_in] = contiguous_inp_shape[i_in + 1] * shared_inp_shape[i_in + 1]
        contiguous_out_shape[out_ndim - 1] = 1
        for i_out in range(out_ndim - 2, -1, -1):
            contiguous_out_shape[i_out] = contiguous_out_shape[i_out + 1] * shared_out_shape[i_out + 1]
    __syncthreads()
    if i < numel:
        i_copy: u64 = i
        for i_in in range(inp_ndim):
            dim_index: u64 = i_copy // contiguous_inp_shape[i_in]
            i_copy %= contiguous_inp_shape[i_in]
            inp_idx += dim_index * shared_inp_strides[i_in]
        i_copy = i
        for i_out in range(out_ndim):
            dim_index: u64 = i_copy // contiguous_out_shape[i_out]
            i_copy %= contiguous_out_shape[i_out]
            out_idx += dim_index * shared_out_strides[i_out]
        out[out_idx] = inp[inp_idx]


@element_wise_module_cuda.compile_fn(
    [
        SpecItem({"T": ctype(dt), "Kernel": f"reshape_copy_kernel_{cname(dt)}"}, f"reshape_copy_{cname(dt)}")
        for dt in all_dtypes
    ],
    pybind=True,
)
def reshape_copy[T: DType, Kernel: Callable](
    inp: list[T],
    inp_strides: list[i64],
    inp_shape: list[i64],
    inp_index: list[i64],
    inp_offset: i64,
    inp_ndim: i64,
    out: list[T],
    out_strides: list[i64],
    out_shape: list[i64],
    out_index: list[i64],
    out_offset: i64,
    out_ndim: i64,
    numel: i64,
) -> None:
    NUM_THREADS: i64 = 1024
    NUM_BLOCKS: i64 = (numel + NUM_THREADS - 1) // NUM_THREADS
    if inp_ndim > 16 or out_ndim > 16:
        printf("Max ndim is 16 for reshape_copy on gpu")
        exit(1)
    Kernel[[[NUM_BLOCKS, NUM_THREADS]]](  # pyrefly: ignore[unsupported-operation]
        inp,
        inp_strides,
        inp_shape,
        inp_offset,
        inp_ndim,
        out,
        out_strides,
        out_shape,
        out_offset,
        out_ndim,
        numel,
    )


element_wise_module_cuda = element_wise_module_cuda.compile("nvcc", ldflags=["-lcudart"])

And the generated CUDA code:

Show CUDA code

#define PY_SSIZE_T_CLEAN
#include <Python.h>
#include <math.h>
#include <cuda_runtime.h>

/* Forward declarations */
__device__ float _prim_add_float(float, float);
__device__ double _prim_add_double(double, double);
__device__ unsigned char _prim_add_unsigned_char(unsigned char, unsigned char);
__device__ char _prim_add_char(char, char);
__device__ unsigned short _prim_add_unsigned_short(unsigned short, unsigned short);
__device__ short _prim_add_short(short, short);
__device__ unsigned int _prim_add_unsigned_int(unsigned int, unsigned int);
__device__ int _prim_add_int(int, int);
__device__ unsigned long long _prim_add_unsigned_long_long(unsigned long long, unsigned long long);
__device__ long long _prim_add_long_long(long long, long long);
__device__ float _prim_sub_float(float, float);
__device__ double _prim_sub_double(double, double);
__device__ unsigned char _prim_sub_unsigned_char(unsigned char, unsigned char);
__device__ char _prim_sub_char(char, char);
__device__ unsigned short _prim_sub_unsigned_short(unsigned short, unsigned short);
__device__ short _prim_sub_short(short, short);
__device__ unsigned int _prim_sub_unsigned_int(unsigned int, unsigned int);
__device__ int _prim_sub_int(int, int);
__device__ unsigned long long _prim_sub_unsigned_long_long(unsigned long long, unsigned long long);
__device__ long long _prim_sub_long_long(long long, long long);
__device__ float _prim_mul_float(float, float);
__device__ double _prim_mul_double(double, double);
__device__ unsigned char _prim_mul_unsigned_char(unsigned char, unsigned char);
__device__ char _prim_mul_char(char, char);
__device__ unsigned short _prim_mul_unsigned_short(unsigned short, unsigned short);
__device__ short _prim_mul_short(short, short);
__device__ unsigned int _prim_mul_unsigned_int(unsigned int, unsigned int);
__device__ int _prim_mul_int(int, int);
__device__ unsigned long long _prim_mul_unsigned_long_long(unsigned long long, unsigned long long);
__device__ long long _prim_mul_long_long(long long, long long);
__device__ float _prim_div_float(float, float);
__device__ double _prim_div_double(double, double);
__device__ unsigned char _prim_div_unsigned_char(unsigned char, unsigned char);
__device__ char _prim_div_char(char, char);
__device__ unsigned short _prim_div_unsigned_short(unsigned short, unsigned short);
__device__ short _prim_div_short(short, short);
__device__ unsigned int _prim_div_unsigned_int(unsigned int, unsigned int);
__device__ int _prim_div_int(int, int);
__device__ unsigned long long _prim_div_unsigned_long_long(unsigned long long, unsigned long long);
__device__ long long _prim_div_long_long(long long, long long);
__device__ float _prim_atan2_float(float, float);
__device__ double _prim_atan2_double(double, double);
__global__ void _add_float(float*, long long, long long, float*, long long, long long, float*, long long, long long, long long);
__global__ void _sub_float(float*, long long, long long, float*, long long, long long, float*, long long, long long, long long);
__global__ void _mul_float(float*, long long, long long, float*, long long, long long, float*, long long, long long, long long);
__global__ void _div_float(float*, long long, long long, float*, long long, long long, float*, long long, long long, long long);
__global__ void _atan2_float(float*, long long, long long, float*, long long, long long, float*, long long, long long, long long);
__global__ void _add_double(double*, long long, long long, double*, long long, long long, double*, long long, long long, long long);
__global__ void _sub_double(double*, long long, long long, double*, long long, long long, double*, long long, long long, long long);
__global__ void _mul_double(double*, long long, long long, double*, long long, long long, double*, long long, long long, long long);
__global__ void _div_double(double*, long long, long long, double*, long long, long long, double*, long long, long long, long long);
__global__ void _atan2_double(double*, long long, long long, double*, long long, long long, double*, long long, long long, long long);
__global__ void _add_unsigned_char(unsigned char*, long long, long long, unsigned char*, long long, long long, unsigned char*, long long, long long, long long);
__global__ void _sub_unsigned_char(unsigned char*, long long, long long, unsigned char*, long long, long long, unsigned char*, long long, long long, long long);
__global__ void _mul_unsigned_char(unsigned char*, long long, long long, unsigned char*, long long, long long, unsigned char*, long long, long long, long long);
__global__ void _div_unsigned_char(unsigned char*, long long, long long, unsigned char*, long long, long long, unsigned char*, long long, long long, long long);
__global__ void _add_char(char*, long long, long long, char*, long long, long long, char*, long long, long long, long long);
__global__ void _sub_char(char*, long long, long long, char*, long long, long long, char*, long long, long long, long long);
__global__ void _mul_char(char*, long long, long long, char*, long long, long long, char*, long long, long long, long long);
__global__ void _div_char(char*, long long, long long, char*, long long, long long, char*, long long, long long, long long);
__global__ void _add_unsigned_short(unsigned short*, long long, long long, unsigned short*, long long, long long, unsigned short*, long long, long long, long long);
__global__ void _sub_unsigned_short(unsigned short*, long long, long long, unsigned short*, long long, long long, unsigned short*, long long, long long, long long);
__global__ void _mul_unsigned_short(unsigned short*, long long, long long, unsigned short*, long long, long long, unsigned short*, long long, long long, long long);
__global__ void _div_unsigned_short(unsigned short*, long long, long long, unsigned short*, long long, long long, unsigned short*, long long, long long, long long);
__global__ void _add_short(short*, long long, long long, short*, long long, long long, short*, long long, long long, long long);
__global__ void _sub_short(short*, long long, long long, short*, long long, long long, short*, long long, long long, long long);
__global__ void _mul_short(short*, long long, long long, short*, long long, long long, short*, long long, long long, long long);
__global__ void _div_short(short*, long long, long long, short*, long long, long long, short*, long long, long long, long long);
__global__ void _add_unsigned_int(unsigned int*, long long, long long, unsigned int*, long long, long long, unsigned int*, long long, long long, long long);
__global__ void _sub_unsigned_int(unsigned int*, long long, long long, unsigned int*, long long, long long, unsigned int*, long long, long long, long long);
__global__ void _mul_unsigned_int(unsigned int*, long long, long long, unsigned int*, long long, long long, unsigned int*, long long, long long, long long);
__global__ void _div_unsigned_int(unsigned int*, long long, long long, unsigned int*, long long, long long, unsigned int*, long long, long long, long long);
__global__ void _add_int(int*, long long, long long, int*, long long, long long, int*, long long, long long, long long);
__global__ void _sub_int(int*, long long, long long, int*, long long, long long, int*, long long, long long, long long);
__global__ void _mul_int(int*, long long, long long, int*, long long, long long, int*, long long, long long, long long);
__global__ void _div_int(int*, long long, long long, int*, long long, long long, int*, long long, long long, long long);
__global__ void _add_unsigned_long_long(unsigned long long*, long long, long long, unsigned long long*, long long, long long, unsigned long long*, long long, long long, long long);
__global__ void _sub_unsigned_long_long(unsigned long long*, long long, long long, unsigned long long*, long long, long long, unsigned long long*, long long, long long, long long);
__global__ void _mul_unsigned_long_long(unsigned long long*, long long, long long, unsigned long long*, long long, long long, unsigned long long*, long long, long long, long long);
__global__ void _div_unsigned_long_long(unsigned long long*, long long, long long, unsigned long long*, long long, long long, unsigned long long*, long long, long long, long long);
__global__ void _add_long_long(long long*, long long, long long, long long*, long long, long long, long long*, long long, long long, long long);
__global__ void _sub_long_long(long long*, long long, long long, long long*, long long, long long, long long*, long long, long long, long long);
__global__ void _mul_long_long(long long*, long long, long long, long long*, long long, long long, long long*, long long, long long, long long);
__global__ void _div_long_long(long long*, long long, long long, long long*, long long, long long, long long*, long long, long long, long long);
void element_wise_binary_float__add(float*, long long, long long, float*, long long, long long, float*, long long, long long, long long);
void element_wise_binary_float__sub(float*, long long, long long, float*, long long, long long, float*, long long, long long, long long);
void element_wise_binary_float__mul(float*, long long, long long, float*, long long, long long, float*, long long, long long, long long);
void element_wise_binary_float__div(float*, long long, long long, float*, long long, long long, float*, long long, long long, long long);
void element_wise_binary_float__atan2(float*, long long, long long, float*, long long, long long, float*, long long, long long, long long);
void element_wise_binary_double__add(double*, long long, long long, double*, long long, long long, double*, long long, long long, long long);
void element_wise_binary_double__sub(double*, long long, long long, double*, long long, long long, double*, long long, long long, long long);
void element_wise_binary_double__mul(double*, long long, long long, double*, long long, long long, double*, long long, long long, long long);
void element_wise_binary_double__div(double*, long long, long long, double*, long long, long long, double*, long long, long long, long long);
void element_wise_binary_double__atan2(double*, long long, long long, double*, long long, long long, double*, long long, long long, long long);
void element_wise_binary_unsigned_char__add(unsigned char*, long long, long long, unsigned char*, long long, long long, unsigned char*, long long, long long, long long);
void element_wise_binary_unsigned_char__sub(unsigned char*, long long, long long, unsigned char*, long long, long long, unsigned char*, long long, long long, long long);
void element_wise_binary_unsigned_char__mul(unsigned char*, long long, long long, unsigned char*, long long, long long, unsigned char*, long long, long long, long long);
void element_wise_binary_unsigned_char__div(unsigned char*, long long, long long, unsigned char*, long long, long long, unsigned char*, long long, long long, long long);
void element_wise_binary_char__add(char*, long long, long long, char*, long long, long long, char*, long long, long long, long long);
void element_wise_binary_char__sub(char*, long long, long long, char*, long long, long long, char*, long long, long long, long long);
void element_wise_binary_char__mul(char*, long long, long long, char*, long long, long long, char*, long long, long long, long long);
void element_wise_binary_char__div(char*, long long, long long, char*, long long, long long, char*, long long, long long, long long);
void element_wise_binary_unsigned_short__add(unsigned short*, long long, long long, unsigned short*, long long, long long, unsigned short*, long long, long long, long long);
void element_wise_binary_unsigned_short__sub(unsigned short*, long long, long long, unsigned short*, long long, long long, unsigned short*, long long, long long, long long);
void element_wise_binary_unsigned_short__mul(unsigned short*, long long, long long, unsigned short*, long long, long long, unsigned short*, long long, long long, long long);
void element_wise_binary_unsigned_short__div(unsigned short*, long long, long long, unsigned short*, long long, long long, unsigned short*, long long, long long, long long);
void element_wise_binary_short__add(short*, long long, long long, short*, long long, long long, short*, long long, long long, long long);
void element_wise_binary_short__sub(short*, long long, long long, short*, long long, long long, short*, long long, long long, long long);
void element_wise_binary_short__mul(short*, long long, long long, short*, long long, long long, short*, long long, long long, long long);
void element_wise_binary_short__div(short*, long long, long long, short*, long long, long long, short*, long long, long long, long long);
void element_wise_binary_unsigned_int__add(unsigned int*, long long, long long, unsigned int*, long long, long long, unsigned int*, long long, long long, long long);
void element_wise_binary_unsigned_int__sub(unsigned int*, long long, long long, unsigned int*, long long, long long, unsigned int*, long long, long long, long long);
void element_wise_binary_unsigned_int__mul(unsigned int*, long long, long long, unsigned int*, long long, long long, unsigned int*, long long, long long, long long);
void element_wise_binary_unsigned_int__div(unsigned int*, long long, long long, unsigned int*, long long, long long, unsigned int*, long long, long long, long long);
void element_wise_binary_int__add(int*, long long, long long, int*, long long, long long, int*, long long, long long, long long);
void element_wise_binary_int__sub(int*, long long, long long, int*, long long, long long, int*, long long, long long, long long);
void element_wise_binary_int__mul(int*, long long, long long, int*, long long, long long, int*, long long, long long, long long);
void element_wise_binary_int__div(int*, long long, long long, int*, long long, long long, int*, long long, long long, long long);
void element_wise_binary_unsigned_long_long__add(unsigned long long*, long long, long long, unsigned long long*, long long, long long, unsigned long long*, long long, long long, long long);
void element_wise_binary_unsigned_long_long__sub(unsigned long long*, long long, long long, unsigned long long*, long long, long long, unsigned long long*, long long, long long, long long);
void element_wise_binary_unsigned_long_long__mul(unsigned long long*, long long, long long, unsigned long long*, long long, long long, unsigned long long*, long long, long long, long long);
void element_wise_binary_unsigned_long_long__div(unsigned long long*, long long, long long, unsigned long long*, long long, long long, unsigned long long*, long long, long long, long long);
void element_wise_binary_long_long__add(long long*, long long, long long, long long*, long long, long long, long long*, long long, long long, long long);
void element_wise_binary_long_long__sub(long long*, long long, long long, long long*, long long, long long, long long*, long long, long long, long long);
void element_wise_binary_long_long__mul(long long*, long long, long long, long long*, long long, long long, long long*, long long, long long, long long);
void element_wise_binary_long_long__div(long long*, long long, long long, long long*, long long, long long, long long*, long long, long long, long long);
__device__ float _primitive_exp_float(float);
__device__ double _primitive_exp_double(double);
__device__ float _primitive_exp2_float(float);
__device__ double _primitive_exp2_double(double);
__device__ float _primitive_log_float(float);
__device__ double _primitive_log_double(double);
__device__ float _primitive_log2_float(float);
__device__ double _primitive_log2_double(double);
__device__ float _primitive_log10_float(float);
__device__ double _primitive_log10_double(double);
__device__ float _primitive_relu_float(float);
__device__ double _primitive_relu_double(double);
__device__ float _primitive_square_float(float);
__device__ double _primitive_square_double(double);
__device__ float _primitive_sqrt_float(float);
__device__ double _primitive_sqrt_double(double);
__device__ float _primitive_sin_float(float);
__device__ double _primitive_sin_double(double);
__device__ float _primitive_cos_float(float);
__device__ double _primitive_cos_double(double);
__device__ float _primitive_tan_float(float);
__device__ double _primitive_tan_double(double);
__device__ float _primitive_asin_float(float);
__device__ double _primitive_asin_double(double);
__device__ float _primitive_acos_float(float);
__device__ double _primitive_acos_double(double);
__device__ float _primitive_atan_float(float);
__device__ double _primitive_atan_double(double);
__device__ float _primitive_sinh_float(float);
__device__ double _primitive_sinh_double(double);
__device__ float _primitive_cosh_float(float);
__device__ double _primitive_cosh_double(double);
__device__ float _primitive_tanh_float(float);
__device__ double _primitive_tanh_double(double);
__global__ void _exp_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _exp2_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _log_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _log2_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _log10_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _relu_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _square_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _sqrt_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _sin_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _cos_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _tan_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _asin_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _acos_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _atan_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _sinh_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _cosh_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _tanh_float(float*, long long, long long, float*, long long, long long, long long);
__global__ void _exp_double(double*, long long, long long, double*, long long, long long, long long);
__global__ void _exp2_double(double*, long long, long long, double*, long long, long long, long long);
__global__ void _log_double(double*, long long, long long, double*, long long, long long, long long);
__global__ void _log2_double(double*, long long, long long, double*, long long, long long, long long);
__global__ void _log10_double(double*, long long, long long, double*, long long, long long, long long);
__global__ void _relu_double(double*, long long, long long, double*, long long, long long, long long);
__global__ void _square_double(double*, long long, long long, double*, long long, long long, long long);
__global__ void _sqrt_double(double*, long long, long long, double*, long long, long long, long long);
__global__ void _sin_double(double*, long long, long long, double*, long long, long long, long long);
__global__ void _cos_double(double*, long long, long long, double*, long long, long long, long long);
__global__ void _tan_double(double*, long long, long long, double*, long long, long long, long long);
__global__ void _asin_double(double*, long long, long long, double*, long long, long long, long long);
__global__ void _acos_double(double*, long long, long long, double*, long long, long long, long long);
__global__ void _atan_double(double*, long long, long long, double*, long long, long long, long long);
__global__ void _sinh_double(double*, long long, long long, double*, long long, long long, long long);
__global__ void _cosh_double(double*, long long, long long, double*, long long, long long, long long);
__global__ void _tanh_double(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_float__exp(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_float__exp2(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_float__log(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_float__log2(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_float__log10(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_float__relu(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_float__square(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_float__sqrt(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_float__sin(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_float__cos(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_float__tan(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_float__asin(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_float__acos(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_float__atan(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_float__sinh(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_float__cosh(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_float__tanh(float*, long long, long long, float*, long long, long long, long long);
void element_wise_unary_double__exp(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_double__exp2(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_double__log(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_double__log2(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_double__log10(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_double__relu(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_double__square(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_double__sqrt(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_double__sin(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_double__cos(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_double__tan(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_double__asin(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_double__acos(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_double__atan(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_double__sinh(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_double__cosh(double*, long long, long long, double*, long long, long long, long long);
void element_wise_unary_double__tanh(double*, long long, long long, double*, long long, long long, long long);
__global__ void arange_kernel_float(float*, long long, long long, long long);
__global__ void arange_kernel_double(double*, long long, long long, long long);
__global__ void arange_kernel_unsigned_char(unsigned char*, long long, long long, long long);
__global__ void arange_kernel_char(char*, long long, long long, long long);
__global__ void arange_kernel_unsigned_short(unsigned short*, long long, long long, long long);
__global__ void arange_kernel_short(short*, long long, long long, long long);
__global__ void arange_kernel_unsigned_int(unsigned int*, long long, long long, long long);
__global__ void arange_kernel_int(int*, long long, long long, long long);
__global__ void arange_kernel_unsigned_long_long(unsigned long long*, long long, long long, long long);
__global__ void arange_kernel_long_long(long long*, long long, long long, long long);
void arange_float(float*, long long, long long, long long);
void arange_double(double*, long long, long long, long long);
void arange_unsigned_char(unsigned char*, long long, long long, long long);
void arange_char(char*, long long, long long, long long);
void arange_unsigned_short(unsigned short*, long long, long long, long long);
void arange_short(short*, long long, long long, long long);
void arange_unsigned_int(unsigned int*, long long, long long, long long);
void arange_int(int*, long long, long long, long long);
void arange_unsigned_long_long(unsigned long long*, long long, long long, long long);
void arange_long_long(long long*, long long, long long, long long);
__global__ void reshape_copy_kernel_float(float*, long long*, long long*, long long, long long, float*, long long*, long long*, long long, long long, long long);
__global__ void reshape_copy_kernel_double(double*, long long*, long long*, long long, long long, double*, long long*, long long*, long long, long long, long long);
__global__ void reshape_copy_kernel_unsigned_char(unsigned char*, long long*, long long*, long long, long long, unsigned char*, long long*, long long*, long long, long long, long long);
__global__ void reshape_copy_kernel_char(char*, long long*, long long*, long long, long long, char*, long long*, long long*, long long, long long, long long);
__global__ void reshape_copy_kernel_unsigned_short(unsigned short*, long long*, long long*, long long, long long, unsigned short*, long long*, long long*, long long, long long, long long);
__global__ void reshape_copy_kernel_short(short*, long long*, long long*, long long, long long, short*, long long*, long long*, long long, long long, long long);
__global__ void reshape_copy_kernel_unsigned_int(unsigned int*, long long*, long long*, long long, long long, unsigned int*, long long*, long long*, long long, long long, long long);
__global__ void reshape_copy_kernel_int(int*, long long*, long long*, long long, long long, int*, long long*, long long*, long long, long long, long long);
__global__ void reshape_copy_kernel_unsigned_long_long(unsigned long long*, long long*, long long*, long long, long long, unsigned long long*, long long*, long long*, long long, long long, long long);
__global__ void reshape_copy_kernel_long_long(long long*, long long*, long long*, long long, long long, long long*, long long*, long long*, long long, long long, long long);
void reshape_copy_float(float*, long long*, long long*, long long*, long long, long long, float*, long long*, long long*, long long*, long long, long long, long long);
void reshape_copy_double(double*, long long*, long long*, long long*, long long, long long, double*, long long*, long long*, long long*, long long, long long, long long);
void reshape_copy_unsigned_char(unsigned char*, long long*, long long*, long long*, long long, long long, unsigned char*, long long*, long long*, long long*, long long, long long, long long);
void reshape_copy_char(char*, long long*, long long*, long long*, long long, long long, char*, long long*, long long*, long long*, long long, long long, long long);
void reshape_copy_unsigned_short(unsigned short*, long long*, long long*, long long*, long long, long long, unsigned short*, long long*, long long*, long long*, long long, long long, long long);
void reshape_copy_short(short*, long long*, long long*, long long*, long long, long long, short*, long long*, long long*, long long*, long long, long long, long long);
void reshape_copy_unsigned_int(unsigned int*, long long*, long long*, long long*, long long, long long, unsigned int*, long long*, long long*, long long*, long long, long long, long long);
void reshape_copy_int(int*, long long*, long long*, long long*, long long, long long, int*, long long*, long long*, long long*, long long, long long, long long);
void reshape_copy_unsigned_long_long(unsigned long long*, long long*, long long*, long long*, long long, long long, unsigned long long*, long long*, long long*, long long*, long long, long long, long long);
void reshape_copy_long_long(long long*, long long*, long long*, long long*, long long, long long, long long*, long long*, long long*, long long*, long long, long long, long long);

/* Transpiled C functions */
__device__ float _prim_add_float(float x, float y) {
    return (x + y);
}

__device__ double _prim_add_double(double x, double y) {
    return (x + y);
}

__device__ unsigned char _prim_add_unsigned_char(unsigned char x, unsigned char y) {
    return (x + y);
}

__device__ char _prim_add_char(char x, char y) {
    return (x + y);
}

__device__ unsigned short _prim_add_unsigned_short(unsigned short x, unsigned short y) {
    return (x + y);
}

__device__ short _prim_add_short(short x, short y) {
    return (x + y);
}

__device__ unsigned int _prim_add_unsigned_int(unsigned int x, unsigned int y) {
    return (x + y);
}

__device__ int _prim_add_int(int x, int y) {
    return (x + y);
}

__device__ unsigned long long _prim_add_unsigned_long_long(unsigned long long x, unsigned long long y) {
    return (x + y);
}

__device__ long long _prim_add_long_long(long long x, long long y) {
    return (x + y);
}

__device__ float _prim_sub_float(float x, float y) {
    return (x - y);
}

__device__ double _prim_sub_double(double x, double y) {
    return (x - y);
}

__device__ unsigned char _prim_sub_unsigned_char(unsigned char x, unsigned char y) {
    return (x - y);
}

__device__ char _prim_sub_char(char x, char y) {
    return (x - y);
}

__device__ unsigned short _prim_sub_unsigned_short(unsigned short x, unsigned short y) {
    return (x - y);
}

__device__ short _prim_sub_short(short x, short y) {
    return (x - y);
}

__device__ unsigned int _prim_sub_unsigned_int(unsigned int x, unsigned int y) {
    return (x - y);
}

__device__ int _prim_sub_int(int x, int y) {
    return (x - y);
}

__device__ unsigned long long _prim_sub_unsigned_long_long(unsigned long long x, unsigned long long y) {
    return (x - y);
}

__device__ long long _prim_sub_long_long(long long x, long long y) {
    return (x - y);
}

__device__ float _prim_mul_float(float x, float y) {
    return (x * y);
}

__device__ double _prim_mul_double(double x, double y) {
    return (x * y);
}

__device__ unsigned char _prim_mul_unsigned_char(unsigned char x, unsigned char y) {
    return (x * y);
}

__device__ char _prim_mul_char(char x, char y) {
    return (x * y);
}

__device__ unsigned short _prim_mul_unsigned_short(unsigned short x, unsigned short y) {
    return (x * y);
}

__device__ short _prim_mul_short(short x, short y) {
    return (x * y);
}

__device__ unsigned int _prim_mul_unsigned_int(unsigned int x, unsigned int y) {
    return (x * y);
}

__device__ int _prim_mul_int(int x, int y) {
    return (x * y);
}

__device__ unsigned long long _prim_mul_unsigned_long_long(unsigned long long x, unsigned long long y) {
    return (x * y);
}

__device__ long long _prim_mul_long_long(long long x, long long y) {
    return (x * y);
}

__device__ float _prim_div_float(float x, float y) {
    return (x / y);
}

__device__ double _prim_div_double(double x, double y) {
    return (x / y);
}

__device__ unsigned char _prim_div_unsigned_char(unsigned char x, unsigned char y) {
    return (x / y);
}

__device__ char _prim_div_char(char x, char y) {
    return (x / y);
}

__device__ unsigned short _prim_div_unsigned_short(unsigned short x, unsigned short y) {
    return (x / y);
}

__device__ short _prim_div_short(short x, short y) {
    return (x / y);
}

__device__ unsigned int _prim_div_unsigned_int(unsigned int x, unsigned int y) {
    return (x / y);
}

__device__ int _prim_div_int(int x, int y) {
    return (x / y);
}

__device__ unsigned long long _prim_div_unsigned_long_long(unsigned long long x, unsigned long long y) {
    return (x / y);
}

__device__ long long _prim_div_long_long(long long x, long long y) {
    return (x / y);
}

__device__ float _prim_atan2_float(float y, float x) {
    return atan2(y, x);
}

__device__ double _prim_atan2_double(double y, double x) {
    return atan2(y, x);
}

__global__ void _add_float(float* a, long long a_off, long long a_stride, float* b, long long b_off, long long b_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_add_float((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _sub_float(float* a, long long a_off, long long a_stride, float* b, long long b_off, long long b_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_sub_float((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _mul_float(float* a, long long a_off, long long a_stride, float* b, long long b_off, long long b_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_mul_float((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _div_float(float* a, long long a_off, long long a_stride, float* b, long long b_off, long long b_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_div_float((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _atan2_float(float* a, long long a_off, long long a_stride, float* b, long long b_off, long long b_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_atan2_float((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _add_double(double* a, long long a_off, long long a_stride, double* b, long long b_off, long long b_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_add_double((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _sub_double(double* a, long long a_off, long long a_stride, double* b, long long b_off, long long b_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_sub_double((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _mul_double(double* a, long long a_off, long long a_stride, double* b, long long b_off, long long b_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_mul_double((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _div_double(double* a, long long a_off, long long a_stride, double* b, long long b_off, long long b_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_div_double((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _atan2_double(double* a, long long a_off, long long a_stride, double* b, long long b_off, long long b_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_atan2_double((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _add_unsigned_char(unsigned char* a, long long a_off, long long a_stride, unsigned char* b, long long b_off, long long b_stride, unsigned char* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_add_unsigned_char((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _sub_unsigned_char(unsigned char* a, long long a_off, long long a_stride, unsigned char* b, long long b_off, long long b_stride, unsigned char* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_sub_unsigned_char((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _mul_unsigned_char(unsigned char* a, long long a_off, long long a_stride, unsigned char* b, long long b_off, long long b_stride, unsigned char* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_mul_unsigned_char((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _div_unsigned_char(unsigned char* a, long long a_off, long long a_stride, unsigned char* b, long long b_off, long long b_stride, unsigned char* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_div_unsigned_char((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _add_char(char* a, long long a_off, long long a_stride, char* b, long long b_off, long long b_stride, char* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_add_char((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _sub_char(char* a, long long a_off, long long a_stride, char* b, long long b_off, long long b_stride, char* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_sub_char((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _mul_char(char* a, long long a_off, long long a_stride, char* b, long long b_off, long long b_stride, char* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_mul_char((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _div_char(char* a, long long a_off, long long a_stride, char* b, long long b_off, long long b_stride, char* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_div_char((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _add_unsigned_short(unsigned short* a, long long a_off, long long a_stride, unsigned short* b, long long b_off, long long b_stride, unsigned short* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_add_unsigned_short((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _sub_unsigned_short(unsigned short* a, long long a_off, long long a_stride, unsigned short* b, long long b_off, long long b_stride, unsigned short* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_sub_unsigned_short((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _mul_unsigned_short(unsigned short* a, long long a_off, long long a_stride, unsigned short* b, long long b_off, long long b_stride, unsigned short* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_mul_unsigned_short((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _div_unsigned_short(unsigned short* a, long long a_off, long long a_stride, unsigned short* b, long long b_off, long long b_stride, unsigned short* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_div_unsigned_short((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _add_short(short* a, long long a_off, long long a_stride, short* b, long long b_off, long long b_stride, short* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_add_short((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _sub_short(short* a, long long a_off, long long a_stride, short* b, long long b_off, long long b_stride, short* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_sub_short((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _mul_short(short* a, long long a_off, long long a_stride, short* b, long long b_off, long long b_stride, short* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_mul_short((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _div_short(short* a, long long a_off, long long a_stride, short* b, long long b_off, long long b_stride, short* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_div_short((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _add_unsigned_int(unsigned int* a, long long a_off, long long a_stride, unsigned int* b, long long b_off, long long b_stride, unsigned int* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_add_unsigned_int((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _sub_unsigned_int(unsigned int* a, long long a_off, long long a_stride, unsigned int* b, long long b_off, long long b_stride, unsigned int* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_sub_unsigned_int((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _mul_unsigned_int(unsigned int* a, long long a_off, long long a_stride, unsigned int* b, long long b_off, long long b_stride, unsigned int* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_mul_unsigned_int((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _div_unsigned_int(unsigned int* a, long long a_off, long long a_stride, unsigned int* b, long long b_off, long long b_stride, unsigned int* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_div_unsigned_int((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _add_int(int* a, long long a_off, long long a_stride, int* b, long long b_off, long long b_stride, int* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_add_int((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _sub_int(int* a, long long a_off, long long a_stride, int* b, long long b_off, long long b_stride, int* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_sub_int((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _mul_int(int* a, long long a_off, long long a_stride, int* b, long long b_off, long long b_stride, int* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_mul_int((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _div_int(int* a, long long a_off, long long a_stride, int* b, long long b_off, long long b_stride, int* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_div_int((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _add_unsigned_long_long(unsigned long long* a, long long a_off, long long a_stride, unsigned long long* b, long long b_off, long long b_stride, unsigned long long* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_add_unsigned_long_long((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _sub_unsigned_long_long(unsigned long long* a, long long a_off, long long a_stride, unsigned long long* b, long long b_off, long long b_stride, unsigned long long* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_sub_unsigned_long_long((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _mul_unsigned_long_long(unsigned long long* a, long long a_off, long long a_stride, unsigned long long* b, long long b_off, long long b_stride, unsigned long long* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_mul_unsigned_long_long((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _div_unsigned_long_long(unsigned long long* a, long long a_off, long long a_stride, unsigned long long* b, long long b_off, long long b_stride, unsigned long long* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_div_unsigned_long_long((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _add_long_long(long long* a, long long a_off, long long a_stride, long long* b, long long b_off, long long b_stride, long long* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_add_long_long((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _sub_long_long(long long* a, long long a_off, long long a_stride, long long* b, long long b_off, long long b_stride, long long* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_sub_long_long((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _mul_long_long(long long* a, long long a_off, long long a_stride, long long* b, long long b_off, long long b_stride, long long* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_mul_long_long((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

__global__ void _div_long_long(long long* a, long long a_off, long long a_stride, long long* b, long long b_off, long long b_stride, long long* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _prim_div_long_long((a[(a_off + (i * a_stride))]), (b[(b_off + (i * b_stride))]));
    }
}

 void element_wise_binary_float__add(float* a, long long a_off, long long a_stride, float* b, long long b_off, long long b_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _add_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_float__sub(float* a, long long a_off, long long a_stride, float* b, long long b_off, long long b_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _sub_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_float__mul(float* a, long long a_off, long long a_stride, float* b, long long b_off, long long b_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _mul_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_float__div(float* a, long long a_off, long long a_stride, float* b, long long b_off, long long b_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _div_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_float__atan2(float* a, long long a_off, long long a_stride, float* b, long long b_off, long long b_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _atan2_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_double__add(double* a, long long a_off, long long a_stride, double* b, long long b_off, long long b_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _add_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_double__sub(double* a, long long a_off, long long a_stride, double* b, long long b_off, long long b_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _sub_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_double__mul(double* a, long long a_off, long long a_stride, double* b, long long b_off, long long b_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _mul_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_double__div(double* a, long long a_off, long long a_stride, double* b, long long b_off, long long b_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _div_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_double__atan2(double* a, long long a_off, long long a_stride, double* b, long long b_off, long long b_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _atan2_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_unsigned_char__add(unsigned char* a, long long a_off, long long a_stride, unsigned char* b, long long b_off, long long b_stride, unsigned char* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _add_unsigned_char<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_unsigned_char__sub(unsigned char* a, long long a_off, long long a_stride, unsigned char* b, long long b_off, long long b_stride, unsigned char* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _sub_unsigned_char<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_unsigned_char__mul(unsigned char* a, long long a_off, long long a_stride, unsigned char* b, long long b_off, long long b_stride, unsigned char* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _mul_unsigned_char<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_unsigned_char__div(unsigned char* a, long long a_off, long long a_stride, unsigned char* b, long long b_off, long long b_stride, unsigned char* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _div_unsigned_char<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_char__add(char* a, long long a_off, long long a_stride, char* b, long long b_off, long long b_stride, char* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _add_char<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_char__sub(char* a, long long a_off, long long a_stride, char* b, long long b_off, long long b_stride, char* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _sub_char<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_char__mul(char* a, long long a_off, long long a_stride, char* b, long long b_off, long long b_stride, char* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _mul_char<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_char__div(char* a, long long a_off, long long a_stride, char* b, long long b_off, long long b_stride, char* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _div_char<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_unsigned_short__add(unsigned short* a, long long a_off, long long a_stride, unsigned short* b, long long b_off, long long b_stride, unsigned short* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _add_unsigned_short<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_unsigned_short__sub(unsigned short* a, long long a_off, long long a_stride, unsigned short* b, long long b_off, long long b_stride, unsigned short* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _sub_unsigned_short<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_unsigned_short__mul(unsigned short* a, long long a_off, long long a_stride, unsigned short* b, long long b_off, long long b_stride, unsigned short* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _mul_unsigned_short<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_unsigned_short__div(unsigned short* a, long long a_off, long long a_stride, unsigned short* b, long long b_off, long long b_stride, unsigned short* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _div_unsigned_short<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_short__add(short* a, long long a_off, long long a_stride, short* b, long long b_off, long long b_stride, short* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _add_short<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_short__sub(short* a, long long a_off, long long a_stride, short* b, long long b_off, long long b_stride, short* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _sub_short<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_short__mul(short* a, long long a_off, long long a_stride, short* b, long long b_off, long long b_stride, short* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _mul_short<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_short__div(short* a, long long a_off, long long a_stride, short* b, long long b_off, long long b_stride, short* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _div_short<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_unsigned_int__add(unsigned int* a, long long a_off, long long a_stride, unsigned int* b, long long b_off, long long b_stride, unsigned int* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _add_unsigned_int<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_unsigned_int__sub(unsigned int* a, long long a_off, long long a_stride, unsigned int* b, long long b_off, long long b_stride, unsigned int* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _sub_unsigned_int<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_unsigned_int__mul(unsigned int* a, long long a_off, long long a_stride, unsigned int* b, long long b_off, long long b_stride, unsigned int* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _mul_unsigned_int<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_unsigned_int__div(unsigned int* a, long long a_off, long long a_stride, unsigned int* b, long long b_off, long long b_stride, unsigned int* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _div_unsigned_int<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_int__add(int* a, long long a_off, long long a_stride, int* b, long long b_off, long long b_stride, int* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _add_int<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_int__sub(int* a, long long a_off, long long a_stride, int* b, long long b_off, long long b_stride, int* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _sub_int<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_int__mul(int* a, long long a_off, long long a_stride, int* b, long long b_off, long long b_stride, int* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _mul_int<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_int__div(int* a, long long a_off, long long a_stride, int* b, long long b_off, long long b_stride, int* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _div_int<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_unsigned_long_long__add(unsigned long long* a, long long a_off, long long a_stride, unsigned long long* b, long long b_off, long long b_stride, unsigned long long* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _add_unsigned_long_long<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_unsigned_long_long__sub(unsigned long long* a, long long a_off, long long a_stride, unsigned long long* b, long long b_off, long long b_stride, unsigned long long* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _sub_unsigned_long_long<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_unsigned_long_long__mul(unsigned long long* a, long long a_off, long long a_stride, unsigned long long* b, long long b_off, long long b_stride, unsigned long long* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _mul_unsigned_long_long<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_unsigned_long_long__div(unsigned long long* a, long long a_off, long long a_stride, unsigned long long* b, long long b_off, long long b_stride, unsigned long long* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _div_unsigned_long_long<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_long_long__add(long long* a, long long a_off, long long a_stride, long long* b, long long b_off, long long b_stride, long long* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _add_long_long<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_long_long__sub(long long* a, long long a_off, long long a_stride, long long* b, long long b_off, long long b_stride, long long* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _sub_long_long<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_long_long__mul(long long* a, long long a_off, long long a_stride, long long* b, long long b_off, long long b_stride, long long* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _mul_long_long<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

 void element_wise_binary_long_long__div(long long* a, long long a_off, long long a_stride, long long* b, long long b_off, long long b_stride, long long* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _div_long_long<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
}

__device__ float _primitive_exp_float(float x) {
    return exp(x);
}

__device__ double _primitive_exp_double(double x) {
    return exp(x);
}

__device__ float _primitive_exp2_float(float x) {
    return exp2(x);
}

__device__ double _primitive_exp2_double(double x) {
    return exp2(x);
}

__device__ float _primitive_log_float(float x) {
    return log(x);
}

__device__ double _primitive_log_double(double x) {
    return log(x);
}

__device__ float _primitive_log2_float(float x) {
    return log2(x);
}

__device__ double _primitive_log2_double(double x) {
    return log2(x);
}

__device__ float _primitive_log10_float(float x) {
    return log10(x);
}

__device__ double _primitive_log10_double(double x) {
    return log10(x);
}

__device__ float _primitive_relu_float(float x) {
    return ((x > 0) ? x : 0);
}

__device__ double _primitive_relu_double(double x) {
    return ((x > 0) ? x : 0);
}

__device__ float _primitive_square_float(float x) {
    return (x * x);
}

__device__ double _primitive_square_double(double x) {
    return (x * x);
}

__device__ float _primitive_sqrt_float(float x) {
    return sqrt(x);
}

__device__ double _primitive_sqrt_double(double x) {
    return sqrt(x);
}

__device__ float _primitive_sin_float(float x) {
    return sin(x);
}

__device__ double _primitive_sin_double(double x) {
    return sin(x);
}

__device__ float _primitive_cos_float(float x) {
    return cos(x);
}

__device__ double _primitive_cos_double(double x) {
    return cos(x);
}

__device__ float _primitive_tan_float(float x) {
    return tan(x);
}

__device__ double _primitive_tan_double(double x) {
    return tan(x);
}

__device__ float _primitive_asin_float(float x) {
    return asin(x);
}

__device__ double _primitive_asin_double(double x) {
    return asin(x);
}

__device__ float _primitive_acos_float(float x) {
    return acos(x);
}

__device__ double _primitive_acos_double(double x) {
    return acos(x);
}

__device__ float _primitive_atan_float(float x) {
    return atan(x);
}

__device__ double _primitive_atan_double(double x) {
    return atan(x);
}

__device__ float _primitive_sinh_float(float x) {
    return sinh(x);
}

__device__ double _primitive_sinh_double(double x) {
    return sinh(x);
}

__device__ float _primitive_cosh_float(float x) {
    return cosh(x);
}

__device__ double _primitive_cosh_double(double x) {
    return cosh(x);
}

__device__ float _primitive_tanh_float(float x) {
    return tanh(x);
}

__device__ double _primitive_tanh_double(double x) {
    return tanh(x);
}

__global__ void _exp_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_exp_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _exp2_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_exp2_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _log_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_log_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _log2_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_log2_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _log10_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_log10_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _relu_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_relu_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _square_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_square_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _sqrt_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_sqrt_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _sin_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_sin_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _cos_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_cos_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _tan_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_tan_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _asin_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_asin_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _acos_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_acos_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _atan_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_atan_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _sinh_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_sinh_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _cosh_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_cosh_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _tanh_float(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_tanh_float((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _exp_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_exp_double((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _exp2_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_exp2_double((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _log_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_log_double((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _log2_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_log2_double((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _log10_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_log10_double((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _relu_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_relu_double((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _square_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_square_double((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _sqrt_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_sqrt_double((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _sin_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_sin_double((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _cos_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_cos_double((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _tan_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_tan_double((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _asin_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_asin_double((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _acos_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_acos_double((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _atan_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_atan_double((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _sinh_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_sinh_double((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _cosh_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_cosh_double((a[(a_off + (i * a_stride))]));
    }
}

__global__ void _tanh_double(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < n)) {
        (c[(c_off + (i * c_stride))]) = _primitive_tanh_double((a[(a_off + (i * a_stride))]));
    }
}

 void element_wise_unary_float__exp(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _exp_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_float__exp2(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _exp2_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_float__log(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _log_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_float__log2(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _log2_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_float__log10(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _log10_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_float__relu(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _relu_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_float__square(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _square_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_float__sqrt(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _sqrt_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_float__sin(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _sin_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_float__cos(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _cos_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_float__tan(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _tan_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_float__asin(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _asin_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_float__acos(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _acos_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_float__atan(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _atan_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_float__sinh(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _sinh_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_float__cosh(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _cosh_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_float__tanh(float* a, long long a_off, long long a_stride, float* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _tanh_float<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__exp(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _exp_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__exp2(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _exp2_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__log(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _log_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__log2(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _log2_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__log10(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _log10_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__relu(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _relu_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__square(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _square_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__sqrt(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _sqrt_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__sin(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _sin_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__cos(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _cos_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__tan(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _tan_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__asin(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _asin_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__acos(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _acos_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__atan(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _atan_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__sinh(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _sinh_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__cosh(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _cosh_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

 void element_wise_unary_double__tanh(double* a, long long a_off, long long a_stride, double* c, long long c_off, long long c_stride, long long n) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((n + NUM_THREADS) - 1) / NUM_THREADS);
    _tanh_double<<<NUM_BLOCKS, NUM_THREADS>>>(a, a_off, a_stride, c, c_off, c_stride, n);
}

__global__ void arange_kernel_float(float* out, long long out_offset, long long out_stride, long long numel) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < numel)) {
        (out[(out_offset + (i * out_stride))]) = i;
    }
}

__global__ void arange_kernel_double(double* out, long long out_offset, long long out_stride, long long numel) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < numel)) {
        (out[(out_offset + (i * out_stride))]) = i;
    }
}

__global__ void arange_kernel_unsigned_char(unsigned char* out, long long out_offset, long long out_stride, long long numel) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < numel)) {
        (out[(out_offset + (i * out_stride))]) = i;
    }
}

__global__ void arange_kernel_char(char* out, long long out_offset, long long out_stride, long long numel) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < numel)) {
        (out[(out_offset + (i * out_stride))]) = i;
    }
}

__global__ void arange_kernel_unsigned_short(unsigned short* out, long long out_offset, long long out_stride, long long numel) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < numel)) {
        (out[(out_offset + (i * out_stride))]) = i;
    }
}

__global__ void arange_kernel_short(short* out, long long out_offset, long long out_stride, long long numel) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < numel)) {
        (out[(out_offset + (i * out_stride))]) = i;
    }
}

__global__ void arange_kernel_unsigned_int(unsigned int* out, long long out_offset, long long out_stride, long long numel) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < numel)) {
        (out[(out_offset + (i * out_stride))]) = i;
    }
}

__global__ void arange_kernel_int(int* out, long long out_offset, long long out_stride, long long numel) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < numel)) {
        (out[(out_offset + (i * out_stride))]) = i;
    }
}

__global__ void arange_kernel_unsigned_long_long(unsigned long long* out, long long out_offset, long long out_stride, long long numel) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < numel)) {
        (out[(out_offset + (i * out_stride))]) = i;
    }
}

__global__ void arange_kernel_long_long(long long* out, long long out_offset, long long out_stride, long long numel) {
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if ((i < numel)) {
        (out[(out_offset + (i * out_stride))]) = i;
    }
}

 void arange_float(float* out, long long out_offset, long long out_stride, long long numel) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    arange_kernel_float<<<NUM_BLOCKS, NUM_THREADS>>>(out, out_offset, out_stride, numel);
}

 void arange_double(double* out, long long out_offset, long long out_stride, long long numel) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    arange_kernel_double<<<NUM_BLOCKS, NUM_THREADS>>>(out, out_offset, out_stride, numel);
}

 void arange_unsigned_char(unsigned char* out, long long out_offset, long long out_stride, long long numel) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    arange_kernel_unsigned_char<<<NUM_BLOCKS, NUM_THREADS>>>(out, out_offset, out_stride, numel);
}

 void arange_char(char* out, long long out_offset, long long out_stride, long long numel) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    arange_kernel_char<<<NUM_BLOCKS, NUM_THREADS>>>(out, out_offset, out_stride, numel);
}

 void arange_unsigned_short(unsigned short* out, long long out_offset, long long out_stride, long long numel) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    arange_kernel_unsigned_short<<<NUM_BLOCKS, NUM_THREADS>>>(out, out_offset, out_stride, numel);
}

 void arange_short(short* out, long long out_offset, long long out_stride, long long numel) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    arange_kernel_short<<<NUM_BLOCKS, NUM_THREADS>>>(out, out_offset, out_stride, numel);
}

 void arange_unsigned_int(unsigned int* out, long long out_offset, long long out_stride, long long numel) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    arange_kernel_unsigned_int<<<NUM_BLOCKS, NUM_THREADS>>>(out, out_offset, out_stride, numel);
}

 void arange_int(int* out, long long out_offset, long long out_stride, long long numel) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    arange_kernel_int<<<NUM_BLOCKS, NUM_THREADS>>>(out, out_offset, out_stride, numel);
}

 void arange_unsigned_long_long(unsigned long long* out, long long out_offset, long long out_stride, long long numel) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    arange_kernel_unsigned_long_long<<<NUM_BLOCKS, NUM_THREADS>>>(out, out_offset, out_stride, numel);
}

 void arange_long_long(long long* out, long long out_offset, long long out_stride, long long numel) {
    int NUM_THREADS = 1024;
    int NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    arange_kernel_long_long<<<NUM_BLOCKS, NUM_THREADS>>>(out, out_offset, out_stride, numel);
}

__global__ void reshape_copy_kernel_float(float* inp, long long* inp_strides, long long* inp_shape, long long inp_offset, long long inp_ndim, float* out, long long* out_strides, long long* out_shape, long long out_offset, long long out_ndim, long long numel) {
    long long inp_idx = inp_offset;
    long long out_idx = out_offset;
    __shared__ long long contiguous_inp_shape[16];
    __shared__ long long shared_inp_shape[16];
    __shared__ long long shared_inp_strides[16];
    __shared__ long long contiguous_out_shape[16];
    __shared__ long long shared_out_shape[16];
    __shared__ long long shared_out_strides[16];
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if (((threadIdx.x) < inp_ndim)) {
        (shared_inp_shape[(threadIdx.x)]) = (inp_shape[(threadIdx.x)]);
        (shared_inp_strides[(threadIdx.x)]) = (inp_strides[(threadIdx.x)]);
    }
    if (((threadIdx.x) < out_ndim)) {
        (shared_out_shape[(threadIdx.x)]) = (out_shape[(threadIdx.x)]);
        (shared_out_strides[(threadIdx.x)]) = (out_strides[(threadIdx.x)]);
    }
    __syncthreads();
    if (((threadIdx.x) == 0)) {
        (contiguous_inp_shape[(inp_ndim - 1)]) = 1;
        for (int i_in = (inp_ndim - 2); i_in > (-1); i_in += (-1)) {
            (contiguous_inp_shape[i_in]) = ((contiguous_inp_shape[(i_in + 1)]) * (shared_inp_shape[(i_in + 1)]));
        }
        (contiguous_out_shape[(out_ndim - 1)]) = 1;
        for (int i_out = (out_ndim - 2); i_out > (-1); i_out += (-1)) {
            (contiguous_out_shape[i_out]) = ((contiguous_out_shape[(i_out + 1)]) * (shared_out_shape[(i_out + 1)]));
        }
    }
    __syncthreads();
    if ((i < numel)) {
        unsigned long long i_copy = i;
        for (int i_in = 0; i_in < inp_ndim; i_in++) {
            unsigned long long dim_index = (i_copy / (contiguous_inp_shape[i_in]));
            i_copy %= (contiguous_inp_shape[i_in]);
            inp_idx += (dim_index * (shared_inp_strides[i_in]));
        }
        i_copy = i;
        for (int i_out = 0; i_out < out_ndim; i_out++) {
            unsigned long long dim_index = (i_copy / (contiguous_out_shape[i_out]));
            i_copy %= (contiguous_out_shape[i_out]);
            out_idx += (dim_index * (shared_out_strides[i_out]));
        }
        (out[out_idx]) = (inp[inp_idx]);
    }
}

__global__ void reshape_copy_kernel_double(double* inp, long long* inp_strides, long long* inp_shape, long long inp_offset, long long inp_ndim, double* out, long long* out_strides, long long* out_shape, long long out_offset, long long out_ndim, long long numel) {
    long long inp_idx = inp_offset;
    long long out_idx = out_offset;
    __shared__ long long contiguous_inp_shape[16];
    __shared__ long long shared_inp_shape[16];
    __shared__ long long shared_inp_strides[16];
    __shared__ long long contiguous_out_shape[16];
    __shared__ long long shared_out_shape[16];
    __shared__ long long shared_out_strides[16];
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if (((threadIdx.x) < inp_ndim)) {
        (shared_inp_shape[(threadIdx.x)]) = (inp_shape[(threadIdx.x)]);
        (shared_inp_strides[(threadIdx.x)]) = (inp_strides[(threadIdx.x)]);
    }
    if (((threadIdx.x) < out_ndim)) {
        (shared_out_shape[(threadIdx.x)]) = (out_shape[(threadIdx.x)]);
        (shared_out_strides[(threadIdx.x)]) = (out_strides[(threadIdx.x)]);
    }
    __syncthreads();
    if (((threadIdx.x) == 0)) {
        (contiguous_inp_shape[(inp_ndim - 1)]) = 1;
        for (int i_in = (inp_ndim - 2); i_in > (-1); i_in += (-1)) {
            (contiguous_inp_shape[i_in]) = ((contiguous_inp_shape[(i_in + 1)]) * (shared_inp_shape[(i_in + 1)]));
        }
        (contiguous_out_shape[(out_ndim - 1)]) = 1;
        for (int i_out = (out_ndim - 2); i_out > (-1); i_out += (-1)) {
            (contiguous_out_shape[i_out]) = ((contiguous_out_shape[(i_out + 1)]) * (shared_out_shape[(i_out + 1)]));
        }
    }
    __syncthreads();
    if ((i < numel)) {
        unsigned long long i_copy = i;
        for (int i_in = 0; i_in < inp_ndim; i_in++) {
            unsigned long long dim_index = (i_copy / (contiguous_inp_shape[i_in]));
            i_copy %= (contiguous_inp_shape[i_in]);
            inp_idx += (dim_index * (shared_inp_strides[i_in]));
        }
        i_copy = i;
        for (int i_out = 0; i_out < out_ndim; i_out++) {
            unsigned long long dim_index = (i_copy / (contiguous_out_shape[i_out]));
            i_copy %= (contiguous_out_shape[i_out]);
            out_idx += (dim_index * (shared_out_strides[i_out]));
        }
        (out[out_idx]) = (inp[inp_idx]);
    }
}

__global__ void reshape_copy_kernel_unsigned_char(unsigned char* inp, long long* inp_strides, long long* inp_shape, long long inp_offset, long long inp_ndim, unsigned char* out, long long* out_strides, long long* out_shape, long long out_offset, long long out_ndim, long long numel) {
    long long inp_idx = inp_offset;
    long long out_idx = out_offset;
    __shared__ long long contiguous_inp_shape[16];
    __shared__ long long shared_inp_shape[16];
    __shared__ long long shared_inp_strides[16];
    __shared__ long long contiguous_out_shape[16];
    __shared__ long long shared_out_shape[16];
    __shared__ long long shared_out_strides[16];
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if (((threadIdx.x) < inp_ndim)) {
        (shared_inp_shape[(threadIdx.x)]) = (inp_shape[(threadIdx.x)]);
        (shared_inp_strides[(threadIdx.x)]) = (inp_strides[(threadIdx.x)]);
    }
    if (((threadIdx.x) < out_ndim)) {
        (shared_out_shape[(threadIdx.x)]) = (out_shape[(threadIdx.x)]);
        (shared_out_strides[(threadIdx.x)]) = (out_strides[(threadIdx.x)]);
    }
    __syncthreads();
    if (((threadIdx.x) == 0)) {
        (contiguous_inp_shape[(inp_ndim - 1)]) = 1;
        for (int i_in = (inp_ndim - 2); i_in > (-1); i_in += (-1)) {
            (contiguous_inp_shape[i_in]) = ((contiguous_inp_shape[(i_in + 1)]) * (shared_inp_shape[(i_in + 1)]));
        }
        (contiguous_out_shape[(out_ndim - 1)]) = 1;
        for (int i_out = (out_ndim - 2); i_out > (-1); i_out += (-1)) {
            (contiguous_out_shape[i_out]) = ((contiguous_out_shape[(i_out + 1)]) * (shared_out_shape[(i_out + 1)]));
        }
    }
    __syncthreads();
    if ((i < numel)) {
        unsigned long long i_copy = i;
        for (int i_in = 0; i_in < inp_ndim; i_in++) {
            unsigned long long dim_index = (i_copy / (contiguous_inp_shape[i_in]));
            i_copy %= (contiguous_inp_shape[i_in]);
            inp_idx += (dim_index * (shared_inp_strides[i_in]));
        }
        i_copy = i;
        for (int i_out = 0; i_out < out_ndim; i_out++) {
            unsigned long long dim_index = (i_copy / (contiguous_out_shape[i_out]));
            i_copy %= (contiguous_out_shape[i_out]);
            out_idx += (dim_index * (shared_out_strides[i_out]));
        }
        (out[out_idx]) = (inp[inp_idx]);
    }
}

__global__ void reshape_copy_kernel_char(char* inp, long long* inp_strides, long long* inp_shape, long long inp_offset, long long inp_ndim, char* out, long long* out_strides, long long* out_shape, long long out_offset, long long out_ndim, long long numel) {
    long long inp_idx = inp_offset;
    long long out_idx = out_offset;
    __shared__ long long contiguous_inp_shape[16];
    __shared__ long long shared_inp_shape[16];
    __shared__ long long shared_inp_strides[16];
    __shared__ long long contiguous_out_shape[16];
    __shared__ long long shared_out_shape[16];
    __shared__ long long shared_out_strides[16];
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if (((threadIdx.x) < inp_ndim)) {
        (shared_inp_shape[(threadIdx.x)]) = (inp_shape[(threadIdx.x)]);
        (shared_inp_strides[(threadIdx.x)]) = (inp_strides[(threadIdx.x)]);
    }
    if (((threadIdx.x) < out_ndim)) {
        (shared_out_shape[(threadIdx.x)]) = (out_shape[(threadIdx.x)]);
        (shared_out_strides[(threadIdx.x)]) = (out_strides[(threadIdx.x)]);
    }
    __syncthreads();
    if (((threadIdx.x) == 0)) {
        (contiguous_inp_shape[(inp_ndim - 1)]) = 1;
        for (int i_in = (inp_ndim - 2); i_in > (-1); i_in += (-1)) {
            (contiguous_inp_shape[i_in]) = ((contiguous_inp_shape[(i_in + 1)]) * (shared_inp_shape[(i_in + 1)]));
        }
        (contiguous_out_shape[(out_ndim - 1)]) = 1;
        for (int i_out = (out_ndim - 2); i_out > (-1); i_out += (-1)) {
            (contiguous_out_shape[i_out]) = ((contiguous_out_shape[(i_out + 1)]) * (shared_out_shape[(i_out + 1)]));
        }
    }
    __syncthreads();
    if ((i < numel)) {
        unsigned long long i_copy = i;
        for (int i_in = 0; i_in < inp_ndim; i_in++) {
            unsigned long long dim_index = (i_copy / (contiguous_inp_shape[i_in]));
            i_copy %= (contiguous_inp_shape[i_in]);
            inp_idx += (dim_index * (shared_inp_strides[i_in]));
        }
        i_copy = i;
        for (int i_out = 0; i_out < out_ndim; i_out++) {
            unsigned long long dim_index = (i_copy / (contiguous_out_shape[i_out]));
            i_copy %= (contiguous_out_shape[i_out]);
            out_idx += (dim_index * (shared_out_strides[i_out]));
        }
        (out[out_idx]) = (inp[inp_idx]);
    }
}

__global__ void reshape_copy_kernel_unsigned_short(unsigned short* inp, long long* inp_strides, long long* inp_shape, long long inp_offset, long long inp_ndim, unsigned short* out, long long* out_strides, long long* out_shape, long long out_offset, long long out_ndim, long long numel) {
    long long inp_idx = inp_offset;
    long long out_idx = out_offset;
    __shared__ long long contiguous_inp_shape[16];
    __shared__ long long shared_inp_shape[16];
    __shared__ long long shared_inp_strides[16];
    __shared__ long long contiguous_out_shape[16];
    __shared__ long long shared_out_shape[16];
    __shared__ long long shared_out_strides[16];
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if (((threadIdx.x) < inp_ndim)) {
        (shared_inp_shape[(threadIdx.x)]) = (inp_shape[(threadIdx.x)]);
        (shared_inp_strides[(threadIdx.x)]) = (inp_strides[(threadIdx.x)]);
    }
    if (((threadIdx.x) < out_ndim)) {
        (shared_out_shape[(threadIdx.x)]) = (out_shape[(threadIdx.x)]);
        (shared_out_strides[(threadIdx.x)]) = (out_strides[(threadIdx.x)]);
    }
    __syncthreads();
    if (((threadIdx.x) == 0)) {
        (contiguous_inp_shape[(inp_ndim - 1)]) = 1;
        for (int i_in = (inp_ndim - 2); i_in > (-1); i_in += (-1)) {
            (contiguous_inp_shape[i_in]) = ((contiguous_inp_shape[(i_in + 1)]) * (shared_inp_shape[(i_in + 1)]));
        }
        (contiguous_out_shape[(out_ndim - 1)]) = 1;
        for (int i_out = (out_ndim - 2); i_out > (-1); i_out += (-1)) {
            (contiguous_out_shape[i_out]) = ((contiguous_out_shape[(i_out + 1)]) * (shared_out_shape[(i_out + 1)]));
        }
    }
    __syncthreads();
    if ((i < numel)) {
        unsigned long long i_copy = i;
        for (int i_in = 0; i_in < inp_ndim; i_in++) {
            unsigned long long dim_index = (i_copy / (contiguous_inp_shape[i_in]));
            i_copy %= (contiguous_inp_shape[i_in]);
            inp_idx += (dim_index * (shared_inp_strides[i_in]));
        }
        i_copy = i;
        for (int i_out = 0; i_out < out_ndim; i_out++) {
            unsigned long long dim_index = (i_copy / (contiguous_out_shape[i_out]));
            i_copy %= (contiguous_out_shape[i_out]);
            out_idx += (dim_index * (shared_out_strides[i_out]));
        }
        (out[out_idx]) = (inp[inp_idx]);
    }
}

__global__ void reshape_copy_kernel_short(short* inp, long long* inp_strides, long long* inp_shape, long long inp_offset, long long inp_ndim, short* out, long long* out_strides, long long* out_shape, long long out_offset, long long out_ndim, long long numel) {
    long long inp_idx = inp_offset;
    long long out_idx = out_offset;
    __shared__ long long contiguous_inp_shape[16];
    __shared__ long long shared_inp_shape[16];
    __shared__ long long shared_inp_strides[16];
    __shared__ long long contiguous_out_shape[16];
    __shared__ long long shared_out_shape[16];
    __shared__ long long shared_out_strides[16];
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if (((threadIdx.x) < inp_ndim)) {
        (shared_inp_shape[(threadIdx.x)]) = (inp_shape[(threadIdx.x)]);
        (shared_inp_strides[(threadIdx.x)]) = (inp_strides[(threadIdx.x)]);
    }
    if (((threadIdx.x) < out_ndim)) {
        (shared_out_shape[(threadIdx.x)]) = (out_shape[(threadIdx.x)]);
        (shared_out_strides[(threadIdx.x)]) = (out_strides[(threadIdx.x)]);
    }
    __syncthreads();
    if (((threadIdx.x) == 0)) {
        (contiguous_inp_shape[(inp_ndim - 1)]) = 1;
        for (int i_in = (inp_ndim - 2); i_in > (-1); i_in += (-1)) {
            (contiguous_inp_shape[i_in]) = ((contiguous_inp_shape[(i_in + 1)]) * (shared_inp_shape[(i_in + 1)]));
        }
        (contiguous_out_shape[(out_ndim - 1)]) = 1;
        for (int i_out = (out_ndim - 2); i_out > (-1); i_out += (-1)) {
            (contiguous_out_shape[i_out]) = ((contiguous_out_shape[(i_out + 1)]) * (shared_out_shape[(i_out + 1)]));
        }
    }
    __syncthreads();
    if ((i < numel)) {
        unsigned long long i_copy = i;
        for (int i_in = 0; i_in < inp_ndim; i_in++) {
            unsigned long long dim_index = (i_copy / (contiguous_inp_shape[i_in]));
            i_copy %= (contiguous_inp_shape[i_in]);
            inp_idx += (dim_index * (shared_inp_strides[i_in]));
        }
        i_copy = i;
        for (int i_out = 0; i_out < out_ndim; i_out++) {
            unsigned long long dim_index = (i_copy / (contiguous_out_shape[i_out]));
            i_copy %= (contiguous_out_shape[i_out]);
            out_idx += (dim_index * (shared_out_strides[i_out]));
        }
        (out[out_idx]) = (inp[inp_idx]);
    }
}

__global__ void reshape_copy_kernel_unsigned_int(unsigned int* inp, long long* inp_strides, long long* inp_shape, long long inp_offset, long long inp_ndim, unsigned int* out, long long* out_strides, long long* out_shape, long long out_offset, long long out_ndim, long long numel) {
    long long inp_idx = inp_offset;
    long long out_idx = out_offset;
    __shared__ long long contiguous_inp_shape[16];
    __shared__ long long shared_inp_shape[16];
    __shared__ long long shared_inp_strides[16];
    __shared__ long long contiguous_out_shape[16];
    __shared__ long long shared_out_shape[16];
    __shared__ long long shared_out_strides[16];
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if (((threadIdx.x) < inp_ndim)) {
        (shared_inp_shape[(threadIdx.x)]) = (inp_shape[(threadIdx.x)]);
        (shared_inp_strides[(threadIdx.x)]) = (inp_strides[(threadIdx.x)]);
    }
    if (((threadIdx.x) < out_ndim)) {
        (shared_out_shape[(threadIdx.x)]) = (out_shape[(threadIdx.x)]);
        (shared_out_strides[(threadIdx.x)]) = (out_strides[(threadIdx.x)]);
    }
    __syncthreads();
    if (((threadIdx.x) == 0)) {
        (contiguous_inp_shape[(inp_ndim - 1)]) = 1;
        for (int i_in = (inp_ndim - 2); i_in > (-1); i_in += (-1)) {
            (contiguous_inp_shape[i_in]) = ((contiguous_inp_shape[(i_in + 1)]) * (shared_inp_shape[(i_in + 1)]));
        }
        (contiguous_out_shape[(out_ndim - 1)]) = 1;
        for (int i_out = (out_ndim - 2); i_out > (-1); i_out += (-1)) {
            (contiguous_out_shape[i_out]) = ((contiguous_out_shape[(i_out + 1)]) * (shared_out_shape[(i_out + 1)]));
        }
    }
    __syncthreads();
    if ((i < numel)) {
        unsigned long long i_copy = i;
        for (int i_in = 0; i_in < inp_ndim; i_in++) {
            unsigned long long dim_index = (i_copy / (contiguous_inp_shape[i_in]));
            i_copy %= (contiguous_inp_shape[i_in]);
            inp_idx += (dim_index * (shared_inp_strides[i_in]));
        }
        i_copy = i;
        for (int i_out = 0; i_out < out_ndim; i_out++) {
            unsigned long long dim_index = (i_copy / (contiguous_out_shape[i_out]));
            i_copy %= (contiguous_out_shape[i_out]);
            out_idx += (dim_index * (shared_out_strides[i_out]));
        }
        (out[out_idx]) = (inp[inp_idx]);
    }
}

__global__ void reshape_copy_kernel_int(int* inp, long long* inp_strides, long long* inp_shape, long long inp_offset, long long inp_ndim, int* out, long long* out_strides, long long* out_shape, long long out_offset, long long out_ndim, long long numel) {
    long long inp_idx = inp_offset;
    long long out_idx = out_offset;
    __shared__ long long contiguous_inp_shape[16];
    __shared__ long long shared_inp_shape[16];
    __shared__ long long shared_inp_strides[16];
    __shared__ long long contiguous_out_shape[16];
    __shared__ long long shared_out_shape[16];
    __shared__ long long shared_out_strides[16];
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if (((threadIdx.x) < inp_ndim)) {
        (shared_inp_shape[(threadIdx.x)]) = (inp_shape[(threadIdx.x)]);
        (shared_inp_strides[(threadIdx.x)]) = (inp_strides[(threadIdx.x)]);
    }
    if (((threadIdx.x) < out_ndim)) {
        (shared_out_shape[(threadIdx.x)]) = (out_shape[(threadIdx.x)]);
        (shared_out_strides[(threadIdx.x)]) = (out_strides[(threadIdx.x)]);
    }
    __syncthreads();
    if (((threadIdx.x) == 0)) {
        (contiguous_inp_shape[(inp_ndim - 1)]) = 1;
        for (int i_in = (inp_ndim - 2); i_in > (-1); i_in += (-1)) {
            (contiguous_inp_shape[i_in]) = ((contiguous_inp_shape[(i_in + 1)]) * (shared_inp_shape[(i_in + 1)]));
        }
        (contiguous_out_shape[(out_ndim - 1)]) = 1;
        for (int i_out = (out_ndim - 2); i_out > (-1); i_out += (-1)) {
            (contiguous_out_shape[i_out]) = ((contiguous_out_shape[(i_out + 1)]) * (shared_out_shape[(i_out + 1)]));
        }
    }
    __syncthreads();
    if ((i < numel)) {
        unsigned long long i_copy = i;
        for (int i_in = 0; i_in < inp_ndim; i_in++) {
            unsigned long long dim_index = (i_copy / (contiguous_inp_shape[i_in]));
            i_copy %= (contiguous_inp_shape[i_in]);
            inp_idx += (dim_index * (shared_inp_strides[i_in]));
        }
        i_copy = i;
        for (int i_out = 0; i_out < out_ndim; i_out++) {
            unsigned long long dim_index = (i_copy / (contiguous_out_shape[i_out]));
            i_copy %= (contiguous_out_shape[i_out]);
            out_idx += (dim_index * (shared_out_strides[i_out]));
        }
        (out[out_idx]) = (inp[inp_idx]);
    }
}

__global__ void reshape_copy_kernel_unsigned_long_long(unsigned long long* inp, long long* inp_strides, long long* inp_shape, long long inp_offset, long long inp_ndim, unsigned long long* out, long long* out_strides, long long* out_shape, long long out_offset, long long out_ndim, long long numel) {
    long long inp_idx = inp_offset;
    long long out_idx = out_offset;
    __shared__ long long contiguous_inp_shape[16];
    __shared__ long long shared_inp_shape[16];
    __shared__ long long shared_inp_strides[16];
    __shared__ long long contiguous_out_shape[16];
    __shared__ long long shared_out_shape[16];
    __shared__ long long shared_out_strides[16];
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if (((threadIdx.x) < inp_ndim)) {
        (shared_inp_shape[(threadIdx.x)]) = (inp_shape[(threadIdx.x)]);
        (shared_inp_strides[(threadIdx.x)]) = (inp_strides[(threadIdx.x)]);
    }
    if (((threadIdx.x) < out_ndim)) {
        (shared_out_shape[(threadIdx.x)]) = (out_shape[(threadIdx.x)]);
        (shared_out_strides[(threadIdx.x)]) = (out_strides[(threadIdx.x)]);
    }
    __syncthreads();
    if (((threadIdx.x) == 0)) {
        (contiguous_inp_shape[(inp_ndim - 1)]) = 1;
        for (int i_in = (inp_ndim - 2); i_in > (-1); i_in += (-1)) {
            (contiguous_inp_shape[i_in]) = ((contiguous_inp_shape[(i_in + 1)]) * (shared_inp_shape[(i_in + 1)]));
        }
        (contiguous_out_shape[(out_ndim - 1)]) = 1;
        for (int i_out = (out_ndim - 2); i_out > (-1); i_out += (-1)) {
            (contiguous_out_shape[i_out]) = ((contiguous_out_shape[(i_out + 1)]) * (shared_out_shape[(i_out + 1)]));
        }
    }
    __syncthreads();
    if ((i < numel)) {
        unsigned long long i_copy = i;
        for (int i_in = 0; i_in < inp_ndim; i_in++) {
            unsigned long long dim_index = (i_copy / (contiguous_inp_shape[i_in]));
            i_copy %= (contiguous_inp_shape[i_in]);
            inp_idx += (dim_index * (shared_inp_strides[i_in]));
        }
        i_copy = i;
        for (int i_out = 0; i_out < out_ndim; i_out++) {
            unsigned long long dim_index = (i_copy / (contiguous_out_shape[i_out]));
            i_copy %= (contiguous_out_shape[i_out]);
            out_idx += (dim_index * (shared_out_strides[i_out]));
        }
        (out[out_idx]) = (inp[inp_idx]);
    }
}

__global__ void reshape_copy_kernel_long_long(long long* inp, long long* inp_strides, long long* inp_shape, long long inp_offset, long long inp_ndim, long long* out, long long* out_strides, long long* out_shape, long long out_offset, long long out_ndim, long long numel) {
    long long inp_idx = inp_offset;
    long long out_idx = out_offset;
    __shared__ long long contiguous_inp_shape[16];
    __shared__ long long shared_inp_shape[16];
    __shared__ long long shared_inp_strides[16];
    __shared__ long long contiguous_out_shape[16];
    __shared__ long long shared_out_shape[16];
    __shared__ long long shared_out_strides[16];
    unsigned long long i = ((threadIdx.x) + ((blockIdx.x) * (blockDim.x)));
    if (((threadIdx.x) < inp_ndim)) {
        (shared_inp_shape[(threadIdx.x)]) = (inp_shape[(threadIdx.x)]);
        (shared_inp_strides[(threadIdx.x)]) = (inp_strides[(threadIdx.x)]);
    }
    if (((threadIdx.x) < out_ndim)) {
        (shared_out_shape[(threadIdx.x)]) = (out_shape[(threadIdx.x)]);
        (shared_out_strides[(threadIdx.x)]) = (out_strides[(threadIdx.x)]);
    }
    __syncthreads();
    if (((threadIdx.x) == 0)) {
        (contiguous_inp_shape[(inp_ndim - 1)]) = 1;
        for (int i_in = (inp_ndim - 2); i_in > (-1); i_in += (-1)) {
            (contiguous_inp_shape[i_in]) = ((contiguous_inp_shape[(i_in + 1)]) * (shared_inp_shape[(i_in + 1)]));
        }
        (contiguous_out_shape[(out_ndim - 1)]) = 1;
        for (int i_out = (out_ndim - 2); i_out > (-1); i_out += (-1)) {
            (contiguous_out_shape[i_out]) = ((contiguous_out_shape[(i_out + 1)]) * (shared_out_shape[(i_out + 1)]));
        }
    }
    __syncthreads();
    if ((i < numel)) {
        unsigned long long i_copy = i;
        for (int i_in = 0; i_in < inp_ndim; i_in++) {
            unsigned long long dim_index = (i_copy / (contiguous_inp_shape[i_in]));
            i_copy %= (contiguous_inp_shape[i_in]);
            inp_idx += (dim_index * (shared_inp_strides[i_in]));
        }
        i_copy = i;
        for (int i_out = 0; i_out < out_ndim; i_out++) {
            unsigned long long dim_index = (i_copy / (contiguous_out_shape[i_out]));
            i_copy %= (contiguous_out_shape[i_out]);
            out_idx += (dim_index * (shared_out_strides[i_out]));
        }
        (out[out_idx]) = (inp[inp_idx]);
    }
}

 void reshape_copy_float(float* inp, long long* inp_strides, long long* inp_shape, long long* inp_index, long long inp_offset, long long inp_ndim, float* out, long long* out_strides, long long* out_shape, long long* out_index, long long out_offset, long long out_ndim, long long numel) {
    long long NUM_THREADS = 1024;
    long long NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    if (((inp_ndim > 16) || (out_ndim > 16))) {
        printf("Max ndim is 16 for reshape_copy on gpu");
        exit(1);
    }
    reshape_copy_kernel_float<<<NUM_BLOCKS, NUM_THREADS>>>(inp, inp_strides, inp_shape, inp_offset, inp_ndim, out, out_strides, out_shape, out_offset, out_ndim, numel);
}

 void reshape_copy_double(double* inp, long long* inp_strides, long long* inp_shape, long long* inp_index, long long inp_offset, long long inp_ndim, double* out, long long* out_strides, long long* out_shape, long long* out_index, long long out_offset, long long out_ndim, long long numel) {
    long long NUM_THREADS = 1024;
    long long NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    if (((inp_ndim > 16) || (out_ndim > 16))) {
        printf("Max ndim is 16 for reshape_copy on gpu");
        exit(1);
    }
    reshape_copy_kernel_double<<<NUM_BLOCKS, NUM_THREADS>>>(inp, inp_strides, inp_shape, inp_offset, inp_ndim, out, out_strides, out_shape, out_offset, out_ndim, numel);
}

 void reshape_copy_unsigned_char(unsigned char* inp, long long* inp_strides, long long* inp_shape, long long* inp_index, long long inp_offset, long long inp_ndim, unsigned char* out, long long* out_strides, long long* out_shape, long long* out_index, long long out_offset, long long out_ndim, long long numel) {
    long long NUM_THREADS = 1024;
    long long NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    if (((inp_ndim > 16) || (out_ndim > 16))) {
        printf("Max ndim is 16 for reshape_copy on gpu");
        exit(1);
    }
    reshape_copy_kernel_unsigned_char<<<NUM_BLOCKS, NUM_THREADS>>>(inp, inp_strides, inp_shape, inp_offset, inp_ndim, out, out_strides, out_shape, out_offset, out_ndim, numel);
}

 void reshape_copy_char(char* inp, long long* inp_strides, long long* inp_shape, long long* inp_index, long long inp_offset, long long inp_ndim, char* out, long long* out_strides, long long* out_shape, long long* out_index, long long out_offset, long long out_ndim, long long numel) {
    long long NUM_THREADS = 1024;
    long long NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    if (((inp_ndim > 16) || (out_ndim > 16))) {
        printf("Max ndim is 16 for reshape_copy on gpu");
        exit(1);
    }
    reshape_copy_kernel_char<<<NUM_BLOCKS, NUM_THREADS>>>(inp, inp_strides, inp_shape, inp_offset, inp_ndim, out, out_strides, out_shape, out_offset, out_ndim, numel);
}

 void reshape_copy_unsigned_short(unsigned short* inp, long long* inp_strides, long long* inp_shape, long long* inp_index, long long inp_offset, long long inp_ndim, unsigned short* out, long long* out_strides, long long* out_shape, long long* out_index, long long out_offset, long long out_ndim, long long numel) {
    long long NUM_THREADS = 1024;
    long long NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    if (((inp_ndim > 16) || (out_ndim > 16))) {
        printf("Max ndim is 16 for reshape_copy on gpu");
        exit(1);
    }
    reshape_copy_kernel_unsigned_short<<<NUM_BLOCKS, NUM_THREADS>>>(inp, inp_strides, inp_shape, inp_offset, inp_ndim, out, out_strides, out_shape, out_offset, out_ndim, numel);
}

 void reshape_copy_short(short* inp, long long* inp_strides, long long* inp_shape, long long* inp_index, long long inp_offset, long long inp_ndim, short* out, long long* out_strides, long long* out_shape, long long* out_index, long long out_offset, long long out_ndim, long long numel) {
    long long NUM_THREADS = 1024;
    long long NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    if (((inp_ndim > 16) || (out_ndim > 16))) {
        printf("Max ndim is 16 for reshape_copy on gpu");
        exit(1);
    }
    reshape_copy_kernel_short<<<NUM_BLOCKS, NUM_THREADS>>>(inp, inp_strides, inp_shape, inp_offset, inp_ndim, out, out_strides, out_shape, out_offset, out_ndim, numel);
}

 void reshape_copy_unsigned_int(unsigned int* inp, long long* inp_strides, long long* inp_shape, long long* inp_index, long long inp_offset, long long inp_ndim, unsigned int* out, long long* out_strides, long long* out_shape, long long* out_index, long long out_offset, long long out_ndim, long long numel) {
    long long NUM_THREADS = 1024;
    long long NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    if (((inp_ndim > 16) || (out_ndim > 16))) {
        printf("Max ndim is 16 for reshape_copy on gpu");
        exit(1);
    }
    reshape_copy_kernel_unsigned_int<<<NUM_BLOCKS, NUM_THREADS>>>(inp, inp_strides, inp_shape, inp_offset, inp_ndim, out, out_strides, out_shape, out_offset, out_ndim, numel);
}

 void reshape_copy_int(int* inp, long long* inp_strides, long long* inp_shape, long long* inp_index, long long inp_offset, long long inp_ndim, int* out, long long* out_strides, long long* out_shape, long long* out_index, long long out_offset, long long out_ndim, long long numel) {
    long long NUM_THREADS = 1024;
    long long NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    if (((inp_ndim > 16) || (out_ndim > 16))) {
        printf("Max ndim is 16 for reshape_copy on gpu");
        exit(1);
    }
    reshape_copy_kernel_int<<<NUM_BLOCKS, NUM_THREADS>>>(inp, inp_strides, inp_shape, inp_offset, inp_ndim, out, out_strides, out_shape, out_offset, out_ndim, numel);
}

 void reshape_copy_unsigned_long_long(unsigned long long* inp, long long* inp_strides, long long* inp_shape, long long* inp_index, long long inp_offset, long long inp_ndim, unsigned long long* out, long long* out_strides, long long* out_shape, long long* out_index, long long out_offset, long long out_ndim, long long numel) {
    long long NUM_THREADS = 1024;
    long long NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    if (((inp_ndim > 16) || (out_ndim > 16))) {
        printf("Max ndim is 16 for reshape_copy on gpu");
        exit(1);
    }
    reshape_copy_kernel_unsigned_long_long<<<NUM_BLOCKS, NUM_THREADS>>>(inp, inp_strides, inp_shape, inp_offset, inp_ndim, out, out_strides, out_shape, out_offset, out_ndim, numel);
}

 void reshape_copy_long_long(long long* inp, long long* inp_strides, long long* inp_shape, long long* inp_index, long long inp_offset, long long inp_ndim, long long* out, long long* out_strides, long long* out_shape, long long* out_index, long long out_offset, long long out_ndim, long long numel) {
    long long NUM_THREADS = 1024;
    long long NUM_BLOCKS = (((numel + NUM_THREADS) - 1) / NUM_THREADS);
    if (((inp_ndim > 16) || (out_ndim > 16))) {
        printf("Max ndim is 16 for reshape_copy on gpu");
        exit(1);
    }
    reshape_copy_kernel_long_long<<<NUM_BLOCKS, NUM_THREADS>>>(inp, inp_strides, inp_shape, inp_offset, inp_ndim, out, out_strides, out_shape, out_offset, out_ndim, numel);
}

/* Python wrapper functions */

static PyObject* element_wise_binary_float__add_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_float__add((float*)a, a_off, a_stride, (float*)b, b_off, b_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_float__sub_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_float__sub((float*)a, a_off, a_stride, (float*)b, b_off, b_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_float__mul_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_float__mul((float*)a, a_off, a_stride, (float*)b, b_off, b_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_float__div_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_float__div((float*)a, a_off, a_stride, (float*)b, b_off, b_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_float__atan2_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_float__atan2((float*)a, a_off, a_stride, (float*)b, b_off, b_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_double__add_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_double__add((double*)a, a_off, a_stride, (double*)b, b_off, b_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_double__sub_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_double__sub((double*)a, a_off, a_stride, (double*)b, b_off, b_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_double__mul_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_double__mul((double*)a, a_off, a_stride, (double*)b, b_off, b_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_double__div_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_double__div((double*)a, a_off, a_stride, (double*)b, b_off, b_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_double__atan2_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_double__atan2((double*)a, a_off, a_stride, (double*)b, b_off, b_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_unsigned_char__add_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_unsigned_char__add((unsigned char*)a, a_off, a_stride, (unsigned char*)b, b_off, b_stride, (unsigned char*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_unsigned_char__sub_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_unsigned_char__sub((unsigned char*)a, a_off, a_stride, (unsigned char*)b, b_off, b_stride, (unsigned char*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_unsigned_char__mul_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_unsigned_char__mul((unsigned char*)a, a_off, a_stride, (unsigned char*)b, b_off, b_stride, (unsigned char*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_unsigned_char__div_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_unsigned_char__div((unsigned char*)a, a_off, a_stride, (unsigned char*)b, b_off, b_stride, (unsigned char*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_char__add_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_char__add((char*)a, a_off, a_stride, (char*)b, b_off, b_stride, (char*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_char__sub_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_char__sub((char*)a, a_off, a_stride, (char*)b, b_off, b_stride, (char*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_char__mul_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_char__mul((char*)a, a_off, a_stride, (char*)b, b_off, b_stride, (char*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_char__div_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_char__div((char*)a, a_off, a_stride, (char*)b, b_off, b_stride, (char*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_unsigned_short__add_wrapper(PyObject* self, PyObject* args) {
    unsigned short* a;
    long long a_off;
    long long a_stride;
    unsigned short* b;
    long long b_off;
    long long b_stride;
    unsigned short* c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_unsigned_short__add(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_unsigned_short__sub_wrapper(PyObject* self, PyObject* args) {
    unsigned short* a;
    long long a_off;
    long long a_stride;
    unsigned short* b;
    long long b_off;
    long long b_stride;
    unsigned short* c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_unsigned_short__sub(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_unsigned_short__mul_wrapper(PyObject* self, PyObject* args) {
    unsigned short* a;
    long long a_off;
    long long a_stride;
    unsigned short* b;
    long long b_off;
    long long b_stride;
    unsigned short* c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_unsigned_short__mul(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_unsigned_short__div_wrapper(PyObject* self, PyObject* args) {
    unsigned short* a;
    long long a_off;
    long long a_stride;
    unsigned short* b;
    long long b_off;
    long long b_stride;
    unsigned short* c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_unsigned_short__div(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_short__add_wrapper(PyObject* self, PyObject* args) {
    short* a;
    long long a_off;
    long long a_stride;
    short* b;
    long long b_off;
    long long b_stride;
    short* c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_short__add(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_short__sub_wrapper(PyObject* self, PyObject* args) {
    short* a;
    long long a_off;
    long long a_stride;
    short* b;
    long long b_off;
    long long b_stride;
    short* c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_short__sub(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_short__mul_wrapper(PyObject* self, PyObject* args) {
    short* a;
    long long a_off;
    long long a_stride;
    short* b;
    long long b_off;
    long long b_stride;
    short* c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_short__mul(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_short__div_wrapper(PyObject* self, PyObject* args) {
    short* a;
    long long a_off;
    long long a_stride;
    short* b;
    long long b_off;
    long long b_stride;
    short* c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_short__div(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_unsigned_int__add_wrapper(PyObject* self, PyObject* args) {
    unsigned int* a;
    long long a_off;
    long long a_stride;
    unsigned int* b;
    long long b_off;
    long long b_stride;
    unsigned int* c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_unsigned_int__add(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_unsigned_int__sub_wrapper(PyObject* self, PyObject* args) {
    unsigned int* a;
    long long a_off;
    long long a_stride;
    unsigned int* b;
    long long b_off;
    long long b_stride;
    unsigned int* c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_unsigned_int__sub(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_unsigned_int__mul_wrapper(PyObject* self, PyObject* args) {
    unsigned int* a;
    long long a_off;
    long long a_stride;
    unsigned int* b;
    long long b_off;
    long long b_stride;
    unsigned int* c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_unsigned_int__mul(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_unsigned_int__div_wrapper(PyObject* self, PyObject* args) {
    unsigned int* a;
    long long a_off;
    long long a_stride;
    unsigned int* b;
    long long b_off;
    long long b_stride;
    unsigned int* c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_unsigned_int__div(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_int__add_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_int__add((int*)a, a_off, a_stride, (int*)b, b_off, b_stride, (int*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_int__sub_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_int__sub((int*)a, a_off, a_stride, (int*)b, b_off, b_stride, (int*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_int__mul_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_int__mul((int*)a, a_off, a_stride, (int*)b, b_off, b_stride, (int*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_int__div_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_int__div((int*)a, a_off, a_stride, (int*)b, b_off, b_stride, (int*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_unsigned_long_long__add_wrapper(PyObject* self, PyObject* args) {
    unsigned long long* a;
    long long a_off;
    long long a_stride;
    unsigned long long* b;
    long long b_off;
    long long b_stride;
    unsigned long long* c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_unsigned_long_long__add(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_unsigned_long_long__sub_wrapper(PyObject* self, PyObject* args) {
    unsigned long long* a;
    long long a_off;
    long long a_stride;
    unsigned long long* b;
    long long b_off;
    long long b_stride;
    unsigned long long* c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_unsigned_long_long__sub(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_unsigned_long_long__mul_wrapper(PyObject* self, PyObject* args) {
    unsigned long long* a;
    long long a_off;
    long long a_stride;
    unsigned long long* b;
    long long b_off;
    long long b_stride;
    unsigned long long* c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_unsigned_long_long__mul(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_unsigned_long_long__div_wrapper(PyObject* self, PyObject* args) {
    unsigned long long* a;
    long long a_off;
    long long a_stride;
    unsigned long long* b;
    long long b_off;
    long long b_stride;
    unsigned long long* c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_unsigned_long_long__div(a, a_off, a_stride, b, b_off, b_stride, c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_long_long__add_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_long_long__add((long long*)a, a_off, a_stride, (long long*)b, b_off, b_stride, (long long*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_long_long__sub_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_long_long__sub((long long*)a, a_off, a_stride, (long long*)b, b_off, b_stride, (long long*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_long_long__mul_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_long_long__mul((long long*)a, a_off, a_stride, (long long*)b, b_off, b_stride, (long long*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_binary_long_long__div_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long b;
    long long b_off;
    long long b_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLKLLL", &a, &a_off, &a_stride, &b, &b_off, &b_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_binary_long_long__div((long long*)a, a_off, a_stride, (long long*)b, b_off, b_stride, (long long*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__exp_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__exp((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__exp2_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__exp2((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__log_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__log((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__log2_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__log2((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__log10_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__log10((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__relu_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__relu((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__square_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__square((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__sqrt_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__sqrt((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__sin_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__sin((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__cos_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__cos((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__tan_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__tan((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__asin_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__asin((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__acos_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__acos((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__atan_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__atan((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__sinh_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__sinh((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__cosh_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__cosh((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_float__tanh_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_float__tanh((float*)a, a_off, a_stride, (float*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__exp_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__exp((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__exp2_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__exp2((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__log_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__log((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__log2_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__log2((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__log10_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__log10((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__relu_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__relu((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__square_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__square((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__sqrt_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__sqrt((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__sin_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__sin((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__cos_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__cos((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__tan_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__tan((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__asin_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__asin((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__acos_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__acos((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__atan_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__atan((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__sinh_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__sinh((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__cosh_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__cosh((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* element_wise_unary_double__tanh_wrapper(PyObject* self, PyObject* args) {
    unsigned long long a;
    long long a_off;
    long long a_stride;
    unsigned long long c;
    long long c_off;
    long long c_stride;
    long long n;
    if (!PyArg_ParseTuple(args, "KLLKLLL", &a, &a_off, &a_stride, &c, &c_off, &c_stride, &n)) {
        return NULL;
    }
    element_wise_unary_double__tanh((double*)a, a_off, a_stride, (double*)c, c_off, c_stride, n);
    Py_RETURN_NONE;
}



static PyObject* arange_float_wrapper(PyObject* self, PyObject* args) {
    unsigned long long out;
    long long out_offset;
    long long out_stride;
    long long numel;
    if (!PyArg_ParseTuple(args, "KLLL", &out, &out_offset, &out_stride, &numel)) {
        return NULL;
    }
    arange_float((float*)out, out_offset, out_stride, numel);
    Py_RETURN_NONE;
}



static PyObject* arange_double_wrapper(PyObject* self, PyObject* args) {
    unsigned long long out;
    long long out_offset;
    long long out_stride;
    long long numel;
    if (!PyArg_ParseTuple(args, "KLLL", &out, &out_offset, &out_stride, &numel)) {
        return NULL;
    }
    arange_double((double*)out, out_offset, out_stride, numel);
    Py_RETURN_NONE;
}



static PyObject* arange_unsigned_char_wrapper(PyObject* self, PyObject* args) {
    unsigned long long out;
    long long out_offset;
    long long out_stride;
    long long numel;
    if (!PyArg_ParseTuple(args, "KLLL", &out, &out_offset, &out_stride, &numel)) {
        return NULL;
    }
    arange_unsigned_char((unsigned char*)out, out_offset, out_stride, numel);
    Py_RETURN_NONE;
}



static PyObject* arange_char_wrapper(PyObject* self, PyObject* args) {
    unsigned long long out;
    long long out_offset;
    long long out_stride;
    long long numel;
    if (!PyArg_ParseTuple(args, "KLLL", &out, &out_offset, &out_stride, &numel)) {
        return NULL;
    }
    arange_char((char*)out, out_offset, out_stride, numel);
    Py_RETURN_NONE;
}



static PyObject* arange_unsigned_short_wrapper(PyObject* self, PyObject* args) {
    unsigned short* out;
    long long out_offset;
    long long out_stride;
    long long numel;
    if (!PyArg_ParseTuple(args, "KLLL", &out, &out_offset, &out_stride, &numel)) {
        return NULL;
    }
    arange_unsigned_short(out, out_offset, out_stride, numel);
    Py_RETURN_NONE;
}



static PyObject* arange_short_wrapper(PyObject* self, PyObject* args) {
    short* out;
    long long out_offset;
    long long out_stride;
    long long numel;
    if (!PyArg_ParseTuple(args, "KLLL", &out, &out_offset, &out_stride, &numel)) {
        return NULL;
    }
    arange_short(out, out_offset, out_stride, numel);
    Py_RETURN_NONE;
}



static PyObject* arange_unsigned_int_wrapper(PyObject* self, PyObject* args) {
    unsigned int* out;
    long long out_offset;
    long long out_stride;
    long long numel;
    if (!PyArg_ParseTuple(args, "KLLL", &out, &out_offset, &out_stride, &numel)) {
        return NULL;
    }
    arange_unsigned_int(out, out_offset, out_stride, numel);
    Py_RETURN_NONE;
}



static PyObject* arange_int_wrapper(PyObject* self, PyObject* args) {
    unsigned long long out;
    long long out_offset;
    long long out_stride;
    long long numel;
    if (!PyArg_ParseTuple(args, "KLLL", &out, &out_offset, &out_stride, &numel)) {
        return NULL;
    }
    arange_int((int*)out, out_offset, out_stride, numel);
    Py_RETURN_NONE;
}



static PyObject* arange_unsigned_long_long_wrapper(PyObject* self, PyObject* args) {
    unsigned long long* out;
    long long out_offset;
    long long out_stride;
    long long numel;
    if (!PyArg_ParseTuple(args, "KLLL", &out, &out_offset, &out_stride, &numel)) {
        return NULL;
    }
    arange_unsigned_long_long(out, out_offset, out_stride, numel);
    Py_RETURN_NONE;
}



static PyObject* arange_long_long_wrapper(PyObject* self, PyObject* args) {
    unsigned long long out;
    long long out_offset;
    long long out_stride;
    long long numel;
    if (!PyArg_ParseTuple(args, "KLLL", &out, &out_offset, &out_stride, &numel)) {
        return NULL;
    }
    arange_long_long((long long*)out, out_offset, out_stride, numel);
    Py_RETURN_NONE;
}



static PyObject* reshape_copy_float_wrapper(PyObject* self, PyObject* args) {
    unsigned long long inp;
    unsigned long long inp_strides;
    unsigned long long inp_shape;
    unsigned long long inp_index;
    long long inp_offset;
    long long inp_ndim;
    unsigned long long out;
    unsigned long long out_strides;
    unsigned long long out_shape;
    unsigned long long out_index;
    long long out_offset;
    long long out_ndim;
    long long numel;
    if (!PyArg_ParseTuple(args, "KKKKLLKKKKLLL", &inp, &inp_strides, &inp_shape, &inp_index, &inp_offset, &inp_ndim, &out, &out_strides, &out_shape, &out_index, &out_offset, &out_ndim, &numel)) {
        return NULL;
    }
    reshape_copy_float((float*)inp, (long long*)inp_strides, (long long*)inp_shape, (long long*)inp_index, inp_offset, inp_ndim, (float*)out, (long long*)out_strides, (long long*)out_shape, (long long*)out_index, out_offset, out_ndim, numel);
    Py_RETURN_NONE;
}



static PyObject* reshape_copy_double_wrapper(PyObject* self, PyObject* args) {
    unsigned long long inp;
    unsigned long long inp_strides;
    unsigned long long inp_shape;
    unsigned long long inp_index;
    long long inp_offset;
    long long inp_ndim;
    unsigned long long out;
    unsigned long long out_strides;
    unsigned long long out_shape;
    unsigned long long out_index;
    long long out_offset;
    long long out_ndim;
    long long numel;
    if (!PyArg_ParseTuple(args, "KKKKLLKKKKLLL", &inp, &inp_strides, &inp_shape, &inp_index, &inp_offset, &inp_ndim, &out, &out_strides, &out_shape, &out_index, &out_offset, &out_ndim, &numel)) {
        return NULL;
    }
    reshape_copy_double((double*)inp, (long long*)inp_strides, (long long*)inp_shape, (long long*)inp_index, inp_offset, inp_ndim, (double*)out, (long long*)out_strides, (long long*)out_shape, (long long*)out_index, out_offset, out_ndim, numel);
    Py_RETURN_NONE;
}



static PyObject* reshape_copy_unsigned_char_wrapper(PyObject* self, PyObject* args) {
    unsigned long long inp;
    unsigned long long inp_strides;
    unsigned long long inp_shape;
    unsigned long long inp_index;
    long long inp_offset;
    long long inp_ndim;
    unsigned long long out;
    unsigned long long out_strides;
    unsigned long long out_shape;
    unsigned long long out_index;
    long long out_offset;
    long long out_ndim;
    long long numel;
    if (!PyArg_ParseTuple(args, "KKKKLLKKKKLLL", &inp, &inp_strides, &inp_shape, &inp_index, &inp_offset, &inp_ndim, &out, &out_strides, &out_shape, &out_index, &out_offset, &out_ndim, &numel)) {
        return NULL;
    }
    reshape_copy_unsigned_char((unsigned char*)inp, (long long*)inp_strides, (long long*)inp_shape, (long long*)inp_index, inp_offset, inp_ndim, (unsigned char*)out, (long long*)out_strides, (long long*)out_shape, (long long*)out_index, out_offset, out_ndim, numel);
    Py_RETURN_NONE;
}



static PyObject* reshape_copy_char_wrapper(PyObject* self, PyObject* args) {
    unsigned long long inp;
    unsigned long long inp_strides;
    unsigned long long inp_shape;
    unsigned long long inp_index;
    long long inp_offset;
    long long inp_ndim;
    unsigned long long out;
    unsigned long long out_strides;
    unsigned long long out_shape;
    unsigned long long out_index;
    long long out_offset;
    long long out_ndim;
    long long numel;
    if (!PyArg_ParseTuple(args, "KKKKLLKKKKLLL", &inp, &inp_strides, &inp_shape, &inp_index, &inp_offset, &inp_ndim, &out, &out_strides, &out_shape, &out_index, &out_offset, &out_ndim, &numel)) {
        return NULL;
    }
    reshape_copy_char((char*)inp, (long long*)inp_strides, (long long*)inp_shape, (long long*)inp_index, inp_offset, inp_ndim, (char*)out, (long long*)out_strides, (long long*)out_shape, (long long*)out_index, out_offset, out_ndim, numel);
    Py_RETURN_NONE;
}



static PyObject* reshape_copy_unsigned_short_wrapper(PyObject* self, PyObject* args) {
    unsigned short* inp;
    unsigned long long inp_strides;
    unsigned long long inp_shape;
    unsigned long long inp_index;
    long long inp_offset;
    long long inp_ndim;
    unsigned short* out;
    unsigned long long out_strides;
    unsigned long long out_shape;
    unsigned long long out_index;
    long long out_offset;
    long long out_ndim;
    long long numel;
    if (!PyArg_ParseTuple(args, "KKKKLLKKKKLLL", &inp, &inp_strides, &inp_shape, &inp_index, &inp_offset, &inp_ndim, &out, &out_strides, &out_shape, &out_index, &out_offset, &out_ndim, &numel)) {
        return NULL;
    }
    reshape_copy_unsigned_short(inp, (long long*)inp_strides, (long long*)inp_shape, (long long*)inp_index, inp_offset, inp_ndim, out, (long long*)out_strides, (long long*)out_shape, (long long*)out_index, out_offset, out_ndim, numel);
    Py_RETURN_NONE;
}



static PyObject* reshape_copy_short_wrapper(PyObject* self, PyObject* args) {
    short* inp;
    unsigned long long inp_strides;
    unsigned long long inp_shape;
    unsigned long long inp_index;
    long long inp_offset;
    long long inp_ndim;
    short* out;
    unsigned long long out_strides;
    unsigned long long out_shape;
    unsigned long long out_index;
    long long out_offset;
    long long out_ndim;
    long long numel;
    if (!PyArg_ParseTuple(args, "KKKKLLKKKKLLL", &inp, &inp_strides, &inp_shape, &inp_index, &inp_offset, &inp_ndim, &out, &out_strides, &out_shape, &out_index, &out_offset, &out_ndim, &numel)) {
        return NULL;
    }
    reshape_copy_short(inp, (long long*)inp_strides, (long long*)inp_shape, (long long*)inp_index, inp_offset, inp_ndim, out, (long long*)out_strides, (long long*)out_shape, (long long*)out_index, out_offset, out_ndim, numel);
    Py_RETURN_NONE;
}



static PyObject* reshape_copy_unsigned_int_wrapper(PyObject* self, PyObject* args) {
    unsigned int* inp;
    unsigned long long inp_strides;
    unsigned long long inp_shape;
    unsigned long long inp_index;
    long long inp_offset;
    long long inp_ndim;
    unsigned int* out;
    unsigned long long out_strides;
    unsigned long long out_shape;
    unsigned long long out_index;
    long long out_offset;
    long long out_ndim;
    long long numel;
    if (!PyArg_ParseTuple(args, "KKKKLLKKKKLLL", &inp, &inp_strides, &inp_shape, &inp_index, &inp_offset, &inp_ndim, &out, &out_strides, &out_shape, &out_index, &out_offset, &out_ndim, &numel)) {
        return NULL;
    }
    reshape_copy_unsigned_int(inp, (long long*)inp_strides, (long long*)inp_shape, (long long*)inp_index, inp_offset, inp_ndim, out, (long long*)out_strides, (long long*)out_shape, (long long*)out_index, out_offset, out_ndim, numel);
    Py_RETURN_NONE;
}



static PyObject* reshape_copy_int_wrapper(PyObject* self, PyObject* args) {
    unsigned long long inp;
    unsigned long long inp_strides;
    unsigned long long inp_shape;
    unsigned long long inp_index;
    long long inp_offset;
    long long inp_ndim;
    unsigned long long out;
    unsigned long long out_strides;
    unsigned long long out_shape;
    unsigned long long out_index;
    long long out_offset;
    long long out_ndim;
    long long numel;
    if (!PyArg_ParseTuple(args, "KKKKLLKKKKLLL", &inp, &inp_strides, &inp_shape, &inp_index, &inp_offset, &inp_ndim, &out, &out_strides, &out_shape, &out_index, &out_offset, &out_ndim, &numel)) {
        return NULL;
    }
    reshape_copy_int((int*)inp, (long long*)inp_strides, (long long*)inp_shape, (long long*)inp_index, inp_offset, inp_ndim, (int*)out, (long long*)out_strides, (long long*)out_shape, (long long*)out_index, out_offset, out_ndim, numel);
    Py_RETURN_NONE;
}



static PyObject* reshape_copy_unsigned_long_long_wrapper(PyObject* self, PyObject* args) {
    unsigned long long* inp;
    unsigned long long inp_strides;
    unsigned long long inp_shape;
    unsigned long long inp_index;
    long long inp_offset;
    long long inp_ndim;
    unsigned long long* out;
    unsigned long long out_strides;
    unsigned long long out_shape;
    unsigned long long out_index;
    long long out_offset;
    long long out_ndim;
    long long numel;
    if (!PyArg_ParseTuple(args, "KKKKLLKKKKLLL", &inp, &inp_strides, &inp_shape, &inp_index, &inp_offset, &inp_ndim, &out, &out_strides, &out_shape, &out_index, &out_offset, &out_ndim, &numel)) {
        return NULL;
    }
    reshape_copy_unsigned_long_long(inp, (long long*)inp_strides, (long long*)inp_shape, (long long*)inp_index, inp_offset, inp_ndim, out, (long long*)out_strides, (long long*)out_shape, (long long*)out_index, out_offset, out_ndim, numel);
    Py_RETURN_NONE;
}



static PyObject* reshape_copy_long_long_wrapper(PyObject* self, PyObject* args) {
    unsigned long long inp;
    unsigned long long inp_strides;
    unsigned long long inp_shape;
    unsigned long long inp_index;
    long long inp_offset;
    long long inp_ndim;
    unsigned long long out;
    unsigned long long out_strides;
    unsigned long long out_shape;
    unsigned long long out_index;
    long long out_offset;
    long long out_ndim;
    long long numel;
    if (!PyArg_ParseTuple(args, "KKKKLLKKKKLLL", &inp, &inp_strides, &inp_shape, &inp_index, &inp_offset, &inp_ndim, &out, &out_strides, &out_shape, &out_index, &out_offset, &out_ndim, &numel)) {
        return NULL;
    }
    reshape_copy_long_long((long long*)inp, (long long*)inp_strides, (long long*)inp_shape, (long long*)inp_index, inp_offset, inp_ndim, (long long*)out, (long long*)out_strides, (long long*)out_shape, (long long*)out_index, out_offset, out_ndim, numel);
    Py_RETURN_NONE;
}


/* Method table */
static PyMethodDef _ndarray_rt_module_methods[] = {
    {"element_wise_binary_float__add", element_wise_binary_float__add_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_float__add"},
    {"element_wise_binary_float__sub", element_wise_binary_float__sub_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_float__sub"},
    {"element_wise_binary_float__mul", element_wise_binary_float__mul_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_float__mul"},
    {"element_wise_binary_float__div", element_wise_binary_float__div_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_float__div"},
    {"element_wise_binary_float__atan2", element_wise_binary_float__atan2_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_float__atan2"},
    {"element_wise_binary_double__add", element_wise_binary_double__add_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_double__add"},
    {"element_wise_binary_double__sub", element_wise_binary_double__sub_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_double__sub"},
    {"element_wise_binary_double__mul", element_wise_binary_double__mul_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_double__mul"},
    {"element_wise_binary_double__div", element_wise_binary_double__div_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_double__div"},
    {"element_wise_binary_double__atan2", element_wise_binary_double__atan2_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_double__atan2"},
    {"element_wise_binary_unsigned_char__add", element_wise_binary_unsigned_char__add_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_unsigned_char__add"},
    {"element_wise_binary_unsigned_char__sub", element_wise_binary_unsigned_char__sub_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_unsigned_char__sub"},
    {"element_wise_binary_unsigned_char__mul", element_wise_binary_unsigned_char__mul_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_unsigned_char__mul"},
    {"element_wise_binary_unsigned_char__div", element_wise_binary_unsigned_char__div_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_unsigned_char__div"},
    {"element_wise_binary_char__add", element_wise_binary_char__add_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_char__add"},
    {"element_wise_binary_char__sub", element_wise_binary_char__sub_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_char__sub"},
    {"element_wise_binary_char__mul", element_wise_binary_char__mul_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_char__mul"},
    {"element_wise_binary_char__div", element_wise_binary_char__div_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_char__div"},
    {"element_wise_binary_unsigned_short__add", element_wise_binary_unsigned_short__add_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_unsigned_short__add"},
    {"element_wise_binary_unsigned_short__sub", element_wise_binary_unsigned_short__sub_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_unsigned_short__sub"},
    {"element_wise_binary_unsigned_short__mul", element_wise_binary_unsigned_short__mul_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_unsigned_short__mul"},
    {"element_wise_binary_unsigned_short__div", element_wise_binary_unsigned_short__div_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_unsigned_short__div"},
    {"element_wise_binary_short__add", element_wise_binary_short__add_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_short__add"},
    {"element_wise_binary_short__sub", element_wise_binary_short__sub_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_short__sub"},
    {"element_wise_binary_short__mul", element_wise_binary_short__mul_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_short__mul"},
    {"element_wise_binary_short__div", element_wise_binary_short__div_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_short__div"},
    {"element_wise_binary_unsigned_int__add", element_wise_binary_unsigned_int__add_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_unsigned_int__add"},
    {"element_wise_binary_unsigned_int__sub", element_wise_binary_unsigned_int__sub_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_unsigned_int__sub"},
    {"element_wise_binary_unsigned_int__mul", element_wise_binary_unsigned_int__mul_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_unsigned_int__mul"},
    {"element_wise_binary_unsigned_int__div", element_wise_binary_unsigned_int__div_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_unsigned_int__div"},
    {"element_wise_binary_int__add", element_wise_binary_int__add_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_int__add"},
    {"element_wise_binary_int__sub", element_wise_binary_int__sub_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_int__sub"},
    {"element_wise_binary_int__mul", element_wise_binary_int__mul_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_int__mul"},
    {"element_wise_binary_int__div", element_wise_binary_int__div_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_int__div"},
    {"element_wise_binary_unsigned_long_long__add", element_wise_binary_unsigned_long_long__add_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_unsigned_long_long__add"},
    {"element_wise_binary_unsigned_long_long__sub", element_wise_binary_unsigned_long_long__sub_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_unsigned_long_long__sub"},
    {"element_wise_binary_unsigned_long_long__mul", element_wise_binary_unsigned_long_long__mul_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_unsigned_long_long__mul"},
    {"element_wise_binary_unsigned_long_long__div", element_wise_binary_unsigned_long_long__div_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_unsigned_long_long__div"},
    {"element_wise_binary_long_long__add", element_wise_binary_long_long__add_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_long_long__add"},
    {"element_wise_binary_long_long__sub", element_wise_binary_long_long__sub_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_long_long__sub"},
    {"element_wise_binary_long_long__mul", element_wise_binary_long_long__mul_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_long_long__mul"},
    {"element_wise_binary_long_long__div", element_wise_binary_long_long__div_wrapper, METH_VARARGS, "Transpiled function element_wise_binary_long_long__div"},
    {"element_wise_unary_float__exp", element_wise_unary_float__exp_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__exp"},
    {"element_wise_unary_float__exp2", element_wise_unary_float__exp2_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__exp2"},
    {"element_wise_unary_float__log", element_wise_unary_float__log_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__log"},
    {"element_wise_unary_float__log2", element_wise_unary_float__log2_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__log2"},
    {"element_wise_unary_float__log10", element_wise_unary_float__log10_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__log10"},
    {"element_wise_unary_float__relu", element_wise_unary_float__relu_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__relu"},
    {"element_wise_unary_float__square", element_wise_unary_float__square_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__square"},
    {"element_wise_unary_float__sqrt", element_wise_unary_float__sqrt_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__sqrt"},
    {"element_wise_unary_float__sin", element_wise_unary_float__sin_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__sin"},
    {"element_wise_unary_float__cos", element_wise_unary_float__cos_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__cos"},
    {"element_wise_unary_float__tan", element_wise_unary_float__tan_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__tan"},
    {"element_wise_unary_float__asin", element_wise_unary_float__asin_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__asin"},
    {"element_wise_unary_float__acos", element_wise_unary_float__acos_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__acos"},
    {"element_wise_unary_float__atan", element_wise_unary_float__atan_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__atan"},
    {"element_wise_unary_float__sinh", element_wise_unary_float__sinh_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__sinh"},
    {"element_wise_unary_float__cosh", element_wise_unary_float__cosh_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__cosh"},
    {"element_wise_unary_float__tanh", element_wise_unary_float__tanh_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_float__tanh"},
    {"element_wise_unary_double__exp", element_wise_unary_double__exp_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__exp"},
    {"element_wise_unary_double__exp2", element_wise_unary_double__exp2_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__exp2"},
    {"element_wise_unary_double__log", element_wise_unary_double__log_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__log"},
    {"element_wise_unary_double__log2", element_wise_unary_double__log2_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__log2"},
    {"element_wise_unary_double__log10", element_wise_unary_double__log10_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__log10"},
    {"element_wise_unary_double__relu", element_wise_unary_double__relu_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__relu"},
    {"element_wise_unary_double__square", element_wise_unary_double__square_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__square"},
    {"element_wise_unary_double__sqrt", element_wise_unary_double__sqrt_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__sqrt"},
    {"element_wise_unary_double__sin", element_wise_unary_double__sin_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__sin"},
    {"element_wise_unary_double__cos", element_wise_unary_double__cos_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__cos"},
    {"element_wise_unary_double__tan", element_wise_unary_double__tan_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__tan"},
    {"element_wise_unary_double__asin", element_wise_unary_double__asin_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__asin"},
    {"element_wise_unary_double__acos", element_wise_unary_double__acos_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__acos"},
    {"element_wise_unary_double__atan", element_wise_unary_double__atan_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__atan"},
    {"element_wise_unary_double__sinh", element_wise_unary_double__sinh_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__sinh"},
    {"element_wise_unary_double__cosh", element_wise_unary_double__cosh_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__cosh"},
    {"element_wise_unary_double__tanh", element_wise_unary_double__tanh_wrapper, METH_VARARGS, "Transpiled function element_wise_unary_double__tanh"},
    {"arange_float", arange_float_wrapper, METH_VARARGS, "Transpiled function arange_float"},
    {"arange_double", arange_double_wrapper, METH_VARARGS, "Transpiled function arange_double"},
    {"arange_unsigned_char", arange_unsigned_char_wrapper, METH_VARARGS, "Transpiled function arange_unsigned_char"},
    {"arange_char", arange_char_wrapper, METH_VARARGS, "Transpiled function arange_char"},
    {"arange_unsigned_short", arange_unsigned_short_wrapper, METH_VARARGS, "Transpiled function arange_unsigned_short"},
    {"arange_short", arange_short_wrapper, METH_VARARGS, "Transpiled function arange_short"},
    {"arange_unsigned_int", arange_unsigned_int_wrapper, METH_VARARGS, "Transpiled function arange_unsigned_int"},
    {"arange_int", arange_int_wrapper, METH_VARARGS, "Transpiled function arange_int"},
    {"arange_unsigned_long_long", arange_unsigned_long_long_wrapper, METH_VARARGS, "Transpiled function arange_unsigned_long_long"},
    {"arange_long_long", arange_long_long_wrapper, METH_VARARGS, "Transpiled function arange_long_long"},
    {"reshape_copy_float", reshape_copy_float_wrapper, METH_VARARGS, "Transpiled function reshape_copy_float"},
    {"reshape_copy_double", reshape_copy_double_wrapper, METH_VARARGS, "Transpiled function reshape_copy_double"},
    {"reshape_copy_unsigned_char", reshape_copy_unsigned_char_wrapper, METH_VARARGS, "Transpiled function reshape_copy_unsigned_char"},
    {"reshape_copy_char", reshape_copy_char_wrapper, METH_VARARGS, "Transpiled function reshape_copy_char"},
    {"reshape_copy_unsigned_short", reshape_copy_unsigned_short_wrapper, METH_VARARGS, "Transpiled function reshape_copy_unsigned_short"},
    {"reshape_copy_short", reshape_copy_short_wrapper, METH_VARARGS, "Transpiled function reshape_copy_short"},
    {"reshape_copy_unsigned_int", reshape_copy_unsigned_int_wrapper, METH_VARARGS, "Transpiled function reshape_copy_unsigned_int"},
    {"reshape_copy_int", reshape_copy_int_wrapper, METH_VARARGS, "Transpiled function reshape_copy_int"},
    {"reshape_copy_unsigned_long_long", reshape_copy_unsigned_long_long_wrapper, METH_VARARGS, "Transpiled function reshape_copy_unsigned_long_long"},
    {"reshape_copy_long_long", reshape_copy_long_long_wrapper, METH_VARARGS, "Transpiled function reshape_copy_long_long"},
    {NULL, NULL, 0, NULL}
};

static struct PyModuleDef _ndarray_rt_module_module = {
    PyModuleDef_HEAD_INIT,
    "_ndarray_rt_module",
    NULL,
    -1,
    _ndarray_rt_module_methods
};

PyMODINIT_FUNC PyInit__ndarray_rt_module(void) {
    return PyModule_Create(&_ndarray_rt_module_module);
}

Some notes on this:

• Some of this Python code, such as the _math_id function and all of the functions that are set to it, are purely for type hinting. The transpiler can accept a function that would otherwise cause a NameError: name 'variable_name' is not defined.

• SpecItem is used to specify what the generic type should string-substitute as, and also specifies the final C function name.

• The __device__ and __global__ functions have pybind=False, only the ones that can be called from host-side Python have pybind=True

• from ._element_wise_cuda_stubs import _ElementWiseModuleClass is actually an automatically generated type stub for PythonModule. It is equivalent to PythonModule but has runtime-generated type stubs for each of the functions with pybind=True. I may remove this in the future if I don’t use it often.

• I have not added define macros yet, but it seems nvcc is smart enough to inline these element-wise operations since these kernels get quite close to the speed of light memory bandwidth. I will investigate this in the next blog post.

This currently suffices for our low-level bridge. When I get to mma PTX instructions, I will likely need to add some more features to support asm volatile, but that is a bit down the line. Also, modules are hashed and cached when compiled into ~/.cache/simplendarray to prevent unnecessary recompilation.

The Buffer

The Buffer classes are necessary to declare that we have a chunk of memory for computation. SimpleNDArray has two buffer classes: Buffer and BufferCuda. Both have the following methods/fields:

• Constructor

• classmethod empty (Buffer currently fills it with all zeros)

• classmethod from_iterable, turns a Python iterable into a contiguous C buffer

• data: always a CPU array. If called on BufferCuda, a d2h copy is done.

• address (int): Memory address of buffer

• typecode (str): The array module typecode to specify the dtype of the buffer

• num_bytes (int): Self-explanatory.

Currently, the constructors do not agree for Buffer and BufferCuda because of the following:

• Buffer’s constructor takes in a Python array and stores it as the data field

• BufferCuda’s constructor takes in a size and dtype and calls cudaMalloc, so BufferCuda.empty just calls the constructor

• The Python array module does not provide the ability to create an empty array so the constructor and empty methods are different

In the future, I hope to clean up these 2 classes. Ideally, I would remove the dependency on the array module and the use of typecodes, as it has cluttered my codebase and will be difficult to use custom types (e.g., bfloat16, float16), which don’t have associated Python typecodes.

To deal with memory management, an invariant I enforced is that for every memory buffer, there will be exactly 1 Buffer/BufferCuda class. This buffer class will have sole ownership of the buffer, and memory views will use the same buffer object. For BufferCuda, cudaFree is called inside the __del__ method for some basic automatic memory management. I will also have to have a more robust solution to remove the possibility of memory leaks or double frees.

I will note that to call functions like cudaMemcpy, cudaMalloc, and cudaFree, the runtime from above was used to bind these to Python! No CUDA files were (directly) written for this.

The Array

The Array class will be a relatively simple wrapper around the Buffer classes with extra metadata. It will contain:

• A data buffer

• A shape, which will be a tuple of ints

• strides, which will be a tuple of ints of the same length as shape

• An offset, indicating where the first element of the array will be in the buffer

We can derive other NumPy/PyTorch array properties from these:

• ndim=len(shape)=len(strides)

• numel()/size = product(shape)

This is the dense strided array layout, which is used in every deep learning/array library today.

Row/Column Major

I will also note that SimpleNDArray will be C/row-major order only, similar to PyTorch. Adding Fortran/column-major support will be complex and unnecessary.

Row major indicates that our indices will start counting from the right. For example, if we have an array of shape (2, 3), our indices to traverse this array in order would be (0, 0), (0, 1), (0, 2), (1, 0), (1, 1), (1, 2).

If we instead use column-major order, the indices would be (0, 0), (1, 0), (0, 1), (1, 1), (0, 2), (1, 2). As it turns out, if the array is contiguous, the memory buffer would have the elements at these indices in these exact orders (depending on whether it is row/column major). But since we have a strided layout, contiguity cannot be assumed.

Strided Layout

I will start this section by explaining what array strides are, if you are unfamiliar. An array dimension’s stride is how far you need to jump to get to the next element in that dimension. For example, in a 1d array, it would be the distances between elements (i.e., the distance between arr[0], arr[1], …). In the 2d array case, we have an (m, n) matrix with m rows and n columns. If it were contiguous, the strides would be (n, 1). This is because assuming row major order, the buffer is laid out in a flattened order as such:

1  2  3  4
5  6  7  8    -> 1 2 3 4 5 6 7 8 9 10 11 12
9  10 11 12

To jump from any element to 1 row below, we will need to traverse 4 elements in flat memory. This is the stride for the first dimension (rows). For the columns, to jump from any element to 1 column to the right, we only need to traverse 1 element in flat memory, so the stride for the second dimension is 1. Overall, in this example, we have a shape of (3, 4) and strides of (4, 1).

However, arrays don’t have to be contiguous like this. We can have overlaps and gaps as long as they can be represented using this strided view.

I will also note that strides (and offsets) are not necessary for an n-dimensional array library. We could assume an invariant that all arrays are contiguous, but this restricts us greatly. Mainly, every single reshape, slice, transpose, and other memory view operation would require an array copy. NumPy and PyTorch (and SimpleNDArray) can perform these memory views in a zero-copy fashion. For example:

>>> from simplendarray import Array
>>> a = Array.arange(4, "d")
>>> a
Array([0.0, 1.0, 2.0, 3.0], shape=(4,), strides=(1,), offset=0)
>>> b = a[::2]
>>> b
Array([0.0, 2.0], shape=(2,), strides=(2,), offset=0)
>>> a.data.address
4310996048
>>> b.data.address
4310996048

As you can see, both a and b’s data buffers share the same memory address because b is a zero-copy memory view of a. b just has double the stride that a has, allowing it to jump by 2 elements for each index increase.

I had ChatGPT make a good ASCII art to show this. I added the bytes because NumPy’s strides are multiplied by the itemsize as the strides are in bytes, not elements.

                 a[0]              a[1]              a[2]              a[3]
                  |<--- 8 bytes --->|<--- 8 bytes --->|<--- 8 bytes --->|
                  v                 v                 v                 v
         +-----------------+-----------------+-----------------+-----------------+
Buffer:  |        0        |        1        |        2        |        3        |
         +-----------------+-----------------+-----------------+-----------------+
                  ^                                   ^
                  |<----------- 16 bytes ------------>|
                b[0]                                 b[1]

a stride = 8 bytes (1 element)
b stride = 16 bytes (2 elements)

By using the same memory buffer but with a different stride, we don’t have to copy the array to do this. Some reshapes can be done by replacing the shape and modifying the stride. A transpose involves reversing both the shapes and the strides. A broadcast can be done by expanding the shape at the broadcast dimension from 1 to n and setting the stride at that dimension to 0. Flipping a dimension is simply negating its stride and increasing the offset of the array by the dimension length. The offset also allows us to start at any element in the dimension. The reason we don’t have one offset per dimension is that calculating the memory location of an element involves the calculation \(\sum_{i=1}^D \text{ndindex}_i \cdot \text{stride}_i\). Adding an offset per dimension would turn this into \(\sum_{i=1}^D (\text{ndindex}_i \cdot \text{stride}_i + \text{offset}_i)=(\sum_{i=1}^D \text{ndindex}_i \cdot \text{stride}_i) + \sum_{i=1}^D \text{offset}_i\). This means we can merge the offsets for all dimensions into a single offset.

Stride tricks like these can eliminate expensive memcpys and speed up algorithms like convolution using im2col.

>>> b.shape = (2, 2)  # broadcast to (2, 2) (repeat b twice)
>>> b.strides = (0, 2)  # Stride of 0 in dim 0
>>> b
Array([[0.0, 2.0], [0.0, 2.0]], shape=(2, 2), strides=(0, 2), offset=0)
>>> b.T
Array([[0.0, 0.0], [2.0, 2.0]], shape=(2, 2), strides=(2, 0), offset=0)
>>> b.T[::-1, :]
Array([[2.0, 2.0], [0.0, 0.0]], shape=(2, 2), strides=(-2, 0), offset=2)

So this is why I chose to use strides and offset.

Thanks to strides and offsets, array indexing [2] is relatively easy (excluding advanced indexing). My implementation currently requires the same number of indices as dimensions, doesn’t do broadcasting, and doesn’t squeeze integer index dimensions, but I will make it match NumPy’s behavior. It currently has support for ints and slices. For each dimension:

• If the index is an int i, the new shape at this dim will be 1 (unlike NumPy/torch, which squeeze the dimension). The new stride doesn’t matter (set to 0). The new offset will have i * stride added on.

• If it is a slice (stop, start, step), then the new shape for this dim will max(0, ceildiv(stop - start, step)). The new stride will be stride * step where stride is the current stride. Finally, the offset will have start * stride added on.

With this, we can implement __getitem__:

Show __getitem__ code

def __getitem__(self, items: tuple[int | slice, ...] | int | slice):
    if not isinstance(items, tuple):
        items = (items,)
    if len(items) != self.ndim:
        raise ValueError("Must index the same number of dimensions as the array")
    new_shape = []
    new_strides = []
    new_offset = self.offset
    for shape, stride, item in zip(self.shape, self.strides, items):
        if isinstance(item, int):
            item = slice(item, item + 1).indices(shape)[0]
            new_shape.append(1)
            new_strides.append(0)
            new_offset += stride * item
        elif isinstance(item, slice):
            start, stop, step = item.indices(shape)
            if step > 0:
                # Num elements in [start, stop) = stop - start
                new_shape.append(max(0, ceildiv(stop - start, step)))
            else:
                new_shape.append(max(0, ceildiv(start - stop, -step)))
            new_strides.append(stride * step)
            new_offset += stride * start
        else:
            raise TypeError(f"Index must be int or slice, got {type(item).__name__}")
    return Array(self.data, tuple(new_shape), tuple(new_strides), new_offset)

To and From Python Iterables

In NumPy and PyTorch, we can say something like torch.tensor([[1, 2], [3, 4]]) and it will give us our (2, 2) array with strides (2, 1). To implement this, I flatten this nested Python Iterable (in this case, a list) while also getting the shape. We also want to ensure this iterable is not jagged (for example, [[1, 2, 3], [4, 5]] is invalid). The code for this is as follows:

Show Python code

type Scalar = int | float | bool
type NestedIterable = Iterable[Scalar] | Iterable[NestedIterable]

def flatten_and_get_shape(
    data: NestedIterable | Scalar,
) -> tuple[list[Scalar], tuple[int, ...]]:  # Returns a tuple of (flat list, shape tuple)
    if isinstance(data, (int, float, bool)):
        return [data], ()  # scalars have shape ()

    if not data:
        return [], (0,)  # empty array/buffer with shape (0,)

    flat: list[Scalar] = []

    cur_dim = list(data)
    first_elem = cur_dim[0]
    if isinstance(first_elem, Iterable):
        first_list = list(first_elem)
        first_flat, first_shape = flatten_and_get_shape(first_list)
    else:
        first_flat, first_shape = [first_elem], ()
    flat.extend(first_flat)
    shape = (len(cur_dim), *first_shape)

    for item in cur_dim[1:]:
        if isinstance(item, Iterable):
            sub_list = list(item)
            sub_flat, sub_shape = flatten_and_get_shape(sub_list)
        else:
            sub_flat, sub_shape = [item], ()
        if sub_shape != first_shape:
            raise ValueError(f"Jagged array: expected shape {first_shape}, got {sub_shape}")
        flat.extend(sub_flat)

    return flat, shape

We can then pass this flat list into a Python array, and we now have our shape. But what about our stride?

For contiguous arrays, the stride will always be a cumulative product of the shape. For example, a contiguous array of shape (3, 4, 5, 6) will have strides (4*5*6, 5 * 6, 6, 1). This is the reversed exclusive cumulative product of the reversed shape:

from itertools import accumulate  # aka associative scan, cumulative op (cumsum, cumprod, generic)
from operator import mul

def contiguous_strides(shape: tuple[int, ...]) -> tuple[int, ...]:
    strides = tuple(accumulate(shape[::-1], mul, initial=1))[-2::-1]
    return strides

With this, and setting offset=0, we can now convert a Python iterable to an n-dimensional Array! But how do we get back to a Python object, or print its contents? We can try converting the data buffer to Python, but that would be out of order and the wrong shape. We instead use a simple recursive algorithm using slicing:

def to_python(self) -> NestedIterable:
    if self.ndim == 0:
        return self.data.data[self.offset]
    nested = []
    for i in range(self.shape[0]):
        # Does [self[0, :, :, ...], self[1, :, :, ...], ..., self[shape[0] - 1, :, :, ...]]
        # Then recursively calls to_python on each of these children, until base case reached
        indexed = self[i, *(slice(None) for _ in range(self.ndim - 1))].squeeze(0)
        nested.append(indexed.to_python())
    return nested

Assuming __getitem__ is implemented correctly, this will work for any array.

Reshaping

Reshaping is the most complex operation to deal with. This is because sometimes reshaping can be done using a zero-copy memory view. Other times, it will require an explicit copy. Here is an example where a copy is necessary:

from simplendarray import Array
Array.arange(4, "f").reshape((2, 2)).T.reshape(-1)

Here are the steps:

• arange: [0, 1, 2, 3]. Shape: (4,), Stride: (1,)

• reshape: [[0, 1], [2, 3]]. Shape: (2, 2), Stride: (2, 1)

• transpose: [[0, 2], [1, 3]]. Shape: (2, 2), Stride: (1, 2)

• reshape: [0, 2, 1, 3]. Shape: (4,), Stride: ?

We are trying to represent [0, 2, 1, 3] in the same 1d flat buffer as [0, 1, 2, 3] with a single stride, which is clearly not possible (we jump 2, then -1, then 2). Therefore, we must copy [0, 2, 1, 3] into a newly allocated buffer and return that. But how can we determine when we can and can’t do a zero-copy reshape?

The way NumPy determines how can be found in their source code[3].

First, they squeeze out any dimensions of length 1. They have no effect on contiguity, and they need special rules.

Next, we must observe that we can partition both the input shape and output shape into subdimensions that iterate together. To explain this, I will use a few examples of (old shape, new shape) pairs.

3    4    5
|    |    |
|____|    |
  |       |
  12      5

  12      5
  |       |
.____.    |
|    |    |
3    4    5

3   6   4   2
|   |   |   |
|___|___|   |
  |   |     |
  12  6     2

For each group pair, the product of the shapes is the same. Iterating over the input group is equivalent to iterating over the output group. We just have to:

• Find these groups sequentially, going from left to right on both input and output shapes

• Check if these groups can be iterated contiguously of each other

Here is the function in SimpleNDArray, transpiled pretty much line for line from NumPy:

Show Python code

def reshape_strides(
    old_shape: tuple[int, ...], new_shape: tuple[int, ...], old_strides: tuple[int, ...]
) -> tuple[int, ...] | None:
    """If the reshape can be done with a zero-copy memory view, return the new strides.
    Otherwise, return None. Based on NumPy's _attempt_nocopy_reshape in _core/src/multiarray/shape.c"""
    # Get rid of the length 1 dimensions
    squeezed = [(x, y) for x, y in zip(old_shape, old_strides) if x != 1]
    old_shape = tuple(p[0] for p in squeezed)
    old_strides = tuple(p[1] for p in squeezed)
    oldnd = len(old_shape)
    newnd = len(new_shape)
    numel = product(old_shape)
    if numel != product(new_shape):
        raise ValueError("Shapes do not share the same number of elements")
    if numel == 0:
        return (0,) * newnd

    oi = 0
    oj = 1
    ni = 0
    nj = 1
    new_strides = [0] * newnd
    while ni < newnd and oi < oldnd:
        np = new_shape[ni]
        op = old_shape[oi]

        while np != op:
            if np < op:
                # Misses trailing 1s, these are handled later
                np *= new_shape[nj]
                nj += 1
            else:
                op *= old_shape[oj]
                oj += 1
        # We have now found our (input, output) group pair which is
        # (old_shape[oi:oj], new_shape[ni:nj])

        # Check whether the original axes can be combined
        for ok in range(oi, oj - 1):
            # C order, the same cumulative product rule for contiguous arrays from before
            if old_strides[ok] != old_shape[ok + 1] * old_strides[ok + 1]:
                # Not contiguous enough
                return None

        # Calculate new strides for all axes currently worked with
        new_strides[nj - 1] = old_strides[oj - 1]
        for nk in range(nj - 1, ni, -1):
            # Again, the same cumulative product rule for contiguous arrays from before
            new_strides[nk - 1] = new_strides[nk] * new_shape[nk]
        ni = nj
        nj += 1
        oi = oj
        oj += 1

    # Set strides corresponding to trailing 1s of the new shape
    if ni >= 1:
        last_stride = new_strides[ni - 1]
    else:
        last_stride = 1
    for nk in range(ni, newnd):
        new_strides[nk] = last_stride

    return tuple(new_strides)

Now we know when we can do a zero-copy memory view and when we need to do a copy. But how do we do a copy?

This requires a custom reshape copy element-wise kernel. How it works is we store an n-dimensional input index and an n-dimensional output index. Over a total of n iterations, where n is the number of elements, we increment both the input and output index once per iteration. Here is the CPU kernel:

Show CPU Kernel code

@element_wise_module.compile_fn(
    types=[SpecItem({"T": ctype(dt)}, f"reshape_copy_{cname(dt)}") for dt in all_dtypes], pybind=True
)
def reshape_copy[T: DType](
    inp: list[T],
    inp_strides: list[i64],
    inp_shape: list[i64],
    inp_index: list[i64],
    inp_offset: i64,
    inp_ndim: i64,
    out: list[T],
    out_strides: list[i64],
    out_shape: list[i64],
    out_index: list[i64],
    out_offset: i64,
    out_ndim: i64,
    numel: i64,
) -> void:
    inp_idx: i64 = inp_offset
    out_idx: i64 = out_offset
    # Set inp_index and out_index to zeros to start off
    for i in range(inp_ndim):
        inp_index[i] = 0
    for i in range(out_ndim):
        out_index[i] = 0
    for i in range(numel):
        out[out_idx] = inp[inp_idx]

        # Now we increment ND index of input
        inp_index[inp_ndim - 1] += 1
        inp_idx += inp_strides[inp_ndim - 1]
        for j in range(inp_ndim - 1, -1, -1):
            if inp_index[j] == inp_shape[j]:
                inp_index[j] = 0
                inp_idx -= inp_strides[j] * inp_shape[j]
                if j > 0:
                    inp_index[j - 1] += 1
                    inp_idx += inp_strides[j - 1]

        # Increment ND index of output
        out_index[out_ndim - 1] += 1
        out_idx += out_strides[out_ndim - 1]
        for j in range(out_ndim - 1, -1, -1):
            if out_index[j] == out_shape[j]:
                out_index[j] = 0
                out_idx -= out_strides[j] * out_shape[j]
                if j > 0:
                    out_index[j - 1] += 1
                    out_idx += out_strides[j - 1]

For example, if we were reshaping a (4, 2) array to a (2, 2, 2) array, this would be doing:

output[0, 0, 0] = input[0, 0]
output[0, 0, 1] = input[0, 1]
output[0, 1, 0] = input[1, 0]
output[0, 1, 1] = input[1, 1]
output[1, 0, 0] = input[2, 0]
output[1, 0, 1] = input[2, 1]
output[1, 1, 0] = input[3, 0]
output[1, 1, 1] = input[3, 1]

Except that we are not doing ND indexing, we are calculating the pointer offsets using the input and output strides using inp_idx and out_idx.

Note

Even though this is an element-wise kernel, the CUDA kernel for this operation does not achieve anywhere close to the speed of light memory bandwidth like the other element-wise kernels because of the overhead of the n-dimensional indexing and pointer offset calculations. I will attempt to improve this in the future.

Note

NumPy and PyTorch have their own abstractions for n-dimensional iteration[4][5]. Both of these flatten adjacent contiguous axes and perform other optimizations to reduce the overhead of n-dimensional iteration.

Kernels and Dispatch

The runtime has given me access to generics, but now we need to dynamically dispatch the correct (kernel, dtype, device) tuple for any arbitrary array.

PyTorch, as a reference, has an extremely complex dispatch engine which covers their (kernel, dtype, device, layout, backend), autograd, autocast, JIT compilation, __torch_function__, __torch_dispatch__, fallback, and more. A great blog post from Edward Yang in 2020 talks about the PyTorch dispatch engine in detail (but is almost certainly out of date now)[6].

I keep it simple. I group similar operations together in PythonModules (element-wise, reduction, mm, conv, etc.). I have one PythonModule per device type (cpu, cuda), and each PythonModule function has type generic functions for all covered dtypes. I then use a dispatch dictionary to get the correct function. It’s simple enough, and it works for my use case.

Conclusion

This is all for part 1 of this series. So far, we have covered:

• The Python->C/CUDA transpiler

• The runtime, which allows for cached compilation

• Buffers and memory management

• Strided arrays and their associated operations

• Dynamic kernel dispatch

These lay the foundations for all the array operations we need to implement. We just define the kernels, add some dispatch logic, and we apply it to arrays.

In part 2, I will cover different kernels such as element-wise, reductions, and matmuls. I will also benchmark my kernels against PyTorch and the theoretical roofline on my testing hardware.

---

[1]

SimpleNDArray requires compilers (nvcc/gcc), the NVIDIA CUDA Toolkit, and drivers to be installed. There are also test dependencies, but those do not count.

[2]

See here

[3]

The C source can be found here

[4]

NumPy NpyIter Docs

[5]

TensorIterator Wiki Page

[6]

Blog is here.