skill_example/references/t4-optimization-guide.md

T4 Turing GPU Optimization Guide

Overview

The NVIDIA T4 GPU, based on the Turing architecture (sm_75), is widely deployed in cloud inference environments. While it lacks many features of newer GPUs (no BF16, limited shared memory), it remains important for cost-effective inference. This guide covers optimization strategies specific to T4 hardware.

T4 Hardware Specifications

Specification	Value
Architecture	Turing (sm_75)
Streaming Multiprocessors (SMs)	40
Memory Bandwidth	320 GB/s (GDDR6)
Shared Memory per SM	64 KB (configurable)
L2 Cache	4 MB
FP32 CUDA Cores	2560
Tensor Cores (2nd gen)	320
Memory	16 GB GDDR6
Memory Bus Width	256-bit
TDP	70W
Max Threads per SM	1024
Max Threads per Block	1024
Warp Size	32
Max Warps per SM	32
Register File per SM	64 KB
Max Registers per Thread	255

Key Limitations Compared to H100/A100

Feature	T4	A100	H100
BF16 Support	No	Yes	Yes
FP16 Tensor Cores	Yes (2nd gen)	Yes (3rd gen)	Yes (4th gen)
Memory Type	GDDR6	HBM2e	HBM3
Bandwidth	320 GB/s	2.0 TB/s	3.35 TB/s
Shared Memory/SM	64 KB	164 KB	192 KB
SMs	40	108	132
Memory	16 GB	40/80 GB	80 GB
TDP	70W	400W	700W
INT8 Tensor Cores	Yes	Yes	Yes
FP8 Support	No	No	Yes
TMA	No	No	Yes

No BF16 Support: FP16 Only

The T4 does not support BF16 (bfloat16). All kernels must use FP16 (float16) or FP32. This is the single most important consideration when porting kernels to T4.

Converting BF16 Kernels to FP16

// BEFORE (H100/A100 -- BF16):
#include <cuda_bf16.h>
__nv_bfloat16 val = __float2bfloat16(x);
float result = __bfloat162float(val);
__nv_bfloat162 pair = __halves2bfloat162(a, b);

// AFTER (T4 -- FP16):
#include <cuda_fp16.h>
half val = __float2half(x);
float result = __half2float(val);
half2 pair = __halves2half2(a, b);

Vectorized FP16 Access on T4

#include <cuda_fp16.h>

// Load 8 FP16 values (128 bits) via float4
__device__ __forceinline__ void load_fp16x8_t4(
    const half* ptr,
    int base_idx,
    float vals[8]
) {
    float4 packed = reinterpret_cast<const float4*>(ptr)[base_idx / 8];
    half2* pairs = reinterpret_cast<half2*>(&packed);

    #pragma unroll
    for (int i = 0; i < 4; i++) {
        vals[2 * i]     = __half2float(__low2half(pairs[i]));
        vals[2 * i + 1] = __half2float(__high2half(pairs[i]));
    }
}

// Store 8 FP16 values (128 bits) via float4
__device__ __forceinline__ void store_fp16x8_t4(
    half* ptr,
    int base_idx,
    const float vals[8]
) {
    float4 packed;
    half2* pairs = reinterpret_cast<half2*>(&packed);

    #pragma unroll
    for (int i = 0; i < 4; i++) {
        pairs[i] = __halves2half2(
            __float2half(vals[2 * i]),
            __float2half(vals[2 * i + 1])
        );
    }

    reinterpret_cast<float4*>(ptr)[base_idx / 8] = packed;
}

FP16 Overflow Considerations

FP16 has a much narrower dynamic range than BF16:

Property	FP16	BF16
Max value	65,504	~3.4e38
Min positive normal	~6.1e-5	~1.2e-38
Precision	~3.3 decimal digits	~3 decimal digits

This means you must be more careful with:

// DANGER: FP16 overflow in attention scores
// If seq_len is large and values are not scaled, scores can overflow 65504
float score = dot_product / sqrtf(head_dim);  // Scale BEFORE exponentiation

// DANGER: FP16 underflow in normalization
// Very small variances can underflow
float variance = sum_sq / n;
float inv_std = rsqrtf(variance + epsilon);  // epsilon prevents 1/0

