SASS Metrics Collection Tutorial

cd sass_metrics
make

./sass_metrics

class SASSMetricsCollector {
private:
    struct SASSMetrics {
        uint64_t totalInstructions;
        uint64_t fp32Instructions;
        uint64_t integerInstructions;
        uint64_t memoryInstructions;
        uint64_t controlInstructions;
        double warpEfficiency;
        double branchEfficiency;
        double threadEfficiency;
    };
    
    std::map<std::string, SASSMetrics> kernelMetrics;
    std::vector<CUpti_EventID> sassEventIds;
    CUpti_EventGroup eventGroup;
 
public:
    void setupSASSMetrics(CUcontext context, CUdevice device);
    void startCollection();
    void stopCollection();
    void analyzeKernelSASS(const std::string& kernelName);
    void generateSASSReport();
};
 
void SASSMetricsCollector::setupSASSMetrics(CUcontext context, CUdevice device) {
    // Set up SASS-level event collection
    CUPTI_CALL(cuptiEventGroupCreate(context, &eventGroup, 0));
    
    // Add SASS instruction counting events
    std::vector<std::string> sassEvents = {
        "inst_executed",
        "inst_fp_32", 
        "inst_integer",
        "inst_compute_ld_st",
        "inst_control",
        "thread_inst_executed",
        "warp_execution_efficiency",
        "branch_efficiency"
    };
    
    for (const auto& eventName : sassEvents) {
        CUpti_EventID eventId;
        CUptiResult result = cuptiEventGetIdFromName(device, eventName.c_str(), &eventId);
        if (result == CUPTI_SUCCESS) {
            CUPTI_CALL(cuptiEventGroupAddEvent(eventGroup, eventId));
            sassEventIds.push_back(eventId);
        }
    }
}

class AssemblyAnalyzer {
private:
    struct InstructionProfile {
        std::string opcode;
        uint64_t executionCount;
        double averageLatency;
        double throughput;
        std::string category;
    };
    
    std::map<std::string, InstructionProfile> instructionProfiles;
 
public:
    void analyzeInstructionMix(const SASSMetrics& metrics) {
        std::cout << "Instruction Mix Analysis:" << std::endl;
        
        uint64_t totalInsts = metrics.totalInstructions;
        
        double fpPercentage = (double)metrics.fp32Instructions / totalInsts * 100;
        double intPercentage = (double)metrics.integerInstructions / totalInsts * 100;
        double memPercentage = (double)metrics.memoryInstructions / totalInsts * 100;
        double ctrlPercentage = (double)metrics.controlInstructions / totalInsts * 100;
        
        std::cout << "  Floating Point: " << fpPercentage << "%" << std::endl;
        std::cout << "  Integer: " << intPercentage << "%" << std::endl;
        std::cout << "  Memory: " << memPercentage << "%" << std::endl;
        std::cout << "  Control: " << ctrlPercentage << "%" << std::endl;
        
        // Analyze instruction balance
        if (memPercentage > 40) {
            std::cout << "  WARNING: Memory-bound kernel detected" << std::endl;
        }
        if (ctrlPercentage > 20) {
            std::cout << "  WARNING: High control overhead detected" << std::endl;
        }
    }
    
    void suggestOptimizations(const SASSMetrics& metrics) {
        std::cout << "Optimization Suggestions:" << std::endl;
        
        if (metrics.warpEfficiency < 0.8) {
            std::cout << "  - Consider optimizing warp utilization" << std::endl;
            std::cout << "  - Check for thread divergence patterns" << std::endl;
        }
        
        if (metrics.branchEfficiency < 0.9) {
            std::cout << "  - Optimize branch patterns to reduce divergence" << std::endl;
            std::cout << "  - Consider predicated execution for simple branches" << std::endl;
        }
        
        double memRatio = (double)metrics.memoryInstructions / metrics.totalInstructions;
        if (memRatio > 0.3) {
            std::cout << "  - Consider memory access optimization" << std::endl;
            std::cout << "  - Evaluate shared memory usage opportunities" << std::endl;
        }
    }
};

class PipelineAnalyzer {
private:
    struct PipelineMetrics {
        double aluUtilization;
        double fpuUtilization;
        double memoryUtilization;
        double controlUtilization;
        std::vector<double> stallReasons;
    };
 
public:
    PipelineMetrics analyzePipelineUtilization(const SASSMetrics& metrics) {
        PipelineMetrics pipelineMetrics;
        
        // Calculate pipeline utilization based on instruction mix
        uint64_t totalInsts = metrics.totalInstructions;
        
        pipelineMetrics.aluUtilization = 
            (double)(metrics.integerInstructions) / totalInsts;
        pipelineMetrics.fpuUtilization = 
            (double)(metrics.fp32Instructions) / totalInsts;
        pipelineMetrics.memoryUtilization = 
            (double)(metrics.memoryInstructions) / totalInsts;
        pipelineMetrics.controlUtilization = 
            (double)(metrics.controlInstructions) / totalInsts;
        
        return pipelineMetrics;
    }
    
