Environment note: Examples and profiling workflows are validated using Docker images with CUDA 13.0.1 (Ubuntu 24.04) and ROCm 7.0.1 (Ubuntu 24.04) for consistent toolchains. Enhanced build system includes profiling integrations.
- Introduction to GPU Performance Optimization
- GPU Performance Analysis Fundamentals
- Profiling Tools and Methodologies
- Memory Optimization Techniques
- Compute Optimization Strategies
- Performance Debugging Methodologies
- Cross-Platform Performance Considerations
- Advanced Optimization Techniques
- Performance Metrics and Validation
- Production Performance Engineering
GPU performance optimization is both an art and a science that requires deep understanding of GPU architecture, systematic analysis methodologies, and iterative refinement techniques. Unlike CPU optimization, GPU optimization must consider the massively parallel nature of GPU execution, complex memory hierarchies, and the interaction between thousands of concurrent threads.
Profile-Guided Optimization: Never optimize without data. Modern GPUs are complex systems with multiple potential bottlenecks, and intuition often leads to incorrect optimization priorities.
Systematic Approach: Performance optimization should follow a methodical process:
- Establish Baseline: Measure current performance comprehensively
- Identify Bottlenecks: Use profiling tools to find actual performance limiters
- Apply Targeted Optimizations: Focus on the most impactful bottlenecks first
- Validate Improvements: Measure the impact of each optimization
- Iterate: Repeat the process as new bottlenecks emerge
Architecture Awareness: Different GPU architectures have different optimization opportunities. Code optimized for one GPU may not perform optimally on another without architecture-specific tuning.
Understanding GPU performance requires modeling the interaction between:
- Compute Resources: SM count, cores per SM, clock frequencies
- Memory Hierarchy: Global, shared, constant, texture memory with different latencies
- Execution Model: Warp scheduling, occupancy, divergence
- Communication: Memory bandwidth, cache behavior, coalescing
The Roofline model provides a framework for understanding performance limits:
Achievable Performance = min(Peak Compute Performance, Memory Bandwidth × Arithmetic Intensity)
Where Arithmetic Intensity = Operations / Bytes Transferred
Key Insights:
- Memory Bound: Performance limited by memory bandwidth (low arithmetic intensity)
- Compute Bound: Performance limited by computational throughput (high arithmetic intensity)
- Balanced: Performance requires optimizing both memory and compute
1. Memory Bottlenecks
- Bandwidth Limited: Not achieving peak memory throughput
- Latency Limited: Irregular access patterns, cache misses
- Capacity Limited: Insufficient memory for optimal algorithms
2. Compute Bottlenecks
- Instruction Throughput: Limited by ALU utilization
- Occupancy Limited: Insufficient parallelism to hide latency
- Divergence Limited: Warp efficiency reduced by branching
3. System Bottlenecks
- Host-Device Transfer: PCIe bandwidth limitations
- Multi-GPU Communication: Inter-GPU bandwidth constraints
- Power/Thermal: Performance throttling due to power limits
Level 1 - Application Metrics
- End-to-end execution time
- Throughput (operations/second, GFLOPS, GB/s)
- Energy efficiency (performance/watt)
Level 2 - Kernel Metrics
- Kernel execution time
- Memory throughput achieved
- Compute utilization
Level 3 - Hardware Metrics
- SM utilization, occupancy
- Memory transaction efficiency
- Cache hit rates, memory coalescing
NVIDIA Nsight Compute (ncu)
- Detailed kernel-level performance analysis
- Hardware performance counters
- Memory and compute throughput analysis
- Optimization recommendations
# Comprehensive kernel analysis
ncu --metrics gpu__time_duration.avg,dram__throughput.avg.pct_of_peak_sustained_elapsed ./program
# Memory analysis focus
ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum,l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum ./program
# Occupancy analysis
ncu --metrics sm__maximum_warps_per_active_cycle_pct ./programNVIDIA Nsight Systems (nsys)
- Timeline-based system analysis
- CUDA API tracing
- Multi-GPU and multi-process analysis
- CPU-GPU interaction analysis
# System-wide timeline analysis
nsys profile --trace=cuda,nvtx,osrt --stats=true --output=profile ./program
# Multi-GPU analysis
nsys profile --trace=cuda,nvtx --stats=true --output=multi_gpu ./programROCProfiler (rocprof)
- HIP API and kernel tracing
- Hardware performance counters
- Memory and compute analysis
# Basic HIP tracing
rocprof --hip-trace --stats ./program
# Hardware counter analysis
rocprof --hsa-trace --stats ./program
# Custom metrics collection
rocprof --input=metrics.txt ./programROCTracer
- Low-level HSA and HIP tracing
- Timeline analysis
- API call overhead analysis
CUDA Events for Timing
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
// Kernel execution
kernelFunction<<<grid, block>>>();
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);HIP Events for Timing
hipEvent_t start, stop;
hipEventCreate(&start);
hipEventCreate(&stop);
hipEventRecord(start);
// Kernel execution
hipLaunchKernelGGL(kernelFunction, grid, block, 0, 0);
hipEventRecord(stop);
hipEventSynchronize(stop);
float milliseconds = 0;
hipEventElapsedTime(&milliseconds, start, stop);NVTX Annotations
#include <nvToolsExt.h>
nvtxRangePushA("Memory Transfer");
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
nvtxRangePop();
nvtxRangePushA("Kernel Execution");
kernelFunction<<<grid, block>>>();
nvtxRangePop();Memory Coalescing Global memory accesses are coalesced when threads in a warp access contiguous memory addresses.
// Coalesced access pattern (good)
__global__ void coalescedAccess(float *data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
data[idx] = data[idx] * 2.0f; // Sequential access
}
}
// Non-coalesced access pattern (poor performance)
__global__ void stridedAccess(float *data, int n, int stride) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx * stride < n) {
data[idx * stride] = data[idx * stride] * 2.0f; // Strided access
}
}Memory Transaction Analysis
- 32-byte transactions: L1 cache line size
- 128-byte transactions: L2 cache and global memory
- Coalescing efficiency: Ratio of useful to total bytes transferred
Bank Conflict Avoidance Shared memory is organized into banks. Simultaneous access to different addresses in the same bank causes conflicts.
// Bank conflict example (avoid)
__shared__ float shared_data[32][32];
// Threads access shared_data[threadIdx.x][0] - all access bank 0
// Bank conflict avoidance with padding
__shared__ float shared_data[32][33]; // +1 padding avoids conflictsShared Memory Tiling Pattern
#define TILE_SIZE 16
__global__ void matrixMulTiled(float *A, float *B, float *C, int width) {
__shared__ float tileA[TILE_SIZE][TILE_SIZE];
__shared__ float tileB[TILE_SIZE][TILE_SIZE];
int row = blockIdx.y * TILE_SIZE + threadIdx.y;
int col = blockIdx.x * TILE_SIZE + threadIdx.x;
float sum = 0.0f;
for (int phase = 0; phase < (width + TILE_SIZE - 1) / TILE_SIZE; phase++) {
// Load tiles into shared memory
if (row < width && phase * TILE_SIZE + threadIdx.x < width) {
tileA[threadIdx.y][threadIdx.x] = A[row * width + phase * TILE_SIZE + threadIdx.x];
} else {
tileA[threadIdx.y][threadIdx.x] = 0.0f;
}
if (col < width && phase * TILE_SIZE + threadIdx.y < width) {
tileB[threadIdx.y][threadIdx.x] = B[(phase * TILE_SIZE + threadIdx.y) * width + col];
} else {
tileB[threadIdx.y][threadIdx.x] = 0.0f;
}
__syncthreads();
// Compute using shared memory
for (int k = 0; k < TILE_SIZE; k++) {
sum += tileA[threadIdx.y][k] * tileB[k][threadIdx.x];
}
__syncthreads();
}
if (row < width && col < width) {
C[row * width + col] = sum;
}
}L1 Cache Optimization
- Temporal locality: Reuse data within short time windows
- Spatial locality: Access nearby memory addresses
- Cache line utilization: Maximize useful data per cache line
L2 Cache Optimization
- Cross-SM data sharing through L2 cache
- Persistent data caching for frequently accessed data
- Avoid cache thrashing with working set management
Constant Memory Usage
__constant__ float const_data[1024];
// Efficient broadcast read pattern
__global__ void useConstantMemory() {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
// All threads read same constant data - broadcasted efficiently
float value = const_data[0];
}Structure of Arrays (SoA) vs Array of Structures (AoS)
// Array of Structures (AoS) - poor coalescing
struct Point { float x, y, z; };
__global__ void processAoS(Point *points, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
points[idx].x += 1.0f; // Non-coalesced: only using 4 bytes of 12-byte struct
}
}
// Structure of Arrays (SoA) - good coalescing
__global__ void processSoA(float *x, float *y, float *z, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
x[idx] += 1.0f; // Coalesced: all threads access contiguous x values
}
}Understanding Occupancy Occupancy = (Active Warps) / (Maximum Possible Warps)
Factors limiting occupancy:
- Registers per thread: Higher register usage reduces occupancy
- Shared memory per block: More shared memory reduces concurrent blocks
- Thread block size: Must be compatible with SM warp schedulers
// Launch bounds to guide occupancy
__global__ void __launch_bounds__(256, 4) // 256 threads/block, min 4 blocks/SM
optimizedKernel() {
// Kernel optimized for specific occupancy target
}Occupancy Analysis Tools
# CUDA Occupancy Calculator
cuda-occupancy-calculator
# Runtime occupancy query
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks, kernel, blockSize, dynamicSMemSize);Divergence Minimization
// Poor: High divergence
__global__ void divergentKernel(int *data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
if (data[idx] % 2 == 0) {
// Only half the threads execute
data[idx] = data[idx] * 2;
} else {
// Other half executes this
data[idx] = data[idx] + 1;
}
}
}
// Better: Reduced divergence using predicated execution
__global__ void optimizedKernel(int *data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
int value = data[idx];
int is_even = (value % 2 == 0);
// All threads execute both operations, results selected by predicate
int even_result = value * 2;
int odd_result = value + 1;
data[idx] = is_even ? even_result : odd_result;
}
}Warp Voting Functions
#include <cooperative_groups.h>
namespace cg = cooperative_groups;
__global__ void warpVotingExample(int *data, int *result, int n) {
auto warp = cg::tiled_partition<32>(cg::this_thread_block());
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
int value = data[idx];
// Check if all threads in warp have positive values
bool all_positive = warp.all(value > 0);
if (all_positive) {
// Warp-optimized operation when all values positive
result[idx] = value * 2;
} else {
// Handle mixed case
result[idx] = (value > 0) ? value * 2 : 0;
}
}
}Mathematical Function Optimization
// Use intrinsic functions for better performance
__device__ float optimizedMath(float x) {
// Fast math intrinsics (less precise but faster)
float result = __sinf(x) * __cosf(x);
result += __expf(x * 0.5f);
return __fdividef(result, x); // Fast divide
}
// Vector loads for memory bandwidth
__global__ void vectorizedLoads(float4 *input, float4 *output, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
float4 data = input[idx]; // 16-byte vectorized load
data.x *= 2.0f;
data.y *= 2.0f;
data.z *= 2.0f;
data.w *= 2.0f;
output[idx] = data; // 16-byte vectorized store
}
}Loop Unrolling
// Manual loop unrolling
__global__ void unrolledLoop(float *data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
idx *= 4; // Process 4 elements per thread
if (idx + 3 < n) {
// Process 4 elements at once
data[idx] = data[idx] * 2.0f;
data[idx + 1] = data[idx + 1] * 2.0f;
data[idx + 2] = data[idx + 2] * 2.0f;
data[idx + 3] = data[idx + 3] * 2.0f;
}
}
// Compiler directive for loop unrolling
__global__ void compilerUnrolled(float *data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
#pragma unroll
for (int i = 0; i < 4; i++) {
if (idx * 4 + i < n) {
data[idx * 4 + i] *= 2.0f;
}
}
}1. Performance Baseline Establishment
// Timing infrastructure
struct PerformanceTimer {
cudaEvent_t start, stop;
PerformanceTimer() {
cudaEventCreate(&start);
cudaEventCreate(&stop);
}
void startTimer() { cudaEventRecord(start); }
float endTimer() {
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float ms;
cudaEventElapsedTime(&ms, start, stop);
return ms;
}
~PerformanceTimer() {
cudaEventDestroy(start);
cudaEventDestroy(stop);
}
};2. Bottleneck Identification Framework
void analyzeKernelPerformance(const char* kernelName) {
// Memory throughput analysis
float achievedBandwidth = (dataSize * 1000.0f) / executionTime / 1e9;
float theoreticalBandwidth = getDeviceMemoryBandwidth();
float memoryEfficiency = achievedBandwidth / theoreticalBandwidth;
// Compute analysis
float achievedGFLOPS = (operationCount * 1000.0f) / executionTime / 1e9;
float theoreticalGFLOPS = getDeviceComputeThroughput();
float computeEfficiency = achievedGFLOPS / theoreticalGFLOPS;
printf("%s Performance Analysis:\n", kernelName);
printf("Memory Efficiency: %.1f%% (%.1f/%.1f GB/s)\n",
memoryEfficiency * 100, achievedBandwidth, theoreticalBandwidth);
printf("Compute Efficiency: %.1f%% (%.1f/%.1f GFLOPS)\n",
computeEfficiency * 100, achievedGFLOPS, theoreticalGFLOPS);
if (memoryEfficiency < 0.8f) {
printf("⚠️ Memory bound - focus on memory optimization\n");
}
if (computeEfficiency < 0.8f) {
printf("⚠️ Compute bound - focus on instruction optimization\n");
}
}3. Performance Regression Testing
class PerformanceRegressionTester {
private:
std::map<std::string, float> baselineMetrics;
float tolerancePercent = 5.0f;
public:
void recordBaseline(const std::string& testName, float metric) {
baselineMetrics[testName] = metric;
}
bool validatePerformance(const std::string& testName, float currentMetric) {
auto it = baselineMetrics.find(testName);
if (it == baselineMetrics.end()) {
printf("No baseline found for %s\n", testName.c_str());
return false;
}
float baseline = it->second;
float changePercent = ((currentMetric - baseline) / baseline) * 100.0f;
printf("%s: %.2f ms (baseline: %.2f ms, change: %.1f%%)\n",
testName.c_str(), currentMetric, baseline, changePercent);
if (changePercent > tolerancePercent) {
printf("❌ Performance regression detected!\n");
return false;
}
if (changePercent < -tolerancePercent) {
printf("✅ Performance improvement detected!\n");
}
return true;
}
};Memory Metrics Interpretation
- Global Memory Throughput: Should approach theoretical peak for memory-bound kernels
- Shared Memory Bank Conflicts: Should be minimized (< 2% conflict rate)
- Cache Hit Rates: L1 cache >80%, L2 cache >60% for cache-friendly access patterns
- Memory Coalescing Efficiency: Should be >90% for optimal global memory access
Compute Metrics Interpretation
- SM Utilization: Should be >70% for compute-intensive kernels
- Achieved Occupancy: Balance with resource usage (registers, shared memory)
- Warp Execution Efficiency: Should be >90% to minimize divergence impact
- Instruction Throughput: Compare against theoretical peak for different instruction types
NVIDIA Architecture Characteristics
- Warp Size: 32 threads
- Shared Memory: Configurable L1/shared memory split
- Compute Units: Streaming Multiprocessors (SMs)
- Memory Hierarchy: More levels, larger caches
- Specialized Units: Tensor Cores, RT Cores (newer architectures)
AMD Architecture Characteristics
- Wavefront Size: 64 threads (some newer: 32)
- Local Data Share (LDS): Similar to shared memory
- Compute Units: Compute Units (CUs)
- Memory Hierarchy: Different cache sizes and behavior
- Specialized Units: Matrix cores, AI accelerators
Generic Memory Coalescing
// Platform-agnostic coalesced access
template<typename T>
__global__ void coalescedProcess(T* data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
// Sequential access pattern works on both platforms
data[idx] = data[idx] * static_cast<T>(2);
}
}Adaptive Block Size Selection
int getOptimalBlockSize(int deviceId) {
hipDeviceProp_t prop;
hipGetDeviceProperties(&prop, deviceId);
if (prop.warpSize == 32) {
// NVIDIA-optimized block sizes
return 256; // Multiple of 32
} else {
// AMD-optimized block sizes
return 256; // Also works well for wavefront size 64
}
}NVIDIA-Specific Optimizations
// Tensor Core utilization (Volta+)
#if __CUDA_ARCH__ >= 700
#include <mma.h>
using namespace nvcuda::wmma;
__global__ void tensorCoreMatMul(half *a, half *b, float *c, int m, int n, int k) {
// Tensor Core fragment declarations
fragment<matrix_a, 16, 16, 16, half, row_major> a_frag;
fragment<matrix_b, 16, 16, 16, half, col_major> b_frag;
fragment<accumulator, 16, 16, 16, float> acc_frag;
// Use Tensor Cores for mixed-precision computation
load_matrix_sync(a_frag, a, 16);
load_matrix_sync(b_frag, b, 16);
fill_fragment(acc_frag, 0.0f);
mma_sync(acc_frag, a_frag, b_frag, acc_frag);
store_matrix_sync(c, acc_frag, 16, mem_row_major);
}
#endifAMD-Specific Optimizations
// Wavefront-specific operations
__global__ void amdOptimized(float* data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int laneId = threadIdx.x % 64; // AMD wavefront size
if (idx < n) {
// AMD-optimized warp shuffle operations
float value = data[idx];
// Wavefront reduction optimized for size 64
for (int offset = 32; offset > 0; offset /= 2) {
value += __shfl_down(value, offset, 64);
}
if (laneId == 0) {
data[blockIdx.x * blockDim.x / 64] = value;
}
}
}Kernel Fusion Benefits
- Reduced memory bandwidth requirements
- Eliminated intermediate memory allocations
- Better data locality and cache utilization
- Reduced kernel launch overhead
// Separate kernels (poor memory efficiency)
__global__ void step1(float *a, float *temp, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) temp[idx] = a[idx] * 2.0f;
}
__global__ void step2(float *temp, float *b, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) b[idx] = temp[idx] + 1.0f;
}
// Fused kernel (better memory efficiency)
__global__ void fusedKernel(float *a, float *b, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
float temp = a[idx] * 2.0f; // No memory round-trip
b[idx] = temp + 1.0f;
}
}Tiling for Cache Optimization
#define TILE_SIZE 32
__global__ void tiledTranspose(float *input, float *output, int width, int height) {
__shared__ float tile[TILE_SIZE][TILE_SIZE + 1]; // +1 to avoid bank conflicts
int x = blockIdx.x * TILE_SIZE + threadIdx.x;
int y = blockIdx.y * TILE_SIZE + threadIdx.y;
// Coalesced load from input
if (x < width && y < height) {
tile[threadIdx.y][threadIdx.x] = input[y * width + x];
}
__syncthreads();
// Coalesced store to output (transposed)
x = blockIdx.y * TILE_SIZE + threadIdx.x;
y = blockIdx.x * TILE_SIZE + threadIdx.y;
if (x < height && y < width) {
output[y * height + x] = tile[threadIdx.x][threadIdx.y];
}
}Cooperative Groups for Fine-Grained Control
#include <cooperative_groups.h>
namespace cg = cooperative_groups;
__global__ void cooperativeKernel(float *data, int n) {
// Different levels of cooperation
auto grid = cg::this_grid();
auto block = cg::this_thread_block();
auto warp = cg::tiled_partition<32>(block);
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
float value = data[idx];
// Warp-level reduction
for (int offset = 16; offset > 0; offset /= 2) {
value += warp.shfl_down(value, offset);
}
// Block-level cooperation
__shared__ float block_sums[32]; // Assuming 1024 threads max
if (warp.thread_rank() == 0) {
block_sums[warp.meta_group_rank()] = value;
}
block.sync();
// Grid-level cooperation (requires cooperative launch)
if (threadIdx.x == 0) {
// Perform block-level operations
float block_sum = 0.0f;
for (int i = 0; i < block.size() / 32; i++) {
block_sum += block_sums[i];
}
// Atomically contribute to global result
atomicAdd(&data[0], block_sum);
}
}
}Performance-Per-Watt Optimization
// Monitor and adapt to thermal conditions
__global__ void thermalAwareKernel(float *data, int n, int iterations) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
float value = data[idx];
// Adaptive work based on thermal state
int actual_iterations = iterations;
#ifdef THERMAL_THROTTLING
// Reduce work under thermal pressure
if (clock() % 1000 == 0) { // Periodically check
actual_iterations = iterations / 2;
}
#endif
for (int i = 0; i < actual_iterations; i++) {
value = sinf(value + i * 0.1f);
}
data[idx] = value;
}
}System-Level Metrics
- End-to-End Latency: Total time from input to output
- Throughput: Operations completed per unit time
- Scalability: Performance improvement with increased resources
- Energy Efficiency: Performance per watt consumed
Application-Level Metrics
- Algorithm Efficiency: Comparison to theoretical optimum
- Memory Efficiency: Achieved vs theoretical memory bandwidth
- Compute Efficiency: Achieved vs theoretical compute throughput
- Load Balance: Work distribution across processing elements
Hardware-Level Metrics
- Occupancy: Active warps vs maximum possible warps
- Utilization: Percentage of time functional units are active
- Cache Performance: Hit rates across different cache levels
- Memory Coalescing: Efficiency of memory access patterns
Performance Variability Management
class PerformanceStatistics {
private:
std::vector<float> measurements;
public:
void addMeasurement(float time) { measurements.push_back(time); }
float getMean() {
return std::accumulate(measurements.begin(), measurements.end(), 0.0f) / measurements.size();
}
float getStandardDeviation() {
float mean = getMean();
float sq_sum = 0.0f;
for (float time : measurements) {
sq_sum += (time - mean) * (time - mean);
}
return sqrt(sq_sum / measurements.size());
}
float getConfidenceInterval(float confidence = 0.95f) {
// Simple approach using standard error
float stddev = getStandardDeviation();
float stderr = stddev / sqrt(measurements.size());
return 1.96 * stderr; // 95% confidence
}
bool isSignificantImprovement(float baseline, float threshold = 0.05f) {
float mean = getMean();
float improvement = (baseline - mean) / baseline;
float ci = getConfidenceInterval();
return (improvement - ci) > threshold;
}
};Continuous Performance Monitoring
class PerformanceMonitor {
private:
std::map<std::string, PerformanceStatistics> testSuite;
public:
void runPerformanceTest(const std::string& testName,
std::function<float()> testFunction,
int iterations = 10) {
PerformanceStatistics stats;
// Warm-up runs
for (int i = 0; i < 3; i++) {
testFunction();
}
// Measurement runs
for (int i = 0; i < iterations; i++) {
float time = testFunction();
stats.addMeasurement(time);
}
testSuite[testName] = stats;
// Report results
printf("%s: %.3f ± %.3f ms (n=%d)\n",
testName.c_str(), stats.getMean(),
stats.getStandardDeviation(), iterations);
}
void generateReport() {
printf("\n=== Performance Report ===\n");
for (const auto& test : testSuite) {
const auto& stats = test.second;
printf("%-30s: %8.3f ± %6.3f ms\n",
test.first.c_str(), stats.getMean(), stats.getStandardDeviation());
}
}
};1. Continuous Performance Integration
- Automated performance testing in CI/CD pipelines
- Performance regression detection and alerts
- Historical performance tracking and trend analysis
2. Production Performance Monitoring
- Real-time performance metrics collection
- Performance anomaly detection
- Adaptive performance optimization based on workload characteristics
3. Performance-Oriented Development Process
- Performance requirements specification
- Performance impact assessment for code changes
- Performance review as part of code review process
Multi-GPU Scaling Strategies
// Adaptive multi-GPU work distribution
class MultiGPUOptimizer {
private:
std::vector<float> gpuPerformance;
std::vector<float> gpuMemoryBandwidth;
public:
std::vector<int> calculateOptimalDistribution(int totalWork, int numGPUs) {
std::vector<int> distribution(numGPUs);
// Calculate total capability
float totalCapability = 0.0f;
for (int i = 0; i < numGPUs; i++) {
totalCapability += gpuPerformance[i];
}
// Distribute work proportionally
int distributedWork = 0;
for (int i = 0; i < numGPUs - 1; i++) {
distribution[i] = (int)(totalWork * gpuPerformance[i] / totalCapability);
distributedWork += distribution[i];
}
distribution[numGPUs - 1] = totalWork - distributedWork;
return distribution;
}
void updatePerformanceMetrics(int gpuId, float performance, float bandwidth) {
if (gpuId < gpuPerformance.size()) {
// Exponential moving average for adaptive learning
float alpha = 0.1f;
gpuPerformance[gpuId] = alpha * performance + (1 - alpha) * gpuPerformance[gpuId];
gpuMemoryBandwidth[gpuId] = alpha * bandwidth + (1 - alpha) * gpuMemoryBandwidth[gpuId];
}
}
};Cost-Benefit Analysis Framework
struct OptimizationROI {
std::string optimizationName;
float developmentTime; // hours
float performanceImprovement; // percentage
float maintenanceOverhead; // ongoing cost factor
float calculateROI(float baselinePerformance, float valuePerImprovement) {
float absoluteImprovement = baselinePerformance * (performanceImprovement / 100.0f);
float totalValue = absoluteImprovement * valuePerImprovement;
float totalCost = developmentTime + maintenanceOverhead;
return (totalValue - totalCost) / totalCost;
}
};Architecture-Adaptive Optimization
// Template-based optimization for different architectures
template<int ARCH_VERSION>
struct ArchOptimizedKernel {
static __global__ void execute(float* data, int n);
};
// Specialization for different architectures
template<>
struct ArchOptimizedKernel<70> { // Volta
static __global__ void execute(float* data, int n) {
// Volta-optimized implementation with Tensor Cores
}
};
template<>
struct ArchOptimizedKernel<80> { // Ampere
static __global__ void execute(float* data, int n) {
// Ampere-optimized implementation with new features
}
};
// Runtime architecture detection and optimization selection
void launchOptimizedKernel(float* data, int n) {
int major, minor;
cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, 0);
cudaDeviceGetAttribute(&minor, cudaDevAttrComputeCapabilityMinor, 0);
int archVersion = major * 10 + minor;
dim3 grid, block;
calculateOptimalLaunchConfig(grid, block, n, archVersion);
switch(archVersion) {
case 70: case 75:
ArchOptimizedKernel<70>::execute<<<grid, block>>>(data, n);
break;
case 80: case 86:
ArchOptimizedKernel<80>::execute<<<grid, block>>>(data, n);
break;
default:
// Fallback generic implementation
genericKernel<<<grid, block>>>(data, n);
}
}GPU performance optimization is a multifaceted discipline requiring deep understanding of hardware architecture, systematic analysis methodologies, and iterative refinement processes. Success in GPU optimization comes from:
- Systematic Approach: Always profile before optimizing, focus on the biggest bottlenecks first
- Architecture Awareness: Understand the target GPU architecture and optimize accordingly
- Holistic View: Consider the entire system, not just individual kernels
- Continuous Validation: Measure and validate all optimization efforts
- Future Adaptability: Design optimization strategies that can evolve with new architectures
The techniques and methodologies covered in this module provide the foundation for achieving optimal GPU performance across a wide range of applications and hardware platforms. Remember that performance optimization is an ongoing process that requires continuous learning and adaptation as GPU architectures continue to evolve.