// SAFE: Always compute reductions in FP32
float sum = 0.0f;  // FP32 accumulator
for (int i = tid; i < n; i += blockDim.x) {
    sum += __half2float(input[i]);  // Convert to FP32 for accumulation
}

Memory-Bound Optimization

The T4's 320 GB/s bandwidth is roughly 6x less than A100 and 10x less than H100. Most deep learning kernels on T4 will be memory-bound.

Bandwidth Utilization Targets

Achieved Bandwidth	Assessment	Action
> 250 GB/s (78%)	Excellent	Kernel is well-optimized
200-250 GB/s (62-78%)	Good	Minor improvements possible
100-200 GB/s (31-62%)	Fair	Check coalescing and vectorization
< 100 GB/s (31%)	Poor	Significant optimization needed

Maximizing Bandwidth on T4

// Rule 1: Always use vectorized loads (float4 = 16 bytes per thread)
float4 data = reinterpret_cast<const float4*>(input)[tid];

// Rule 2: Ensure coalesced access
// Good: consecutive threads access consecutive addresses
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = input[idx];

// Rule 3: Minimize global memory round-trips by fusing operations
// BAD: Two kernel launches = two memory round-trips
//   rmsnorm(x, weight, temp);
//   gelu(temp, output);
// GOOD: Single fused kernel
//   fused_rmsnorm_gelu(x, weight, output);

// Rule 4: Use shared memory for data reuse
__shared__ float smem[TILE_SIZE];
// Load once from global memory, use multiple times from shared

Arithmetic Intensity

To determine if a kernel is memory-bound or compute-bound on T4:

Arithmetic Intensity = FLOPs / Bytes Transferred

T4 compute-to-memory ratio:
  FP32: 8.1 TFLOPS / 320 GB/s = 25.3 FLOP/byte
  FP16: 65 TFLOPS / 320 GB/s = 203 FLOP/byte (with Tensor Cores)

If your kernel's arithmetic intensity < 25.3 (FP32) or < 203 (FP16):
  => Memory-bound: Focus on bandwidth optimization
Otherwise:
  => Compute-bound: Focus on Tensor Core utilization

Common operations and their arithmetic intensity:

Operation	Approximate AI	Bound on T4
RMSNorm	~5 FLOP/byte	Memory-bound
GELU	~3 FLOP/byte	Memory-bound
Softmax	~5 FLOP/byte	Memory-bound
RoPE	~4 FLOP/byte	Memory-bound
MatMul (large)	~100+ FLOP/byte	Compute-bound
Attention (small seq)	~10 FLOP/byte	Memory-bound

Smaller Tile Sizes

The T4's 64 KB shared memory per SM (compared to 164 KB on A100 or 192 KB on H100) requires smaller tile sizes.

Tile Size Recommendations

Operation	H100 Tile	A100 Tile	T4 Tile
MatMul	128x128	128x128	64x64
Attention	128x64	128x64	64x32
Reduction	N/A	N/A	N/A
Conv2D tile	64x64	64x64	32x32

Shared Memory Budget

// T4: 64 KB shared memory max
// Common configurations:

// For reduction kernels (RMSNorm, LayerNorm):
// Only need warp reduction buffer -- minimal shared memory
__shared__ float warp_results[32];  // 128 bytes -- plenty of room

// For tiled matmul:
// Two tiles in shared memory
// 64x64 tile of FP16 = 64 * 64 * 2 = 8 KB per tile
// Two tiles = 16 KB -- fits easily
__shared__ half tile_a[64][64 + 1];  // +1 for bank conflict avoidance
__shared__ half tile_b[64][64 + 1];

// For attention with larger tiles:
// 64x32 Q tile + 32x64 K tile + 32x64 V tile
// = 64*32*2 + 32*64*2 + 32*64*2 = 4 + 4 + 4 = 12 KB
// Plus score buffer: 64*32*4 = 8 KB
// Total: ~20 KB -- fits in 64 KB

// WARNING: Do NOT use tile sizes designed for H100/A100
// __shared__ half tile[128][128];  // = 32 KB per tile -- too large for T4!

