7_Conv3d_ReLU_LeakyReLU_GELU_Sigmoid_BiasAdd • optimized_thread_block_mapping_kernel_baseimport torch
import torch.nn as nn
import torch.nn.functional as F
def module_fn(
x: torch.Tensor,
conv_weight: torch.Tensor,
conv_bias: torch.Tensor,
bias: torch.Tensor,
) -> torch.Tensor:
"""
Applies 3D convolution followed by ReLU, LeakyReLU, GELU, Sigmoid activations and bias addition.
Args:
x (torch.Tensor): Input tensor of shape (batch_size, in_channels, depth, height, width)
conv_weight (torch.Tensor): 3D convolution weight tensor of shape
(out_channels, in_channels, kernel_size, kernel_size, kernel_size)
conv_bias (torch.Tensor): Bias tensor for 3D convolution of shape (out_channels)
bias (torch.Tensor): Bias tensor for addition of shape (out_channels, 1, 1, 1)
Returns:
torch.Tensor: Output tensor after applying convolution and activations
"""
x = F.conv3d(x, conv_weight, bias=conv_bias)
x = F.relu(x)
x = F.leaky_relu(x, negative_slope=0.01)
x = F.gelu(x)
x = torch.sigmoid(x)
x = x + bias
return x
class Model(nn.Module):
"""
Model that performs a 3D convolution, applies ReLU, LeakyReLU, GELU, Sigmoid activations, and bias in sequence.
"""
def __init__(self, in_channels, out_channels, kernel_size, bias_shape):
super(Model, self).__init__()
conv = nn.Conv3d(in_channels, out_channels, kernel_size)
self.bias = nn.Parameter(torch.randn(bias_shape) * 0.02)
self.conv_weight = nn.Parameter(conv.weight)
self.conv_bias = nn.Parameter(conv.bias)
self.bias = self.bias
def forward(self, x, fn=module_fn):
return fn(x, self.conv_weight, self.conv_bias, self.bias)
batch_size = 128
in_channels = 3
out_channels = 16
depth, height, width = 16, 32, 32
kernel_size = 3
bias_shape = (out_channels, 1, 1, 1)
def get_inputs():
return [torch.randn(batch_size, in_channels, depth, height, width)]
def get_init_inputs():
return [in_channels, out_channels, kernel_size, bias_shape]
import torch
import torch.nn as nn
class Model(nn.Module):
"""
Model that performs a 3D convolution, applies ReLU, LeakyReLU, GELU, Sigmoid activations, and bias in sequence.
"""
def __init__(self, in_channels, out_channels, kernel_size, bias_shape):
super(Model, self).__init__()
self.conv = nn.Conv3d(in_channels, out_channels, kernel_size)
self.bias = nn.Parameter(torch.randn(bias_shape) * 0.02)
def forward(self, x):
x = self.conv(x)
x = torch.relu(x)
x = torch.nn.functional.leaky_relu(x, negative_slope=0.01)
x = torch.nn.functional.gelu(x)
x = torch.sigmoid(x)
x = x + self.bias
return x
batch_size = 128
in_channels = 3
out_channels = 16
depth, height, width = 16, 32, 32
kernel_size = 3
bias_shape = (out_channels, 1, 1, 1)
def get_inputs():
return [torch.randn(batch_size, in_channels, depth, height, width)]
def get_init_inputs():
return [in_channels, out_channels, kernel_size, bias_shape]#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>
#define CHECK_CUDA(x) TORCH_CHECK(x.is_cuda(), #x " must be a CUDA tensor")
#define CHECK_CONTIGUOUS(x) TORCH_CHECK(x.is_contiguous(), #x " must be contiguous")
#define CHECK_INPUT(x) CHECK_CUDA(x); CHECK_CONTIGUOUS(x)
#define BLOCK_SIZE 256
#define WARP_SIZE 32
__device__ __forceinline__ float4 load_float4(const float* addr) {
float4 val;
val = *reinterpret_cast<const float4*>(addr);
return val;
}
__device__ __forceinline__ void store_float4(float* addr, float4 val) {
*reinterpret_cast<float4*>(addr) = val;
}
__device__ __forceinline__ float process_value(float val, const float* bias, int bias_idx) {
// ReLU
val = fmaxf(0.0f, val);
// LeakyReLU
val = fmaxf(0.01f * val, val);
// GELU
const float sqrt_2_over_pi = sqrtf(2.0f / M_PI);
val = 0.5f * val * (1.0f + tanhf(sqrt_2_over_pi * (val + 0.044715f * powf(val, 3.0f))));
// Sigmoid
val = 1.0f / (1.0f + expf(-val));
// Add bias
val += __ldg(&bias[bias_idx]);
return val;
}
__global__ void apply_activations_and_bias_kernel(
float* __restrict__ output, const float* __restrict__ bias,
int batch_size, int out_channels, int depth, int height, int width
) {
const int spatial_size = depth * height * width;
const int total_elements = batch_size * out_channels * spatial_size;
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
const int elements_per_thread = 4;
const int stride = gridDim.x * blockDim.x * elements_per_thread;
for (int idx = tid * elements_per_thread; idx < total_elements; idx += stride) {
if (idx + 3 < total_elements) {
// Load 4 elements
float4 data = load_float4(&output[idx]);
// Calculate bias index for the current position
int base_idx = idx / spatial_size;
int bias_idx = base_idx % out_channels;
// Process each component
data.x = process_value(data.x, bias, bias_idx);
data.y = process_value(data.y, bias, bias_idx);
data.z = process_value(data.z, bias, bias_idx);
data.w = process_value(data.w, bias, bias_idx);
// Store results back
store_float4(&output[idx], data);
} else {
// Handle remaining elements
for (int i = 0; i < 4 && idx + i < total_elements; ++i) {
int curr_idx = idx + i;
float val = output[curr_idx];
int bias_idx = (curr_idx / spatial_size) % out_channels;
output[curr_idx] = process_value(val, bias, bias_idx);
}
}
}
}
torch::Tensor module_fn_cuda(
torch::Tensor x,
torch::Tensor conv_weight,
torch::Tensor conv_bias,
torch::Tensor bias
) {
CHECK_INPUT(x);
CHECK_INPUT(conv_weight);
CHECK_INPUT(conv_bias);
CHECK_INPUT(bias);
auto output = torch::conv3d(x, conv_weight, conv_bias);
int batch_size = output.size(0);
int out_channels = output.size(1);
int depth = output.size(2);
int height = output.size(3);
int width = output.size(4);
int total_vectors = (batch_size * out_channels * depth * height * width + 3) / 4;
int blocks = (total_vectors + BLOCK_SIZE - 1) / BLOCK_SIZE;
apply_activations_and_bias_kernel<<<blocks, BLOCK_SIZE>>>(
output.data_ptr<float>(), bias.data_ptr<float>(),
batch_size, out_channels, depth, height, width
);
return output;
}
PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
m.def("forward", &module_fn_cuda, "CUDA implementation of module_fn");
}| Metric | Value | Unit | Variance | Samples |
|---|---|---|---|---|
| Executed Ipc Active | 3.068 | inst/cycle | 0.000 | 5 |
| Executed Ipc Elapsed | 2.968 | inst/cycle | 0.000 | 5 |
| Issue Slots Busy | 76.862 | % | 0.007 | 5 |
| Issued Ipc Active | 3.076 | inst/cycle | 0.000 | 5 |
| SM Busy | 77.212 | % | 0.007 | 5 |
| Memory Throughput | 2389519037799.892 | byte/second | 78076366010807107584.000 | 5 |
| Mem Busy | 40.434 | % | 0.028 | 5 |
| Max Bandwidth | 71.322 | % | 0.069 | 5 |
| L1/TEX Hit Rate | 43.458 | % | 0.013 | 5 |
| L2 Hit Rate | 50.436 | % | 0.005 | 5 |
| Mem Pipes Busy | 15.136 | % | 0.004 | 5 |
| Warp Cycles Per Issued Instruction | 18.642 | cycle | 0.000 | 5 |
| Warp Cycles Per Executed Instruction | 18.662 | cycle | 0.000 | 5 |
| Avg. Active Threads Per Warp | 27.850 | 0.000 | 5 | |
| Avg. Not Predicated Off Threads Per Warp | 25.140 | 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 | 8.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 | 89.978 | % | 0.010 | 5 |
| Achieved Active Warps Per SM | 57.586 | warp | 0.004 | 5 |
| Rule | Description |
|---|---|
| INF HighPipeUtilization | ALU is the highest-utilized pipeline (42.6%) based on active cycles, taking into account the rates of its different instructions. It executes integer and logic operations. It is well-utilized, but should not be a bottleneck. |
| INF Occupancy | This kernel's theoretical occupancy is not impacted by any block limit. |
| Operation / Metric | Value | Unit |
|---|---|---|
| aten::conv3d | ||
| CPU Time | 591687.79 | μs |
| Device Time | 3845126.60 | μs |
| Self CPU Time | 9573.46 | μ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::convolution | ||
| CPU Time | 582114.33 | μs |
| Device Time | 3845126.60 | μs |
| Self CPU Time | 13034.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::_convolution | ||
| CPU Time | 569079.44 | μs |
| Device Time | 3845126.60 | μs |
| Self CPU Time | 26279.66 | μ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::cudnn_convolution | ||
| CPU Time | 469830.20 | μs |
| Device Time | 3337495.91 | μs |
| Self CPU Time | 164946.33 | μs |
| Self Device Time | 3337495.91 | μ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_fprop_implicit_gemm_indexed_f32f32_f32f32_f32_nchwkcrs_nchw_tilesize32x32x8_stage3_warpsize1x2x1_g1_ffma_aligna4_alignc4_execute_kernel__5x_cudnn | ||
| CPU Time | 0.00 | μs |
| Device Time | 3337494.44 | μs |
| Self CPU Time | 0.00 | μs |
| Self Device Time | 3337494.44 | μ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 | 3586242.76 | μs |
| Device Time | 71527.45 | μs |
| Self CPU Time | 3586242.76 | μs |
| Self Device Time | 71527.45 | μs |
| CPU Memory Usage | 0 | B |
| Device Memory Usage | 0 | B |
| Self CPU Memory Usage | 0 | B |
| Self Device Memory Usage | 0 | B |
45290 warnings generated when compiling for host. Suppressed 45324 warnings (45277 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.