import torch
import torch.nn as nn
import torch.nn.functional as F
def module_fn(
x: torch.Tensor,
i2h_weight: torch.Tensor,
i2h_bias: torch.Tensor,
h2o_weight: torch.Tensor,
h2o_bias: torch.Tensor,
hidden: torch.Tensor,
) -> torch.Tensor:
"""
Vanilla RNN forward pass
Args:
x: Input tensor of shape (batch_size, input_size)
i2h_weight: Weight tensor for input-to-hidden layer
i2h_bias: Bias tensor for input-to-hidden layer
h2o_weight: Weight tensor for hidden-to-output layer
h2o_bias: Bias tensor for hidden-to-output layer
hidden: Hidden state tensor
Returns:
Output tensor of shape (batch_size, output_size)
"""
hidden = hidden.to(x.device)
combined = torch.cat((x, hidden), dim=1)
hidden = torch.tanh(F.linear(combined, i2h_weight, i2h_bias))
output = F.linear(hidden, h2o_weight, h2o_bias)
return output
class Model(nn.Module):
def __init__(self, input_size: int, hidden_size: int, output_size: int):
"""
Initialize the Vanilla RNN model.
:param input_size: The number of input features (int).
:param hidden_size: The size of the hidden state (int).
:param output_size: The number of output features (int).
"""
super(Model, self).__init__()
self.input_size = input_size
self.hidden_size = hidden_size
self.output_size = output_size
self.hidden = nn.Parameter(torch.randn((batch_size, hidden_size)))
# Extract parameters from linear layers
i2h = nn.Linear(input_size + hidden_size, hidden_size)
self.i2h_weight = nn.Parameter(i2h.weight.data.clone())
self.i2h_bias = nn.Parameter(i2h.bias.data.clone())
h2o = nn.Linear(hidden_size, output_size)
self.h2o_weight = nn.Parameter(h2o.weight.data.clone())
self.h2o_bias = nn.Parameter(h2o.bias.data.clone())
def forward(self, x: torch.Tensor, fn=module_fn) -> torch.Tensor:
return fn(
x,
self.i2h_weight,
self.i2h_bias,
self.h2o_weight,
self.h2o_bias,
self.hidden,
)
batch_size = 8
input_size = 1024
hidden_size = 256
output_size = 128
sequence_length = 256
def get_inputs():
return [torch.randn(batch_size, input_size)]
def get_init_inputs():
return [input_size, hidden_size, output_size]
import torch
import torch.nn as nn
class Model(nn.Module):
def __init__(self, input_size: int, hidden_size: int, output_size: int):
"""
Initialize the Vanilla RNN model.
:param input_size: The number of input features (int).
:param hidden_size: The size of the hidden state (int).
:param output_size: The number of output features (int).
"""
super(Model, self).__init__()
self.input_size = input_size
self.hidden_size = hidden_size
self.output_size = output_size
self.hidden = torch.randn((batch_size, hidden_size))
# Define the RNN cell components (input to hidden, hidden to hidden, and hidden to output)
self.i2h = nn.Linear(input_size + hidden_size, hidden_size) # Input to hidden
self.h2o = nn.Linear(hidden_size, output_size) # Hidden to output
self.tanh = nn.Tanh() # Activation function for hidden state
def forward(self, x: torch.Tensor) -> torch.Tensor:
"""
Forward pass of the Vanilla RNN.
:param x: Input tensor of shape (batch_size, input_size).
:param hidden: Hidden state tensor of shape (batch_size, hidden_size).
:return: Output tensor of shape (batch_size, output_size), and the new hidden state.
"""
self.hidden = self.hidden.to(x.device)
combined = torch.cat((x, self.hidden), dim=1) # Concatenate input and hidden state
self.hidden = self.tanh(self.i2h(combined)) # Update hidden state
output = self.h2o(self.hidden) # Compute output
return output
batch_size = 8
input_size = 1024
hidden_size = 256
output_size = 128
sequence_length = 256
def get_inputs():
return [torch.randn(batch_size, input_size)]
def get_init_inputs():
return [input_size, hidden_size, output_size]
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#define BLOCK_SIZE 256
__global__ void balanced_concat_kernel(
const float* __restrict__ x,
const float* __restrict__ hidden,
float* __restrict__ combined,
int batch_size,
int x_size,
int hidden_size
) {
const int total_cols = x_size + hidden_size;
const int total_elements = batch_size * total_cols;
for (int idx = blockIdx.x * blockDim.x + threadIdx.x;
idx < total_elements;
idx += blockDim.x * gridDim.x) {
int row = idx / total_cols;
int col = idx % total_cols;
if (col < x_size) {
combined[idx] = x[row * x_size + col];
} else {
combined[idx] = hidden[row * hidden_size + (col - x_size)];
}
}
}
torch::Tensor module_fn_cuda(
torch::Tensor x,
torch::Tensor i2h_weight,
torch::Tensor i2h_bias,
torch::Tensor h2o_weight,
torch::Tensor h2o_bias,
torch::Tensor hidden
) {
x = x.contiguous().cuda();
i2h_weight = i2h_weight.contiguous().cuda();
i2h_bias = i2h_bias.contiguous().cuda();
h2o_weight = h2o_weight.contiguous().cuda();
h2o_bias = h2o_bias.contiguous().cuda();
hidden = hidden.contiguous().cuda();
const int batch_size = x.size(0);
const int x_size = x.size(1);
const int hidden_size = hidden.size(1);
auto options = torch::TensorOptions().dtype(x.dtype()).device(x.device());
torch::Tensor combined = torch::empty({batch_size, x_size + hidden_size}, options);
const int total_elements = batch_size * (x_size + hidden_size);
const int blocks = (total_elements + BLOCK_SIZE - 1) / BLOCK_SIZE;
balanced_concat_kernel<<<blocks, BLOCK_SIZE>>>(
x.data_ptr<float>(),
hidden.data_ptr<float>(),
combined.data_ptr<float>(),
batch_size,
x_size,
hidden_size
);
torch::Tensor hidden_new = torch::tanh(torch::addmm(i2h_bias, combined, i2h_weight.t()));
torch::Tensor output = torch::addmm(h2o_bias, hidden_new, h2o_weight.t());
return output;
}
PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
m.def("forward", &module_fn_cuda, "Module forward (CUDA)");
}
Metric | Value | Unit | Variance | Samples |
---|---|---|---|---|
Executed Ipc Active | 0.304 | inst/cycle | 0.001 | 5 |
Executed Ipc Elapsed | 0.030 | inst/cycle | 0.000 | 5 |
Issue Slots Busy | 8.100 | % | 0.391 | 5 |
Issued Ipc Active | 0.324 | inst/cycle | 0.001 | 5 |
SM Busy | 8.100 | % | 0.391 | 5 |
Memory Throughput | 13992682349.122 | byte/second | 176910433180597920.000 | 5 |
Mem Busy | 10.260 | % | 0.147 | 5 |
Max Bandwidth | 5.500 | % | 0.032 | 5 |
L1/TEX Hit Rate | 0.000 | % | 0.000 | 5 |
L2 Hit Rate | 97.316 | % | 0.318 | 5 |
Mem Pipes Busy | 0.878 | % | 0.001 | 5 |
Warp Cycles Per Issued Instruction | 23.398 | cycle | 0.338 | 5 |
Warp Cycles Per Executed Instruction | 24.930 | cycle | 0.386 | 5 |
Avg. Active Threads Per Warp | 32.000 | 0.000 | 5 | |
Avg. Not Predicated Off Threads Per Warp | 25.920 | 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 | 10.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 | 12.228 | % | 0.002 | 5 |
Achieved Active Warps Per SM | 7.828 | warp | 0.001 | 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 (12.2%) 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 | 632341.49 | μs |
Device Time | 67.04 | μs |
Self CPU Time | 34799.80 | μ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 | 597541.69 | μs |
Device Time | 67.04 | μs |
Self CPU Time | 126.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::empty_strided | ||
CPU Time | 596999.58 | μs |
Device Time | 0.00 | μs |
Self CPU Time | 147.30 | μ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 | 594232.30 | μs |
Device Time | 0.00 | μs |
Self CPU Time | 594232.30 | μ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::addmm | ||
CPU Time | 807161.42 | μs |
Device Time | 344749.79 | μs |
Self CPU Time | 448262.21 | μs |
Self Device Time | 344749.79 | μs |
CPU Memory Usage | 0 | B |
Device Memory Usage | 0 | B |
Self CPU Memory Usage | 0 | B |
Self Device Memory Usage | 0 | B |
sm80_xmma_gemm_f32f32_f32f32_f32_tn_n_tilesize32x32x8_stage3_warpsize1x2x1_ffma_aligna4_alignc4_execute_kernel__51_cublas | ||
CPU Time | 0.00 | μs |
Device Time | 166649.01 | μs |
Self CPU Time | 0.00 | μs |
Self Device Time | 166649.01 | μ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 | 116409.78 | μs |
Device Time | 1099210.64 | μs |
Self CPU Time | 25804.63 | μ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 | 90606.85 | μs |
Device Time | 1099210.64 | μs |
Self CPU Time | 30116.53 | μs |
Self Device Time | 1099210.64 | μ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 | 1099210.64 | μs |
Self CPU Time | 0.00 | μs |
Self Device Time | 1099210.64 | μ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.