Dynamic Shared Memory on T4

// Query available shared memory
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
printf("Shared memory per block: %zu bytes\n", prop.sharedMemPerBlock);
printf("Shared memory per SM: %zu bytes\n", prop.sharedMemPerMultiprocessor);

// T4 output:
// Shared memory per block: 49152 bytes (48 KB default)
// Shared memory per SM: 65536 bytes (64 KB)

// To use more than 48 KB, request explicitly:
cudaFuncSetAttribute(
    my_kernel,
    cudaFuncAttributeMaxDynamicSharedMemorySize,
    64 * 1024  // Request up to 64 KB
);

INT8 Quantization

The T4 has strong INT8 Tensor Core support, making quantization an effective optimization strategy.

INT8 Quantization for Inference

import torch
from torch.quantization import quantize_dynamic

# Dynamic quantization (easiest)
model_int8 = quantize_dynamic(
    model,
    {torch.nn.Linear},
    dtype=torch.qint8,
)

# Or using bitsandbytes for LLM quantization
import bitsandbytes as bnb

# 8-bit linear layer
linear_8bit = bnb.nn.Linear8bitLt(
    input_features=4096,
    output_features=4096,
    bias=False,
    has_fp16_weights=False,
)

INT8 Kernel Example

#include <cuda_runtime.h>
#include <cstdint>

// INT8 GEMV kernel for T4 (simplified)
__global__ void int8_gemv_kernel(
    const int8_t* __restrict__ weight,  // [out_features, in_features]
    const int8_t* __restrict__ input,   // [in_features]
    float* __restrict__ output,          // [out_features]
    const float* __restrict__ scale_w,   // [out_features] per-channel scale
    const float scale_x,                 // input scale
    const int in_features,
    const int out_features
) {
    const int row = blockIdx.x * blockDim.x + threadIdx.x;
    if (row >= out_features) return;

    const int8_t* w_row = weight + row * in_features;

    // Accumulate dot product in INT32
    int32_t acc = 0;
    for (int i = 0; i < in_features; i += 4) {
        // Load 4 INT8 values at once (32 bits)
        int32_t w4 = reinterpret_cast<const int32_t*>(w_row)[i / 4];
        int32_t x4 = reinterpret_cast<const int32_t*>(input)[i / 4];

        int8_t* w_bytes = reinterpret_cast<int8_t*>(&w4);
        int8_t* x_bytes = reinterpret_cast<int8_t*>(&x4);

        acc += (int32_t)w_bytes[0] * (int32_t)x_bytes[0];
        acc += (int32_t)w_bytes[1] * (int32_t)x_bytes[1];
        acc += (int32_t)w_bytes[2] * (int32_t)x_bytes[2];
        acc += (int32_t)w_bytes[3] * (int32_t)x_bytes[3];
    }

    // Dequantize to float
    output[row] = (float)acc * scale_w[row] * scale_x;
}

TensorRT Integration

TensorRT is highly optimized for T4 and often provides the best inference performance.

Converting a Model to TensorRT

import torch
import tensorrt as trt

# Export model to ONNX first
torch.onnx.export(
    model,
    dummy_input,
    "model.onnx",
    opset_version=17,
    input_names=["input"],
    output_names=["output"],
    dynamic_axes={
        "input": {0: "batch_size", 1: "seq_len"},
        "output": {0: "batch_size", 1: "seq_len"},
    },
)

# Build TensorRT engine
logger = trt.Logger(trt.Logger.WARNING)
builder = trt.Builder(logger)
network = builder.create_network(
    1 << int(trt.NetworkDefinitionCreationFlag.EXPLICIT_BATCH)
)
parser = trt.OnnxParser(network, logger)

with open("model.onnx", "rb") as f:
    parser.parse(f.read())

config = builder.create_builder_config()
config.set_memory_pool_limit(trt.MemoryPoolType.WORKSPACE, 1 << 30)  # 1 GB

# Enable FP16 (T4 supports this)
config.set_flag(trt.BuilderFlag.FP16)

# Enable INT8 for maximum performance
config.set_flag(trt.BuilderFlag.INT8)