    void identifyBottlenecks(const PipelineMetrics& pipeline) {
        std::cout << "Pipeline Bottleneck Analysis:" << std::endl;
        
        double maxUtilization = std::max({
            pipeline.aluUtilization,
            pipeline.fpuUtilization, 
            pipeline.memoryUtilization,
            pipeline.controlUtilization
        });
        
        if (maxUtilization == pipeline.memoryUtilization) {
            std::cout << "  Primary bottleneck: Memory pipeline" << std::endl;
            std::cout << "  Suggestions: Optimize memory access patterns, use shared memory" << std::endl;
        } else if (maxUtilization == pipeline.fpuUtilization) {
            std::cout << "  Primary bottleneck: FP pipeline" << std::endl;
            std::cout << "  Suggestions: Balance compute with memory operations" << std::endl;
        } else if (maxUtilization == pipeline.controlUtilization) {
            std::cout << "  Primary bottleneck: Control pipeline" << std::endl;
            std::cout << "  Suggestions: Reduce branch divergence, simplify control flow" << std::endl;
        }
    }
};

class PerformanceModel {
private:
    struct ModelParameters {
        double peakALUThroughput;
        double peakFPUThroughput;
        double memoryBandwidth;
        double instructionIssueRate;
    };
    
    ModelParameters getGPUParameters(int deviceId) {
        ModelParameters params;
        
        cudaDeviceProp prop;
        cudaGetDeviceProperties(&prop, deviceId);
        
        // Estimate peak throughput based on GPU architecture
        params.peakALUThroughput = prop.multiProcessorCount * prop.clockRate * 2; // Simplified
        params.peakFPUThroughput = prop.multiProcessorCount * prop.clockRate;
        params.memoryBandwidth = prop.memoryBusWidth * prop.memoryClockRate * 2 / 8; // GB/s
        params.instructionIssueRate = prop.multiProcessorCount * prop.clockRate;
        
        return params;
    }
 
public:
    void validatePerformanceModel(const SASSMetrics& metrics, int deviceId) {
        ModelParameters params = getGPUParameters(deviceId);
        
        // Calculate achieved vs theoretical performance
        double achievedALUUtilization = 
            (double)metrics.integerInstructions / metrics.totalInstructions;
        double achievedFPUUtilization = 
            (double)metrics.fp32Instructions / metrics.totalInstructions;
        
        std::cout << "Performance Model Validation:" << std::endl;
        std::cout << "  ALU Utilization: " << achievedALUUtilization * 100 << "%" << std::endl;
        std::cout << "  FPU Utilization: " << achievedFPUUtilization * 100 << "%" << std::endl;
        
        // Compare with theoretical limits
        double efficiencyGap = 1.0 - std::max(achievedALUUtilization, achievedFPUUtilization);
        if (efficiencyGap > 0.3) {
            std::cout << "  Significant efficiency gap detected: " << efficiencyGap * 100 << "%" << std::endl;
        }
    }
};

void analyzeMatrixMultiplication() {
    SASSMetricsCollector collector;
    AssemblyAnalyzer analyzer;
    
    // Set up metrics collection
    collector.setupSASSMetrics(context, device);
    
    // Launch matrix multiplication kernel
    collector.startCollection();
    
    dim3 grid(32, 32);
    dim3 block(16, 16);
    matrixMulKernel<<<grid, block>>>(d_A, d_B, d_C, N);
    
    cudaDeviceSynchronize();
    collector.stopCollection();
    
    // Analyze results
    SASSMetrics metrics = collector.getMetrics("matrixMulKernel");
    analyzer.analyzeInstructionMix(metrics);
    analyzer.suggestOptimizations(metrics);
    
    std::cout << "Matrix Multiplication SASS Analysis:" << std::endl;
    std::cout << "  Compute Intensity: " << calculateComputeIntensity(metrics) << std::endl;
    std::cout << "  Memory Efficiency: " << calculateMemoryEfficiency(metrics) << std::endl;
}

