diff --git a/skills/rocprofiler-pc-sampling/.federated.json b/skills/rocprofiler-pc-sampling/.federated.json index 16f05e4..8cd8f6f 100644 --- a/skills/rocprofiler-pc-sampling/.federated.json +++ b/skills/rocprofiler-pc-sampling/.federated.json @@ -2,9 +2,9 @@ "source": "rocm-rocm-systems", "repo": "ROCm/rocm-systems", "ref": "develop", - "commit": "62872f06058460d0a7fc2a1c629bc182550f8427", + "commit": "8a45f2c24fa5a56d4bba1ac8d08e5731e1dc89c5", "path": "projects/rocprofiler-sdk/skills/pc-sampling", "license": "MIT", - "content_hash": "8d2eee99e01b1aa20d257c17249990d01417b3998e8821688857529830ccb6b6", - "imported_at": "2026-10-16T19:55:17Z" + "content_hash": "4519f7b689cd7f494e274bda6558fcc1953a564bc76aa7ddca1eb8d843cb7f84", + "imported_at": "2026-10-28T06:35:53Z" } diff --git a/skills/rocprofiler-pc-sampling/SKILL.md b/skills/rocprofiler-pc-sampling/SKILL.md index 006d7b3..c4c1e80 100644 --- a/skills/rocprofiler-pc-sampling/SKILL.md +++ b/skills/rocprofiler-pc-sampling/SKILL.md @@ -16,7 +16,7 @@ description: Profiles a target application using rocprofv3 with program counter To run rocprofv3 with PC sampling: - explicitly enable pc sampling with the flag `--pc-sampling-beta-enabled`. -- specify the type of pc sampling method with `--pc-sampling-method`: `host_trap` or `stochastic` +- specify the pc sampling method with `--pc-sampling-method`: `host_trap` or `stochastic` - specify the interval unit for the sampling with `--pc-sampling-unit`: `time` for host_trap and `cycles` for stochastic - specify the sampling interval with `--pc-sampling-interval`: an integer representing microseconds if you specified `time` or cycles if you specified `cycles`. - add the `--kernel-trace` flag to track kernel launches for attribution @@ -90,4 +90,4 @@ A brief overview of stall types and their causes is available in this skill's re The rocm-systems repository has several PC sampling documentation files in this folder: projects/rocprofiler-sdk/source/docs/how-to/. You may also access them here: https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/develop/how-to/cdna3-cdna4-pc-sampling.html -Only one PC sampling configurations is supported at a time, so if exactly one valid configuration is reported then PC sampling is probably being actively used by another process. +Only one PC sampling configuration can be used at a time, so if exactly one valid configuration is reported then PC sampling is probably being actively used by another process. If this happens, report it to the user and inform them that results may be affected. diff --git a/skills/rocprofiler-pc-sampling/evals/evals.json b/skills/rocprofiler-pc-sampling/evals/evals.json index e0c9388..a740698 100644 --- a/skills/rocprofiler-pc-sampling/evals/evals.json +++ b/skills/rocprofiler-pc-sampling/evals/evals.json @@ -3,35 +3,35 @@ { "id": "explicit-ask-for-pc-sampling", "skill_should_trigger": true, - "prompt": "use pc sampling to profile my HIP application. my application is demo_program.hip.cpp in the skill's evals/resources/ folder", + "prompt": "use pc sampling to profile my HIP application. my application is myHipProgram.hip.cpp in the skill's evals/resources/ folder", "expected_behavior": ["calls rocprofv3 with pc sampling flags, runs the analysis script, and informs the user of the results"] }, { "id": "ask-to-find-slow-gpu-code", "skill_should_trigger": true, - "prompt": "figure out slowest parts of my gpu code and suggest solutions. my application is demo_program.hip.cpp in the skill's evals/resources/ folder", + "prompt": "figure out slowest parts of my gpu code and suggest solutions. my application is myHipProgram.hip.cpp in the skill's evals/resources/ folder", "expected_behavior": ["uses rocprofv3 with pc sampling flags, analyzes the results to determine bottlenecks, and suggests solutions"] }, { "id": "ask-to-characterize-performance", "skill_should_trigger": true, - "prompt": "analyze the performance characteristics of my HIP program. my application is demo_program.hip.cpp in the skill's evals/resources/ folder", + "prompt": "analyze the performance characteristics of my HIP program. my application is myHipProgram.hip.cpp in the skill's evals/resources/ folder", "expected_behavior": ["uses rocprofv3 with pc sampling flags, analyzes the results, and informs the user"] }, { "id": "vague-performance-question", "skill_should_trigger": false, - "prompt": "how can i make my application faster. my application is demo_program.hip.cpp in the skill's evals/resources/ folder" + "prompt": "how can i make my application faster. my application is myHipProgram.hip.cpp in the skill's evals/resources/ folder" }, { "id": "vague-gpu-improvement-question", "skill_should_trigger": false, - "prompt": "how could i improve this gpu program. my application is demo_program.hip.cpp in the skill's evals/resources/ folder" + "prompt": "how could i improve this gpu program. my application is myHipProgram.hip.cpp in the skill's evals/resources/ folder" }, { "id": "cuda-profiling-request", "skill_should_trigger": false, - "prompt": "perform pc sampling on my cuda application. my application is cuda_demo.cu in the skill's evals/resources/ folder" + "prompt": "perform pc sampling on my cuda application. my application is myCudaProgram.cu in the skill's evals/resources/ folder" } ] } diff --git a/skills/rocprofiler-pc-sampling/evals/resources/cuda_demo.cu b/skills/rocprofiler-pc-sampling/evals/resources/myCudaProgram.cu similarity index 73% rename from skills/rocprofiler-pc-sampling/evals/resources/cuda_demo.cu rename to skills/rocprofiler-pc-sampling/evals/resources/myCudaProgram.cu index bd7c62f..798f705 100644 --- a/skills/rocprofiler-pc-sampling/evals/resources/cuda_demo.cu +++ b/skills/rocprofiler-pc-sampling/evals/resources/myCudaProgram.cu @@ -1,10 +1,30 @@ -// cuda_demo.cu +// MIT License // +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + // Build: -// nvcc -O3 -g cuda_demo.cu -o cuda_demo +// nvcc -O3 -g myCudaProgram.cu -o myCudaProgram // // Run: -// ./cuda_demo [repeat] # repeat defaults to 1 +// ./myCudaProgram [repeat] # repeat defaults to 1 #include @@ -23,8 +43,9 @@ } \ } while (0) +// Same problem sizes as the HIP version (myHipProgram.hip.cpp). constexpr int kN = 8192; -constexpr int kTransposeReps = 36; +constexpr int kTransposeReps = 35; constexpr int kBlock = 256; constexpr int kDivThreads = 1 << 20; constexpr int kDivIters = 58000; @@ -32,22 +53,12 @@ constexpr int kAtomicThreads = 1 << 20; constexpr int kAtomicBins = 4; constexpr int kAtomicReps = 14; -__global__ __launch_bounds__(kBlock) void k_naive_transpose(float* __restrict__ a, - float* __restrict__ b, int n, - int reps) { +__global__ __launch_bounds__(kBlock) void k_naive_transpose(const float* __restrict__ in, + float* __restrict__ out, int n) { const int col = blockIdx.x * blockDim.x + threadIdx.x; const int row = blockIdx.y * blockDim.y + threadIdx.y; if (col >= n || row >= n) return; - - for (int i = 0; i < reps; ++i) { - if ((i & 1) == 0) { - b[static_cast(col) * n + row] = - a[static_cast(row) * n + col] + 1.0f; - } else { - a[static_cast(col) * n + row] = - b[static_cast(row) * n + col] + 1.0f; - } - } + out[static_cast(col) * n + row] = in[static_cast(row) * n + col] + 1.0f; } __global__ __launch_bounds__(kBlock) void k_integer_div(uint32_t* __restrict__ out, @@ -95,7 +106,7 @@ int main(int argc, char** argv) { cudaDeviceProp prop; CUDA_CHECK(cudaGetDeviceProperties(&prop, 0)); - std::printf("Device: %s (SM %d.%d), %d SMs\n", prop.name, prop.major, prop.minor, + std::printf("Device: %s (sm_%d%d), %d SMs\n", prop.name, prop.major, prop.minor, prop.multiProcessorCount); const size_t elems = static_cast(kN) * kN; @@ -119,7 +130,12 @@ int main(int argc, char** argv) { Timer timer; for (int rep = 0; rep < repeat; ++rep) { timer.start(); - k_naive_transpose<<>>(d_in, d_out, kN, kTransposeReps); + for (int i = 0; i < kTransposeReps; ++i) { + if ((i & 1) == 0) + k_naive_transpose<<>>(d_in, d_out, kN); + else + k_naive_transpose<<>>(d_out, d_in, kN); + } transpose_ms += timer.stop(); timer.start(); diff --git a/skills/rocprofiler-pc-sampling/evals/resources/demo_program.hip.cpp b/skills/rocprofiler-pc-sampling/evals/resources/myHipProgram.hip.cpp similarity index 81% rename from skills/rocprofiler-pc-sampling/evals/resources/demo_program.hip.cpp rename to skills/rocprofiler-pc-sampling/evals/resources/myHipProgram.hip.cpp index 4e277cc..5530ff5 100644 --- a/skills/rocprofiler-pc-sampling/evals/resources/demo_program.hip.cpp +++ b/skills/rocprofiler-pc-sampling/evals/resources/myHipProgram.hip.cpp @@ -1,8 +1,30 @@ +// MIT License +// +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + // Build: -// hipcc -O3 -g demo_program.hip.cpp -o demo_program +// hipcc -O3 -g myHipProgram.hip.cpp -o myHipProgram // // Run: -// ./demo_program [repeat] # repeat defaults to 1 (~100 ms of kernel time) +// ./myHipProgram [repeat] # repeat defaults to 1 (~100 ms of kernel time) #include