engine = builder.build_serialized_network(network, config)

TensorRT with Custom Plugins

If you have custom CUDA kernels, you can wrap them as TensorRT plugins:

#include "NvInferPlugin.h"

class RMSNormPlugin : public nvinfer1::IPluginV2DynamicExt {
public:
    // ... TensorRT plugin interface implementation
    // This allows your custom RMSNorm to be used within TensorRT graphs

    int enqueue(
        const nvinfer1::PluginTensorDesc* inputDesc,
        const nvinfer1::PluginTensorDesc* outputDesc,
        const void* const* inputs,
        void* const* outputs,
        void* workspace,
        cudaStream_t stream
    ) noexcept override {
        // Launch your custom kernel here
        const int hidden_size = inputDesc[0].dims.d[inputDesc[0].dims.nbDims - 1];
        const int num_rows = /* compute total rows */;

        rmsnorm_kernel<256><<<num_rows, 256, 0, stream>>>(
            static_cast<const half*>(inputs[0]),
            static_cast<const half*>(inputs[1]),
            static_cast<half*>(outputs[0]),
            hidden_size,
            1e-6f
        );
        return 0;
    }
};

16 GB Memory Strategies

The T4's 16 GB memory is a significant constraint for modern models. Here are strategies to work within this limit.

Model Size Estimates

Model	FP32	FP16	INT8	INT4
LLaMA-7B	28 GB	14 GB	7 GB	3.5 GB
LLaMA-13B	52 GB	26 GB	13 GB	6.5 GB
Mistral-7B	28 GB	14 GB	7 GB	3.5 GB
SD 1.5 UNet	3.4 GB	1.7 GB	0.85 GB	0.43 GB
SDXL UNet	10 GB	5 GB	2.5 GB	1.25 GB
SD3 Transformer	8 GB	4 GB	2 GB	1 GB

Memory Optimization Techniques

import torch

# 1. Use FP16 (required on T4 -- no BF16!)
model = model.half()  # Convert to FP16

# 2. Use quantization
from transformers import BitsAndBytesConfig

quantization_config = BitsAndBytesConfig(
    load_in_8bit=True,  # INT8 quantization
    # Or for 4-bit:
    # load_in_4bit=True,
    # bnb_4bit_compute_dtype=torch.float16,
)

model = AutoModelForCausalLM.from_pretrained(
    "meta-llama/Llama-2-7b-hf",
    quantization_config=quantization_config,
    device_map="auto",
)

# 3. Use CPU offloading for diffusers
pipe = StableDiffusionPipeline.from_pretrained(
    "runwayml/stable-diffusion-v1-5",
    torch_dtype=torch.float16,
)
pipe.enable_model_cpu_offload()  # Moves parts to CPU when not in use

# 4. Enable attention slicing to reduce peak memory
pipe.enable_attention_slicing(slice_size=1)

# 5. Use VAE tiling for high-resolution image generation
pipe.enable_vae_tiling()

# 6. Gradient checkpointing (for fine-tuning)
model.gradient_checkpointing_enable()

Memory-Efficient Kernel Design

// Design kernels that minimize temporary allocations

// BAD: Allocates a large temporary buffer
torch::Tensor bad_kernel(torch::Tensor input) {
    auto temp = torch::empty_like(input);  // Doubles memory usage!
    // ... compute into temp ...
    auto output = torch::empty_like(input);
    // ... compute into output from temp ...
    return output;
}

// GOOD: In-place computation or single output buffer
torch::Tensor good_kernel(torch::Tensor input) {
    auto output = torch::empty_like(input);
    // Compute directly into output, using registers and shared memory
    // for intermediate values instead of global memory temp buffers
    return output;
}

Occupancy Tuning for 40 SMs

The T4 has only 40 SMs with a maximum of 1024 threads per SM (32 warps). This is significantly less than A100 (108 SMs, 2048 threads/SM) or H100 (132 SMs, 2048 threads/SM).

Grid Sizing for T4

// T4: 40 SMs, max 1024 threads/SM, max 32 warps/SM
int num_sms = 40;

// Minimum: 1 block per SM
int min_grid = 40;

