53_Gemm_Scaling_Hardtanh_GELU • 53gemm_scaling_hardtanh_gelu_fast_baseimport torch
import torch.nn as nn
import torch.nn.functional as F
def module_fn(
x: torch.Tensor,
scaling_factor: float,
hardtanh_min: float,
hardtanh_max: float,
weight: torch.Tensor,
bias: torch.Tensor,
) -> torch.Tensor:
"""
Applies GEMM, scaling, hardtanh and GELU activation.
Args:
x (torch.Tensor): Input tensor of shape (batch_size, in_features)
scaling_factor (float): Factor to scale the GEMM output
hardtanh_min (float): Minimum value for hardtanh
hardtanh_max (float): Maximum value for hardtanh
weight (torch.Tensor): Weight matrix of shape (out_features, in_features)
bias (torch.Tensor): Bias vector of shape (out_features)
Returns:
torch.Tensor: Output tensor after applying GEMM, scaling, hardtanh and GELU,
with shape (batch_size, out_features)
"""
x = F.linear(x, weight, bias)
x = x * scaling_factor
x = F.hardtanh(x, min_val=hardtanh_min, max_val=hardtanh_max)
x = F.gelu(x)
return x
class Model(nn.Module):
"""
Model that performs a GEMM, scaling, hardtanh, and GELU activation.
"""
def __init__(
self, in_features, out_features, scaling_factor, hardtanh_min, hardtanh_max
):
super(Model, self).__init__()
gemm = nn.Linear(in_features, out_features)
self.weight = nn.Parameter(gemm.weight)
self.bias = nn.Parameter(gemm.bias)
self.scaling_factor = scaling_factor
self.hardtanh_min = hardtanh_min
self.hardtanh_max = hardtanh_max
def forward(self, x, fn=module_fn):
return fn(
x,
self.scaling_factor,
self.hardtanh_min,
self.hardtanh_max,
self.weight,
self.bias,
)
batch_size = 128
in_features = 1024
out_features = 512
scaling_factor = 0.5
hardtanh_min = -2
hardtanh_max = 2
def get_inputs():
return [torch.randn(batch_size, in_features)]
def get_init_inputs():
return [in_features, out_features, scaling_factor, hardtanh_min, hardtanh_max]
import torch
import torch.nn as nn
class Model(nn.Module):
"""
Model that performs a GEMM, scaling, hardtanh, and GELU activation.
"""
def __init__(self, in_features, out_features, scaling_factor, hardtanh_min, hardtanh_max):
super(Model, self).__init__()
self.gemm = nn.Linear(in_features, out_features)
self.scaling_factor = scaling_factor
self.hardtanh = nn.Hardtanh(min_val=hardtanh_min, max_val=hardtanh_max)
self.gelu = nn.GELU()
def forward(self, x):
x = self.gemm(x)
x = x * self.scaling_factor
x = self.hardtanh(x)
x = self.gelu(x)
return x
batch_size = 128
in_features = 1024
out_features = 512
scaling_factor = 0.5
hardtanh_min = -2
hardtanh_max = 2
def get_inputs():
return [torch.randn(batch_size, in_features)]
def get_init_inputs():
return [in_features, out_features, scaling_factor, hardtanh_min, hardtanh_max]#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
// This kernel applies scaling, Hardtanh and GELU activation using a grid-stride loop.
// Note: No shared memory is used so __syncthreads() is not required.
template <typename scalar_t>
__global__ void fused_activation_kernel_fast(
scalar_t* x,
scalar_t scaling_factor,
scalar_t hardtanh_min,
scalar_t hardtanh_max,
int64_t numel) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;
for (int i = idx; i < numel; i += stride) {
scalar_t val = x[i];
// Scaling
val = val * scaling_factor;
// Hardtanh
val = min(max(val, hardtanh_min), hardtanh_max);
// GELU approximation
const scalar_t c = static_cast<scalar_t>(0.044715);
const scalar_t sqrt_2_over_pi = static_cast<scalar_t>(0.7978845608028654); // sqrt(2.0 / pi)
scalar_t x_cube = val * val * val;
scalar_t tanh_arg = sqrt_2_over_pi * (val + c * x_cube);
scalar_t tanh_res = tanh(tanh_arg);
val = static_cast<scalar_t>(0.5) * val * (static_cast<scalar_t>(1.0) + tanh_res);
x[i] = val;
}
}
void fused_activation_cuda(
torch::Tensor& x,
double scaling_factor,
double hardtanh_min,
double hardtanh_max) {
const auto numel = x.numel();
const int threads = 1024;
const int blocks = (numel + threads - 1) / threads;
AT_DISPATCH_FLOATING_TYPES(x.scalar_type(), "fused_activation_cuda", ([&] {
fused_activation_kernel_fast<scalar_t><<<blocks, threads>>>(
x.data_ptr<scalar_t>(),
static_cast<scalar_t>(scaling_factor),
static_cast<scalar_t>(hardtanh_min),
static_cast<scalar_t>(hardtanh_max),
numel);
}));
}
torch::Tensor module_fn_forward(
torch::Tensor x,
double scaling_factor,
double hardtanh_min,
double hardtanh_max,
torch::Tensor weight,
torch::Tensor bias) {
// Ensure inputs are contiguous and on CUDA
x = x.contiguous().cuda();
weight = weight.contiguous().cuda();
bias = bias.contiguous().cuda();
// Linear transformation: x = x @ weight.T + bias
auto xw = torch::matmul(x, weight.t()) + bias;
// Apply fused activation functions
fused_activation_cuda(xw, scaling_factor, hardtanh_min, hardtanh_max);
return xw;
}
PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
m.def("forward", &module_fn_forward, "Module function forward (CUDA)");
}
| Metric | Value | Unit | Variance | Samples |
|---|---|---|---|---|
| Executed Ipc Active | 0.846 | inst/cycle | 0.004 | 5 |
| Executed Ipc Elapsed | 0.156 | inst/cycle | 0.000 | 5 |
| Issue Slots Busy | 23.014 | % | 2.623 | 5 |
| Issued Ipc Active | 0.922 | inst/cycle | 0.004 | 5 |
| SM Busy | 23.014 | % | 2.623 | 5 |
| Memory Throughput | 78061840308.370 | byte/second | 5854716052622974976.000 | 5 |
| Mem Busy | 10.948 | % | 0.104 | 5 |
| Max Bandwidth | 6.982 | % | 0.044 | 5 |
| L1/TEX Hit Rate | 50.000 | % | 0.000 | 5 |
| L2 Hit Rate | 81.850 | % | 0.164 | 5 |
| Mem Pipes Busy | 1.738 | % | 0.003 | 5 |
| Warp Cycles Per Issued Instruction | 29.400 | cycle | 0.027 | 5 |
| Warp Cycles Per Executed Instruction | 32.048 | cycle | 0.034 | 5 |
| Avg. Active Threads Per Warp | 30.140 | 0.000 | 5 | |
| Avg. Not Predicated Off Threads Per Warp | 28.950 | 0.000 | 5 | |
| Max Active Clusters | 0.000 | cluster | 0.000 | 5 |
| Max Cluster Size | 8.000 | block | 0.000 | 5 |
| Overall GPU Occupancy | 0.000 | % | 0.000 | 5 |
| Cluster Occupancy | 0.000 | % | 0.000 | 5 |
| Block Limit SM | 32.000 | block | 0.000 | 5 |
| Block Limit Registers | 4.000 | block | 0.000 | 5 |
| Block Limit Shared Mem | 8.000 | block | 0.000 | 5 |
| Block Limit Warps | 2.000 | block | 0.000 | 5 |
| Theoretical Active Warps per SM | 64.000 | warp | 0.000 | 5 |
| Theoretical Occupancy | 100.000 | % | 0.000 | 5 |
| Achieved Occupancy | 45.026 | % | 0.778 | 5 |
| Achieved Active Warps Per SM | 28.818 | warp | 0.317 | 5 |
| Rule | Description |
|---|---|
| WRN HighPipeUtilization | All compute pipelines are under-utilized. Either this kernel is very small or it doesn't issue enough warps per scheduler. Check the Launch Statistics and Scheduler Statistics sections for further details. |
| INF CPIStall | Check the Warp Stall Sampling (All Cycles) table for the top stall locations in your source based on sampling data. The Kernel Profiling Guide (https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html#metrics-reference) provides more details on each stall reason. |
| WRN Occupancy | This kernel's theoretical occupancy is not impacted by any block limit. The difference between calculated theoretical (100.0%) and measured achieved occupancy (45.4%) can be the result of warp scheduling overheads or workload imbalances during the kernel execution. Load imbalances can occur between warps within a block as well as across blocks of the same kernel. See the CUDA Best Practices Guide (https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#occupancy) for more details on optimizing occupancy. |
| Operation / Metric | Value | Unit |
|---|---|---|
| aten::to | ||
| CPU Time | 326637.84 | μs |
| Device Time | 178.50 | μs |
| Self CPU Time | 10002.86 | μs |
| Self Device Time | 0.00 | μs |
| CPU Memory Usage | 0 | B |
| Device Memory Usage | 0 | B |
| Self CPU Memory Usage | 0 | B |
| Self Device Memory Usage | 0 | B |
| aten::_to_copy | ||
| CPU Time | 316634.98 | μs |
| Device Time | 178.50 | μs |
| Self CPU Time | 120.75 | μs |
| Self Device Time | 0.00 | μs |
| CPU Memory Usage | 0 | B |
| Device Memory Usage | 0 | B |
| Self CPU Memory Usage | 0 | B |
| Self Device Memory Usage | 0 | B |
| aten::matmul | ||
| CPU Time | 309310.31 | μs |
| Device Time | 125462.86 | μs |
| Self CPU Time | 9084.69 | μs |
| Self Device Time | 0.00 | μs |
| CPU Memory Usage | 0 | B |
| Device Memory Usage | 0 | B |
| Self CPU Memory Usage | 0 | B |
| Self Device Memory Usage | 0 | B |
| aten::mm | ||
| CPU Time | 300225.62 | μs |
| Device Time | 125462.86 | μs |
| Self CPU Time | 175472.98 | μs |
| Self Device Time | 125462.86 | μs |
| CPU Memory Usage | 0 | B |
| Device Memory Usage | 0 | B |
| Self CPU Memory Usage | 0 | B |
| Self Device Memory Usage | 0 | B |
| cudaLaunchKernel | ||
| CPU Time | 380349.81 | μs |
| Device Time | 12800.51 | μs |
| Self CPU Time | 380349.81 | μs |
| Self Device Time | 12800.51 | μs |
| CPU Memory Usage | 0 | B |
| Device Memory Usage | 0 | B |
| Self CPU Memory Usage | 0 | B |
| Self Device Memory Usage | 0 | B |
| aten::zero_ | ||
| CPU Time | 336896.41 | μs |
| Device Time | 676752.16 | μs |
| Self CPU Time | 16038.89 | μs |
| Self Device Time | 0.00 | μs |
| CPU Memory Usage | 0 | B |
| Device Memory Usage | 0 | B |
| Self CPU Memory Usage | 0 | B |
| Self Device Memory Usage | 0 | B |
| aten::fill_ | ||
| CPU Time | 320859.50 | μs |
| Device Time | 676752.16 | μs |
| Self CPU Time | 19070.45 | μs |
| Self Device Time | 676752.16 | μs |
| CPU Memory Usage | 0 | B |
| Device Memory Usage | 0 | B |
| Self CPU Memory Usage | 0 | B |
| Self Device Memory Usage | 0 | B |
| void at::native::vectorized_elementwise_kernel<4, at::native::FillFunctor<int>, at::detail::Array<char*, 1> >(int, at::native::FillFunctor<int>, at::detail::Array<char*, 1>) | ||
| CPU Time | 0.00 | μs |
| Device Time | 676752.16 | μs |
| Self CPU Time | 0.00 | μs |
| Self Device Time | 676752.16 | μs |
| CPU Memory Usage | 0 | B |
| Device Memory Usage | 0 | B |
| Self CPU Memory Usage | 0 | B |
| Self Device Memory Usage | 0 | B |
45286 warnings generated when compiling for host. Suppressed 45326 warnings (45279 in non-user code, 47 NOLINT). Use -header-filter=.* to display errors from all non-system headers. Use -system-headers to display errors from system headers as well.