Repository navigation
[AIROCMLIR-901] Enable the Benchmark and Report Performance nightly stage - #546
Muhamed-Husic wants to merge 4 commits into
Conversation
| add_library(ck-common INTERFACE IMPORTED) | ||
| target_link_libraries(ck-common INTERFACE ${LIBS}) | ||
| target_compile_options(ck-common INTERFACE "-std=c++17") | ||
| get_target_property(CK_INCLUDE_DIRS composable_kernel::device_gemm_operations |
There was a problem hiding this comment.
This file is under mlir/utils/performance/, which docs/PR_REVIEW_CHECKLIST.md lists under "rocMLIR back-port check -> Likely needs a back-port (shared with rocMLIR)". The FP8-mode detection added here is not Triton-specific -- it reads CK's own ck/config.h and would apply verbatim to ROCm/rocMLIR, which carries the same ck-benchmark-driver. Per the checklist's Verdict subsection, the PR description must contain one of: (a) a link to a parallel rocMLIR PR, (b) a one-line note explaining why the divergence is intentional, or (c) confirmation that the file no longer exists / has been refactored on the rocMLIR side. The description currently describes the change but does not address back-porting, so please add the note (or open the parallel PR). The Jenkinsfile half of this PR does not trigger the rule -- mlir/utils/jenkins/Jenkinsfile is not in the shared path list (only mlir/utils/jenkins/static-checks/ is), so this CMakeLists is the only shared path here.
| ${ckDtypesCmakeOptions(CHIP)} | ||
| ${ckFp8CmakeOptions(CHIP)} | ||
| """, | ||
| 'device_gemm_operations') |
There was a problem hiding this comment.
The CK build is now narrowed to the device_gemm_operations target, and installCKGemmOnly (line 404) copies only libdevice_gemm_operations.a plus the composable_kerneldevice_gemm_operationsTargets*.cmake exports. But line 1347 then builds the ck-benchmark-driver target, which is a custom target depending on both ck-gemm-benchmark-driver and ck-atten-benchmark-driver (mlir/utils/performance/ck-benchmark-driver/CMakeLists.txt:59-60). The attention driver includes ck/library/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute.hpp and .../batched_gemm_bias_softmax_gemm_permute.hpp (ck-atten-benchmark-driver.cpp:7-8), whose add_device_*_instances definitions do not live in libdevice_gemm_operations.a. Two plausible failure modes: undefined references when linking ck-atten-benchmark-driver, or a CMake configure error if the imported composable_kernel::device_gemm_operations target's INTERFACE_LINK_LIBRARIES names a CK target whose export file was not copied. Either would be swallowed by the catchError(buildResult: null) at line 1325, so the nightly would stay green with the CK column permanently empty. The Test Plan states the CK comparison was not run locally, so this path is currently unverified. Please either build/install the attention op library as well, or narrow the build target at line 1347 to ck-gemm-benchmark-driver (which is all the --external-gemm-library CK run at line 1357 actually needs), and confirm against a nightly before merge.
There was a problem hiding this comment.
Neither failure mode applies. CK builds both attention instance sets the driver uses (plain and bias batched_gemm_softmax_gemm_permute, from one directory) into device_gemm_operations, and CK's package config only loads the export file for the requested component, whose link interface is just Threads::Threads;hip::device. This is the same build and install sequence rocMLIR uses. Since catchError hides failures in this stage, I'll need to check the "Test MLIR vs CK" log and the CK report in the Jenkins run.
| stage("Create performance reports") { | ||
| // perfRunner.py names its CSVs after the chip it runs on, which | ||
| // need not match the CHIP matrix value. | ||
| def reportChip = get_gpu_architecture() |
There was a problem hiding this comment.
get_gpu_architecture() (line 623) returns the string 'N/A' when rocminfo fails or its output does not match the pattern. Every report command below, plus postProcessPerfRes(reportChip) at line 1381, then uses that literal as the filename stem, so a detection failure surfaces as a confusing "no such file: N/A_mlir_vs_miopen_perf.csv" several steps later instead of at the point of failure. Add an explicit guard right after this line -- e.g. if (reportChip == 'N/A') { error "Could not detect the GPU chip on ${env.NODE_NAME} for ${CHIP} reports" } -- so the stage fails with a diagnostic naming the node. The same applies to ckChip at line 1360.
|
|
||
| stage("Copy tuning database") { | ||
| // A job that has not picked up the new parameters reports null. | ||
| def tuningDBJob = params.tuningDBJob ?: 'MLIR/rocmlirTriton-weekly' |
There was a problem hiding this comment.
The default weekly job name 'MLIR/rocmlirTriton-weekly' is spelled here and again as the tuningDBJob parameter default at line 1409. The PR description already flags this as a two-place edit once the real job is created, which is exactly the maintenance hazard worth removing now. Hoist it to a single @Field constant near DOCKER_HUB_CREDS (line 17), e.g. @Field final String DEFAULT_TUNING_DB_JOB = 'MLIR/rocmlirTriton-weekly', and reference it from both the parameter default and this ?: fallback so the name is defined once.
There was a problem hiding this comment.
Verdict: COMMENT · Findings: 4 (0 Critical, 2 Major, 2 Minor)
Scope
Enables the nightly "Benchmark and Report Performance" stage in mlir/utils/jenkins/Jenkinsfile by moving its per-chip body into a new top-level runBenchmarkMatrixRow(CHIP) (keeping the pipeline {} block under the JVM 64 KB method limit), reordering "Build MLIR" ahead of the artifact copies, adding tuningDBJob/tuningDBBuild parameters, re-enabling CKBranch/checkCK, narrowing get_gpu_architecture() to the bare chip name, and restricting the CK comparison build to CK's GEMM library. mlir/utils/performance/ck-benchmark-driver/CMakeLists.txt gains FP8-mode detection from CK's ck/config.h.
The structural work reads well. Moving the body out of pipeline {}, switching the perf/split helpers to shStrict (so a failing perfRunner.py is not masked by tee), and fixing the Build-MLIR-before-copy ordering are all clear improvements over the commented-out copy — the old ordering really would have had cmake.sh delete the tuning DBs it had just copied in.
Findings
mlir/utils/performance/ck-benchmark-driver/CMakeLists.txt:11— Major: shared-with-rocMLIR path with no back-port note in the PR description (see "rocMLIR back-port check" indocs/PR_REVIEW_CHECKLIST.md).mlir/utils/jenkins/Jenkinsfile:1342— Major: building/installing only CK'sdevice_gemm_operationswhile still building theck-benchmark-driverumbrella target, which also buildsck-atten-benchmark-driver.mlir/utils/jenkins/Jenkinsfile:1370— Minor:get_gpu_architecture()'s'N/A'fallback is used unguarded as the report filename stem.mlir/utils/jenkins/Jenkinsfile:1248— Minor: the default weekly job name is duplicated between the parameter default and the in-function fallback.
Notes
- "Copy tuning database" uses
copyArtifacts optional: truefollowed by a baresh 'ls -l build/tuning-date build/*.tsv', so the real failure the operator sees is anlsexit code rather than "no tuning DB found in<job>". Since the PR already expects this to be the nightly failure mode until the weekly job exists, consider replacing thelswith an expliciterrormessage naming the job and build selector. postProcessPerfReslists${chip}_MLIR_vs_CK.htmlinpublishHTML(allowMissing: false, ...), but the CK stage is skipped ongfx1100(in the matrix) and is wrapped incatchError(buildResult: null)elsewhere. Worth confirming on the first nightly that a missing CK report does not fail or blank the publish step. This is inherited from the rocMLIR copy, so out of scope for inline review.- Chip classification in
ckFp8CmakeOptions/ckDtypesCmakeOptionsusesgfxprefix tests. That is acceptable here (Groovy has norock::/amd_arch_dbaccess, and the flags describe CK's own FP8 ABI, not a rock hardware feature), and thestartsWith('gfx12')grouping is correct for OCP FP8. Flagging only so the divergence from the checklist's "don't classify by gfx prefix" rule is a conscious one. supportsCKBenchmark's!= 'gfx906'guard cannot fire today: the matrix axis atJenkinsfile:1986has nogfx906row. Harmless, but it is currently unreachable.
CI status
No failing or cancelled checks on 4bcc4ff. py-checks and this review job were still in progress at review time; detect passed and notify-fork-pr was skipped.
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## users/bpetkovi/AIROCMLIR-899-enable-weekly-ci #546 +/- ##
================================================================================
Coverage ? 84.66%
================================================================================
Files ? 208
Lines ? 36263
Branches ? 6549
================================================================================
Hits ? 30702
Misses ? 3333
Partials ? 2228
Flags with carried forward coverage won't be shown. Click here to find out more. 🚀 New features to boost your workflow:
|
42afd25 to
06c4c3a
Compare
5aa4a44 to
662fb3b
Compare
pabloantoniom
left a comment
There was a problem hiding this comment.
High-level question:
You say in the description that nightly will fail at "Copy tuning database" until the weekly job exists. I see that it would be executed in "Benchmark and Report Performance" stage, before "MIGraphX". I think if one fails, Jenkins will skip the rest? So it would mean that such failure would also make Jenkins skip the MIGraphX stage, so we'd lose nightly MIGraphX coverage in the meantime?
If that is the case, maybe we could we make the copy non-fatal (not sure how). Or maybe only merge this until the weekly job has a green run?
You are right, until the weekly job has a successful run the copy would fail and Jenkins would skip MIGraphX stage. I made the copy non-fatal as you suggested, so if no tuning DBs can be copied, the stage marks the build unstable and the benchmarks run untuned, so MIGraphX still runs. With this, we don't have to wait for the first weekly run anymore, before merging this PR, and once the weekly job has a green run, the next nightly picks up its DBs automatically. |
9a5833b to
fc17725
Compare
231d26b to
862096b
Compare
Depends on #361
Motivation
"Benchmark and Report Performance" is the last rocMLIR nightly stage still missing in rocmlirTriton. The stage was already in our Jenkinsfile, commented out, as an older copy of rocMLIR's. This PR enables it and ports the fixes rocMLIR has made to the stage since our copy was taken.
Technical Details
runBenchmarkMatrixRow(CHIP), to keep thepipeline {}block under the JVM's 64 KB method limit.get_gpu_architecture()now returns the bare chip name (e.g.gfx942instead ofgfx942:sramecc+:xnack-); the MIGraphX stage uses the same function. For the CK comparison, only CK's GEMM library is built, andck-benchmark-driver/CMakeLists.txtdetects CK's FP8 mode.build/, becausecmake.shwipesbuild/. In the old order, the copied tuning DBs would have been deleted and every benchmark would have run untuned, with only a warning in the log.tuningDBJobandtuningDBBuild, choose which Jenkins job and build the tuning DBs are copied from. They replace the hard-coded/MLIR/mlir-weekly, which is rocMLIR's weekly job, left over from the copied stage. The default for rocmlirTriton is explained in Notes.classifyBuildFailure.Dependencies
shStrict). It also needs its-archfix for the fusion benchmarks. I'll rebase onto develop once [AIROCMLIR-899] Enable weekly CI stages (parameter sweeps + tuning) #361 merges.Test Plan
Test Result
Notes
MLIR/rocmlirTriton-weeklyis a placeholder, andtuningDBJobdefaults to it, but no rocmlirTriton weekly job exists yet. I chose the name to match the existing naming, since our nightly job isMLIR/rocmlirTriton-nightly, and rocMLIR pairsMLIR/mlir-nightly-allwithMLIR/mlir-weekly.Submission Checklist