// Recommended: 2-4 blocks per SM
int grid_2x = 80;
int grid_4x = 160;

// For memory-bound kernels: more blocks help hide latency
int grid_8x = 320;

Block Size Recommendations for T4

Kernel Type	Block Size	Blocks/SM	Warps/SM	Occupancy
Element-wise	256	4	32	100%
Row reduction	256	4	32	100%
Row reduction	128	8	32	100%
Tiled matmul	128	4	16	50%
Complex kernel	64	8	16	50%

Register Pressure on T4

// T4: 64 KB registers per SM = 16,384 32-bit registers
// Max 255 registers per thread

// With 1024 threads/SM: 16 registers per thread (if fully occupied)
// With 512 threads/SM: 32 registers per thread (50% occupancy)
// With 256 threads/SM: 64 registers per thread (25% occupancy)

// Use __launch_bounds__ to help the compiler
__global__ __launch_bounds__(256, 4)  // 256 threads, aim for 4 blocks/SM
void t4_optimized_kernel(...) {
    // With 4 blocks of 256 threads = 1024 threads/SM
    // Budget: 16 registers per thread for 100% occupancy
    // Realistically: 32-48 registers with some occupancy trade-off
}

Cloud Instance Types

Provider	Instance	GPUs	vCPUs	RAM	Price (approx)
AWS	g4dn.xlarge	1x T4	4	16 GB	$0.526/hr
AWS	g4dn.2xlarge	1x T4	8	32 GB	$0.752/hr
AWS	g4dn.12xlarge	4x T4	48	192 GB	$3.912/hr
GCP	n1-standard-4 + T4	1x T4	4	15 GB	$0.35/hr
GCP	n1-standard-8 + T4	1x T4	8	30 GB	$0.45/hr
Azure	NC4as_T4_v3	1x T4	4	28 GB	$0.526/hr
Azure	NC8as_T4_v3	1x T4	8	56 GB	$0.752/hr

Choosing the Right Instance

• g4dn.xlarge: Best for single-model inference, batch size 1
• g4dn.2xlarge: Good for inference with larger batch sizes or CPU-heavy preprocessing
• g4dn.12xlarge: Multi-model serving or batch processing

build.toml Configuration for T4

[build]
cuda-version = "12.4"
cuda-capabilities = ["7.5"]  # T4 = sm_75

[kernel.rmsnorm]
src = ["src/rmsnorm_fp16.cu"]  # FP16 version!

[kernel.gelu]
src = ["src/gelu_fp16.cu"]

[kernel.rope]
src = ["src/rope_fp16.cu"]

Multi-Target Configuration

[build]
cuda-version = "12.4"
# Support T4, A100, and H100
cuda-capabilities = ["7.5", "8.0", "9.0"]

Conditional Compilation for T4

// Use preprocessor to handle architecture differences
#if __CUDA_ARCH__ >= 800
    // A100/H100: Use BF16
    #include <cuda_bf16.h>
    using compute_t = __nv_bfloat16;
    using compute_t2 = __nv_bfloat162;
    #define FLOAT_TO_COMPUTE __float2bfloat16
    #define COMPUTE_TO_FLOAT __bfloat162float
#else
    // T4: Use FP16
    #include <cuda_fp16.h>
    using compute_t = half;
    using compute_t2 = half2;
    #define FLOAT_TO_COMPUTE __float2half
    #define COMPUTE_TO_FLOAT __half2float
#endif

T4 RMSNorm Example (FP16)

#include <cuda_fp16.h>
#include <cuda_runtime.h>
#include <torch/extension.h>

