import torch
import torch.nn as nn
import torch.nn.functional as F
def module_fn(x: torch.Tensor) -> torch.Tensor:
"""
Applies Sigmoid activation to the input tensor.
Args:
x (torch.Tensor): Input tensor of any shape.
Returns:
torch.Tensor: Output tensor with Sigmoid applied, same shape as input.
"""
return torch.sigmoid(x)
class Model(nn.Module):
"""
Simple model that performs a Sigmoid activation.
"""
def __init__(self):
super(Model, self).__init__()
def forward(self, x: torch.Tensor, fn=module_fn) -> torch.Tensor:
return fn(x)
batch_size = 16
dim = 16384
def get_inputs():
x = torch.randn(batch_size, dim)
return [x]
def get_init_inputs():
return [] # No special initialization inputs needed
import torch
import torch.nn as nn
class Model(nn.Module):
"""
Simple model that performs a Sigmoid activation.
"""
def __init__(self):
super(Model, self).__init__()
def forward(self, x: torch.Tensor) -> torch.Tensor:
"""
Applies Sigmoid activation to the input tensor.
Args:
x (torch.Tensor): Input tensor of any shape.
Returns:
torch.Tensor: Output tensor with Sigmoid applied, same shape as input.
"""
return torch.sigmoid(x)
batch_size = 16
dim = 16384
def get_inputs():
x = torch.randn(batch_size, dim)
return [x]
def get_init_inputs():
return [] # No special initialization inputs needed#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
template <typename scalar_t>
__forceinline__ __device__ float sigmoid_compute(float x) {
return 1.0f / (1.0f + expf(-x));
}
template <typename scalar_t>
__forceinline__ __device__ void process_element(const scalar_t* __restrict__ input,
scalar_t* __restrict__ output,
const int idx) {
float val = static_cast<float>(input[idx]);
output[idx] = static_cast<scalar_t>(sigmoid_compute<scalar_t>(val));
}
template <typename scalar_t>
__global__ void sigmoid_kernel(const scalar_t* __restrict__ input,
scalar_t* __restrict__ output,
const int64_t size) {
const int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < size) {
process_element<scalar_t>(input, output, idx);
}
}
torch::Tensor forward(torch::Tensor input) {
auto output = torch::empty_like(input);
const int64_t size = input.numel();
constexpr int threads = 256;
const int blocks = (size + threads - 1) / threads;
AT_DISPATCH_FLOATING_TYPES(input.scalar_type(), "sigmoid_kernel", [&] {
const auto* input_data = input.data_ptr<scalar_t>();
auto* output_data = output.data_ptr<scalar_t>();
sigmoid_kernel<scalar_t><<<blocks, threads>>>(input_data, output_data, size);
});
return output;
}
PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
m.def("forward", &forward, "Sigmoid forward (CUDA)");
}| Metric | Value | Unit | Variance | Samples |
|---|---|---|---|---|
| Executed Ipc Active | 0.720 | inst/cycle | 0.000 | 5 |
| Executed Ipc Elapsed | 0.280 | inst/cycle | 0.000 | 5 |
| Issue Slots Busy | 19.406 | % | 0.018 | 5 |
| Issued Ipc Active | 0.774 | inst/cycle | 0.000 | 5 |
| SM Busy | 19.406 | % | 0.018 | 5 |
| Memory Throughput | 287763869848.960 | byte/second | 8671476223554697216.000 | 5 |
| Mem Busy | 13.600 | % | 0.021 | 5 |
| Max Bandwidth | 12.606 | % | 0.020 | 5 |
| L1/TEX Hit Rate | 0.000 | % | 0.000 | 5 |
| L2 Hit Rate | 66.932 | % | 0.001 | 5 |
| Mem Pipes Busy | 7.552 | % | 0.005 | 5 |
| Warp Cycles Per Issued Instruction | 61.696 | cycle | 4.621 | 5 |
| Warp Cycles Per Executed Instruction | 66.404 | cycle | 5.352 | 5 |
| Avg. Active Threads Per Warp | 32.000 | 0.000 | 5 | |
| Avg. Not Predicated Off Threads Per Warp | 30.770 | 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 | 16.000 | block | 0.000 | 5 |
| Block Limit Shared Mem | 32.000 | block | 0.000 | 5 |
| Block Limit Warps | 8.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 | 74.572 | % | 0.128 | 5 |
| Achieved Active Warps Per SM | 47.724 | warp | 0.053 | 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 (74.9%) 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 | 612880.44 | μs |
| Device Time | 40.26 | μs |
| Self CPU Time | 40.08 | μ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 | 612840.36 | μs |
| Device Time | 40.26 | μs |
| Self CPU Time | 78.20 | μ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::empty_strided | ||
| CPU Time | 628558.68 | μs |
| Device Time | 0.00 | μs |
| Self CPU Time | 16170.62 | μ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 |
| cudaDeviceGetStreamPriorityRange | ||
| CPU Time | 611624.09 | μs |
| Device Time | 0.00 | μs |
| Self CPU Time | 611624.09 | μ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 |
| cudaLaunchKernel | ||
| CPU Time | 424706.29 | μs |
| Device Time | 19488.54 | μs |
| Self CPU Time | 424706.29 | μs |
| Self Device Time | 19488.54 | μs |
| CPU Memory Usage | 0 | B |
| Device Memory Usage | 0 | B |
| Self CPU Memory Usage | 0 | B |
| Self Device Memory Usage | 0 | B |
| void sigmoid_kernel<float>(float const*, float*, long) | ||
| CPU Time | 0.00 | μs |
| Device Time | 22043.94 | μs |
| Self CPU Time | 0.00 | μs |
| Self Device Time | 22043.94 | μs |
| CPU Memory Usage | 0 | B |
| Device Memory Usage | 0 | B |
| Self CPU Memory Usage | 0 | B |
| Self Device Memory Usage | 0 | B |
| cudaEventRecord | ||
| CPU Time | 15839.61 | μs |
| Device Time | 37508.34 | μs |
| Self CPU Time | 15839.61 | μs |
| Self Device Time | 37508.34 | μ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 | 59302.81 | μs |
| Device Time | 556654.52 | μs |
| Self CPU Time | 11426.34 | μ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 | 47878.12 | μs |
| Device Time | 556654.52 | μs |
| Self CPU Time | 13851.18 | μs |
| Self Device Time | 556654.52 | μ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 | 556654.52 | μs |
| Self CPU Time | 0.00 | μs |
| Self Device Time | 556654.52 | μs |
| CPU Memory Usage | 0 | B |
| Device Memory Usage | 0 | B |
| Self CPU Memory Usage | 0 | B |
| Self Device Memory Usage | 0 | B |
45279 warnings generated when compiling for host. Suppressed 45321 warnings (45274 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.