diff --git a/.github/federation.json b/.github/federation.json index 4b8cfd9..60f3bf1 100644 --- a/.github/federation.json +++ b/.github/federation.json @@ -100,6 +100,9 @@ "path": "projects/rocprofiler-sdk/skills/pc-sampling", "as": "rocprofiler-pc-sampling" }, + { + "path": "projects/rocprofiler-compute/skills/rocprof-compute-speed-of-light" + }, { "path": "projects/rocprofiler-systems/skills/choosing-a-profiling-preset", "as": "rocprofiler-systems-choosing-a-profiling-preset" diff --git a/skills/rocprof-compute-speed-of-light/.federated.json b/skills/rocprof-compute-speed-of-light/.federated.json new file mode 100644 index 0000000..49c5be7 --- /dev/null +++ b/skills/rocprof-compute-speed-of-light/.federated.json @@ -0,0 +1,10 @@ +{ + "source": "rocm-rocm-systems", + "repo": "ROCm/rocm-systems", + "ref": "users/vedithal-amd/rocprof-compute-skills-install-check", + "commit": "ebf465e112ae1ca99a2e6e45de06c0ad84388439", + "path": "projects/rocprofiler-compute/skills/rocprof-compute-speed-of-light", + "license": "MIT", + "content_hash": "846aeabf5b7cb88a07c742bd4e80646957a10ecb58485888d2e2513b6781638f", + "imported_at": "2026-11-09T18:28:59Z" +} diff --git a/skills/rocprof-compute-speed-of-light/SKILL.md b/skills/rocprof-compute-speed-of-light/SKILL.md new file mode 100644 index 0000000..c6e582d --- /dev/null +++ b/skills/rocprof-compute-speed-of-light/SKILL.md @@ -0,0 +1,131 @@ +--- +name: rocprof-compute-speed-of-light +description: Reads the rocprof-compute System Speed-of-Light panel to show how close an AMD GPU kernel runs to hardware peak for compute throughput, memory bandwidth, and utilization. Use when the user asks how close to peak a kernel is, what percent of peak it reaches, whether a measured percentage is good enough or normal, whether a utilization number should be higher, or asks for an overall efficiency summary. Not for CUDA tools, Windows, per-instruction analysis, or a datasheet or theoretical-peak question when there is no profile to read. +--- + +# Read System Speed-of-Light + +Speed-of-Light is the first panel to read after profiling. It puts every +headline metric next to the hardware peak for the current GPU, so one table +says whether a kernel is near any limit at all. + +Run `rocprof-compute analyze --help` before choosing flags. + +## Before you start + +If `command -v rocprof-compute` finds the tool, skip this section. + +Otherwise, check that ROCm and a GPU are present: + +```bash +rocminfo +amd-smi list +``` + +If either command fails, ROCm is not installed or no GPU is visible. Stop and +point the user to "Install the ROCm Core SDK" in the +[install guide](https://github.com/ROCm/rocm-systems/blob/ebf465e112ae1ca99a2e6e45de06c0ad84388439/projects/rocprofiler-compute/docs/install/core-install.rst). Installing rocprof-compute +alone does not help. + +Check the GPU against +[compatible accelerators](https://github.com/ROCm/rocm-systems/blob/ebf465e112ae1ca99a2e6e45de06c0ad84388439/projects/rocprofiler-compute/docs/reference/compatible-accelerators.rst). +The table lists product families, so match the `Marketing Name` that `rocminfo` +prints, and use the `gfx` name to confirm. If the GPU is not supported, stop and +tell the user. + +Then follow "Install the analyze mode dependencies" in the +[install guide](https://github.com/ROCm/rocm-systems/blob/ebf465e112ae1ca99a2e6e45de06c0ad84388439/projects/rocprofiler-compute/docs/install/core-install.rst), which installs +rocprof-compute in a virtual environment of its own. Use the nightly index so the +version matches these skills, and confirm with +`rocprof-compute analyze --verify-deps`. Call that environment's +`rocprof-compute` by its full path, so the application keeps running with its own +Python. + +## 1. Get the panel + +```bash +rocprof-compute analyze --path ./workloads// --list-stats +rocprof-compute analyze --path ./workloads// -k -b sol +``` + +`--path` points at the directory containing `profiling_config.yaml`. Always +pass `-k` with one kernel id from `--list-stats`; a Speed-of-Light table +averaged over unrelated kernels means nothing. + +If there is no workload yet, or the panel is missing or empty, collect +Speed-of-Light alone: + +```bash +rocprof-compute profile --name -b sol -- +``` + +### Collect only what the question needs + +Never run a full profile. It replays the application once per counter pass and +also runs the roofline benchmarks. + +- List what this GPU can collect with `rocprof-compute --list-blocks ` + and `rocprof-compute --list-metrics `. Pass `-b` only the block + aliases, block ids, or metric ids (for example `2.1.0`) the question needs. + Aliases differ by architecture. +- `--set ` collects a predefined topic in one pass. + `rocprof-compute profile --list-sets` lists the sets. +- Add `-k ` when the user names the kernels that matter. +- Add `--no-roof` unless the question is about roofline. A `-b` list without + block 4 (`roof`) already skips it. +- `--iteration-multiplexing kernel` spreads counters over repeated dispatches of + each kernel, so it only helps when every profiled kernel is dispatched many + times. Pair it with `-k`. See Iteration multiplexing in + [profile mode](https://github.com/ROCm/rocm-systems/blob/ebf465e112ae1ca99a2e6e45de06c0ad84388439/projects/rocprofiler-compute/docs/how-to/profile/mode.rst). +- Tell the user which analysis panels the narrower profile leaves empty. + +## 2. Read it + +The table has `Metric`, `Avg`, `Unit`, `Peak`, and `Percent of Peak` columns. +The peak is already computed for the detected architecture. Compare `Avg` +against the `Peak` in the same row and never calculate a peak yourself. + +`Percent of Peak` is `N/A` for metrics that are already percentages, such as +the utilization rows. Read those as the percentage they are. + +What the shape of the table tells you: + +| Reading | Meaning | Next step | +|---|---|---| +| A memory row near peak | bandwidth-limited | the `rocprof-compute-memory` skill | +| An MFMA or VALU row near peak | compute-limited | the `rocprof-compute-roofline` skill | +| Every row far below peak | not limited by throughput | occupancy and launch limits, via the `rocprof-compute-kernel-bottleneck` skill | +| MFMA at zero on a GEMM | matrix pipeline unused | check the kernel's instruction mix with `-b cu_ins` | + +A row at zero is a real reading only if the kernel was expected to use that +unit. MFMA at zero is normal for a kernel with no matrix instructions. + +Do not apply remembered thresholds. "Above 70% is saturated" is not a rule the +tool encodes, and the correct comparison is always against the printed peak, +the neighbouring kernels, and a known-good baseline for the same workload. + +## 3. Understand the metrics + +Each metric's definition and the counters behind it are in +[system-speed-of-light.rst](https://github.com/ROCm/rocm-systems/blob/ebf465e112ae1ca99a2e6e45de06c0ad84388439/projects/rocprofiler-compute/docs/conceptual/cdna/system-speed-of-light.rst) +for CDNA accelerators and +[the RDNA equivalent](https://github.com/ROCm/rocm-systems/blob/ebf465e112ae1ca99a2e6e45de06c0ad84388439/projects/rocprofiler-compute/docs/conceptual/rdna/system-speed-of-light.rst) +for client APUs. The whole metric hierarchy is in the +[performance model](https://github.com/ROCm/rocm-systems/blob/ebf465e112ae1ca99a2e6e45de06c0ad84388439/projects/rocprofiler-compute/docs/conceptual/performance-model.rst). + +## 4. Compute Speed-of-Light + +The compute pipeline has its own Speed-of-Light table with the same shape: + +```bash +rocprof-compute analyze --path ./workloads// -k -b cu_pipe +``` + +Read it when the system-level table shows compute near peak and you need to +know which pipeline is responsible. + +## 5. Normalization + +`-n` changes the unit for the entire report, not one metric. Leave it at +`per_kernel` for Speed-of-Light. The peaks are expressed to match the default, +and switching units makes the peak comparison harder to read, not easier. diff --git a/skills/rocprof-compute-speed-of-light/evals/evals.json b/skills/rocprof-compute-speed-of-light/evals/evals.json new file mode 100644 index 0000000..628f8c8 --- /dev/null +++ b/skills/rocprof-compute-speed-of-light/evals/evals.json @@ -0,0 +1,72 @@ +{ + "evaluations": [ + { + "id": "sol-percent-of-peak", + "skill_should_trigger": true, + "prompt": "how close to peak is this kernel running in my rocprof-compute workload?\n\nThe application is ./vcopy.cpp.", + "workspace": "evals/samples", + "expected_behavior": [ + "analyzes the sol block for a single kernel id", + "reads the Percent of Peak column from the report", + "compares against the Peak value the report prints rather than a remembered threshold" + ], + "unexpected_behavior": [ + "calculates a theoretical peak instead of using the one in the report" + ], + "logs_contain": ["-b sol"] + }, + { + "id": "sol-overall-efficiency", + "skill_should_trigger": true, + "prompt": "give me an overall efficiency summary of this rocprofiler-compute capture\n\nThe application is ./vcopy.cpp.", + "workspace": "evals/samples", + "expected_behavior": [ + "selects one kernel with -k before reading Speed-of-Light", + "summarizes using the printed peak and percent of peak values" + ], + "unexpected_behavior": [ + "reports a Speed-of-Light table averaged across unrelated kernels" + ], + "logs_contain": ["-b sol"] + }, + { + "id": "sol-mfma-zero", + "skill_should_trigger": true, + "prompt": "my rocprof-compute report shows MFMA FLOPs at zero, is that a bug?", + "expected_behavior": [ + "explains that zero is expected when the kernel issues no matrix instructions", + "suggests confirming the instruction mix before treating it as a problem" + ], + "unexpected_behavior": [ + "treats zero MFMA as a defect without checking whether the kernel should use matrix instructions" + ] + }, + { + "id": "sol-threshold-pushback", + "skill_should_trigger": true, + "prompt": "my kernel is at 62% of peak memory bandwidth, is that good enough?", + "expected_behavior": [ + "declines to apply a fixed pass or fail cutoff", + "suggests comparing against neighbouring kernels or a known-good baseline" + ], + "unexpected_behavior": [ + "states a universal threshold such as 70% as if the tool defined one" + ] + }, + { + "id": "sol-instruction-level", + "skill_should_trigger": false, + "prompt": "which individual instruction in this kernel stalls the most?" + }, + { + "id": "sol-cuda-peak", + "skill_should_trigger": false, + "prompt": "what percent of peak is my CUDA kernel hitting in Nsight Compute?" + }, + { + "id": "sol-hardware-spec-lookup", + "skill_should_trigger": false, + "prompt": "what is the theoretical peak FP64 throughput of an MI300X?" + } + ] +} diff --git a/skills/rocprof-compute-speed-of-light/evals/machine.yml b/skills/rocprof-compute-speed-of-light/evals/machine.yml new file mode 100644 index 0000000..78cd197 --- /dev/null +++ b/skills/rocprof-compute-speed-of-light/evals/machine.yml @@ -0,0 +1 @@ +os: [Linux] diff --git a/skills/rocprof-compute-speed-of-light/evals/samples/vcopy.cpp b/skills/rocprof-compute-speed-of-light/evals/samples/vcopy.cpp new file mode 100644 index 0000000..1e3071b --- /dev/null +++ b/skills/rocprof-compute-speed-of-light/evals/samples/vcopy.cpp @@ -0,0 +1,230 @@ +// Copyright (c) Advanced Micro Devices, Inc. +// SPDX-License-Identifier: MIT + +// Keep in sync with the original under projects/rocprofiler-compute/: +// sample/vcopy.cpp + +#include "hip/hip_runtime.h" +#include +#include +#include +#include +#include +#include +using namespace std; + +#define HIP_ASSERT(x) (assert((x)==hipSuccess)) + +// Unlike HIP_ASSERT, this reports the failing call and error string, and stays +// active when NDEBUG is defined. +#define HIP_CHECK(x) \ + do { \ + hipError_t _hip_check_err = (x); \ + if (_hip_check_err != hipSuccess) { \ + fflush(stdout); /* keep the error in sequence when stdout is a pipe */ \ + fprintf(stderr, "HIP error at %s:%d in '%s': %s (%d)\n", __FILE__, \ + __LINE__, #x, hipGetErrorString(_hip_check_err), \ + (int)_hip_check_err); \ + exit(EXIT_FAILURE); \ + } \ + } while (0) + +// Cap on how many individual mismatches are printed. A failed launch leaves +// every element mismatched, which is millions of lines for a typical -n. +#define MAX_REPORTED_ERRORS 10 + +// HIP kernel. Each thread takes care of one element of c +__global__ void vecCopy(double *a, double *b, double *c, int n, int stride) { + // Get our global thread ID + int id = blockIdx.x*blockDim.x+threadIdx.x; + if (id < n) + c[id] = a[id]; +} + +// Duplicate of vecCopy kernel. Included for testing purposes +__global__ void vecCopy_2(double *a, double *b, double *c, int n, int stride) { + // Get our global thread ID + int id = blockIdx.x*blockDim.x+threadIdx.x; + if (id < n) + c[id] = a[id]; +} + +void usage() { + std::cout << "Usage: vcopy [OPTIONS]\n"; + std::cout << "Required:\n"; + std::cout << " -n/--numThreads Set the num of threads\n"; + std::cout << " -b/--blockSize Set the block size\n"; + std::cout << "Optional:\n"; + std::cout << " -d/--dev Set the device ID [Default: 0]\n"; + std::cout << " -i/--iter Set the num of iterations [Default: 1]\n"; + std::cout << " -h/--help Display this help message\n"; + exit(1); + return; +} + +int main(int argc, char* argv[]) { + // Size of vectors + int n; //64 MB + int blockSize, gridSize; + + // Launch multiple kernels + bool multiKernel = false; + + // Host input vectors + double *h_a; + double *h_b; + //Host output vector + double *h_c; + //Host output vector for verification + double *h_verify_c; + + // Device input vectors + double *d_a; + double *d_b; + // Device output vector + double *d_c; + + int stride = 1; + int devId = 0; + int numIter = 1; + + for (int i = 0; i < argc; i++){ + std::string arg = argv[i]; + if ((arg == "--blockSize" || arg == "-b") && i+1 < argc) + blockSize = std::atoi(argv[i+1]); + + else if ((arg == "--vec" || arg == "-n") && i+1 < argc) + n = std::atoi(argv[i+1]); + + else if ((arg == "--device" || arg == "-d") && i+1 < argc) + devId = std::atoi(argv[i+1]); + + else if ((arg == "--iter" || arg == "-i") && i+1 < argc) + numIter = std::atoi(argv[i+1]); + + else if (arg == "--multikernel") + multiKernel = true; + + else if (arg == "--help" || arg == "-h") + usage(); + } + + if (blockSize == 0) + usage(); + + if (n == 0) + usage(); + + + int numGpuDevices; + HIP_ASSERT(hipGetDeviceCount(&numGpuDevices)); + if(devId >= numGpuDevices) + devId = 0; + HIP_ASSERT(hipSetDevice(devId)); + + printf("vcopy testing on GCD %d\n", devId); + + assert(n > 0); + assert(blockSize > 0); + + // Size, in bytes, of each vector + size_t bytes = n*sizeof(double)*stride; + + // Allocate memory for each vector on host + h_a = (double*)malloc(bytes); + h_b = (double*)malloc(bytes); + h_c = (double*)malloc(bytes); + h_verify_c = (double*)malloc(bytes); + + printf("Finished allocating vectors on the CPU\n"); + + // Allocate memory for each vector on GPU + HIP_ASSERT(hipMalloc(&d_a, bytes)); + HIP_ASSERT(hipMalloc(&d_b, bytes)); + HIP_ASSERT(hipMalloc(&d_c, bytes)); + + printf("Finished allocating vectors on the GPU\n"); + + // Initialize vectors on host + for(int i = 0; i < n; i++) { + h_a[i] = i; + h_b[i] = i; + } + + // Copy host vectors to device + HIP_ASSERT(hipMemcpy(d_a, h_a, bytes, hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(d_b, h_b, bytes, hipMemcpyHostToDevice)); + + printf("Finished copying vectors to the GPU\n"); + + // Number of thread blocks in grid + gridSize = (int)ceil((float)n/blockSize); + int tot_waves = (blockSize*gridSize)/64; + float num_bytes_kb = ((sizeof(double))*n)/(1024); + float num_bytes_wave = (1.0*num_bytes_kb)/(1.0*tot_waves); + + printf("sw thinks it moved %f KB per wave \n", (2.0*num_bytes_wave)); + printf("Total threads: %d, Grid Size: %d block Size:%d, Wavefronts:%d:\n", n, gridSize, blockSize, tot_waves); + printf("Launching the kernel on the GPU\n"); + + // Execute the kernel + for(int i = 0; i < numIter; i++){ + hipLaunchKernelGGL(vecCopy, dim3(gridSize), dim3(blockSize), 0, 0, d_a, d_b, d_c, n, stride); + // Catches launch-time failures (e.g. no kernel image for this device); + // hipDeviceSynchronize() catches failures during execution. + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + printf("Finished executing kernel\n"); + // Optionally, launch a second kernel. Only here for testing purposes + if (multiKernel){ + hipLaunchKernelGGL(vecCopy_2, dim3(gridSize), dim3(blockSize), 0, 0, d_a, d_b, d_c, n, stride); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + printf("Finished executing kernel\n"); + } + } + + // Copy array back to host + HIP_ASSERT(hipMemcpy( h_c, d_c, bytes, hipMemcpyDeviceToHost)); + printf("Finished copying the output vector from the GPU to the CPU\n"); + + // Compute for CPU + for(int i=0; i 1e-5) { + if (numErrors < MAX_REPORTED_ERRORS) + printf("Error at position i %d, Expected: %f, Found: %f \n", i, + h_verify_c[i*stride], h_c[i*stride]); + numErrors++; + } + } + if (numErrors > MAX_REPORTED_ERRORS) + printf("... %d further mismatches not shown\n", numErrors - MAX_REPORTED_ERRORS); + if (numErrors > 0) + printf("Verification FAILED: %d of %d elements mismatched\n", numErrors, n); + //printf("Printing few elements from the output vector\n"); + for(int i = 0; i < 20; i++) { + //printf("Output[%d]:%f\n",i, h_c[i]); + } + + printf("Releasing GPU memory\n"); + + // Release device memory + HIP_ASSERT(hipFree(d_a)); + HIP_ASSERT(hipFree(d_b)); + HIP_ASSERT(hipFree(d_c)); + + // Release host memory + printf("Releasing CPU memory\n"); + free(h_a); + free(h_b); + free(h_c); + + return numErrors > 0 ? EXIT_FAILURE : EXIT_SUCCESS; +} diff --git a/skills/rocprof-compute-speed-of-light/skill-card.md b/skills/rocprof-compute-speed-of-light/skill-card.md new file mode 100644 index 0000000..eac7e0e --- /dev/null +++ b/skills/rocprof-compute-speed-of-light/skill-card.md @@ -0,0 +1,13 @@ +# Skill Card + +## Description + +Reads the rocprof-compute System Speed-of-Light panel to show how close an AMD GPU kernel runs to hardware peak. + +## Owner + +rocprofiler-compute, a project in [rocm-systems](https://github.com/ROCm/rocm-systems/) + +## License + +MIT