void analyzeReductionKernel() {
    SASSMetricsCollector collector;
    
    collector.setupSASSMetrics(context, device);
    collector.startCollection();
    
    // Launch reduction kernel
    reductionKernel<<<grid, block, sharedMemSize>>>(d_input, d_output, N);
    cudaDeviceSynchronize();
    
    collector.stopCollection();
    
    SASSMetrics metrics = collector.getMetrics("reductionKernel");
    
    std::cout << "Reduction Kernel Analysis:" << std::endl;
    std::cout << "  Shared Memory Usage Efficiency: " << 
        analyzeSharedMemoryUsage(metrics) << std::endl;
    std::cout << "  Synchronization Overhead: " << 
        analyzeSynchronizationOverhead(metrics) << std::endl;
    std::cout << "  Warp Divergence: " << 
        (100.0 - metrics.branchEfficiency) << "%" << std::endl;
}

class CompilerAnalyzer {
public:
    void compareOptimizationLevels() {
        std::vector<std::string> optimizationFlags = {"-O0", "-O1", "-O2", "-O3"};
        
        for (const auto& flag : optimizationFlags) {
            std::cout << "Testing with " << flag << ":" << std::endl;
            
            // Recompile with different optimization level
            recompileKernel(flag);
            
            // Collect SASS metrics
            SASSMetrics metrics = collectMetricsForOptimization(flag);
            
            std::cout << "  Instructions: " << metrics.totalInstructions << std::endl;
            std::cout << "  Efficiency: " << metrics.warpEfficiency << std::endl;
            
            analyzeOptimizationImpact(metrics, flag);
        }
    }
    
private:
    void analyzeOptimizationImpact(const SASSMetrics& metrics, const std::string& flag) {
        // Analyze how compiler optimizations affect SASS metrics
        if (flag == "-O3" && metrics.controlInstructions > baseline.controlInstructions) {
            std::cout << "    Warning: High optimization may have increased control overhead" << std::endl;
        }
        
        if (metrics.warpEfficiency < 0.8) {
            std::cout << "    Consider different optimization strategies" << std::endl;
        }
    }
};

class CodeCorrelator {
public:
    void correlateCUDAtoSASS(const std::string& cudaFile, const std::string& sassFile) {
        // Parse CUDA source code
        auto cudaLines = parseCUDASource(cudaFile);
        
        // Parse SASS assembly
        auto sassInstructions = parseSASSAssembly(sassFile);
        
        // Correlate performance metrics with source lines
        for (const auto& [lineNum, sassInsts] : sassInstructions) {
            if (lineNum < cudaLines.size()) {
                std::cout << "CUDA Line " << lineNum << ": " << cudaLines[lineNum] << std::endl;
                std::cout << "  SASS Instructions: " << sassInsts.size() << std::endl;
                
                for (const auto& inst : sassInsts) {
                    std::cout << "    " << inst.opcode << " (executed " 
                             << inst.executionCount << " times)" << std::endl;
                }
                
                // Identify performance hotspots
                if (getTotalExecutionCount(sassInsts) > hotspotThreshold) {
                    std::cout << "  >>> PERFORMANCE HOTSPOT DETECTED <<<" << std::endl;
                }
            }
        }
    }
};

class SASSDebugger {
public:
    void validateMetricSupport(CUdevice device) {
        uint32_t numEvents;
        CUPTI_CALL(cuptiDeviceGetNumEventDomains(device, &numEvents));
        
        std::cout << "Available event domains: " << numEvents << std::endl;
        
        // Check specific SASS events
        std::vector<std::string> requiredEvents = {
            "inst_executed", "inst_fp_32", "warp_execution_efficiency"
        };
        
        for (const auto& eventName : requiredEvents) {
            CUpti_EventID eventId;
            CUptiResult result = cuptiEventGetIdFromName(device, eventName.c_str(), &eventId);
            
            if (result == CUPTI_SUCCESS) {
                std::cout << "✓ " << eventName << " supported" << std::endl;
            } else {
                std::cout << "✗ " << eventName << " not supported" << std::endl;
            }
        }
    }
    
    void checkEventGroupCompatibility(const std::vector<CUpti_EventID>& events, CUcontext context) {
        // Test if events can be collected together
        CUpti_EventGroup testGroup;
        CUptiResult result = cuptiEventGroupCreate(context, &testGroup, 0);
        
        for (auto eventId : events) {
            result = cuptiEventGroupAddEvent(testGroup, eventId);
            if (result != CUPTI_SUCCESS) {
                std::cout << "Event " << eventId << " cannot be added to group" << std::endl;
            }
        }
        
        cuptiEventGroupDestroy(testGroup);
    }
};