You write custom CUDA kernels to replace the pytorch operators in the given architecture to get speedups. You have complete freedom to choose the set of operators you want to replace. You may make the decision to replace some operators with custom CUDA kernels and leave others unchanged. You may replace multiple operators with custom implementations, consider operator fusion opportunities (combining multiple operators into a single kernel, for example, combining matmul+relu), or algorithmic changes (such as online softmax). You are only limited by your imagination.

Here's an example to show you the syntax of inline embedding custom CUDA operators in torch: The example given architecture is:

python
import torch
import torch.nn as nn
import torch.nn.functional as F

class Model(nn.Module):
def init(self) -> None:
super().init()

def forward(self, a, b):
    return a + b
def get_inputs():
# randomly generate input tensors based on the model architecture
a = torch.randn(1, 128).cuda()
b = torch.randn(1, 128).cuda()
return [a, b]

def get_init_inputs():
# randomly generate tensors required for initialization based on the model architecture
return []



The example new arch with custom CUDA kernels looks like this:

python
import torch
from torch.utils.cpp_extension import load_inline
relu_source = “”"
#include <torch/extension.h>
#include <cuda_runtime.h>

global void relu_kernel(const float* x, float* y, int size) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < size) {
y[idx] = fmaxf(x[idx], 0.f);
}
}

torch::Tensor relu_cuda(torch::Tensor x) {
auto size = x.numel();
auto y = torch::empty_like(x);
const int block_size = 256;
int num_blocks = (size + block_size - 1) / block_size;
relu_kernel<<<num_blocks, block_size>>>(x.data_ptr<float>(), y.data_ptr<float>(), size);
return y;
}
“”"

relu_cpp_source = “”"
torch::Tensor relu_cuda(torch::Tensor x);
“”"

Compile the inline CUDA code
relu = load_inline(
name=“relu”,
cpp_sources=relu_cpp_source,
cuda_sources=relu_source,
functions=[“relu_cuda”],
verbose=True
)

class ModelNew(torch.nn.Module):
def init(self):
super(ModelNew, self).init()
self.relu = relu # The module containing the kernel

def forward(self, x):
    return self.relu.relu_cuda(x)
def get_inputs():
# randomly generate input tensors based on the model architecture
a = torch.randn(1, 128).cuda()
b = torch.randn(1, 128).cuda()
return [a, b]

def get_init_inputs():
# randomly generate tensors required for initialization based on the model architecture
return []



You are given the following architecture:

python
import torch
import torch.nn as nn

class Model(nn.Module):
“”"
Bray-Curtis Distance implementation.
Computes the Bray-Curtis distance between two sets of vectors.
“”"
def init(self):
super(Model, self).init()

def forward(self, x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
    """
    Compute Bray-Curtis distance between x and y.

    Args:
        x (torch.Tensor): First set of vectors [batch_size, feature_dim]
        y (torch.Tensor): Second set of vectors [batch_size, feature_dim]

    Returns:
        torch.Tensor: Bray-Curtis distances [batch_size]
    """
    # Input validation
    if x.shape != y.shape:
        raise ValueError(f"Input tensors must have the same shape, got {x.shape} and {y.shape}")
    
    if x.dim() != 2:
        raise ValueError(f"Input tensors must be 2D, got {x.dim()}D")
    
    # Compute Bray-Curtis distance: Σ|x_i - y_i| / Σ(|x_i| + |y_i|)
    # Step 1: Compute absolute differences
    diff = torch.abs(x - y)
    
    # Step 2: Compute numerator (sum of absolute differences)
    numerator = torch.sum(diff, dim=1)
    
    # Step 3: Compute denominator (sum of absolute values)
    denominator = torch.sum(torch.abs(x) + torch.abs(y), dim=1)
    
    # Step 4: Handle division by zero
    # When denominator is 0 (both vectors are all zeros), distance is 0
    distance = torch.where(denominator > 0, numerator / denominator, torch.zeros_like(numerator))
    
    return distance
batch_size = 256
feature_dim = 512

def get_inputs():
# Generate two sets of positive vectors (ecological data is typically non-negative)
x = torch.abs(torch.randn(batch_size, feature_dim)) + 0.1
y = torch.abs(torch.randn(batch_size, feature_dim)) + 0.1
return [x, y]

def get_init_inputs():
return [] # No special initialization inputs needed