template<int BLOCK_SIZE>
__global__ void rmsnorm_fp16_t4(
    const half* __restrict__ input,
    const half* __restrict__ weight,
    half* __restrict__ output,
    const int hidden_size,
    const float epsilon
) {
    const int row = blockIdx.x;
    const int tid = threadIdx.x;

    const half* x = input + row * hidden_size;
    half* out = output + row * hidden_size;

    // Compute sum of squares with FP32 accumulation
    float sum_sq = 0.0f;

    // Vectorized: process 8 FP16 values per iteration
    const int num_vecs = hidden_size / 8;
    for (int v = tid; v < num_vecs; v += BLOCK_SIZE) {
        float4 packed = reinterpret_cast<const float4*>(x)[v];
        half2* pairs = reinterpret_cast<half2*>(&packed);

        #pragma unroll
        for (int j = 0; j < 4; j++) {
            float lo = __half2float(__low2half(pairs[j]));
            float hi = __half2float(__high2half(pairs[j]));
            sum_sq += lo * lo + hi * hi;
        }
    }

    // Warp reduction
    #pragma unroll
    for (int offset = 16; offset > 0; offset >>= 1) {
        sum_sq += __shfl_xor_sync(0xffffffff, sum_sq, offset);
    }

    // Block reduction
    __shared__ float warp_sums[BLOCK_SIZE / 32];
    int warp_id = tid / 32;
    int lane_id = tid % 32;
    if (lane_id == 0) warp_sums[warp_id] = sum_sq;
    __syncthreads();

    __shared__ float rms_scale;
    if (warp_id == 0) {
        float val = (lane_id < BLOCK_SIZE / 32) ? warp_sums[lane_id] : 0.0f;
        for (int offset = 16; offset > 0; offset >>= 1)
            val += __shfl_xor_sync(0xffffffff, val, offset);
        if (lane_id == 0)
            rms_scale = rsqrtf(val / hidden_size + epsilon);
    }
    __syncthreads();

    float scale = rms_scale;

    // Vectorized normalization
    for (int v = tid; v < num_vecs; v += BLOCK_SIZE) {
        float4 in_packed = reinterpret_cast<const float4*>(x)[v];
        float4 w_packed = reinterpret_cast<const float4*>(weight)[v];

        half2* in_pairs = reinterpret_cast<half2*>(&in_packed);
        half2* w_pairs = reinterpret_cast<half2*>(&w_packed);

        float4 out_packed;
        half2* out_pairs = reinterpret_cast<half2*>(&out_packed);

        #pragma unroll
        for (int j = 0; j < 4; j++) {
            float x0 = __half2float(__low2half(in_pairs[j]));
            float x1 = __half2float(__high2half(in_pairs[j]));
            float w0 = __half2float(__low2half(w_pairs[j]));
            float w1 = __half2float(__high2half(w_pairs[j]));

            out_pairs[j] = __halves2half2(
                __float2half(x0 * scale * w0),
                __float2half(x1 * scale * w1)
            );
        }

        reinterpret_cast<float4*>(out)[v] = out_packed;
    }
}

torch::Tensor rmsnorm_forward(torch::Tensor input, torch::Tensor weight, double eps) {
    TORCH_CHECK(input.dtype() == torch::kFloat16, "T4 requires FP16 input");
    TORCH_CHECK(input.is_contiguous());

    int hidden_size = input.size(-1);
    auto input_2d = input.view({-1, hidden_size});
    int num_rows = input_2d.size(0);
    auto output = torch::empty_like(input_2d);

    constexpr int BLOCK_SIZE = 256;

    rmsnorm_fp16_t4<BLOCK_SIZE><<<num_rows, BLOCK_SIZE>>>(
        reinterpret_cast<const half*>(input_2d.data_ptr<at::Half>()),
        reinterpret_cast<const half*>(weight.data_ptr<at::Half>()),
        reinterpret_cast<half*>(output.data_ptr<at::Half>()),
        hidden_size,
        static_cast<float>(eps)
    );

    return output.view_as(input);
}

Summary

When targeting the T4:

1. Use FP16 only -- BF16 is not supported on Turing (sm_75)
2. Optimize for 320 GB/s bandwidth -- most kernels will be memory-bound
3. Use smaller tile sizes -- 64 KB shared memory limits tile dimensions
4. Leverage INT8 Tensor Cores for quantized inference
5. Consider TensorRT for production inference workloads
6. Plan for 16 GB memory -- use quantization, offloading, and attention slicing
7. Tune for 40 SMs -- grid sizes should be multiples of 40
8. Guard against FP16 overflow -- especially in attention score computation
9. Use conditional compilation (__CUDA_ARCH__) to support T4 alongside A100/H100