diff --git a/mlir/utils/jenkins/Jenkinsfile b/mlir/utils/jenkins/Jenkinsfile index 72d5423d0574..5d8297761088 100644 --- a/mlir/utils/jenkins/Jenkinsfile +++ b/mlir/utils/jenkins/Jenkinsfile @@ -1,1406 +1,19 @@ +// rocMLIR Jenkinsfile β€” pipeline definition only. +// Helper methods live in mlir/utils/jenkins/helpers/*.groovy and are loaded in +// the Bootstrap stage. They are then accessed through the namespaced handles +// declared at the top of this file (e.g. nodeUtils.dockerArgs(), +// ciLogic.setHeartbeat()). // ON CHANGING THESE, ALSO CHANGE Jenkinsfile.downstream // used for private CI -import com.cloudbees.groovy.cps.NonCPS -import groovy.transform.Field -import hudson.plugins.git.extensions.impl.CheckoutOption -import hudson.plugins.git.extensions.impl.CloneOption -import java.util.concurrent.ConcurrentHashMap -// ConcurrentHashMap helps when we need to write variables in parallel -// one instance for the whole run -@Field -ConcurrentHashMap DOCKER_ARGS_BY_NODE = new ConcurrentHashMap<>() -// Jenkins Git plugin defaults to 10 minutes per command. Use 2h for fetches and -// checkouts that can exceed that on a slow network. -@Field -final int GIT_SCM_TIMEOUT_MINUTES = 120 -// Max automatic re-kicks for a nightly/weekly build failing for transient reasons. -@Field -final int MAX_REKICK_ATTEMPTS = 2 -// Characters from the end of the log scanned to classify the failure; the decisive cause sits at the end. -@Field -final int FAILURE_LOG_TAIL_CHARS = 1000000 - -// Run `script` through bash with errexit + pipefail. Use this whenever a -// command pipes through tee/awk/grep/etc. so failures in the upstream command -// are not masked by the pipeline's last exit code. Plain `sh` runs under -// /bin/sh -xe (errexit but no pipefail); a #!/bin/bash shebang bypasses -// Jenkins's default flags, so we re-enable both explicitly here. -def shStrict(String script) { - // When running inside withHealthyNode (REKICK_ROW_LOG set), mirror this step's output to a - // per-row log so the retry handler can classify transient failures (e.g. GPU hang) that only - // appear in stdout. pipefail keeps the real command's exit code from being masked by tee. - if (env.REKICK_ROW_LOG) { - sh "#!/bin/bash\nset -eo pipefail\n{\n${script}\n} 2>&1 | tee -a \"${env.REKICK_ROW_LOG}\"" - } else { - sh "#!/bin/bash\nset -eo pipefail\n${script}" - } -} - -void buildProject(String target, String cmakeOpts) { - timeout(time: 60, activity: true, unit: 'MINUTES') { - // Configure with the CMake plugin (unchanged: same source/build dir resolution as before). - cmakeBuild generator: 'Ninja',\ - buildDir: 'build',\ - buildType: 'RelWithDebInfo',\ - installation: 'InSearchPath',\ - cmakeArgs: """-DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ - -DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang - ${cmakeOpts}""" - // Build via shStrict (was the plugin's `steps: [[args: target]]`, i.e. `ninja ` in - // build dir) so build output is mirrored to the per-row log and build-time OOM - // (ninja exit 137) can be classified as a per-server transient in withHealthyNode. - // `ninja -C ` is CWD-independent, matching the plugin's workspace-relative build dir. - shStrict "ninja -C ${env.WORKSPACE}/build ${target}" - } -} - -// Lightweight Git probe: verifies auth + network + ref exists -void gitHealthCheck() { - // Check if git installed - sh "git --version" - - // Check if git commands are healthy - String repo = scm?.userRemoteConfigs?.getAt(0)?.url - String cred = scm?.userRemoteConfigs?.getAt(0)?.credentialsId - String ref = env.CHANGE_ID ? "refs/pull/${env.CHANGE_ID}/head" - : env.BRANCH_NAME ? "refs/heads/${env.BRANCH_NAME}" - : "HEAD" - - if (!repo || !cred) { - error "[healthcheck] SCM not configured (repo='${repo}', cred='${cred}')" - } - echo "[healthcheck] Probing git: repo=${repo}, ref=${ref}" - - timeout(time: 2, unit: 'MINUTES') { - withCredentials([usernamePassword(credentialsId: cred, - usernameVariable: 'GIT_USER', - passwordVariable: 'GIT_PASS')]) { - withEnv(["REPO=${repo}", "REF=${ref}"]) { - sh ''' - set -eu - ASK="$(mktemp)"; trap 'rm -f "$ASK"' EXIT - printf '#!/bin/sh\nprintf %s "$GIT_PASS"\n' > "$ASK" - chmod +x "$ASK" - GIT_ASKPASS="$ASK" \ - git -c credential.username="$GIT_USER" \ - ls-remote --exit-code "$REPO" "$REF" >/dev/null - ''' - } - } - } - echo "[healthcheck] Git OK" -} - -@NonCPS -Map scmWithGitTimeout(Object baseScm) { - List extensions = [] - boolean hasCloneOption = false - boolean hasCheckoutOption = false - - (baseScm.extensions ?: []).each { ext -> - if (ext instanceof CloneOption) { - extensions << [ - $class: 'CloneOption', - depth: ext.depth ?: 0, - shallow: ext.shallow ?: false, - noTags: ext.noTags ?: false, - reference: ext.reference ?: '', - honorRefspec: ext.honorRefspec ?: false, - timeout: gitTimeoutAtLeast(ext.timeout) - ] - hasCloneOption = true - } else if (ext instanceof CheckoutOption) { - extensions << [$class: 'CheckoutOption', timeout: gitTimeoutAtLeast(ext.timeout)] - hasCheckoutOption = true - } else { - extensions << ext - } - } - - if (!hasCloneOption) { - extensions << [$class: 'CloneOption', timeout: GIT_SCM_TIMEOUT_MINUTES] - } - if (!hasCheckoutOption) { - extensions << [$class: 'CheckoutOption', timeout: GIT_SCM_TIMEOUT_MINUTES] - } - - Map checkoutScm = [ - $class: 'GitSCM', - branches: baseScm.branches, - doGenerateSubmoduleConfigurations: baseScm.doGenerateSubmoduleConfigurations ?: false, - extensions: extensions, - submoduleCfg: baseScm.submoduleCfg ?: [], - userRemoteConfigs: baseScm.userRemoteConfigs - ] - if (baseScm.gitTool) { - checkoutScm.gitTool = baseScm.gitTool - } - if (baseScm.browser) { - checkoutScm.browser = baseScm.browser - } - return checkoutScm -} - -@NonCPS -int gitTimeoutAtLeast(Integer timeout) { - int currentTimeout = timeout ?: 0 - return Math.max(currentTimeout, GIT_SCM_TIMEOUT_MINUTES) -} - -String scmCheckoutRetryContext(Object err) { - String msg = "${err}".toLowerCase() - try { - def logLines = currentBuild?.rawBuild?.getLog(500) ?: [] - msg = msg + '\n' + logLines.join('\n').toLowerCase() - } catch (ignored) { - // Fall back to the exception text; retry classification should not mask the checkout failure. - } - return msg -} - -boolean isRetriableScmCheckoutError(String msg) { - return [ - "connection reset by peer", - "curl 18", - "transfer closed with outstanding read data remaining", - "bytes of body are still expected", - "unexpected disconnect while reading sideband packet", - "bad pack header", - "git-remote-https died of signal 15", - "early eof", - "invalid index-pack output" - ].any { msg.contains(it) } -} - -// Retry checkout without shallow clone if GitSCM chokes on a specific SHA -void robustScmCheckout() { - int maxAttempts = 2 - for (int attempt = 1; attempt <= maxAttempts; attempt++) { - try { - // This inner 'try' handles the "reference is not a tree" fallback - try { - echo "[SCM] Attempting checkout (${attempt}/${maxAttempts})..." - checkout(scmWithGitTimeout(scm)) - echo "[SCM] Checkout successful" - // If checkout succeeds, exit the function immediately - return - } catch (err) { - def msg = "${err}".toLowerCase() - if (!msg.contains("reference is not a tree") && !msg.contains("could not checkout")) { - // If it's not a known transient error, re-throw it to be caught by the outer block - throw err - } - - // This is the fallback logic for the "reference is not a tree" error - echo "[SCM] Default checkout failed: ${err}. Retrying ONCE with robust deep clone" - String repo = scm?.userRemoteConfigs?.getAt(0)?.url - String cred = scm?.userRemoteConfigs?.getAt(0)?.credentialsId - String ref = env.CHANGE_ID ? "refs/pull/${env.CHANGE_ID}/head" - : env.BRANCH_NAME ? "refs/heads/${env.BRANCH_NAME}" - : "HEAD" - - def deepScm = [ - $class: 'GitSCM', - userRemoteConfigs: [[url: repo, credentialsId: cred, refspec: "+${ref}:${ref}"]], - branches: [[name: ref]], - doGenerateSubmoduleConfigurations: false, - extensions: [ - [$class: 'CloneOption', depth: 0, shallow: false, noTags: false, honorRefspec: true, timeout: GIT_SCM_TIMEOUT_MINUTES], - [$class: 'CheckoutOption', timeout: GIT_SCM_TIMEOUT_MINUTES] - ] - ] - checkout(deepScm) - echo "[SCM] Deep clone checkout successful." - // If the deep clone succeeds, exit the function - return - } - } catch (err) { - // This outer 'catch' block is specifically for retrying network errors - def msg = scmCheckoutRetryContext(err) - if (isRetriableScmCheckoutError(msg) && attempt < maxAttempts) { - echo "[SCM] Attempt ${attempt}/${maxAttempts} failed due to a transient git fetch error." - echo "[SCM] Waiting 2 minutes before retrying..." - sleep(time: 2, unit: 'MINUTES') - // The loop will now continue to the next attempt. - } else { - // This is either not a network error, or it was the final attempt. Fail the build - echo "[SCM] Unrecoverable SCM error after ${attempt} attempt(s)." - throw err - } - } - } -} - -Map externalGitScm(String url, String branch) { - return [ - $class: 'GitSCM', - branches: [[name: "*/${branch}"]], - doGenerateSubmoduleConfigurations: false, - extensions: [ - [ - $class: 'CloneOption', - depth: 0, - shallow: false, - noTags: false, - reference: '', - honorRefspec: false, - timeout: GIT_SCM_TIMEOUT_MINUTES - ], - [$class: 'CheckoutOption', timeout: GIT_SCM_TIMEOUT_MINUTES] - ], - submoduleCfg: [], - userRemoteConfigs: [[url: url]] - ] -} - -void robustExternalCheckout(String url, String branch) { - int maxAttempts = 2 - for (int attempt = 1; attempt <= maxAttempts; attempt++) { - try { - // Discard a partial pack before retrying the clone. - deleteDir() - checkout( - changelog: false, - poll: false, - scm: externalGitScm(url, branch) - ) - return - } catch (err) { - String context = scmCheckoutRetryContext(err) - if (attempt == maxAttempts || !isRetriableScmCheckoutError(context)) { - throw err - } - - echo "[SCM] External checkout attempt ${attempt}/${maxAttempts} failed due to a transient git fetch error." - echo "[SCM] Waiting 2 minutes before retrying..." - sleep(time: 2, unit: 'MINUTES') - } - } -} - -def resetGPUs() { - // Abort this if runs longer than 10 minutes - timeout(time: 10, unit: 'MINUTES') { - // Run the reset, but don't fail the build if anything is wrong - def rc = sh( - script: ''' - reset_all_gpus() { - echo "Scanning GPUs..." - GPU_IDS=$(rocm-smi | awk '/^[0-9]+[[:space:]]+[0-9]+[[:space:]]+0x/ { print $1 }') - if [ -z "$GPU_IDS" ]; then - echo "WARNING: No GPUs found to reset." - return 0 - fi - for id in $GPU_IDS; do - echo "Resetting GPU ID: $id" - if ! rocm-smi --gpureset -d $id; then - echo "WARNING: Unable to reset GPU $id" - fi - sleep 2 - done - return 0 - } - reset_all_gpus - ''', - returnStatus: true - ) - if (rc != 0) { - echo "WARNING: reset_all_gpus exited with code ${rc}, but continuing anyway" - } - } -} - -def advancedNodeCheck(Map params) { - script { - echo "Jenkins-side PATH = '${env.PATH}'" - } - boolean doCleanWs = params.doCleanWs - boolean doGPUcheck = params.doGPUcheck - - if (doCleanWs) { - timeout(time: 15, unit: 'MINUTES', activity: true) { - cleanWs() - } - } - - resetGPUs() - - timeout(time: 5, unit: 'MINUTES', activity: true) { sh 'docker version' } - - ['ls -la /dev/kfd', 'ls -la /dev/dri'].each { cmd -> - timeout(time: 5, unit: 'MINUTES', activity: true) { sh cmd } - } - - String nodeSpecMessage = "\nNode specification:\n" - timeout(time: 5, unit: 'MINUTES', activity: true) { - nodeSpecMessage += "\nOS info:\n" + sh(script: 'sudo dkms status', returnStdout: true).trim() + '\n' - } - echo nodeSpecMessage - - if (env.NODE_LABELS && !env.NODE_LABELS.contains('build-only')) { - timeout(time: 5, unit: 'MINUTES', activity: true) { sh 'rocminfo' } - timeout(time: 5, unit: 'MINUTES', activity: true) { sh 'rocm-smi' } - timeout(time: 5, unit: 'MINUTES', activity: true) { sh 'cat /opt/rocm/.info/version' } - if (doGPUcheck) { - timeout(time: 5, unit: 'MINUTES', activity: true) { - def n = sh(script: "lspci | grep -e 'controller' -e 'accelerators' | grep 'AMD/ATI' | wc -l", - returnStdout: true).trim().toInteger() - if (n == 0) { - error "No GPUs detected on ${env.NODE_NAME}" - } - echo "Number of GPUs on ${env.NODE_NAME}: ${n}" - } - } - } else { - echo 'Skipping GPU checks…' - } -} - -def checkNodeHealth(Map opts = [:]) { - advancedNodeCheck( - doCleanWs: opts.get('doCleanWs', true), - doGPUcheck: opts.get('doGPUcheck', true) - ) -} - -// Classifies build failure from console log. Returns [reason:, codepath:, stage:] (empty string = not found). -// Add new scenarios here by matching log patterns (order = first match wins). -Map classifyBuildFailure(String logText) { - def reason = '' - def codepath = '' - def stage = '' - def failureList = '' - def failureListLabel = '' - def failedTestsSnippet = '' - if (!logText) return [reason: reason, codepath: codepath, stage: stage, failureList: failureList, failureListLabel: failureListLabel, failedTestsSnippet: failedTestsSnippet] - - // Scenario 1: Tuning failed - errors detected in tuning log (Tune rocMLIR) - if (!reason && logText.contains('Tuning failed: Detected errors in tuning log')) { - reason = 'Tune rocMLIR: errors in tuning log (check logs for details)' - } - - // Scenario 2: SCM checkout failed (max retries, clone error, or channel error) - if (!reason && (logText.contains('ERROR: Checkout failed') || logText.contains('Maximum checkout retry attempts reached') || logText.contains("ERROR: Error cloning remote repo"))) { - reason = 'SCM checkout failed (max retries or agent/channel error)' - } - - // Scenario 3: Parameter sweeps - failing configurations discovered. - // 3a: Conv/perf sweeps (parameterSweeps.py) use "*** Summary of failures ***". - if (!reason && logText.contains('*** Summary of failures ***')) { - reason = 'Parameter sweeps: failing configurations discovered' - def summaryStart = logText.indexOf('*** Summary of failures ***') - def summaryEnd = logText.indexOf('Passed:', summaryStart) - if (summaryEnd < 0) summaryEnd = logText.indexOf('script returned exit code', summaryStart) - if (summaryEnd < 0) summaryEnd = logText.length() - failureList = logText.substring(summaryStart, summaryEnd).trim() - if (failureList.length() > 2000) failureList = failureList.substring(0, 2000) + '\n... (truncated)' - } - // 3b: Attention sweeps (attentionSweeps.py) use "Failing Configurations". - if (!reason && logText.contains('Failing Configurations')) { - reason = 'Attention parameter sweeps: failing configurations discovered' - def headerPos = logText.lastIndexOf('Failing Configurations') - def configStart = logText.indexOf('\n', headerPos) - configStart = (configStart >= 0) ? configStart + 1 : headerPos - def summaryEnd = logText.indexOf('Passed:', configStart) - if (summaryEnd < 0) summaryEnd = logText.indexOf('script returned exit code', configStart) - if (summaryEnd < 0) summaryEnd = Math.min(configStart + 3000, logText.length()) - def snippet = logText.substring(configStart, summaryEnd).trim() - snippet = snippet.replaceAll(/\[\d{4}-\d{2}-\d{2}T[\d:.]+Z\]\s*/, '') - // Append the Passed/Invalid/Failed summary line if present. - def statsLineEnd = logText.indexOf('\n', summaryEnd) - if (statsLineEnd < 0) statsLineEnd = logText.length() - def statsLine = logText.substring(summaryEnd, statsLineEnd).trim() - .replaceAll(/\[\d{4}-\d{2}-\d{2}T[\d:.]+Z\]\s*/, '') - if (statsLine) snippet = snippet + '\n\n' + statsLine - if (snippet.length() > 2000) snippet = snippet.substring(0, 2000) + '\n... (truncated)' - failureList = snippet - } - - // Scenario 4: HIP no device (hipErrorNoDevice) - if (!reason && logText.contains('RuntimeError: hipError_t.hipErrorNoDevice')) { - reason = 'HIP: no device (hipErrorNoDevice)' - } - - // Scenario 5: One or more tests failed (Failed Tests (N): ...) - if (!reason && logText.contains('Failed Tests (')) { - reason = 'One or more tests failed' - def failedStart = logText.indexOf('Failed Tests (') - def failedEnd = logText.indexOf('Testing Time:', failedStart) - if (failedEnd < 0) failedEnd = logText.indexOf('Total Discovered Tests:', failedStart) - if (failedEnd < 0) failedEnd = Math.min(failedStart + 2000, logText.length()) - failedTestsSnippet = logText.substring(failedStart, failedEnd).trim() - if (failedTestsSnippet.length() > 2000) failedTestsSnippet = failedTestsSnippet.substring(0, 2000) + '\n... (truncated)' - } - - // Scenario 6: MIGraphX CMake configuration failed. - // Match by context around "Configuring incomplete" (MIGraphX path or composable_kernel_host) so we don't rely on stage order in interleaved logs. - def cmakeConfigErrorPos = logText.lastIndexOf('Configuring incomplete, errors occurred!') - if (!reason && cmakeConfigErrorPos >= 0) { - def ctxStart = Math.max(0, cmakeConfigErrorPos - 4000) - def ctxAround = logText.substring(ctxStart, Math.min(logText.length(), cmakeConfigErrorPos + 500)) - if (ctxAround.contains('MIGraphX') || ctxAround.contains('composable_kernel_host') || ctxAround.contains('Findcomposable_kernel_host')) { - reason = 'MIGraphX: CMake configuration failed' - // Extract the last "CMake Error" block before "Configuring incomplete" as a snippet. - def cmakeErrorStart = logText.lastIndexOf('CMake Error', cmakeConfigErrorPos) - if (cmakeErrorStart >= 0) { - def snippet = logText.substring(cmakeErrorStart, cmakeConfigErrorPos).trim() - snippet = snippet.replaceAll(/\[\d{4}-\d{2}-\d{2}T[\d:.]+Z\]\s*/, '') - if (snippet.length() > 2000) snippet = snippet.substring(0, 2000) + '\n... (truncated)' - failureList = snippet - failureListLabel = 'CMake error:' - } - } - } - - // Scenario 7: Agent flapping (node repeatedly offline/online). - // Checked last: agent disconnect messages often appear as a side effect of pod termination after a real build error. - if (!reason) { - def flappingMatch = logText =~ /(\S+)\s+seems to be removed or offline.*will wait for.*come back online/ - if (flappingMatch.find()) { - reason = "Agent flapping: ${flappingMatch.group(1)} went offline/online repeatedly" - } - } - - if (!reason) reason = 'Could not match a known error pattern. See build log for details.' - - // Failure anchor: position in log where this failure was detected (used to extract stage/CODEPATH from the failing branch, not from later branches). - def failureAnchor = -1 - - // Prefer detecting the anchor directly from log patterns instead of the human-facing reason text. - def scmAnchor = Math.max(logText.lastIndexOf('Maximum checkout retry attempts reached'), - logText.lastIndexOf('[SCM] Checkout failed on')) - if (scmAnchor < 0) scmAnchor = logText.lastIndexOf("ERROR: Error cloning remote repo") - if (scmAnchor < 0) scmAnchor = logText.lastIndexOf('ERROR: Checkout failed') - - if (scmAnchor >= 0) { - failureAnchor = scmAnchor - } else { - def tuneAnchor = logText.lastIndexOf('Tuning failed: Detected errors in tuning log') - if (tuneAnchor >= 0) { - failureAnchor = tuneAnchor - } else { - def sweepsAnchor = logText.indexOf('*** Summary of failures ***') - if (sweepsAnchor < 0) sweepsAnchor = logText.indexOf('Failing Configurations') - if (sweepsAnchor >= 0) { - failureAnchor = sweepsAnchor - } else { - def hipNoDeviceAnchor = logText.lastIndexOf('hipErrorNoDevice') - if (hipNoDeviceAnchor >= 0) { - failureAnchor = hipNoDeviceAnchor - } else { - def testsFailedAnchor = logText.lastIndexOf('Failed Tests (') - if (testsFailedAnchor >= 0) { - failureAnchor = testsFailedAnchor - } else { - def migraphxAnchor = logText.lastIndexOf('Configuring incomplete, errors occurred!') - if (migraphxAnchor >= 0) { - failureAnchor = migraphxAnchor - } else { - def agentFlappingAnchor = logText.lastIndexOf('seems to be removed or offline') - if (agentFlappingAnchor >= 0) { - failureAnchor = agentFlappingAnchor - } - } - } - } - } - } - } - - def searchStart = (failureAnchor >= 0) ? Math.max(0, failureAnchor - 8000) : 0 - def searchEnd = (failureAnchor >= 0) ? Math.min(logText.length(), failureAnchor + 500) : logText.length() - def contextWindow = (failureAnchor >= 0) ? logText.substring(searchStart, searchEnd) : logText - def logBeforeAnchor = (failureAnchor > 0) ? logText.substring(0, failureAnchor) : '' - - // CODEPATH: prefer "Failed in branch Matrix - CODEPATH = 'X'" near the failure; else any CODEPATH in context window; else global. - def branchMatch = contextWindow =~ /Failed in branch Matrix - CODEPATH = ['"](\w+)['"]/ - if (branchMatch.find()) { - codepath = branchMatch.group(1) - } else { - def cpMatch = contextWindow =~ /CODEPATH\s*=\s*['"]?(\w+)['"]?|Running\s+(\w+)\s+on\s+\S+/ - if (cpMatch.find()) codepath = cpMatch[0][1] ?: cpMatch[0][2] ?: '' - } - if (!codepath) { - def cpMatch = logText =~ /CODEPATH\s*=\s*['"]?(\w+)['"]?|Running\s+(\w+)\s+on\s+\S+/ - if (cpMatch.find()) codepath = cpMatch[0][1] ?: cpMatch[0][2] ?: '' - } - - // Stage: last stage name that appears *before* the failure anchor (so we report the stage that was running when it failed). - def stageNames = ['SCM Checkout', 'Build and Test', 'Parameter sweeps', 'Parameter Sweep', 'Tune MLIR kernels', 'Tune rocMLIR', 'Code coverage', 'Archive performance DB', 'MIGraphX', 'Build and Verify MIGraphX with MLIR'] - def stageSearchText = (logBeforeAnchor.length() > 0) ? logBeforeAnchor : logText - def stageIdx = -1 - for (def name in stageNames) { - def idx = stageSearchText.lastIndexOf(name) - if (idx >= 0 && idx > stageIdx) { stage = name; stageIdx = idx } - } - - return [reason: reason, codepath: codepath, stage: stage, failureList: failureList, failureListLabel: failureListLabel, failedTestsSnippet: failedTestsSnippet] -} - -// Genuine logic/test failures that must never be re-kicked. Shared veto for both the whole-job -// classifier (isTransientHardwareFailure) and the per-server one (isPerServerTransient). -List realTestFailureSignals() { - return [ - 'failed tests (', // lit - 'error: no match found', // FileCheck - 'filecheck error', - '*** summary of failures ***', // conv/perf sweeps - 'failing configurations', // attention sweeps - 'tuning failed: detected errors', - 'invalid mlir created', // MIGraphX - ] -} - -// True if a FAILED build is transient (re-kickable), not a genuine code/test failure. Scans the -// log tail only, since the decisive cause is at the end. -boolean isTransientHardwareFailure(String logText) { - if (!logText) return false - String tail = logText.length() > FAILURE_LOG_TAIL_CHARS - ? logText.substring(logText.length() - FAILURE_LOG_TAIL_CHARS) : logText - tail = tail.toLowerCase() - - // Hardware loss wins even over the veto below: a dead GPU also makes lit report spurious failures. - def hardwareSignals = [ - 'hiperror_t.hiperrornodevice', - 'unable to reset gpu', - 'unsupported hip gpu architecture: n/a', - 'no performance report found for n/a', - 'no healthy node found', - '[withhealthynode] transient', - ] - if (hardwareSignals.any { tail.contains(it) }) return true - - // Genuine logic/test failures: never re-kick. - if (realTestFailureSignals().any { tail.contains(it) }) return false - if (tail =~ /no performance report found for gfx/) return false - - // Transient infra/connection. After the veto, since failFast and dying agents emit some of - // these as collateral of a real failure. - def abortLikeSignals = [ - 'seems to be removed or offline', - 'agentofflineexception', - 'issue with creating launcher for agent', - 'failed to run image', - 'outofmemoryerror', - 'interruptedexception', - 'ninja exited with error code 137', // OOM - 'closedchannelexception', - 'requestabortedexception', - 'broken pipe', - 'script returned exit code -1', - 'script returned exit code -2', - 'maximum checkout retry attempts reached', - 'error cloning remote repo', - 'error fetching remote repo', - 'gpu hang', - 'hw exception by gpu', - ] - if (abortLikeSignals.any { tail.contains(it) }) return true - return isRetriableScmCheckoutError(tail) -} - -// Group-1 transients that can be retried on a fresh node in-pipeline (per matrix row), as opposed -// to whole-job transients like "no healthy node found" (handled by the post-block re-kick). `text` -// is the thrown exception plus the row's console tail. Case-insensitive. Deliberately excludes -// "no healthy node found"/"[withHealthyNode] transient" (nothing to retry on), "InterruptedException" -// (failFast collateral), and all genuine test-failure markers. -boolean isPerServerTransient(String text) { - if (!text) return false - String t = text.toLowerCase() - - // GPU lost/hung on this node; these surface in the test stdout, not in the thrown exception. - // Pre-veto: a dead GPU also makes lit report spurious test failures, so it wins over realSignals. - def gpuSignals = [ - 'hiperror_t.hiperrornodevice', - 'unable to reset gpu', - 'unsupported hip gpu architecture: n/a', - 'no performance report found for n/a', - 'gpu hang', - 'hw exception by gpu', - ] - if (gpuSignals.any { t.contains(it) }) return true - - // Veto: genuine test failures are never per-server retried (mirror the whole-job classifier). - if (realTestFailureSignals().any { t.contains(it) }) return false - if (t =~ /no performance report found for gfx/) return false - - // Node/agent died mid-run, or docker/OOM on this node; these surface in the exception. - def nodeSignals = [ - 'seems to be removed or offline', - 'agentofflineexception', - 'issue with creating launcher for agent', - 'closedchannelexception', - 'requestabortedexception', - 'broken pipe', - 'script returned exit code -1', - 'script returned exit code -2', - 'failed to run image', - 'outofmemoryerror', - 'ninja exited with error code 137', - 'maximum checkout retry attempts reached', - 'error cloning remote repo', - 'error fetching remote repo', - ] - if (nodeSignals.any { t.contains(it) }) return true - return isRetriableScmCheckoutError(t) -} - -// Forwards all current parameters unchanged, with the incremented re-kick attempt counter. -List rekickParameters(int nextAttempt, String reason) { - return [ - booleanParam(name: 'nightly', value: params.nightly), - booleanParam(name: 'canXdlops', value: params.canXdlops), - booleanParam(name: 'weekly', value: params.weekly), - string(name: 'MIGraphXBranch', value: params.MIGraphXBranch), - string(name: 'CKBranch', value: params.CKBranch), - booleanParam(name: 'sharedLib', value: params.sharedLib), - booleanParam(name: 'staticLib', value: params.staticLib), - booleanParam(name: 'checkMIGraphX', value: params.checkMIGraphX), - booleanParam(name: 'checkCK', value: params.checkCK), - booleanParam(name: 'runCodeCoverage', value: params.runCodeCoverage), - booleanParam(name: 'runCoverageHtml', value: params.runCoverageHtml), - string(name: 'codecovCredentialsId', value: params.codecovCredentialsId), - string(name: 'codepath', value: params.codepath), - booleanParam(name: 'disableGfx103x', value: params.disableGfx103x), - booleanParam(name: 'disableGfx110x', value: params.disableGfx110x), - booleanParam(name: 'disableGfx120x', value: params.disableGfx120x), - booleanParam(name: 'disable90a', value: params.disable90a), - booleanParam(name: 'disable908', value: params.disable908), - booleanParam(name: 'disable942', value: params.disable942), - booleanParam(name: 'disable950', value: params.disable950), - booleanParam(name: 'ignoreExternalLinting', value: params.ignoreExternalLinting), - string(name: 'weeklyTasks', value: params.weeklyTasks), - booleanParam(name: 'rekickEnabled', value: params.rekickEnabled), - string(name: 'rekickAttempt', value: nextAttempt.toString()), - string(name: 'rekickReason', value: reason ?: ''), - ] -} - -// Parse "Aborted by USERNAME" from console log (Jenkins writes this when a user aborts the build). -def parseAbortedByFromLog(String logText) { - if (!logText) return '' - def m = logText =~ /Aborted by ([^\r\n]+)/ - return m.find() ? m.group(1).trim() : '' -} - -// Sends a Teams adaptive card for build result (webhook URL from Jenkins credential 'CI_MONITORING_TEAMS'). -// statusMessage: full phrase e.g. "Build 42 completed successfully". color: Adaptive Card color ("good"=green, "warning"=yellow, "attention"=red). -// runType: "nightly" or "weekly" (subtitle line). blueOceanUrl: Blue Ocean pipeline URL. jobUrl: classic Jenkins job/build URL. -// failureDetails: optional Map [reason:, codepath:, stage:, failureList:] β€” when set, adds Stage/CODEPATH/Details and optionally a code block for failureList. -void sendTeamsBuildNotification(String buildNumber, String statusMessage, String color, String runType, String blueOceanUrl, String jobUrl, Map failureDetails = null) { - try { - def subtitle = (runType == 'nightly') ? 'MLIR Nightly πŸŒ™' : 'MLIR Weekly πŸ“…' - def timestamp = new Date().format('yyyy-MM-dd HH:mm z') - def escapeJson = { String s -> (s ?: '').replace('\\', '\\\\').replace('"', '\\"').replace('\n', ' ') } - def escapeJsonMultiline = { String s -> (s ?: '').replace('\\', '\\\\').replace('"', '\\"').replace('\n', '\\n').replace('\r', '') } - def detailBlocks = '' - if (failureDetails) { - if (failureDetails.abortedBy != null) { - def ab = escapeJson(failureDetails.abortedBy) - detailBlocks += ",{\"type\":\"RichTextBlock\",\"inlines\":[{\"type\":\"TextRun\",\"text\":\"Aborted by: \",\"weight\":\"bolder\"},{\"type\":\"TextRun\",\"text\":\"${ab}\"}]}" - } - if (failureDetails.reason) { - def r = escapeJson(failureDetails.reason) - def c = failureDetails.codepath ? escapeJson(failureDetails.codepath) : 'β€”' - def t = failureDetails.stage ? escapeJson(failureDetails.stage) : 'β€”' - detailBlocks += ",{\"type\":\"RichTextBlock\",\"inlines\":[{\"type\":\"TextRun\",\"text\":\"Stage: \",\"weight\":\"bolder\"},{\"type\":\"TextRun\",\"text\":\"${t}\"}]},{\"type\":\"RichTextBlock\",\"inlines\":[{\"type\":\"TextRun\",\"text\":\"CODEPATH: \",\"weight\":\"bolder\"},{\"type\":\"TextRun\",\"text\":\"${c}\"}]},{\"type\":\"RichTextBlock\",\"inlines\":[{\"type\":\"TextRun\",\"text\":\"Details: \",\"weight\":\"bolder\"},{\"type\":\"TextRun\",\"text\":\"${r}\"}],\"wrap\":true}" - } - if (failureDetails.failureList) { - def flLabel = failureDetails.failureListLabel ?: 'Failing configs:' - def fl = escapeJsonMultiline(failureDetails.failureList) - detailBlocks += ",{\"type\":\"RichTextBlock\",\"inlines\":[{\"type\":\"TextRun\",\"text\":\"${escapeJson(flLabel)}\",\"weight\":\"bolder\"}]},{\"type\":\"TextBlock\",\"text\":\"${fl}\",\"wrap\":true,\"fontType\":\"monospace\",\"size\":\"small\",\"separator\":true}" - } - if (failureDetails.failedTestsSnippet) { - def fts = escapeJsonMultiline(failureDetails.failedTestsSnippet) - detailBlocks += ",{\"type\":\"RichTextBlock\",\"inlines\":[{\"type\":\"TextRun\",\"text\":\"Failed tests:\",\"weight\":\"bolder\"}]},{\"type\":\"TextBlock\",\"text\":\"${fts}\",\"wrap\":true,\"fontType\":\"monospace\",\"size\":\"small\",\"separator\":true}" - } - if (failureDetails.rekick) { - def rk = escapeJson(failureDetails.rekick) - detailBlocks += ",{\"type\":\"RichTextBlock\",\"inlines\":[{\"type\":\"TextRun\",\"text\":\"Auto re-kick: \",\"weight\":\"bolder\"},{\"type\":\"TextRun\",\"text\":\"${rk}\"}],\"wrap\":true}" - } - } - def payload = """ -{"attachments":[{"contentType":"application/vnd.microsoft.card.adaptive","content":{"type":"AdaptiveCard","\$schema":"http://adaptivecards.io/schemas/adaptive-card.json","version":"1.4","body":[{"type":"TextBlock","text":"CI Update","weight":"bolder","size":"extraLarge","separator":true},{"type":"TextBlock","text":"${subtitle}"},{"type":"TextBlock","text":"${statusMessage}","color":"${color}"}${detailBlocks},{"type":"TextBlock","text":"Finished: ${timestamp}","size":"small","isSubtle":true}],"actions":[{"type":"Action.OpenUrl","url":"${blueOceanUrl}","title":"Open Blue Ocean 🌊"},{"type":"Action.OpenUrl","url":"${jobUrl}","title":"Open Job πŸ—οΈ"}]}}]} -""" - writeFile file: 'teams-payload.json', text: payload.trim(), encoding: 'UTF-8' - withCredentials([ - string(credentialsId: 'CI_MONITORING_TEAMS', variable: 'WEBHOOK_URL'), - string(credentialsId: 'MLIR_CI_CHANNEL', variable: 'WEBHOOK_URL_MLIR') - ]) { - ['CI_MONITORING_TEAMS': 'WEBHOOK_URL', 'MLIR_CI_CHANNEL': 'WEBHOOK_URL_MLIR'].each { name, envVar -> - def resp = sh(script: "curl -s -w '\\n%{http_code}' -X POST \"\$${envVar}\" -H 'Content-Type: application/json; charset=utf-8' -d @teams-payload.json", returnStdout: true).trim() - def lines = resp.split('\n') - def code = lines[-1] - def body = lines.length > 1 ? lines[0..-2].join('\n') : '' - echo "Teams webhook (${name}) response: HTTP ${code}${body ? ' body=' + body : ''}" - if (code != '200' && code != '202') { - echo "Teams notification (${name}) may have failed (expected 200/202, got ${code})" - } - } - } - } catch (e) { - echo "Teams notification skipped or failed: ${e}" - } -} - -Map dockerArgs() { - echo "Getting Docker args from ${env.NODE_NAME}..." - def run = { cmd -> sh(script: cmd, returnStdout: true).trim() } - // discover devices - String renderFlags = run("ls -1 /dev/dri/renderD* 2>/dev/null || true") - .split() - .collect { "--device=${it}" } - .join(' ') - // /dev/kfd appears only on GPU-enabled nodes - boolean haveKfd = sh(script: '[ -e /dev/kfd ]', returnStatus: true) == 0 - String kfdFlg = haveKfd ? '--device=/dev/kfd' : '' - - // Get the GIDs of the render and video groups - String renderGid = run("getent group render | cut -d':' -f3") - String videoGid = run("getent group video | cut -d':' -f3") - - String args = """ - ${kfdFlg} \ - ${renderFlags} \ - --group-add ${renderGid} --group-add ${videoGid} - """.trim().replaceAll(/\s+/, ' ') - - DOCKER_ARGS_BY_NODE[env.NODE_NAME] = args - echo "Received Docker args for ${env.NODE_NAME}: ${args}" - return DOCKER_ARGS_BY_NODE // ConcurrentHashMap -} - - -void buildCK(String cmakeOpts, String buildTarget = '') { - sh '[ ! -d build ] || rm -rf build' - // CK's Unix Makefiles do not expose device_gemm_operations as a directly - // buildable target, while Ninja handles this target graph correctly. - cmakeBuild generator: (buildTarget ? 'Ninja' : 'Unix Makefiles'),\ - buildDir: 'build',\ - buildType: 'Release',\ - installation: 'InSearchPath',\ - cmakeArgs: """-DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ - -DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang - ${cmakeOpts} - """ - if (buildTarget) { - shStrict "cmake --build build --target ${buildTarget} --parallel \$(nproc)" - } else { - shStrict 'cd build; make -j $(nproc)' - } -} - -void installCKGemmOnly(String installDir) { - sh """#!/usr/bin/env bash - set -euo pipefail - - install_dir="${installDir}" - cmake_dir="\${install_dir}/lib/cmake/composable_kernel" - mkdir -p "\${install_dir}/include/ck" "\${install_dir}/lib" "\${cmake_dir}" - - cp -R include/ck/. "\${install_dir}/include/ck/" - cp -R library/include/ck/. "\${install_dir}/include/ck/" - cp build/include/ck/config.h build/include/ck/version.h "\${install_dir}/include/ck/" - cp build/lib/libdevice_gemm_operations.a "\${install_dir}/lib/" - cp build/composable_kernelConfig.cmake \ - build/composable_kernelConfigVersion.cmake \ - "\${cmake_dir}/" - - mapfile -t gemm_export_files < <(find build -type f -name 'composable_kerneldevice_gemm_operationsTargets*.cmake' -print) - if [ "\${#gemm_export_files[@]}" -eq 0 ]; then - echo "Could not find CK device_gemm_operations CMake export files" - exit 1 - fi - cp "\${gemm_export_files[@]}" "\${cmake_dir}/" - """ -} - -void buildMIGraphX(String cmakeOpts) { - sh '[ ! -d build ] || rm -rf build' - cmakeBuild generator: 'Unix Makefiles',\ - buildDir: 'build',\ - buildType: 'Release',\ - installation: 'InSearchPath',\ - cmakeArgs: """-DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ - -DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang - -DMIGRAPHX_USE_COMPOSABLEKERNEL=OFF - ${cmakeOpts} - """ - shStrict 'cd build; make -j $(nproc)' -} - -void getAndBuildMIGraphX(String cmakeOpts) { - robustExternalCheckout('https://github.com/ROCm/AMDMIGraphX.git', params.MIGraphXBranch) - buildMIGraphX(cmakeOpts) -} - -void getAndBuildCK(String cmakeOpts, String buildTarget = '') { - robustExternalCheckout('https://github.com/ROCm/composable_kernel.git', params.CKBranch) - buildCK(cmakeOpts, buildTarget) -} - -void showEnv() { - echo "$env.NODE_NAME" - sh 'cat /etc/os-release' - sh 'ulimit -a' - // Ignore rocm-smi failures in ixt-sjc2-05 - sh '/opt/rocm/bin/rocm-smi || true' - sh '/opt/rocm/bin/rocm_agent_enumerator' - sh 'id' - sh 'printenv' -} - -String dockerImage() { - // If this is being changed please change Dockerfile.migraphx-ci's base image as well - return 'rocm/mlir:rocm7.2-latest' -} - -String dockerImageCIMIGraphX() { - return 'rocm/mlir-migraphx-ci:rocm7.2-latest' -} - -def retryDockerOperation(String description, Closure operation) { - int attempt = 0 - def result = null - retry(10) { - attempt += 1 - try { - result = operation() - } catch (err) { - echo "[Docker retry] ${description} failed on attempt ${attempt}/10 on ${env.NODE_NAME}: ${err}" - if (attempt < 10) { - echo "[Docker retry] Waiting 5 seconds before retrying ${description}" - sleep(time: 5, unit: 'SECONDS') - } - throw err - } - } - return result -} - -// For when the docker image is in a private repo -void explicitDockerLogin() { - withCredentials([usernamePassword(credentialsId: 'DOCKER_HUB_CREDS', - usernameVariable: 'D_USER', - passwordVariable: 'D_PASS')]) { - retryDockerOperation('docker login to DockerHub') { - sh ''' - set +x - printf "%s\n" "$D_PASS" | docker login -u "$D_USER" --password-stdin - ''' - } - } -} - -def pullDockerImage(String imageName) { - def img = docker.image(imageName) - retryDockerOperation("docker pull ${imageName}") { - img?.pull() - } - return img -} - -void preMergeCheck(String codepath) { - // Only do static check on mfma codepath during PR CI - if ( (params.nightly == false) && (codepath == "mfma") ) { - echo "Performing Static Test (preMergeCheck)" - sh ''' - if [ ! -f ./build/compile_commands.json ]; then - echo "No compile commands, bailing." - exit 1 - fi - if [ ! -f ./compile_commands.json ]; then - ln -s build/compile_commands.json compile_commands.json - fi - ''' - def targetBranch = env.CHANGE_TARGET - if (!targetBranch) { - targetBranch = "develop" - } - if (params.ignoreExternalLinting == true) { - sh "python3 ./mlir/utils/jenkins/static-checks/premerge-checks.py --base-commit=origin/${targetBranch} --ignore-external" - } - else { - sh "python3 ./mlir/utils/jenkins/static-checks/premerge-checks.py --base-commit=origin/${targetBranch}" - } - } else { - echo "Static Test step skipped" - } -} - -void preMergeCheckPackage(String codepath) { - // Only do static check on mfma codepath during PR CI - if ( (params.nightly == false) && (codepath == "mfma") ) { - echo "Checking if the fat library target list is accurate" - dir('build') { - sh '../mlir/utils/jenkins/static-checks/get_fat_library_deps_list.pl > ./librockcompiler_deps.cmake.new' - } - sh 'diff -up mlir/tools/rocmlir-lib/librockcompiler_deps.cmake ./build/librockcompiler_deps.cmake.new' - } else { - echo "Skipping fat library target list check" - } -} - -void splitConfigFile(String inputFilePath, String outputFilePath, int run, int totalSplits = 5) { - shStrict """ - lines=\$(grep -Ev '(^\\s*\$|^\\s*#)' ${inputFilePath} | wc -l) - lines_per_chunk=\$(((lines + ${totalSplits} - 1) / ${totalSplits})) - start_line=\$((lines_per_chunk * (${run} - 1) + 1)) - end_line=\$((lines_per_chunk * ${run})) - - grep -Ev '(^\\s*\$|^\\s*#)' ${inputFilePath} | sed -n "\${start_line},\${end_line}p" | tee ${outputFilePath} - """ -} - -void postProcessPerfRes(String chip) { - publishHTML (target: [ - allowMissing: false, - alwaysLinkToLastBuild: false, - keepAll: true, - reportDir: 'build/reports', - reportFiles: "${chip}_MLIR_Performance_Changes.html,${chip}_MLIR_vs_MIOpen.html,${chip}_MLIR_Performance_Changes_Gemm.html,${chip}_MLIR_vs_hipBLASLt.html,${chip}_MLIR_vs_CK.html,${chip}_conv_fusion.html,${chip}_gemm_fusion.html", - reportName: "Performance report for ${chip}" - ]) - - if (fileExists("build/${chip}_mlir_vs_miopen_perf_for_plot.csv")) { - plot csvFileName: "${chip}_plot-nightly-perf-results-000001.csv",\ - csvSeries: [[file: "build/${chip}_mlir_vs_miopen_perf_for_plot.csv", displayTableFlag: false]],\ - title: "Test performance summary ${chip}, Conv",\ - yaxis: 'TFlops',\ - style: 'line',\ - group: 'Performance plots' - } - if (fileExists("build/${chip}_mlir_vs_hipblaslt_perf_for_plot.csv")) { - plot csvFileName: "${chip}_plot-nightly-perf-results-gemm-000001.csv",\ - csvSeries: [[file: "build/${chip}_mlir_vs_hipblaslt_perf_for_plot.csv", displayTableFlag: false]],\ - title: "Test performance summary ${chip}, GEMM",\ - yaxis: 'TFlops',\ - style: 'line',\ - group: 'Performance plots' - } - // Save results for future comparison - archiveArtifacts artifacts: 'build/*_mlir_*.csv,build/perf-run-date', allowEmptyArchive: true, onlyIfSuccessful: true -} - -// Get the base GPU chip name as reported by the runtime (e.g. gfx1200, gfx942). -def get_gpu_architecture() { - try { - def result = sh(script: 'rocminfo', returnStdout: true).trim() - def arch_pattern = /Name:\s+amdgcn-amd-amdhsa--(gfx[0-9a-z]+)/ - def matches = (result =~ arch_pattern) - if (matches) { - return matches[0][1] - } - return 'N/A' - } catch (Exception e) { - echo "Error getting GPU architecture name: ${e}" - return 'N/A' - } -} - -//makes sure multiple builds are not triggered for branch indexing -def resetBuild() { - if (currentBuild.getPreviousBuild() == null - || currentBuild.getPreviousBuild().getBuildCauses().toString().contains('BranchIndexingCause')) { - def buildNumber = BUILD_NUMBER as int; - if (buildNumber > 1) - milestone(buildNumber - 1); - milestone(buildNumber) - } -} - -void setHeartbeat() { - script { - System.setProperty("org.jenkinsci.plugins.durabletask.BourneShellScript.HEARTBEAT_CHECK_INTERVAL", "86400"); - } -} - -String getLabelFromCodepath(String codepath) { - echo "codepath is ${codepath}" - String label = '' - if (codepath == "mfma") { - label = 'mlir && (gfx942 || gfx908 || gfx90a)' - } else if (codepath == "gfx950") { - if (params.weekly) { - label = 'mlir && linux-mi350-8' - } else { - label = 'mlir && linux-mi350-1' - } - } else if (codepath == "gfx103x") { - // For non-performance related testing, use both workstations (gfx1030w) - // and server nodes (gfx1030) - label = 'mlir && ( gfx1030w || gfx1030 )' - } else if (codepath == "vanilla"){ - label = 'mlir' - } else if (codepath == "gfx110x") { - if (params.nightly || params.weekly) { - label = 'mlir && gfx1100' - } else { - label = 'mlir && ( gfx1100 || gfx1101 )' - } - } else if (codepath == "gfx120x") { - if (params.nightly || params.weekly) { - label = 'mlir && gfx1201' - } else { - label = 'mlir && ( gfx1200 || gfx1201 )' - } - } else { - echo "${codepath} is not supported" - label = 'wrongLabel' - } - echo "label is ${label}" - return label -} - -String getLabelFromChip(String chip) { - switch (chip) { - case "gfx906": - return getLabelFromCodepath("vanilla") - case "gfx908": - return "mlir && gfx908" - case "gfx90a": - return "mlir && gfx90a" - case "gfx942": - return "mlir && gfx942" - case "gfx950": - if (params.weekly) { - return "mlir && linux-mi350-8" - } else { - return "mlir && linux-mi350-1" - } - case "gfx1030": - // For [Tune MLIR Kernels] and [Performance report] stages, - // fix the vm-5 workstation for testing - return "mlir && vm-5" - case "gfx1100": - return "mlir && gfx1100" - case "gfx1101": - return "mlir && gfx1101" - case "gfx1200": - return "mlir && gfx1200" - case "gfx1201": - return "mlir && gfx1201" - } -} - -int setLitWorkerCount() { - int limit_lit_workers = 8 - def gpu_arch = get_gpu_architecture() - if (gpu_arch.contains('gfx908') || gpu_arch.contains('gfx90a')) { - limit_lit_workers = 20 - } else if (gpu_arch.contains('gfx942')) { - limit_lit_workers = 64 - } - return limit_lit_workers -} - -void build_fixedE2ETests(String codepath) { - // Limit the number of lit workers for gfx908, gfx90a to (8, 30) on CI as a workaround for issue #1845 and #1841 - int limit_lit_workers = setLitWorkerCount() - buildProject("check-mlir-build-only check-rocmlir-build-only${params.nightly ? ' hipblaslt-benchmark-driver' : ''}", """ - -DROCMLIR_DRIVER_PR_E2E_TEST_ENABLED=${params.nightly ? '0' : '1'} - -DROCMLIR_DRIVER_E2E_TEST_ENABLED=${params.nightly ? '1' : '0'} - -DROCK_E2E_TEST_ENABLED=${params.nightly ? '1' : '0'} - -DROCMLIR_DRIVER_TEST_GPU_VALIDATION=1 - -DROCMLIR_ENABLE_BENCHMARKS=${params.nightly ? 'hipblaslt' : ''} - -DLLVM_LIT_ARGS='-v --time-tests --timeout=3600 --max-failures=1 -j ${limit_lit_workers}' - -DCMAKE_EXPORT_COMPILE_COMMANDS=1 - """) -} - -void check_randomE2ETests(String codepath) { - // Limit the number of lit workers for gfx908, gfx90a to (8, 30) on CI as a workaround for issue #1845 and #1841 - int limit_lit_workers = setLitWorkerCount() - // Configure and build the E2E deps without running the tests, then run the GPU tests via - // shStrict so their stdout is mirrored to the per-row log (withHealthyNode classifies GPU - // hangs there and retries only this node). Running check-rocmlir directly through cmakeBuild - // would bypass shStrict and force a whole-job re-kick instead. - buildProject('check-rocmlir-build-only', """ - -DROCMLIR_DRIVER_PR_E2E_TEST_ENABLED=0 - -DROCMLIR_DRIVER_E2E_TEST_ENABLED=1 - -DROCK_E2E_TEST_ENABLED=1 - -DROCMLIR_DRIVER_RANDOM_DATA_SEED=1 - -DROCMLIR_DRIVER_TEST_GPU_VALIDATION=0 - -DLLVM_LIT_ARGS='-v --time-tests --timeout=3600 --max-failures=1 -j ${limit_lit_workers}' - -DCMAKE_EXPORT_COMPILE_COMMANDS=1 - """) - timeout(time: 60, activity: true, unit: 'MINUTES') { - shStrict 'cd build; ninja check-rocmlir' - } -} - -void parameterSweep(String CONFIG, String sweepType = "default") { - int limit_lit_workers = setLitWorkerCount() - timeout(time: 300, activity: true, unit: 'MINUTES') { - dir('build') { - if (sweepType == "attention" || sweepType == "gemm_gemm") { - String accelCodepath = "auto" - if (CONFIG == "mfma" || CONFIG == "gfx950") { - accelCodepath = "mfma" - } else if (CONFIG == "gfx103x" || CONFIG == "gfx110x" || CONFIG == "gfx120x") { - accelCodepath = "wmma" - } - shStrict """python3 ./bin/attentionSweeps.py --op ${sweepType} -j ${limit_lit_workers} --codepath ${accelCodepath} --log-failures --debug-fails""" - } else { - shStrict """python3 ./bin/parameterSweeps.py -j ${limit_lit_workers} ${CONFIG} --log-failures""" - } - } - } -} - -boolean shouldRunFromCodepath(String codepath) { - // Run vanilla on public CI - if ((codepath == "vanilla") && (params.canXdlops == false)) { - return true - } - // Run mfma on private CI - if ((codepath == "mfma") && params.canXdlops) { - return true - } - if (codepath == "gfx950" && params.canXdlops && params.disable950 == false) { - return true - } - // Run gfx103x on private nightly or weekly CI if it is not disabled - if (params.canXdlops && (params.disableGfx103x == false) && (codepath == "gfx103x") && - (params.nightly || params.weekly)) { - return true - } - // Run gfx110x on private CI if it is not disabled - if (params.canXdlops && (params.disableGfx110x == false) && (codepath == "gfx110x")) { - return true - } - // Run gfx120x on private CI if it is not disabled - if (params.canXdlops && (params.disableGfx120x == false) && (codepath == "gfx120x")) { - return true - } - return false -} - -boolean shouldRunFromChip(String chip) { - switch (chip) { - default: - return shouldRunFromCodepath("vanilla") - case "gfx90a": - // Special case because all our "vanilla" hosts are gfx90a. - return params.disable90a == false && - (shouldRunFromCodepath("mfma") || shouldRunFromCodepath("vanilla")) - case "gfx908": - return params.disable908 == false && shouldRunFromCodepath("mfma") - case "gfx942": - return params.disable942 == false && shouldRunFromCodepath("mfma") - case "gfx950": - return params.disable950 == false && shouldRunFromCodepath("gfx950") - case "gfx1030": - return shouldRunFromCodepath("gfx103x") - case "gfx1100": - return shouldRunFromCodepath("gfx110x") - case "gfx1200": - case "gfx1201": - return shouldRunFromCodepath("gfx120x") - } -} - -void archivePerfDB() { - // Note: add additional architectures here - dir ('build/perfDB') { - def architectures = params.canXdlops ? ['gfx908', 'gfx90a', 'gfx942', 'gfx950', 'gfx1100', 'gfx1201'] : ['vanilla'] - for (arch in architectures) { - try { - unstash name: "MLIR-PerfDB-${arch}" - } catch (Exception e) { - echo "No stash found for MLIR-PerfDB-${arch}, skipping." - } - } - sh 'date --utc +%Y-%m-%d >tuning-date' - } - archiveArtifacts artifacts: 'build/perfDB/**',\ - onlyIfSuccessful: true -} - -boolean shouldRunBuildAndTest(String codepath) { - // When default codepath is selected, we test mfma, gfx103x, gfx110x and gfx120x on - // private CI and vanilla on public CI - if (params.codepath == "default" && shouldRunFromCodepath(codepath)) - return true - - // When a particular codepath is selected, we only test the codepath - // on private CI - if (params.codepath == codepath && params.canXdlops) { - if (params.codepath == "mfma") return true - if (params.codepath == "vanilla") return true - if (params.codepath == "gfx950" && params.disable950 == false) return true - if (params.codepath == "gfx103x" && params.disableGfx103x == false) return true - if (params.codepath == "gfx110x" && params.disableGfx110x == false) return true - if (params.codepath == "gfx120x" && params.disableGfx120x == false) return true - return false - } -} - -boolean isNotGfx11x(String chip) { - return "${chip}" != 'gfx1100' && "${chip}" != 'gfx1101' -} - -boolean supportsCKBenchmark(String chip) { - // CK does not generate a device_gemm_operations target for gfx906. - return isNotGfx11x(chip) && "${chip}" != 'gfx906' -} - -String ckFp8CmakeOptions(String chip) { - // CK auto-derives CK_USE_*_FP8 from GPU_TARGETS, so we only force the - // numeric macro values via CMAKE_CXX_FLAGS to keep CK's `#if` checks correct. - if ("${chip}".startsWith('gfx94')) { - return ''' - -DCMAKE_CXX_FLAGS="-O3 -DCK_USE_FNUZ_FP8=1" - ''' - } - if ("${chip}" == 'gfx950' || "${chip}".startsWith('gfx12')) { - return ''' - -DCMAKE_CXX_FLAGS="-O3 -DCK_USE_OCP_FP8=1 -DCK_TILE_USE_OCP_FP8=1" - ''' - } - return ''' - -DCMAKE_CXX_FLAGS="-O3" - ''' -} - -String ckDtypesCmakeOptions(String chip) { - // The MLIR vs CK perf configs exercise f16/f32 GEMM, with int8 available - // for the CK GEMM driver. Restricting DTYPES avoids building unused CK - // operation instances on older arches such as gfx908. - if ("${chip}".startsWith('gfx94') || "${chip}" == 'gfx950' || "${chip}".startsWith('gfx12')) { - return ''' - -DDTYPES="fp8;bf8;fp16;fp32;int8" - ''' - } - return ''' - -DDTYPES="fp16;fp32;int8" - ''' -} - -void collectCoverageData(String profdata, String cov, String cpath) { - // Runs `ninja check-rocmlir` (GPU E2E), so use shStrict to mirror output to the per-row log. - shStrict """ - rm -f *.profraw - # Arbitrarily 150 GB; we typically see 125 GB of *.profraw. - if [ `df --output=avail -k . | tail -1l` -lt 153600000 ]; then - echo Not enough free disk space for profiling. - exit 1 - fi - ninja check-rocmlir - # Profile processing. - ${profdata} merge -sparse ./*.profraw -o ./coverage.profdata - rm -f build/*.profraw - ${cov} report --object ./bin/rocmlir-opt --object ./bin/rocmlir-driver \ - --object ./bin/rocmlir-gen --instr-profile ./coverage.profdata \ - --ignore-filename-regex='external/llvm-project|mlir/test/|mlir/unittests/' > ./coverage_${cpath}.report - cat ./coverage_${cpath}.report - ${cov} export --object ./bin/rocmlir-opt --object ./bin/rocmlir-driver \ - --object ./bin/rocmlir-gen --instr-profile ./coverage.profdata \ - --ignore-filename-regex='external/llvm-project|mlir/test/|mlir/unittests/' --format=lcov \ - --compilation-dir ${WORKSPACE} > ./coverage_${cpath}.lcov - """ -} - -// Produce the HTML coverage report -void produceCoverageHtml(String cov, String cpath) { - sh """ - ${cov} show --object ./bin/rocmlir-opt --object ./bin/rocmlir-driver \ - --object ./bin/rocmlir-gen --instr-profile ./coverage.profdata \ - --ignore-filename-regex='external/llvm-project|mlir/test/|mlir/unittests/' -Xdemangler=llvm-cxxfilt \ - --format=html > ./coverage_${cpath}.html - """ -} - -// Run the body on a node that passes the supplied healthChecks() block -// The health check is retried on fresh executors; the body is not retried. -// This function also retries the main 'body' if it fails due to a recoverable node-related issue (e.g., agent disconnect). -def withHealthyNode(String baseLabel, Closure healthChecks, Closure body, int maxAttempts = 3) { - def blacklist = [] // nodes and pods that already failed the check - int attempt = 0 - boolean done = false - - while (!done && attempt < maxAttempts) { - attempt += 1 - - // Build a dynamic label that excludes everything that failed before - def expr = new StringBuilder(baseLabel) - blacklist.each { expr.append(' && !').append(it) } - - echo "[withHealthyNode] attempt #${attempt}: looking for '${expr}'" - node(expr.toString()) { - // Retry ONLY the health-check. We don't want to retry the actual stages - try { - stage("Health checks on ${env.NODE_NAME}") { - echo 'Cleaning up old Docker images...' - def pruneStatus = sh(script: 'docker image prune -af --filter "until=720h"', returnStatus: true) - if (pruneStatus != 0) { - echo "[withHealthyNode] WARNING: Docker image prune failed with exit code ${pruneStatus}. Continuing health check." - } - healthChecks() - gitHealthCheck() - } - } catch (Exception err) { - echo "[withHealthyNode] ❌ ${env.NODE_NAME} rejected: ${err}" - blacklist << env.NODE_NAME - // return exits the node {} block here, not the whole function. Some groovy magic - return - } - stage("Node selected") { - // Health-check passed. Do real work - echo "[withHealthyNode] βœ… using ${env.NODE_NAME}" - } - // Per-row console log: shStrict mirrors output here so we can classify transient - // failures (e.g. GPU hang) that only appear in stdout, not in the thrown exception. - String rowLog = "${env.WORKSPACE}/.rekick-row.log" - try { - withEnv(["REKICK_ROW_LOG=${rowLog}"]) { - body() - } - // If body succeeds, we're done with the loop - done = true - } catch (Exception err) { - String rowText = '' - try { - if (fileExists(rowLog)) { - rowText = readFile(rowLog) - if (rowText.length() > FAILURE_LOG_TAIL_CHARS) { - rowText = rowText.substring(rowText.length() - FAILURE_LOG_TAIL_CHARS) - } - } - } catch (Exception ignored) { } - - if (isPerServerTransient("${err}\n${rowText}")) { - // Group-1 transient on this node: blacklist it and retry the same arch on a - // fresh node. The while loop continues (done still false); if attempts run out - // this becomes "no healthy node found", which the post-block re-kicks whole-job. - echo "[withHealthyNode] Per-server transient on ${env.NODE_NAME}. Blacklisting the node and retrying.." - echo "[withHealthyNode] Error was: ${err}" - blacklist << env.NODE_NAME - return - } - // Real failure (or a whole-job transient like no-healthy-node): fail immediately. - echo "[withHealthyNode] Execution failed with a non-recoverable error on ${env.NODE_NAME}" - echo "[withHealthyNode] Error was: ${err}" - throw err - } finally { - // Clean here (moved out of the matrix bodies) so the per-row log above survives - // until it has been classified. Never let cleanup mask the body's exception. - try { - cleanWs() - } catch (Exception cleanErr) { - echo "[withHealthyNode] cleanWs failed: ${cleanErr}" - } - } - } - } - - if (!done) { - // In-stage breadcrumb: the post block reads the log before the final "error" line is printed. - echo "[withHealthyNode] TRANSIENT: no healthy node for '${baseLabel}' after ${maxAttempts} attempts" - error "No healthy node found for '${baseLabel}' after ${maxAttempts} attempts" - } -} +// Per-domain helper handles, populated by the Bootstrap stage's load() calls. +// Declared at script scope so subsequent stages can reference them. +def scmUtils +def nodeUtils +def buildUtils +def testUtils +def reportUtils +def ciLogic pipeline { agent none @@ -1480,9 +93,50 @@ pipeline { description: 'Internal: reason of the transient failure that triggered the re-kick.') } stages { + stage("Bootstrap") { + // Loads per-domain helper scripts from the workspace and wires + // cross-helper dependencies. Must be the first stage so every + // subsequent stage can reach the helpers via their namespaced + // handles declared at the top of this file. + // + // Needs a real executor because checkout scm + load require a + // workspace. The matrix stages still allocate their own agents + // through nodeUtils.withHealthyNode(); pipeline-level agent stays + // 'none' so Bootstrap is the only stage that consumes an mlir + // executor up-front. + agent { label 'mlir' } + steps { + checkout scm + script { + def base = 'mlir/utils/jenkins/helpers' + scmUtils = load "${base}/scmUtils.groovy" + nodeUtils = load "${base}/nodeUtils.groovy" + buildUtils = load "${base}/buildUtils.groovy" + testUtils = load "${base}/testUtils.groovy" + reportUtils = load "${base}/reportUtils.groovy" + ciLogic = load "${base}/ciLogic.groovy" + + // Wire cross-helper references (helpers calling helpers + // in another file). These mirror the bare-name calls that + // used to work when every helper lived in the Jenkinsfile. + nodeUtils.scmUtils = scmUtils + buildUtils.scmUtils = scmUtils + testUtils.nodeUtils = nodeUtils + testUtils.buildUtils = buildUtils + // ciLogic owns the matrix-row orchestrators extracted + // from the Jenkinsfile, so it needs handles to every + // other helper they call. + ciLogic.scmUtils = scmUtils + ciLogic.nodeUtils = nodeUtils + ciLogic.buildUtils = buildUtils + ciLogic.testUtils = testUtils + ciLogic.reportUtils = reportUtils + } + } + } stage("Set System Property") { steps { - setHeartbeat() + script { ciLogic.setHeartbeat() } } } stage("Kill old PR builds") { @@ -1491,7 +145,7 @@ pipeline { equals expected: false, actual: params.nightly; } steps { - resetBuild() + script { ciLogic.resetBuild() } } } stage('Build and Test') { @@ -1511,131 +165,10 @@ pipeline { stages { stage('Matrix row orchestration') { when { - expression { shouldRunBuildAndTest(CODEPATH) } + expression { ciLogic.shouldRunBuildAndTest(CODEPATH) } } steps { - script { - // Prepare node - withHealthyNode( - getLabelFromCodepath(CODEPATH), - { - checkNodeHealth([doCleanWs: true]) - }, - { - stage("SCM Checkout") { - try { - robustScmCheckout() - } catch (e) { - error "[SCM] Checkout failed on ${env.NODE_NAME}: ${e}" - } - } - try { - String args = '' - def img = null - stage("Prepare Docker environment") { - // Fill in the docker args from the node - dockerArgs() - - args = DOCKER_ARGS_BY_NODE[env.NODE_NAME] - // Check these args - echo "Running ${CODEPATH} on ${env.NODE_NAME} with: ${args}" - explicitDockerLogin() - img = pullDockerImage(dockerImage()) - } - // Spin up ONE container and stay in it for all substages - img.inside(args) { - withEnv([ - "HOME=${env.WORKSPACE}", - "PATH=/opt/rocm/llvm/bin:${env.PATH}" - ]) { - - if (params.sharedLib) { - stage('Shared Library: fixed E2E') { - echo "codepath is ${CODEPATH}" - echo "Container environment:" - showEnv() - - build_fixedE2ETests("${CODEPATH}") - preMergeCheck("${CODEPATH}") - timeout(time: 60, activity: true, unit: 'MINUTES') { - shStrict 'cd build; ninja check-mlir check-rocmlir' - } - } - } - - if (params.sharedLib && params.nightly) { - stage('Shared Library: random E2E') { - check_randomE2ETests("${CODEPATH}") - } - } - - if (params.sharedLib && !params.nightly) { - stage('Tune selected rocMLIR configs') { - buildProject('ci-performance-scripts', '') - dir('build') { - timeout(time: 60, activity: true, unit: 'MINUTES') { - shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ - --op gemm \ - -c ../mlir/utils/jenkins/ci-configs/selected-gemm-configs \ - -o tuning_gemm.tsv - [ -f tuning_gemm.tsv ]""" - shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ - --op conv \ - -c ../mlir/utils/jenkins/ci-configs/selected-conv-configs \ - -o tuning_conv.tsv - [ -f tuning_conv.tsv ]""" - shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ - --op attention \ - -c ../mlir/utils/jenkins/ci-configs/selected-attention-configs \ - -o tuning_attention.tsv - [ -f tuning_attention.tsv ]""" - shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ - --op gemm_gemm \ - -c ../mlir/utils/jenkins/ci-configs/selected-gemmgemm-configs \ - -o tuning_gemmgemm.tsv - [ -f tuning_gemmgemm.tsv ]""" - shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ - --op gemm --tuning-space quick \ - -c ../mlir/utils/jenkins/ci-configs/selected-gemm-configs \ - -o quick_tuning_gemm.tsv - [ -f quick_tuning_gemm.tsv ]""" - shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ - --op conv --tuning-space quick \ - -c ../mlir/utils/jenkins/ci-configs/selected-conv-configs \ - -o quick_tuning_conv.tsv - [ -f quick_tuning_conv.tsv ]""" - shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ - --op attention --tuning-space quick \ - -c ../mlir/utils/jenkins/ci-configs/selected-attention-configs \ - -o quick_tuning_attention.tsv - [ -f quick_tuning_attention.tsv ]""" - } - } - } - } - - if (params.staticLib && !params.nightly) { - stage('Static Lib: build packages') { - sh 'rm -f build/CMakeCache.txt' - buildProject('package', '-DBUILD_FAT_LIBROCKCOMPILER=ON') - preMergeCheckPackage("${CODEPATH}") - echo "Running tests on the newly-built static library" - dir ('build') { - shStrict 'ninja check-rocmlir' - } - } - } - } - } - } - // Post block in scripted Jenkins works as try {} catch {}. - // Workspace cleanup is centralized in withHealthyNode so the per-row log survives. - catch (e) { - throw e - } - } - ) - } + script { ciLogic.runBuildAndTestMatrixRow(CODEPATH) } } } } @@ -1662,69 +195,10 @@ pipeline { stages { stage('Matrix row orchestration') { when { - expression { shouldRunFromCodepath(CODEPATH) } + expression { ciLogic.shouldRunFromCodepath(CODEPATH) } } steps { - script { - // Prepare node - withHealthyNode( - getLabelFromCodepath(CODEPATH), - { - checkNodeHealth() - }, - { - stage("SCM Checkout") { - try { - robustScmCheckout() - } catch (e) { - error "[SCM] Checkout failed on ${env.NODE_NAME}: ${e}" - } - } - try { - String args = '' - def img = null - stage("Prepare Docker environment") { - // Fill in the docker args from the node - dockerArgs() - - args = DOCKER_ARGS_BY_NODE[env.NODE_NAME] - // Check these args - echo "Running ${CODEPATH} on ${env.NODE_NAME} with: ${args}" - explicitDockerLogin() - img = pullDockerImage(dockerImage()) - } - // Spin up ONE container and stay in it for all substages - img.inside(args) { - // The only way the env variables worked with all other changes - withEnv([ - "HOME=${env.WORKSPACE}" - ]) { - stage("Prepare Performance Scripts") { - echo "codepath is ${CODEPATH}" - echo "Container environment:" - showEnv() - setHeartbeat() - buildProject('check-rocmlir-build-only ci-performance-scripts', '') - } - - stage("Parameter Sweep") { - parameterSweep("conv_structure") - parameterSweep("perf_config") - parameterSweep(CODEPATH, "attention") - parameterSweep(CODEPATH, "gemm_gemm") - archiveArtifacts artifacts: 'build/failing_attn_configs.txt,build/failing_gemmgemm_configs.txt,build/failing_conv_configs.txt', allowEmptyArchive: true - } - } - } - } - // Post block in scripted Jenkins works as try {} catch {}. - // Workspace cleanup is centralized in withHealthyNode so the per-row log survives. - catch (e) { - throw e - } - } - ) - } + script { ciLogic.runParameterSweepsMatrixRow(CODEPATH) } } } } @@ -1751,148 +225,10 @@ pipeline { stages { stage('Matrix row orchestration') { when { - expression { shouldRunFromChip(CHIP) } + expression { ciLogic.shouldRunFromChip(CHIP) } } steps { - script { - // Prepare node - withHealthyNode( - getLabelFromChip(CHIP), - { - checkNodeHealth() - }, - { - stage("SCM Checkout") { - try { - robustScmCheckout() - } catch (e) { - error "[SCM] Checkout failed on ${env.NODE_NAME}: ${e}" - } - } - try { - String args = '' - def img = null - stage("Prepare Docker environment") { - // Fill in the docker args from the node - dockerArgs() - - args = DOCKER_ARGS_BY_NODE[env.NODE_NAME] - // Check these args - echo "Running ${CHIP} on ${env.NODE_NAME} with: ${args}" - explicitDockerLogin() - img = pullDockerImage(dockerImage()) - } - // Spin up ONE container and stay in it for all substages - img.inside(args) { - // The only way the env variables worked with all other changes - withEnv([ - "HOME=${env.WORKSPACE}", - "PATH=/opt/rocm/llvm/bin:${env.PATH}" - ]) { - if (CHIP == "gfx90a") { - stage("Set System Property on Lockhart nodes") { - showEnv() - setHeartbeat() - } - } - - stage("Tune rocMLIR") { - buildProject('check-rocmlir-build-only ci-performance-scripts', '') - dir('build') { - def tuningLog = "tune_rocmlir_${CHIP}.log" - shStrict """echo "=== Tuning rocMLIR for ${CHIP} ===" | tee ${tuningLog}""" - // Tune gemms with default datatypes, fail if the tuning DB is not created - // (Includes int8xint8->int8 for performance comparisons against CK.) - shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ - --op gemm \ - -c ../mlir/utils/performance/configs/tier1-gemm-configs \ - -o mlir_tuning_${CHIP}.tsv 2>&1 | tee -a ${tuningLog} - [ -f mlir_tuning_${CHIP}.tsv ]""" - // Tune resnet50 and unet configs - shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ - --op conv \ - -c ../mlir/utils/performance/configs/tier1-conv-configs \ - -o mlir_tuning_${CHIP}.tsv 2>&1 | tee -a ${tuningLog}""" - // Tune attention configs - shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ - --op attention \ - -c ../mlir/utils/performance/configs/tier1-attention-configs \ - -o mlir_tuning_${CHIP}.tsv 2>&1 | tee -a ${tuningLog}""" - // Tune gemm_gemm configs - shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ - --op gemm_gemm \ - -c ../mlir/utils/performance/configs/tier1-gemmgemm-configs \ - -o mlir_tuning_${CHIP}.tsv 2>&1 | tee -a ${tuningLog}""" - // Quick tuning - shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ - --op gemm --tuning-space quick \ - -c ../mlir/utils/performance/configs/tier1-gemm-configs \ - -o mlir_quick_tuning_${CHIP}.tsv 2>&1 | tee -a ${tuningLog} - [ -f mlir_quick_tuning_${CHIP}.tsv ]""" - shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ - --op conv --tuning-space quick \ - -c ../mlir/utils/performance/configs/tier1-conv-configs \ - -o mlir_quick_tuning_${CHIP}.tsv 2>&1 | tee -a ${tuningLog}""" - shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ - --op attention --tuning-space quick \ - -c ../mlir/utils/performance/configs/tier1-attention-configs \ - -o mlir_quick_tuning_${CHIP}.tsv 2>&1 | tee -a ${tuningLog}""" - shStrict """echo "=== Tuning rocMLIR for ${CHIP} completed ===" | tee -a ${tuningLog}""" - // Check for errors in the tuning log - script { - def tuneLog = readFile(tuningLog).split('\n') - // Find errors that are not part of a warning line - def errors = tuneLog.findAll { it =~ /(?i)error/ && !(it =~ /(?i)\bWARNING\b.*error/) } - - if (errors) { - currentBuild.result = 'FAILURE' - echo "Detected ${errors.size()} error(s) in tuning log:" - errors.each { echo "ERROR LINE: ${it}" } - error("Tuning failed: Detected errors in tuning log") - } else { - echo "No errors found in tuning log" - } - } - } - } - - stage("Tune Fusion") { - dir('build') { - // Tune resnet50 - shStrict """python3 ./bin/tuningRunner.py --abort-on-error --op fusion --test-dir ../mlir/test/fusion/resnet50-e2e/ -o tuning_fusion_${CHIP}.tsv""" - - // Tune bert - shStrict """python3 ./bin/tuningRunner.py --abort-on-error --op fusion --test-dir ../mlir/test/xmir/bert-torch-tosa-e2e/ -o tuning_fusion_${CHIP}.tsv""" - } - sh 'rm -f build/CMakeCache.txt' - } - - stage("Stash Databases") { - // Save user database for nightly jobs - dir ('build') { - stash name: "MLIR-PerfDB-${params.canXdlops ? CHIP : 'vanilla'}", includes: "*.tsv" - } - } - } - } - } - // Post block in scripted Jenkins works as try {} catch {} finally {} - catch (e) { - throw e - } - finally { - // Publish per-arch tuning DBs as soon as this branch finishes so artifacts published without waiting for other parallel branches - try { - archiveArtifacts artifacts: "build/*.tsv", - allowEmptyArchive: true, onlyIfSuccessful: false - } catch (Exception archiveErr) { - echo "[CI] archiveArtifacts of tuning DBs failed: ${archiveErr.message}" - } - // Workspace cleanup is centralized in withHealthyNode so the per-row log survives. - } - } - ) - } + script { ciLogic.runTuneMatrixRow(CHIP) } } } } @@ -1913,7 +249,9 @@ pipeline { skipDefaultCheckout() } steps { - archivePerfDB() + script { + reportUtils.archivePerfDB() + } } post { always { @@ -1937,225 +275,10 @@ pipeline { stages { stage('Matrix row orchestration') { when { - expression { shouldRunFromChip(CHIP) } + expression { ciLogic.shouldRunFromChip(CHIP) } } steps { - script { - // Prepare node - withHealthyNode( - getLabelFromChip(CHIP), - { - checkNodeHealth() - }, - { - stage("SCM Checkout") { - try { - robustScmCheckout() - } catch (e) { - error "[SCM] Checkout failed on ${env.NODE_NAME}: ${e}" - } - } - try { - String args = '' - def img = null - stage("Prepare Docker environment") { - // Fill in the docker args from the node - dockerArgs() - - args = DOCKER_ARGS_BY_NODE[env.NODE_NAME] - // Check these args - echo "Running ${CHIP} on ${env.NODE_NAME} with: ${args}" - explicitDockerLogin() - img = pullDockerImage(dockerImage()) - } - // Spin up ONE container and stay in it for all substages - img.inside(args) { - // The only way the env variables worked with all other changes - withEnv([ - "HOME=${env.WORKSPACE}", - "PATH=/opt/rocm/llvm/bin:${env.PATH}" - ]) { - stage("Copy tuning database") { - echo "chip is ${CHIP}" - echo "Container environment:" - showEnv() - copyArtifacts filter: 'build/perfDB/**',\ - optional: true,\ - flatten: true,\ - projectName: "/MLIR/mlir-weekly",\ - selector: lastSuccessful(),\ - target: 'build' - sh 'ls build' - sh 'cat build/tuning-date' - } - - stage("Build MLIR") { - // Clean up build settings to disable static library and allow ROCm testing - buildProject( - 'check-rocmlir-build-only ci-performance-scripts hipblaslt-benchmark-driver', - '-DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ ' + - '-DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang ' + - '-DROCMLIR_ENABLE_BENCHMARKS=hipblaslt' - ) - } - - stage("Copy earlier performance results") { - copyArtifacts filter: 'build/*.csv,build/perf-run-date',\ - optional: true,\ - flatten: true,\ - projectName: "/${JOB_NAME}",\ - selector: lastSuccessful(),\ - target: 'build/oldData' - } - - stage("Test MLIR vs MIOpen/hipBLASLt") { - dir('build') { - def convInput = "${WORKSPACE}/mlir/utils/performance/configs/tier1-conv-configs" - def convToUse = "${WORKSPACE}/build/tier1-conv-configs" - def gemmInput = "${WORKSPACE}/mlir/utils/performance/configs/tier1-gemm-configs" - def gemmToUse = "${WORKSPACE}/build/tier1-gemm-configs" - script { - if (params.nightly) { - def runIndex = ((env.BUILD_NUMBER as int) - 1) % 5 + 1 - splitConfigFile(convInput, convToUse, runIndex) - splitConfigFile(gemmInput, gemmToUse, runIndex) - } - } - sh 'date --utc +%Y-%m-%d > perf-run-date' - sh 'ls -l /dev/kfd' - sh 'ls -l /dev/dri' - // Run MLIR vs MIOpen perf benchmarks. - shStrict """python3 ./bin/perfRunner.py --op=conv --batch-all \ - --configs-file=${convToUse} \ - --tuning-db=${WORKSPACE}/build/mlir_tuning_${CHIP}.tsv \ - --quick-tuning-db=${WORKSPACE}/build/mlir_quick_tuning_${CHIP}.tsv""" - // Run MLIR vs hipBLASLt perf benchmarks - shStrict """python3 ./bin/perfRunner.py --op=gemm --batch-all \ - --configs-file=${gemmToUse} \ - --tuning-db=${WORKSPACE}/build/mlir_tuning_${CHIP}.tsv \ - --quick-tuning-db=${WORKSPACE}/build/mlir_quick_tuning_${CHIP}.tsv""" - } - } - - stage("Test Fusion") { - dir('build') { - // Run fusion resnet50 perf benchmarks - shStrict """python3 ./bin/perfRunner.py --op=fusion --test-dir=${WORKSPACE}/mlir/test/fusion/resnet50-e2e/ --tuning-db=${WORKSPACE}/build/tuning_fusion_${CHIP}.tsv""" - // Run bert perf benchmarks - shStrict """python3 ./bin/perfRunner.py --op fusion --test-dir=${WORKSPACE}/mlir/test/xmir/bert-torch-tosa-e2e/ --tuning-db=${WORKSPACE}/build/tuning_fusion_${CHIP}.tsv""" - } - } - - if (isNotGfx11x(CHIP)) { - stage("Test Attention") { - dir('build') { - def attnInput = "${WORKSPACE}/mlir/utils/performance/configs/tier1-attention-configs" - def attnToUse = "${WORKSPACE}/build/tier1-attention-configs" - script { - if (params.nightly) { - def runIndex = ((env.BUILD_NUMBER as int) - 1) % 5 + 1 - splitConfigFile(attnInput, attnToUse, runIndex) - } - } - // Run attention benchmarks - shStrict """python3 ./bin/perfRunner.py --op=attention -b \ - --configs-file=${attnToUse} \ - --tuning-db=${WORKSPACE}/build/mlir_tuning_${CHIP}.tsv""" - } - } - - stage("Test Gemm+Gemm") { - dir('build') { - def gemmGemmInput = "${WORKSPACE}/mlir/utils/performance/configs/tier1-gemmgemm-configs" - def gemmGemmToUse = "${WORKSPACE}/build/tier1-gemmgemm-configs" - script { - if (params.nightly) { - def runIndex = ((env.BUILD_NUMBER as int) - 1) % 5 + 1 - splitConfigFile(gemmGemmInput, gemmGemmToUse, runIndex) - } - } - shStrict """python3 ./bin/perfRunner.py --op=gemm_gemm -b \ - --configs-file=${gemmGemmToUse} \ - --tuning-db=${WORKSPACE}/build/mlir_tuning_${CHIP}.tsv""" - } - } - } - - if (params.checkCK && supportsCKBenchmark(CHIP)) { - stage("Test MLIR vs CK") { - catchError (buildResult: null) { // This is an optional stage - def ckInstallDir = "${WORKSPACE}/composable_kernel/build/CKInstallDir" - dir('composable_kernel') { - sh 'rm -rf composable_kernel' - getAndBuildCK(''' - -DGPU_TARGETS=${CHIP} - -DCMAKE_PREFIX_PATH="/opt/rocm" - -DCMAKE_INSTALL_PREFIX=''' + ckInstallDir + ''' - -DCMAKE_BUILD_TYPE=Release - -DBUILD_TESTING=OFF - -DBUILD_CK_EXAMPLES=OFF - -DBUILD_CK_TUTORIALS=OFF - -DBUILD_CK_PROFILER=OFF - -DENABLE_CLANG_CPP_CHECKS=OFF - ''' + ckDtypesCmakeOptions(CHIP) + ''' - ''' + ckFp8CmakeOptions(CHIP) + ''' - ''', - 'device_gemm_operations') - installCKGemmOnly(ckInstallDir) - sh 'echo `git rev-parse HEAD`' - } - sh 'rm -f build/CMakeCache.txt' - buildProject("ck-benchmark-driver", - '''-DCMAKE_PREFIX_PATH=''' + ckInstallDir + ''' - -DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ - -DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang - -DROCMLIR_ENABLE_BENCHMARKS=ck''') - - - dir('build') { - def gemmInput = "${WORKSPACE}/mlir/utils/performance/configs/tier1-gemm-configs" - def gemmToUse = "${WORKSPACE}/build/tier1-gemm-configs" - script { - if (params.nightly) { - def runIndex = ((env.BUILD_NUMBER as int) - 1) % 5 + 1 - splitConfigFile(gemmInput, gemmToUse, runIndex) - } - } - shStrict """python3 ./bin/perfRunner.py --op=gemm --batch-all \ - --configs-file=${gemmToUse} \ - --tuning-db=${WORKSPACE}/build/mlir_tuning_${CHIP}.tsv --data-type f32 f16 i8_i8 --external-gemm-library CK""" - def ckChip = get_gpu_architecture() - sh "python3 ./bin/createPerformanceReports.py ${ckChip} CK" - } - } - } - } - - stage("Create performance reports") { - dir('build') { - sh 'ls -l' - def reportChip = get_gpu_architecture() - echo "Detected GPU chip for reports: ${reportChip} (CHIP matrix value: ${CHIP})" - sh "python3 ./bin/createPerformanceReports.py ${reportChip} MIOpen" - sh "python3 ./bin/createPerformanceReports.py ${reportChip} hipBLASLt" - sh "python3 ./bin/createFusionPerformanceReports.py ${reportChip}" - sh "python3 ./bin/perfRegressionReport.py ${reportChip}" - sh "python3 ./bin/perfRegressionReport.py ${reportChip} ./oldData/${reportChip}_mlir_vs_hipblaslt_perf.csv ./${reportChip}_mlir_vs_hipblaslt_perf.csv" - sh 'mkdir -p reports && cp ./*.html reports' - } - postProcessPerfRes(get_gpu_architecture()) - } - } - } - } - // Post block in scripted Jenkins works as try {} catch {}. - // Workspace cleanup is centralized in withHealthyNode so the per-row log survives. - catch (e) { - throw e - } - } - ) - } + script { ciLogic.runBenchmarkMatrixRow(CHIP) } } } } @@ -2180,113 +303,10 @@ pipeline { stages { stage('Matrix row orchestration') { when { - expression { shouldRunFromCodepath(CODEPATH) } + expression { ciLogic.shouldRunFromCodepath(CODEPATH) } } steps { - script { - // Prepare node - withHealthyNode( - getLabelFromCodepath(CODEPATH), - { - checkNodeHealth() - }, - { - stage("SCM Checkout") { - try { - robustScmCheckout() - } catch (e) { - error "[SCM] Checkout failed on ${env.NODE_NAME}: ${e}" - } - } - try { - String args = '' - def img = null - stage("Prepare Docker environment") { - // Fill in the docker args from the node - dockerArgs() - - args = DOCKER_ARGS_BY_NODE[env.NODE_NAME] - // Check these args - echo "Running ${CODEPATH} on ${env.NODE_NAME} with: ${args}" - - // Explicit docker login since this repo is private - explicitDockerLogin() - img = pullDockerImage(dockerImageCIMIGraphX()) - } - // Spin up ONE container and stay in it for all substages - img.inside(args) { - // The only way the env variables worked with all other changes - withEnv([ - "HOME=${env.WORKSPACE}", - "PYTHONPATH=${env.WORKSPACE}/MIGraphX/build/lib:${env.PYTHONPATH}" - ]) { - stage("Install MIGraphX Dependencies") { - echo "codepath is ${CODEPATH}" - echo "Container environment:" - showEnv() - // Package and install current checkout of rocMLIR as MIGraphX dependency (builds the fat lib, can OOM). - shStrict 'cget -p ${WORKSPACE}/MIGraphXDeps install ${WORKSPACE} -DBUILD_FAT_LIBROCKCOMPILER=On -DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ -DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang' - } - - stage("Build and Verify MIGraphX with MLIR") { - def gpu_arch = get_gpu_architecture() - sh 'rm -rf MIGraphX' - dir('MIGraphX') { - getAndBuildMIGraphX(""" - -DCMAKE_PREFIX_PATH='${WORKSPACE}/MIGraphXDeps;/MIGraphXDeps;/opt/rocm' - -DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ - -DGPU_TARGETS="${gpu_arch}" - """) - } - } - - stage("Verify MIGraphX with MLIR") { - // f32 attention is unsupported on RDNA (no f32 WMMA), and --int8 - // does not quantize attention ops leaving them in f32. Exclude - // attention from MLIR ops on RDNA for int8 - def mlirOps = 'convolution,fused,dot,attention' - def mlirOpsInt8 = (CODEPATH == 'gfx103x' || CODEPATH == 'gfx120x') - ? 'convolution,fused,dot' - : 'convolution,fused,dot,attention' - - dir('MIGraphX/build') { - timeout(time: 120, activity: true, unit: 'MINUTES') { - // run test_verify for accuracy and run MLIR related unit-tests - withEnv(["MIGRAPHX_MLIR_USE_SPECIFIC_OPS=${mlirOps}", 'MIGRAPHX_ENABLE_MLIR_INPUT_FUSION=1', 'MIGRAPHX_ENABLE_MLIR_REDUCE_FUSION=1', 'MIGRAPHX_MLIR_ENABLE_SPLITK=1', 'MIGRAPHX_ENABLE_EXTRA_MLIR=1', 'MIGRAPHX_DISABLE_LAYERNORM_FUSION=1', 'MIGRAPHX_ENABLE_SPLIT_REDUCE=1']) { - shStrict 'make -j$(nproc) test_verify test_gpu_mlir test_gpu_fuse_mlir' - // Verify ResNet50, Bert, Gpt2 with fp16 - shStrict './bin/migraphx-driver verify --gpu --onnx /MIGraphXDeps/resnet50-v1-7.onnx --fp16' - shStrict './bin/migraphx-driver verify --gpu --onnx /MIGraphXDeps/bert_base_cased_1.onnx --fill1 input_ids --input-dim @input_ids 1 384 --fp16' - shStrict './bin/migraphx-driver verify --gpu --onnx /MIGraphXDeps/distilgpt2_1.onnx --fill1 input_ids --input-dim @input_ids 1 384 --fp16' - } - // int8 runs: exclude attention on RDNA to avoid f32 WMMA failure - withEnv(["MIGRAPHX_MLIR_USE_SPECIFIC_OPS=${mlirOpsInt8}", 'MIGRAPHX_ENABLE_MLIR_INPUT_FUSION=1', 'MIGRAPHX_ENABLE_MLIR_REDUCE_FUSION=1', 'MIGRAPHX_MLIR_ENABLE_SPLITK=1', 'MIGRAPHX_ENABLE_EXTRA_MLIR=1', 'MIGRAPHX_DISABLE_LAYERNORM_FUSION=1', 'MIGRAPHX_ENABLE_SPLIT_REDUCE=1']) { - shStrict './bin/migraphx-driver verify --gpu --onnx /MIGraphXDeps/resnet50-v1-7.onnx --int8' - shStrict './bin/migraphx-driver verify --gpu --onnx /MIGraphXDeps/bert_base_cased_1.onnx --fill1 input_ids --input-dim @input_ids 1 384 --int8' - shStrict './bin/migraphx-driver verify --gpu --onnx /MIGraphXDeps/distilgpt2_1.onnx --fill1 input_ids --input-dim @input_ids 1 384 --int8' - } - } - } - //Accuracy_checker will compare outputs from MIGraphX and onnx runtime - dir('MIGraphX/tools/accuracy') { - withEnv(["MIGRAPHX_MLIR_USE_SPECIFIC_OPS=${mlirOpsInt8}", 'MIGRAPHX_ENABLE_MLIR_INPUT_FUSION=1', 'MIGRAPHX_ENABLE_MLIR_REDUCE_FUSION=1', 'MIGRAPHX_MLIR_ENABLE_SPLITK=1', 'MIGRAPHX_ENABLE_EXTRA_MLIR=1', 'MIGRAPHX_DISABLE_LAYERNORM_FUSION=1', 'MIGRAPHX_ENABLE_SPLIT_REDUCE=1']) { - shStrict 'python3 accuracy_checker.py --onnx /MIGraphXDeps/resnet50-v1-7.onnx' - shStrict 'python3 accuracy_checker.py --fill1 --onnx /MIGraphXDeps/bert_base_cased_1.onnx --input-dim input_ids:1,384' - shStrict 'python3 accuracy_checker.py --fill1 --onnx /MIGraphXDeps/distilgpt2_1.onnx --input-dim input_ids:1,384' - } - } - } - } - } - } - // Post block in scripted Jenkins works as try {} catch {}. - // Workspace cleanup is centralized in withHealthyNode so the per-row log survives. - catch (e) { - throw e - } - } - ) - } + script { ciLogic.runMIGraphXMatrixRow(CODEPATH) } } } } @@ -2311,120 +331,7 @@ pipeline { steps { // Do not fail the build on code coverage catchError(buildResult: 'SUCCESS', stageResult: 'FAILURE') { - script { - // Prepare node - withHealthyNode( - getLabelFromCodepath(CODEPATH), - { - checkNodeHealth() - }, - { - stage("SCM Checkout") { - try { - robustScmCheckout() - } catch (e) { - error "[SCM] Checkout failed on ${env.NODE_NAME}: ${e}" - } - } - try { - String args = '' - def img = null - stage("Prepare Docker environment") { - // Fill in the docker args from the node - dockerArgs() - - args = DOCKER_ARGS_BY_NODE[env.NODE_NAME] - // Check these args - echo "Running ${CODEPATH} on ${env.NODE_NAME} with: ${args}" - explicitDockerLogin() - img = pullDockerImage(dockerImage()) - } - // Spin up ONE container and stay in it for all substages - img.inside(args) { - // The only way the env variables worked with all other changes - withEnv([ - "HOME=${env.WORKSPACE}", - "PYTHONPATH=${env.WORKSPACE}/MIGraphX/build/lib:${env.PYTHONPATH}", - // Note the %m to avoid issues with threads and dynamic libraries. - "LLVM_PROFILE_FILE=${env.WORKSPACE}/build/%m-%p.profraw", - "LLVM_PROFDATA=/opt/rocm/llvm/bin/llvm-profdata", - "LLVM_COV=/opt/rocm/llvm/bin/llvm-cov" - ]) { - stage ("body") { - echo "Container environment:" - showEnv() - // Build with profiling on, and just code-generation tests. - try { - // Wall-clock timeout (no `activity:`): the coverage stage has long - // silent phases (lit buffers progress; `llvm-profdata merge` on - // ~125 GB of *.profraw produces no output for several minutes), - // which makes activity-based timeouts fire spuriously and the - // Codecov upload never run. 180 min covers ~60 min tests + - // profdata merge + three llvm-cov calls + upload with margin. - timeout(time: 180, unit: 'MINUTES') { - sh 'rm -f build/CMakeCache.txt' - sh 'rm -f build/*.profraw' - buildProject('check-rocmlir-build-only', - '-DBUILD_FAT_LIBROCKCOMPILER=ON -DCMAKE_BUILD_TYPE=debug -DLLVM_BUILD_INSTRUMENTED_COVERAGE=ON') - dir ('build') { - // Run tests. - collectCoverageData("${LLVM_PROFDATA}", "${LLVM_COV}", "${CODEPATH}") - // Interpolate ${CODEPATH} via Groovy at the call site (bare CODEPATH is empty inside nested closures); shell-side vars are escaped with \$. - withCredentials([string(credentialsId: params.codecovCredentialsId ?: 'codecov-token-rocmlir', - variable: 'CODECOV_TOKEN')]) { - def uploadStatus = sh(script: """ - curl -Os https://uploader.codecov.io/latest/linux/codecov && chmod +x ./codecov - proxy_opt="" - if [ -n "\${http_proxy}" ]; then - proxy_opt="-U \${http_proxy}" - fi - set +e - ./codecov -t "\${CODECOV_TOKEN}" --flags "${CODEPATH}" -f ./coverage_${CODEPATH}.lcov \${proxy_opt} -Z - codecov_exit=\$? - set -e - echo "Codecov upload exit code: \${codecov_exit}" - exit \${codecov_exit} - """, returnStatus: true) - if (uploadStatus != 0) { - // Yellow stage but green build: don't block PR merge on Codecov hiccups. - catchError(buildResult: 'SUCCESS', stageResult: 'UNSTABLE', - message: "Codecov upload failed (exit ${uploadStatus})") { - error("Codecov upload failed (exit code ${uploadStatus}). Check that credential '${params.codecovCredentialsId ?: 'codecov-token-rocmlir'}' contains a valid token from codecov.io and that the repo is linked.") - } - } - } - // Best-effort HTML coverage report: runs AFTER the Codecov upload so a slow/timed-out llvm-cov show cannot prevent the LCOV upload, and is bounded by its own timeout so it cannot block archiveArtifacts. Gated by the runCoverageHtml parameter for builds that only need the Codecov upload. - if (params.runCoverageHtml) { - catchError(buildResult: 'SUCCESS', stageResult: 'UNSTABLE', message: 'Skipped HTML coverage report (llvm-cov show was slow or failed)') { - timeout(time: 45, unit: 'MINUTES') { - produceCoverageHtml("${LLVM_COV}", "${CODEPATH}") - } - } - } - } - } - } catch (Exception e) { - // Yellow stage but green build, same rationale as above. - catchError(buildResult: 'SUCCESS', stageResult: 'UNSTABLE', - message: 'Code coverage stage had an error or timeout') { - error("Code coverage stage had an error or timeout: ${e}") - } - } finally { - // Always archive whatever was produced, even on UNSTABLE / timeout / exception. - archiveArtifacts artifacts: 'build/coverage*.report, build/coverage*.lcov, build/coverage*.html', allowEmptyArchive: true - } - } - } - } - } - // Post block in scripted Jenkins works as try {} catch {}. - // Workspace cleanup is centralized in withHealthyNode so the per-row log survives. - catch (e) { - throw e - } - } - ) - } + script { ciLogic.runCodeCoverageMatrixRow(CODEPATH) } } } } @@ -2434,95 +341,7 @@ pipeline { } post { always { - script { - def result = currentBuild?.currentResult ?: 'UNKNOWN' - def statusMessage = 'Unknown' - def color = 'warning' - if (result == 'SUCCESS') { - statusMessage = "Build ${env.BUILD_NUMBER ?: '?'} completed successfully" - color = 'good' // green - } else if (result == 'FAILURE') { - statusMessage = "Build ${env.BUILD_NUMBER ?: '?'} failed" - color = 'attention' // red - } else if (result == 'ABORTED') { - statusMessage = "Build ${env.BUILD_NUMBER ?: '?'} was aborted" - color = 'warning' // yellow - } else { - statusMessage = "Build ${env.BUILD_NUMBER ?: '?'} ${result}" - color = 'warning' - } - def buildNum = env.BUILD_NUMBER ?: '?' - def buildUrl = env.BUILD_URL ?: '' - def jobName = env.JOB_NAME ?: '' - def isOfficialNightlyOrWeekly = (jobName == 'MLIR/mlir-nightly-all' || jobName == 'MLIR/mlir-weekly') - def failureDetails = null - def logText = '' - if ((result == 'FAILURE' || result == 'ABORTED') && (params.nightly || params.weekly) && isOfficialNightlyOrWeekly) { - try { - // Use large limit so early failures (e.g. SCM checkout) are included; getLog(N) may return last N lines on some setups. - def logLines = currentBuild.rawBuild.getLog(100000) - logText = logLines.join('\n') - failureDetails = classifyBuildFailure(logText) - } catch (e) { - echo "Could not classify failure: ${e}" - } - } - if (result == 'ABORTED') { - if (failureDetails == null) failureDetails = [:] - failureDetails.abortedBy = parseAbortedByFromLog(logText) ?: 'β€”' - } - // For FAILURE/ABORTED, ensure we always have a details map so the card shows Stage/CODEPATH/Details (with fallback if classification failed). - if ((result == 'FAILURE' || result == 'ABORTED') && failureDetails == null) { - failureDetails = [reason: 'Could not match a known error pattern. See build log for details.', codepath: '', stage: ''] - } - if (failureDetails != null && !failureDetails.reason) { - failureDetails.reason = 'Could not match a known error pattern. See build log for details.' - } - // Auto re-kick transient failures (official nightly/weekly only). - def rekickAttempt = (params.rekickAttempt?.toString()?.isInteger()) ? Math.max(0, params.rekickAttempt as int) : 0 - def willRekick = (result == 'FAILURE') && params.rekickEnabled && isOfficialNightlyOrWeekly && (params.nightly || params.weekly) && - (rekickAttempt < MAX_REKICK_ATTEMPTS) && isTransientHardwareFailure(logText) - def transientReason = (failureDetails?.reason) ?: 'transient failure' - - // Schedule the re-kick before notifying so we can suppress this run's card on success. - boolean rekickScheduled = false - if (willRekick) { - try { - echo "[re-kick] Transient failure on ${jobName} #${buildNum}; re-kicking (attempt ${rekickAttempt + 1}/${MAX_REKICK_ATTEMPTS}), suppressing this run's notification." - // Absolute path: a relative job name resolves against the current job's folder (MLIR/). - build job: "/${jobName}", wait: false, propagate: false, parameters: rekickParameters(rekickAttempt + 1, transientReason) - rekickScheduled = true - } catch (e) { - echo "[re-kick] Could not schedule re-kick: ${e}" - } - } - - // Exactly one card per logical build: skip it only when we successfully handed off to a re-kick. - if (!(willRekick && rekickScheduled) && (params.nightly || params.weekly) && isOfficialNightlyOrWeekly && buildUrl && jobName) { - if (willRekick && !rekickScheduled) { - if (failureDetails == null) failureDetails = [:] - failureDetails.rekick = "Transient failure detected, but the re-kick could not be scheduled β€” see build log." - } else if (rekickAttempt > 0) { - if (failureDetails == null) failureDetails = [:] - failureDetails.rekick = (result == 'SUCCESS') - ? "Auto re-kicked ${rekickAttempt}Γ— after a transient failure (${params.rekickReason ?: 'β€”'}); passed on retry." - : "Auto re-kicked ${rekickAttempt}Γ— after transient failure(s) (last: ${params.rekickReason ?: 'β€”'}); this run is not being re-kicked." - } - def runType = params.nightly ? 'nightly' : 'weekly' - def jenkinsBase = buildUrl.replaceFirst('/job/.*', '') - def jobNameEncoded = jobName.replace('/', '%2F') - def jobShortName = jobName.contains('/') ? jobName.split('/').last() : jobName - def blueOceanUrl = "${jenkinsBase}/blue/organizations/jenkins/${jobNameEncoded}/detail/${jobShortName}/${buildNum}/pipeline" - def jobUrl = buildUrl - try { - node('build-only') { - sendTeamsBuildNotification(buildNum, statusMessage, color, runType, blueOceanUrl, jobUrl, failureDetails) - } - } catch (e) { - echo "Could not send Teams notification: ${e}" - } - } - } + script { ciLogic.handlePostBuildNotification() } } } } diff --git a/mlir/utils/jenkins/helpers/buildUtils.groovy b/mlir/utils/jenkins/helpers/buildUtils.groovy new file mode 100644 index 000000000000..a4c0087de2d4 --- /dev/null +++ b/mlir/utils/jenkins/helpers/buildUtils.groovy @@ -0,0 +1,145 @@ +// Build helpers: rocMLIR build, plus building/checking out CK and MIGraphX. +// Loaded by Jenkinsfile's Bootstrap stage; consumed as buildUtils.(). +// ON CHANGING THESE, ALSO CHANGE Jenkinsfile.downstream + +import groovy.transform.Field + +// Cross-helper handle, populated by Jenkinsfile's Bootstrap stage: +// buildUtils.scmUtils = scmUtils +@Field def scmUtils + +// Run `script` through bash with errexit + pipefail. Use this whenever a +// command pipes through tee/awk/grep/etc. so failures in the upstream command +// are not masked by the pipeline's last exit code. Plain `sh` runs under +// /bin/sh -xe (errexit but no pipefail); a #!/bin/bash shebang bypasses +// Jenkins's default flags, so we re-enable both explicitly here. +def shStrict(String script) { + // When running inside withHealthyNode (REKICK_ROW_LOG set), mirror this step's output to a + // per-row log so the retry handler can classify transient failures (e.g. GPU hang) that only + // appear in stdout. pipefail keeps the real command's exit code from being masked by tee. + if (env.REKICK_ROW_LOG) { + sh "#!/bin/bash\nset -eo pipefail\n{\n${script}\n} 2>&1 | tee -a \"${env.REKICK_ROW_LOG}\"" + } else { + sh "#!/bin/bash\nset -eo pipefail\n${script}" + } +} + +void buildProject(String target, String cmakeOpts) { + timeout(time: 60, activity: true, unit: 'MINUTES') { + // Configure with the CMake plugin (unchanged: same source/build dir resolution as before). + cmakeBuild generator: 'Ninja',\ + buildDir: 'build',\ + buildType: 'RelWithDebInfo',\ + installation: 'InSearchPath',\ + cmakeArgs: """-DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ + -DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang + ${cmakeOpts}""" + // Build via shStrict (was the plugin's `steps: [[args: target]]`, i.e. `ninja ` in + // build dir) so build output is mirrored to the per-row log and build-time OOM + // (ninja exit 137) can be classified as a per-server transient in withHealthyNode. + // `ninja -C ` is CWD-independent, matching the plugin's workspace-relative build dir. + shStrict "ninja -C ${env.WORKSPACE}/build ${target}" + } +} + +void buildCK(String cmakeOpts, String buildTarget = '') { + sh '[ ! -d build ] || rm -rf build' + // CK's Unix Makefiles do not expose device_gemm_operations as a directly + // buildable target, while Ninja handles this target graph correctly. + cmakeBuild generator: (buildTarget ? 'Ninja' : 'Unix Makefiles'),\ + buildDir: 'build',\ + buildType: 'Release',\ + installation: 'InSearchPath',\ + cmakeArgs: """-DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ + -DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang + ${cmakeOpts} + """ + if (buildTarget) { + shStrict "cmake --build build --target ${buildTarget} --parallel \$(nproc)" + } else { + shStrict 'cd build; make -j $(nproc)' + } +} + +void installCKGemmOnly(String installDir) { + sh """#!/usr/bin/env bash + set -euo pipefail + + install_dir="${installDir}" + cmake_dir="\${install_dir}/lib/cmake/composable_kernel" + mkdir -p "\${install_dir}/include/ck" "\${install_dir}/lib" "\${cmake_dir}" + + cp -R include/ck/. "\${install_dir}/include/ck/" + cp -R library/include/ck/. "\${install_dir}/include/ck/" + cp build/include/ck/config.h build/include/ck/version.h "\${install_dir}/include/ck/" + cp build/lib/libdevice_gemm_operations.a "\${install_dir}/lib/" + cp build/composable_kernelConfig.cmake \ + build/composable_kernelConfigVersion.cmake \ + "\${cmake_dir}/" + + mapfile -t gemm_export_files < <(find build -type f -name 'composable_kerneldevice_gemm_operationsTargets*.cmake' -print) + if [ "\${#gemm_export_files[@]}" -eq 0 ]; then + echo "Could not find CK device_gemm_operations CMake export files" + exit 1 + fi + cp "\${gemm_export_files[@]}" "\${cmake_dir}/" + """ +} + +void buildMIGraphX(String cmakeOpts) { + sh '[ ! -d build ] || rm -rf build' + cmakeBuild generator: 'Unix Makefiles',\ + buildDir: 'build',\ + buildType: 'Release',\ + installation: 'InSearchPath',\ + cmakeArgs: """-DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ + -DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang + -DMIGRAPHX_USE_COMPOSABLEKERNEL=OFF + ${cmakeOpts} + """ + shStrict 'cd build; make -j $(nproc)' +} + +void getAndBuildMIGraphX(String cmakeOpts) { + scmUtils.robustExternalCheckout('https://github.com/ROCm/AMDMIGraphX.git', params.MIGraphXBranch) + buildMIGraphX(cmakeOpts) +} + +void getAndBuildCK(String cmakeOpts, String buildTarget = '') { + scmUtils.robustExternalCheckout('https://github.com/ROCm/composable_kernel.git', params.CKBranch) + buildCK(cmakeOpts, buildTarget) +} + +String ckFp8CmakeOptions(String chip) { + // CK auto-derives CK_USE_*_FP8 from GPU_TARGETS, so we only force the + // numeric macro values via CMAKE_CXX_FLAGS to keep CK's `#if` checks correct. + if ("${chip}".startsWith('gfx94')) { + return ''' + -DCMAKE_CXX_FLAGS="-O3 -DCK_USE_FNUZ_FP8=1" + ''' + } + if ("${chip}" == 'gfx950' || "${chip}".startsWith('gfx12')) { + return ''' + -DCMAKE_CXX_FLAGS="-O3 -DCK_USE_OCP_FP8=1 -DCK_TILE_USE_OCP_FP8=1" + ''' + } + return ''' + -DCMAKE_CXX_FLAGS="-O3" + ''' +} + +String ckDtypesCmakeOptions(String chip) { + // The MLIR vs CK perf configs exercise f16/f32 GEMM, with int8 available + // for the CK GEMM driver. Restricting DTYPES avoids building unused CK + // operation instances on older arches such as gfx908. + if ("${chip}".startsWith('gfx94') || "${chip}" == 'gfx950' || "${chip}".startsWith('gfx12')) { + return ''' + -DDTYPES="fp8;bf8;fp16;fp32;int8" + ''' + } + return ''' + -DDTYPES="fp16;fp32;int8" + ''' +} + +return this diff --git a/mlir/utils/jenkins/helpers/ciLogic.groovy b/mlir/utils/jenkins/helpers/ciLogic.groovy new file mode 100644 index 000000000000..3c9cf3548cf1 --- /dev/null +++ b/mlir/utils/jenkins/helpers/ciLogic.groovy @@ -0,0 +1,1404 @@ +// CI flow helpers: heartbeat, build resets, label resolution, codepath/chip +// gating, config-file splitting, build-failure classification, the Teams +// notification card, and the per-stage matrix-row orchestrators. +// Loaded by Jenkinsfile's Bootstrap stage; consumed as ciLogic.(). +// ON CHANGING THESE, ALSO CHANGE Jenkinsfile.downstream + +import groovy.transform.Field + +// Cross-helper handles, populated by Jenkinsfile's Bootstrap stage: +// ciLogic.scmUtils = scmUtils +// ciLogic.nodeUtils = nodeUtils +// ciLogic.buildUtils = buildUtils +// ciLogic.testUtils = testUtils +// ciLogic.reportUtils = reportUtils +// Used by the matrix-row orchestrators (runBuildAndTestMatrixRow, etc.) which +// were extracted from the Jenkinsfile to keep the main pipeline body within +// the JVM's 64KB CPS bytecode limit. +@Field def scmUtils +@Field def nodeUtils +@Field def buildUtils +@Field def testUtils +@Field def reportUtils + +// Max automatic re-kicks for a nightly/weekly build failing for transient reasons. +@Field final int MAX_REKICK_ATTEMPTS = 2 +// Characters from the end of the log scanned to classify the failure; the decisive cause sits at the end. +@Field final int FAILURE_LOG_TAIL_CHARS = 1000000 + +//makes sure multiple builds are not triggered for branch indexing +def resetBuild() { + if (currentBuild.getPreviousBuild() == null + || currentBuild.getPreviousBuild().getBuildCauses().toString().contains('BranchIndexingCause')) { + def buildNumber = BUILD_NUMBER as int; + if (buildNumber > 1) + milestone(buildNumber - 1); + milestone(buildNumber) + } +} + +void setHeartbeat() { + script { + System.setProperty("org.jenkinsci.plugins.durabletask.BourneShellScript.HEARTBEAT_CHECK_INTERVAL", "86400"); + } +} + +String getLabelFromCodepath(String codepath) { + echo "codepath is ${codepath}" + String label = '' + if (codepath == "mfma") { + label = 'mlir && (gfx942 || gfx908 || gfx90a)' + } else if (codepath == "gfx950") { + if (params.weekly) { + label = 'mlir && linux-mi350-8' + } else { + label = 'mlir && linux-mi350-1' + } + } else if (codepath == "gfx103x") { + // For non-performance related testing, use both workstations (gfx1030w) + // and server nodes (gfx1030) + label = 'mlir && ( gfx1030w || gfx1030 )' + } else if (codepath == "vanilla"){ + label = 'mlir' + } else if (codepath == "gfx110x") { + if (params.nightly || params.weekly) { + label = 'mlir && gfx1100' + } else { + label = 'mlir && ( gfx1100 || gfx1101 )' + } + } else if (codepath == "gfx120x") { + if (params.nightly || params.weekly) { + label = 'mlir && gfx1201' + } else { + label = 'mlir && ( gfx1200 || gfx1201 )' + } + } else { + echo "${codepath} is not supported" + label = 'wrongLabel' + } + echo "label is ${label}" + return label +} + +String getLabelFromChip(String chip) { + switch (chip) { + case "gfx906": + return getLabelFromCodepath("vanilla") + case "gfx908": + return "mlir && gfx908" + case "gfx90a": + return "mlir && gfx90a" + case "gfx942": + return "mlir && gfx942" + case "gfx950": + if (params.weekly) { + return "mlir && linux-mi350-8" + } else { + return "mlir && linux-mi350-1" + } + case "gfx1030": + // For [Tune MLIR Kernels] and [Performance report] stages, + // fix the vm-5 workstation for testing + return "mlir && vm-5" + case "gfx1100": + return "mlir && gfx1100" + case "gfx1101": + return "mlir && gfx1101" + case "gfx1200": + return "mlir && gfx1200" + case "gfx1201": + return "mlir && gfx1201" + } +} + +boolean shouldRunFromCodepath(String codepath) { + // Run vanilla on public CI + if ((codepath == "vanilla") && (params.canXdlops == false)) { + return true + } + // Run mfma on private CI + if ((codepath == "mfma") && params.canXdlops) { + return true + } + if (codepath == "gfx950" && params.canXdlops && params.disable950 == false) { + return true + } + // Run gfx103x on private nightly or weekly CI if it is not disabled + if (params.canXdlops && (params.disableGfx103x == false) && (codepath == "gfx103x") && + (params.nightly || params.weekly)) { + return true + } + // Run gfx110x on private CI if it is not disabled + if (params.canXdlops && (params.disableGfx110x == false) && (codepath == "gfx110x")) { + return true + } + // Run gfx120x on private CI if it is not disabled + if (params.canXdlops && (params.disableGfx120x == false) && (codepath == "gfx120x")) { + return true + } + return false +} + +boolean shouldRunFromChip(String chip) { + switch (chip) { + default: + return shouldRunFromCodepath("vanilla") + case "gfx90a": + // Special case because all our "vanilla" hosts are gfx90a. + return params.disable90a == false && + (shouldRunFromCodepath("mfma") || shouldRunFromCodepath("vanilla")) + case "gfx908": + return params.disable908 == false && shouldRunFromCodepath("mfma") + case "gfx942": + return params.disable942 == false && shouldRunFromCodepath("mfma") + case "gfx950": + return params.disable950 == false && shouldRunFromCodepath("gfx950") + case "gfx1030": + return shouldRunFromCodepath("gfx103x") + case "gfx1100": + return shouldRunFromCodepath("gfx110x") + case "gfx1200": + case "gfx1201": + return shouldRunFromCodepath("gfx120x") + } +} + +boolean shouldRunBuildAndTest(String codepath) { + // When default codepath is selected, we test mfma, gfx103x, gfx110x and gfx120x on + // private CI and vanilla on public CI + if (params.codepath == "default" && shouldRunFromCodepath(codepath)) + return true + + // When a particular codepath is selected, we only test the codepath + // on private CI + if (params.codepath == codepath && params.canXdlops) { + if (params.codepath == "mfma") return true + if (params.codepath == "vanilla") return true + if (params.codepath == "gfx950" && params.disable950 == false) return true + if (params.codepath == "gfx103x" && params.disableGfx103x == false) return true + if (params.codepath == "gfx110x" && params.disableGfx110x == false) return true + if (params.codepath == "gfx120x" && params.disableGfx120x == false) return true + return false + } +} + +boolean isNotGfx11x(String chip) { + return "${chip}" != 'gfx1100' && "${chip}" != 'gfx1101' +} + +boolean supportsCKBenchmark(String chip) { + // CK does not generate a device_gemm_operations target for gfx906. + return isNotGfx11x(chip) && "${chip}" != 'gfx906' +} + +void splitConfigFile(String inputFilePath, String outputFilePath, int run, int totalSplits = 5) { + buildUtils.shStrict """ + lines=\$(grep -Ev '(^\\s*\$|^\\s*#)' ${inputFilePath} | wc -l) + lines_per_chunk=\$(((lines + ${totalSplits} - 1) / ${totalSplits})) + start_line=\$((lines_per_chunk * (${run} - 1) + 1)) + end_line=\$((lines_per_chunk * ${run})) + + grep -Ev '(^\\s*\$|^\\s*#)' ${inputFilePath} | sed -n "\${start_line},\${end_line}p" | tee ${outputFilePath} + """ +} + +// Classifies build failure from console log. Returns [reason:, codepath:, stage:] (empty string = not found). +// Add new scenarios here by matching log patterns (order = first match wins). +Map classifyBuildFailure(String logText) { + def reason = '' + def codepath = '' + def stage = '' + def failureList = '' + def failureListLabel = '' + def failedTestsSnippet = '' + if (!logText) return [reason: reason, codepath: codepath, stage: stage, failureList: failureList, failureListLabel: failureListLabel, failedTestsSnippet: failedTestsSnippet] + + // Scenario 1: Tuning failed - errors detected in tuning log (Tune rocMLIR) + if (!reason && logText.contains('Tuning failed: Detected errors in tuning log')) { + reason = 'Tune rocMLIR: errors in tuning log (check logs for details)' + } + + // Scenario 2: SCM checkout failed (max retries, clone error, or channel error) + if (!reason && (logText.contains('ERROR: Checkout failed') || logText.contains('Maximum checkout retry attempts reached') || logText.contains("ERROR: Error cloning remote repo"))) { + reason = 'SCM checkout failed (max retries or agent/channel error)' + } + + // Scenario 3: Parameter sweeps - failing configurations discovered. + // 3a: Conv/perf sweeps (parameterSweeps.py) use "*** Summary of failures ***". + if (!reason && logText.contains('*** Summary of failures ***')) { + reason = 'Parameter sweeps: failing configurations discovered' + def summaryStart = logText.indexOf('*** Summary of failures ***') + def summaryEnd = logText.indexOf('Passed:', summaryStart) + if (summaryEnd < 0) summaryEnd = logText.indexOf('script returned exit code', summaryStart) + if (summaryEnd < 0) summaryEnd = logText.length() + failureList = logText.substring(summaryStart, summaryEnd).trim() + if (failureList.length() > 2000) failureList = failureList.substring(0, 2000) + '\n... (truncated)' + } + // 3b: Attention sweeps (attentionSweeps.py) use "Failing Configurations". + if (!reason && logText.contains('Failing Configurations')) { + reason = 'Attention parameter sweeps: failing configurations discovered' + def headerPos = logText.lastIndexOf('Failing Configurations') + def configStart = logText.indexOf('\n', headerPos) + configStart = (configStart >= 0) ? configStart + 1 : headerPos + def summaryEnd = logText.indexOf('Passed:', configStart) + if (summaryEnd < 0) summaryEnd = logText.indexOf('script returned exit code', configStart) + if (summaryEnd < 0) summaryEnd = Math.min(configStart + 3000, logText.length()) + def snippet = logText.substring(configStart, summaryEnd).trim() + snippet = snippet.replaceAll(/\[\d{4}-\d{2}-\d{2}T[\d:.]+Z\]\s*/, '') + // Append the Passed/Invalid/Failed summary line if present. + def statsLineEnd = logText.indexOf('\n', summaryEnd) + if (statsLineEnd < 0) statsLineEnd = logText.length() + def statsLine = logText.substring(summaryEnd, statsLineEnd).trim() + .replaceAll(/\[\d{4}-\d{2}-\d{2}T[\d:.]+Z\]\s*/, '') + if (statsLine) snippet = snippet + '\n\n' + statsLine + if (snippet.length() > 2000) snippet = snippet.substring(0, 2000) + '\n... (truncated)' + failureList = snippet + } + + // Scenario 4: HIP no device (hipErrorNoDevice) + if (!reason && logText.contains('RuntimeError: hipError_t.hipErrorNoDevice')) { + reason = 'HIP: no device (hipErrorNoDevice)' + } + + // Scenario 5: One or more tests failed (Failed Tests (N): ...) + if (!reason && logText.contains('Failed Tests (')) { + reason = 'One or more tests failed' + def failedStart = logText.indexOf('Failed Tests (') + def failedEnd = logText.indexOf('Testing Time:', failedStart) + if (failedEnd < 0) failedEnd = logText.indexOf('Total Discovered Tests:', failedStart) + if (failedEnd < 0) failedEnd = Math.min(failedStart + 2000, logText.length()) + failedTestsSnippet = logText.substring(failedStart, failedEnd).trim() + if (failedTestsSnippet.length() > 2000) failedTestsSnippet = failedTestsSnippet.substring(0, 2000) + '\n... (truncated)' + } + + // Scenario 6: MIGraphX CMake configuration failed. + // Match by context around "Configuring incomplete" (MIGraphX path or composable_kernel_host) so we don't rely on stage order in interleaved logs. + def cmakeConfigErrorPos = logText.lastIndexOf('Configuring incomplete, errors occurred!') + if (!reason && cmakeConfigErrorPos >= 0) { + def ctxStart = Math.max(0, cmakeConfigErrorPos - 4000) + def ctxAround = logText.substring(ctxStart, Math.min(logText.length(), cmakeConfigErrorPos + 500)) + if (ctxAround.contains('MIGraphX') || ctxAround.contains('composable_kernel_host') || ctxAround.contains('Findcomposable_kernel_host')) { + reason = 'MIGraphX: CMake configuration failed' + // Extract the last "CMake Error" block before "Configuring incomplete" as a snippet. + def cmakeErrorStart = logText.lastIndexOf('CMake Error', cmakeConfigErrorPos) + if (cmakeErrorStart >= 0) { + def snippet = logText.substring(cmakeErrorStart, cmakeConfigErrorPos).trim() + snippet = snippet.replaceAll(/\[\d{4}-\d{2}-\d{2}T[\d:.]+Z\]\s*/, '') + if (snippet.length() > 2000) snippet = snippet.substring(0, 2000) + '\n... (truncated)' + failureList = snippet + failureListLabel = 'CMake error:' + } + } + } + + // Scenario 7: Agent flapping (node repeatedly offline/online). + // Checked last: agent disconnect messages often appear as a side effect of pod termination after a real build error. + if (!reason) { + def flappingMatch = logText =~ /(\S+)\s+seems to be removed or offline.*will wait for.*come back online/ + if (flappingMatch.find()) { + reason = "Agent flapping: ${flappingMatch.group(1)} went offline/online repeatedly" + } + } + + if (!reason) reason = 'Could not match a known error pattern. See build log for details.' + + // Failure anchor: position in log where this failure was detected (used to extract stage/CODEPATH from the failing branch, not from later branches). + def failureAnchor = -1 + + // Prefer detecting the anchor directly from log patterns instead of the human-facing reason text. + def scmAnchor = Math.max(logText.lastIndexOf('Maximum checkout retry attempts reached'), + logText.lastIndexOf('[SCM] Checkout failed on')) + if (scmAnchor < 0) scmAnchor = logText.lastIndexOf("ERROR: Error cloning remote repo") + if (scmAnchor < 0) scmAnchor = logText.lastIndexOf('ERROR: Checkout failed') + + if (scmAnchor >= 0) { + failureAnchor = scmAnchor + } else { + def tuneAnchor = logText.lastIndexOf('Tuning failed: Detected errors in tuning log') + if (tuneAnchor >= 0) { + failureAnchor = tuneAnchor + } else { + def sweepsAnchor = logText.indexOf('*** Summary of failures ***') + if (sweepsAnchor < 0) sweepsAnchor = logText.indexOf('Failing Configurations') + if (sweepsAnchor >= 0) { + failureAnchor = sweepsAnchor + } else { + def hipNoDeviceAnchor = logText.lastIndexOf('hipErrorNoDevice') + if (hipNoDeviceAnchor >= 0) { + failureAnchor = hipNoDeviceAnchor + } else { + def testsFailedAnchor = logText.lastIndexOf('Failed Tests (') + if (testsFailedAnchor >= 0) { + failureAnchor = testsFailedAnchor + } else { + def migraphxAnchor = logText.lastIndexOf('Configuring incomplete, errors occurred!') + if (migraphxAnchor >= 0) { + failureAnchor = migraphxAnchor + } else { + def agentFlappingAnchor = logText.lastIndexOf('seems to be removed or offline') + if (agentFlappingAnchor >= 0) { + failureAnchor = agentFlappingAnchor + } + } + } + } + } + } + } + + def searchStart = (failureAnchor >= 0) ? Math.max(0, failureAnchor - 8000) : 0 + def searchEnd = (failureAnchor >= 0) ? Math.min(logText.length(), failureAnchor + 500) : logText.length() + def contextWindow = (failureAnchor >= 0) ? logText.substring(searchStart, searchEnd) : logText + def logBeforeAnchor = (failureAnchor > 0) ? logText.substring(0, failureAnchor) : '' + + // CODEPATH: prefer "Failed in branch Matrix - CODEPATH = 'X'" near the failure; else any CODEPATH in context window; else global. + def branchMatch = contextWindow =~ /Failed in branch Matrix - CODEPATH = ['"](\w+)['"]/ + if (branchMatch.find()) { + codepath = branchMatch.group(1) + } else { + def cpMatch = contextWindow =~ /CODEPATH\s*=\s*['"]?(\w+)['"]?|Running\s+(\w+)\s+on\s+\S+/ + if (cpMatch.find()) codepath = cpMatch[0][1] ?: cpMatch[0][2] ?: '' + } + if (!codepath) { + def cpMatch = logText =~ /CODEPATH\s*=\s*['"]?(\w+)['"]?|Running\s+(\w+)\s+on\s+\S+/ + if (cpMatch.find()) codepath = cpMatch[0][1] ?: cpMatch[0][2] ?: '' + } + + // Stage: last stage name that appears *before* the failure anchor (so we report the stage that was running when it failed). + def stageNames = ['SCM Checkout', 'Build and Test', 'Parameter sweeps', 'Parameter Sweep', 'Tune MLIR kernels', 'Tune rocMLIR', 'Code coverage', 'Archive performance DB', 'MIGraphX', 'Build and Verify MIGraphX with MLIR'] + def stageSearchText = (logBeforeAnchor.length() > 0) ? logBeforeAnchor : logText + def stageIdx = -1 + for (def name in stageNames) { + def idx = stageSearchText.lastIndexOf(name) + if (idx >= 0 && idx > stageIdx) { stage = name; stageIdx = idx } + } + + return [reason: reason, codepath: codepath, stage: stage, failureList: failureList, failureListLabel: failureListLabel, failedTestsSnippet: failedTestsSnippet] +} + +// Parse "Aborted by USERNAME" from console log (Jenkins writes this when a user aborts the build). +// True if a FAILED build is transient (re-kickable), not a genuine code/test failure. Scans the +// log tail only, since the decisive cause is at the end. Shares the genuine-failure veto with the +// per-server classifier via nodeUtils.realTestFailureSignals(). +boolean isTransientHardwareFailure(String logText) { + if (!logText) return false + String tail = logText.length() > FAILURE_LOG_TAIL_CHARS + ? logText.substring(logText.length() - FAILURE_LOG_TAIL_CHARS) : logText + tail = tail.toLowerCase() + + // Hardware loss wins even over the veto below: a dead GPU also makes lit report spurious failures. + def hardwareSignals = [ + 'hiperror_t.hiperrornodevice', + 'unable to reset gpu', + 'unsupported hip gpu architecture: n/a', + 'no performance report found for n/a', + 'no healthy node found', + '[withhealthynode] transient', + ] + if (hardwareSignals.any { tail.contains(it) }) return true + + // Genuine logic/test failures: never re-kick. + if (nodeUtils.realTestFailureSignals().any { tail.contains(it) }) return false + if (tail =~ /no performance report found for gfx/) return false + + // Transient infra/connection. After the veto, since failFast and dying agents emit some of + // these as collateral of a real failure. + def abortLikeSignals = [ + 'seems to be removed or offline', + 'agentofflineexception', + 'issue with creating launcher for agent', + 'failed to run image', + 'outofmemoryerror', + 'interruptedexception', + 'ninja exited with error code 137', // OOM + 'closedchannelexception', + 'requestabortedexception', + 'broken pipe', + 'script returned exit code -1', + 'script returned exit code -2', + 'maximum checkout retry attempts reached', + 'error cloning remote repo', + 'error fetching remote repo', + 'gpu hang', + 'hw exception by gpu', + ] + if (abortLikeSignals.any { tail.contains(it) }) return true + return scmUtils.isRetriableScmCheckoutError(tail) +} + +// Forwards all current parameters unchanged, with the incremented re-kick attempt counter. +List rekickParameters(int nextAttempt, String reason) { + return [ + booleanParam(name: 'nightly', value: params.nightly), + booleanParam(name: 'canXdlops', value: params.canXdlops), + booleanParam(name: 'weekly', value: params.weekly), + string(name: 'MIGraphXBranch', value: params.MIGraphXBranch), + string(name: 'CKBranch', value: params.CKBranch), + booleanParam(name: 'sharedLib', value: params.sharedLib), + booleanParam(name: 'staticLib', value: params.staticLib), + booleanParam(name: 'checkMIGraphX', value: params.checkMIGraphX), + booleanParam(name: 'checkCK', value: params.checkCK), + booleanParam(name: 'runCodeCoverage', value: params.runCodeCoverage), + booleanParam(name: 'runCoverageHtml', value: params.runCoverageHtml), + string(name: 'codecovCredentialsId', value: params.codecovCredentialsId), + string(name: 'codepath', value: params.codepath), + booleanParam(name: 'disableGfx103x', value: params.disableGfx103x), + booleanParam(name: 'disableGfx110x', value: params.disableGfx110x), + booleanParam(name: 'disableGfx120x', value: params.disableGfx120x), + booleanParam(name: 'disable90a', value: params.disable90a), + booleanParam(name: 'disable908', value: params.disable908), + booleanParam(name: 'disable942', value: params.disable942), + booleanParam(name: 'disable950', value: params.disable950), + booleanParam(name: 'ignoreExternalLinting', value: params.ignoreExternalLinting), + string(name: 'weeklyTasks', value: params.weeklyTasks), + booleanParam(name: 'rekickEnabled', value: params.rekickEnabled), + string(name: 'rekickAttempt', value: nextAttempt.toString()), + string(name: 'rekickReason', value: reason ?: ''), + ] +} + +def parseAbortedByFromLog(String logText) { + if (!logText) return '' + def m = logText =~ /Aborted by ([^\r\n]+)/ + return m.find() ? m.group(1).trim() : '' +} + +// Sends a Teams adaptive card for build result (webhook URL from Jenkins credential 'CI_MONITORING_TEAMS'). +// statusMessage: full phrase e.g. "Build 42 completed successfully". color: Adaptive Card color ("good"=green, "warning"=yellow, "attention"=red). +// runType: "nightly" or "weekly" (subtitle line). blueOceanUrl: Blue Ocean pipeline URL. jobUrl: classic Jenkins job/build URL. +// failureDetails: optional Map [reason:, codepath:, stage:, failureList:] β€” when set, adds Stage/CODEPATH/Details and optionally a code block for failureList. +void sendTeamsBuildNotification(String buildNumber, String statusMessage, String color, String runType, String blueOceanUrl, String jobUrl, Map failureDetails = null) { + try { + def subtitle = (runType == 'nightly') ? 'MLIR Nightly πŸŒ™' : 'MLIR Weekly πŸ“…' + def timestamp = new Date().format('yyyy-MM-dd HH:mm z') + def escapeJson = { String s -> (s ?: '').replace('\\', '\\\\').replace('"', '\\"').replace('\n', ' ') } + def escapeJsonMultiline = { String s -> (s ?: '').replace('\\', '\\\\').replace('"', '\\"').replace('\n', '\\n').replace('\r', '') } + def detailBlocks = '' + if (failureDetails) { + if (failureDetails.abortedBy != null) { + def ab = escapeJson(failureDetails.abortedBy) + detailBlocks += ",{\"type\":\"RichTextBlock\",\"inlines\":[{\"type\":\"TextRun\",\"text\":\"Aborted by: \",\"weight\":\"bolder\"},{\"type\":\"TextRun\",\"text\":\"${ab}\"}]}" + } + if (failureDetails.reason) { + def r = escapeJson(failureDetails.reason) + def c = failureDetails.codepath ? escapeJson(failureDetails.codepath) : 'β€”' + def t = failureDetails.stage ? escapeJson(failureDetails.stage) : 'β€”' + detailBlocks += ",{\"type\":\"RichTextBlock\",\"inlines\":[{\"type\":\"TextRun\",\"text\":\"Stage: \",\"weight\":\"bolder\"},{\"type\":\"TextRun\",\"text\":\"${t}\"}]},{\"type\":\"RichTextBlock\",\"inlines\":[{\"type\":\"TextRun\",\"text\":\"CODEPATH: \",\"weight\":\"bolder\"},{\"type\":\"TextRun\",\"text\":\"${c}\"}]},{\"type\":\"RichTextBlock\",\"inlines\":[{\"type\":\"TextRun\",\"text\":\"Details: \",\"weight\":\"bolder\"},{\"type\":\"TextRun\",\"text\":\"${r}\"}],\"wrap\":true}" + } + if (failureDetails.failureList) { + def flLabel = failureDetails.failureListLabel ?: 'Failing configs:' + def fl = escapeJsonMultiline(failureDetails.failureList) + detailBlocks += ",{\"type\":\"RichTextBlock\",\"inlines\":[{\"type\":\"TextRun\",\"text\":\"${escapeJson(flLabel)}\",\"weight\":\"bolder\"}]},{\"type\":\"TextBlock\",\"text\":\"${fl}\",\"wrap\":true,\"fontType\":\"monospace\",\"size\":\"small\",\"separator\":true}" + } + if (failureDetails.failedTestsSnippet) { + def fts = escapeJsonMultiline(failureDetails.failedTestsSnippet) + detailBlocks += ",{\"type\":\"RichTextBlock\",\"inlines\":[{\"type\":\"TextRun\",\"text\":\"Failed tests:\",\"weight\":\"bolder\"}]},{\"type\":\"TextBlock\",\"text\":\"${fts}\",\"wrap\":true,\"fontType\":\"monospace\",\"size\":\"small\",\"separator\":true}" + } + if (failureDetails.rekick) { + def rk = escapeJson(failureDetails.rekick) + detailBlocks += ",{\"type\":\"RichTextBlock\",\"inlines\":[{\"type\":\"TextRun\",\"text\":\"Auto re-kick: \",\"weight\":\"bolder\"},{\"type\":\"TextRun\",\"text\":\"${rk}\"}],\"wrap\":true}" + } + } + def payload = """ +{"attachments":[{"contentType":"application/vnd.microsoft.card.adaptive","content":{"type":"AdaptiveCard","\$schema":"http://adaptivecards.io/schemas/adaptive-card.json","version":"1.4","body":[{"type":"TextBlock","text":"CI Update","weight":"bolder","size":"extraLarge","separator":true},{"type":"TextBlock","text":"${subtitle}"},{"type":"TextBlock","text":"${statusMessage}","color":"${color}"}${detailBlocks},{"type":"TextBlock","text":"Finished: ${timestamp}","size":"small","isSubtle":true}],"actions":[{"type":"Action.OpenUrl","url":"${blueOceanUrl}","title":"Open Blue Ocean 🌊"},{"type":"Action.OpenUrl","url":"${jobUrl}","title":"Open Job πŸ—οΈ"}]}}]} +""" + writeFile file: 'teams-payload.json', text: payload.trim(), encoding: 'UTF-8' + withCredentials([ + string(credentialsId: 'CI_MONITORING_TEAMS', variable: 'WEBHOOK_URL'), + string(credentialsId: 'MLIR_CI_CHANNEL', variable: 'WEBHOOK_URL_MLIR') + ]) { + ['CI_MONITORING_TEAMS': 'WEBHOOK_URL', 'MLIR_CI_CHANNEL': 'WEBHOOK_URL_MLIR'].each { name, envVar -> + def resp = sh(script: "curl -s -w '\\n%{http_code}' -X POST \"\$${envVar}\" -H 'Content-Type: application/json; charset=utf-8' -d @teams-payload.json", returnStdout: true).trim() + def lines = resp.split('\n') + def code = lines[-1] + def body = lines.length > 1 ? lines[0..-2].join('\n') : '' + echo "Teams webhook (${name}) response: HTTP ${code}${body ? ' body=' + body : ''}" + if (code != '200' && code != '202') { + echo "Teams notification (${name}) may have failed (expected 200/202, got ${code})" + } + } + } + } catch (e) { + echo "Teams notification skipped or failed: ${e}" + } +} + +// Post-build entrypoint: classifies the result, optionally enriches with +// failure details from the build log, and sends a Teams notification on +// official nightly/weekly runs. Extracted from Jenkinsfile's post.always +// block to keep the main pipeline body within JVM's 64KB CPS bytecode limit. +void handlePostBuildNotification() { + def result = currentBuild?.currentResult ?: 'UNKNOWN' + def statusMessage = 'Unknown' + def color = 'warning' + if (result == 'SUCCESS') { + statusMessage = "Build ${env.BUILD_NUMBER ?: '?'} completed successfully" + color = 'good' // green + } else if (result == 'FAILURE') { + statusMessage = "Build ${env.BUILD_NUMBER ?: '?'} failed" + color = 'attention' // red + } else if (result == 'ABORTED') { + statusMessage = "Build ${env.BUILD_NUMBER ?: '?'} was aborted" + color = 'warning' // yellow + } else { + statusMessage = "Build ${env.BUILD_NUMBER ?: '?'} ${result}" + color = 'warning' + } + def buildNum = env.BUILD_NUMBER ?: '?' + def buildUrl = env.BUILD_URL ?: '' + def jobName = env.JOB_NAME ?: '' + def isOfficialNightlyOrWeekly = (jobName == 'MLIR/mlir-nightly-all' || jobName == 'MLIR/mlir-weekly') + def failureDetails = null + def logText = '' + if ((result == 'FAILURE' || result == 'ABORTED') && (params.nightly || params.weekly) && isOfficialNightlyOrWeekly) { + try { + // Use large limit so early failures (e.g. SCM checkout) are included; getLog(N) may return last N lines on some setups. + def logLines = currentBuild.rawBuild.getLog(100000) + logText = logLines.join('\n') + failureDetails = classifyBuildFailure(logText) + } catch (e) { + echo "Could not classify failure: ${e}" + } + } + if (result == 'ABORTED') { + if (failureDetails == null) failureDetails = [:] + failureDetails.abortedBy = parseAbortedByFromLog(logText) ?: 'β€”' + } + // For FAILURE/ABORTED, ensure we always have a details map so the card shows Stage/CODEPATH/Details (with fallback if classification failed). + if ((result == 'FAILURE' || result == 'ABORTED') && failureDetails == null) { + failureDetails = [reason: 'Could not match a known error pattern. See build log for details.', codepath: '', stage: ''] + } + if (failureDetails != null && !failureDetails.reason) { + failureDetails.reason = 'Could not match a known error pattern. See build log for details.' + } + // Auto re-kick transient failures (official nightly/weekly only). + def rekickAttempt = (params.rekickAttempt?.toString()?.isInteger()) ? Math.max(0, params.rekickAttempt as int) : 0 + def willRekick = (result == 'FAILURE') && params.rekickEnabled && isOfficialNightlyOrWeekly && (params.nightly || params.weekly) && + (rekickAttempt < MAX_REKICK_ATTEMPTS) && isTransientHardwareFailure(logText) + def transientReason = (failureDetails?.reason) ?: 'transient failure' + + // Schedule the re-kick before notifying so we can suppress this run's card on success. + boolean rekickScheduled = false + if (willRekick) { + try { + echo "[re-kick] Transient failure on ${jobName} #${buildNum}; re-kicking (attempt ${rekickAttempt + 1}/${MAX_REKICK_ATTEMPTS}), suppressing this run's notification." + // Absolute path: a relative job name resolves against the current job's folder (MLIR/). + build job: "/${jobName}", wait: false, propagate: false, parameters: rekickParameters(rekickAttempt + 1, transientReason) + rekickScheduled = true + } catch (e) { + echo "[re-kick] Could not schedule re-kick: ${e}" + } + } + + // Exactly one card per logical build: skip it only when we successfully handed off to a re-kick. + if (!(willRekick && rekickScheduled) && (params.nightly || params.weekly) && isOfficialNightlyOrWeekly && buildUrl && jobName) { + if (willRekick && !rekickScheduled) { + if (failureDetails == null) failureDetails = [:] + failureDetails.rekick = "Transient failure detected, but the re-kick could not be scheduled β€” see build log." + } else if (rekickAttempt > 0) { + if (failureDetails == null) failureDetails = [:] + failureDetails.rekick = (result == 'SUCCESS') + ? "Auto re-kicked ${rekickAttempt}Γ— after a transient failure (${params.rekickReason ?: 'β€”'}); passed on retry." + : "Auto re-kicked ${rekickAttempt}Γ— after transient failure(s) (last: ${params.rekickReason ?: 'β€”'}); this run is not being re-kicked." + } + def runType = params.nightly ? 'nightly' : 'weekly' + def jenkinsBase = buildUrl.replaceFirst('/job/.*', '') + def jobNameEncoded = jobName.replace('/', '%2F') + def jobShortName = jobName.contains('/') ? jobName.split('/').last() : jobName + def blueOceanUrl = "${jenkinsBase}/blue/organizations/jenkins/${jobNameEncoded}/detail/${jobShortName}/${buildNum}/pipeline" + def jobUrl = buildUrl + try { + node('build-only') { + sendTeamsBuildNotification(buildNum, statusMessage, color, runType, blueOceanUrl, jobUrl, failureDetails) + } + } catch (e) { + echo "Could not send Teams notification: ${e}" + } + } +} + +// runBuildAndTestMatrixRow: Matrix-row body for the "Build and Test" stage. +// Extracted verbatim from Jenkinsfile to keep the main pipeline body +// within JVM's 64KB CPS bytecode limit. +def runBuildAndTestMatrixRow(String CODEPATH) { + // Prepare node + nodeUtils.withHealthyNode( + getLabelFromCodepath(CODEPATH), + { + nodeUtils.checkNodeHealth([doCleanWs: true]) + }, + { + stage("SCM Checkout") { + try { + scmUtils.robustScmCheckout() + } catch (e) { + error "[SCM] Checkout failed on ${env.NODE_NAME}: ${e}" + } + } + try { + String args = '' + def img = null + stage("Prepare Docker environment") { + // Fill in the docker args from the node + nodeUtils.dockerArgs() + + args = nodeUtils.DOCKER_ARGS_BY_NODE[env.NODE_NAME] + // Check these args + echo "Running ${CODEPATH} on ${env.NODE_NAME} with: ${args}" + nodeUtils.explicitDockerLogin() + img = nodeUtils.pullDockerImage(nodeUtils.dockerImage()) + } + // Spin up ONE container and stay in it for all substages + img.inside(args) { + withEnv([ + "HOME=${env.WORKSPACE}", + "PATH=/opt/rocm/llvm/bin:${env.PATH}" + ]) { + + if (params.sharedLib) { + stage('Shared Library: fixed E2E') { + echo "codepath is ${CODEPATH}" + echo "Container environment:" + nodeUtils.showEnv() + + testUtils.build_fixedE2ETests("${CODEPATH}") + testUtils.preMergeCheck("${CODEPATH}") + timeout(time: 60, activity: true, unit: 'MINUTES') { + buildUtils.shStrict 'cd build; ninja check-mlir check-rocmlir' + } + } + } + + if (params.sharedLib && params.nightly) { + stage('Shared Library: random E2E') { + testUtils.check_randomE2ETests("${CODEPATH}") + } + } + + if (params.sharedLib && !params.nightly) { + stage('Tune selected rocMLIR configs') { + buildUtils.buildProject('ci-performance-scripts', '') + dir('build') { + timeout(time: 60, activity: true, unit: 'MINUTES') { + buildUtils.shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ + --op gemm \ + -c ../mlir/utils/jenkins/ci-configs/selected-gemm-configs \ + -o tuning_gemm.tsv + [ -f tuning_gemm.tsv ]""" + buildUtils.shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ + --op conv \ + -c ../mlir/utils/jenkins/ci-configs/selected-conv-configs \ + -o tuning_conv.tsv + [ -f tuning_conv.tsv ]""" + buildUtils.shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ + --op attention \ + -c ../mlir/utils/jenkins/ci-configs/selected-attention-configs \ + -o tuning_attention.tsv + [ -f tuning_attention.tsv ]""" + buildUtils.shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ + --op gemm_gemm \ + -c ../mlir/utils/jenkins/ci-configs/selected-gemmgemm-configs \ + -o tuning_gemmgemm.tsv + [ -f tuning_gemmgemm.tsv ]""" + buildUtils.shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ + --op gemm --tuning-space quick \ + -c ../mlir/utils/jenkins/ci-configs/selected-gemm-configs \ + -o quick_tuning_gemm.tsv + [ -f quick_tuning_gemm.tsv ]""" + buildUtils.shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ + --op conv --tuning-space quick \ + -c ../mlir/utils/jenkins/ci-configs/selected-conv-configs \ + -o quick_tuning_conv.tsv + [ -f quick_tuning_conv.tsv ]""" + buildUtils.shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ + --op attention --tuning-space quick \ + -c ../mlir/utils/jenkins/ci-configs/selected-attention-configs \ + -o quick_tuning_attention.tsv + [ -f quick_tuning_attention.tsv ]""" + } + } + } + } + + if (params.staticLib && !params.nightly) { + stage('Static Lib: build packages') { + sh 'rm -f build/CMakeCache.txt' + buildUtils.buildProject('package', '-DBUILD_FAT_LIBROCKCOMPILER=ON') + testUtils.preMergeCheckPackage("${CODEPATH}") + echo "Running tests on the newly-built static library" + dir ('build') { + buildUtils.shStrict 'ninja check-rocmlir' + } + } + } + } + } + } + // Post block in scripted Jenkins works as try {} catch {}. + // Workspace cleanup is centralized in withHealthyNode so the per-row log survives. + catch (e) { + throw e + } + } + ) +} + + +// runParameterSweepsMatrixRow: Matrix-row body for the "Parameter sweeps" stage. +// Extracted verbatim from Jenkinsfile to keep the main pipeline body +// within JVM's 64KB CPS bytecode limit. +def runParameterSweepsMatrixRow(String CODEPATH) { + // Prepare node + nodeUtils.withHealthyNode( + getLabelFromCodepath(CODEPATH), + { + nodeUtils.checkNodeHealth() + }, + { + stage("SCM Checkout") { + try { + scmUtils.robustScmCheckout() + } catch (e) { + error "[SCM] Checkout failed on ${env.NODE_NAME}: ${e}" + } + } + try { + String args = '' + def img = null + stage("Prepare Docker environment") { + // Fill in the docker args from the node + nodeUtils.dockerArgs() + + args = nodeUtils.DOCKER_ARGS_BY_NODE[env.NODE_NAME] + // Check these args + echo "Running ${CODEPATH} on ${env.NODE_NAME} with: ${args}" + nodeUtils.explicitDockerLogin() + img = nodeUtils.pullDockerImage(nodeUtils.dockerImage()) + } + // Spin up ONE container and stay in it for all substages + img.inside(args) { + // The only way the env variables worked with all other changes + withEnv([ + "HOME=${env.WORKSPACE}" + ]) { + stage("Prepare Performance Scripts") { + echo "codepath is ${CODEPATH}" + echo "Container environment:" + nodeUtils.showEnv() + setHeartbeat() + buildUtils.buildProject('check-rocmlir-build-only ci-performance-scripts', '') + } + + stage("Parameter Sweep") { + testUtils.parameterSweep("conv_structure") + testUtils.parameterSweep("perf_config") + testUtils.parameterSweep(CODEPATH, "attention") + testUtils.parameterSweep(CODEPATH, "gemm_gemm") + archiveArtifacts artifacts: 'build/failing_attn_configs.txt,build/failing_gemmgemm_configs.txt,build/failing_conv_configs.txt', allowEmptyArchive: true + } + } + } + } + // Post block in scripted Jenkins works as try {} catch {}. + // Workspace cleanup is centralized in withHealthyNode so the per-row log survives. + catch (e) { + throw e + } + } + ) +} + + +// runTuneMatrixRow: Matrix-row body for the "Tune MLIR kernels" stage. +// Extracted verbatim from Jenkinsfile to keep the main pipeline body +// within JVM's 64KB CPS bytecode limit. +def runTuneMatrixRow(String CHIP) { + // Prepare node + nodeUtils.withHealthyNode( + getLabelFromChip(CHIP), + { + nodeUtils.checkNodeHealth() + }, + { + stage("SCM Checkout") { + try { + scmUtils.robustScmCheckout() + } catch (e) { + error "[SCM] Checkout failed on ${env.NODE_NAME}: ${e}" + } + } + try { + String args = '' + def img = null + stage("Prepare Docker environment") { + // Fill in the docker args from the node + nodeUtils.dockerArgs() + + args = nodeUtils.DOCKER_ARGS_BY_NODE[env.NODE_NAME] + // Check these args + echo "Running ${CHIP} on ${env.NODE_NAME} with: ${args}" + nodeUtils.explicitDockerLogin() + img = nodeUtils.pullDockerImage(nodeUtils.dockerImage()) + } + // Spin up ONE container and stay in it for all substages + img.inside(args) { + // The only way the env variables worked with all other changes + withEnv([ + "HOME=${env.WORKSPACE}", + "PATH=/opt/rocm/llvm/bin:${env.PATH}" + ]) { + if (CHIP == "gfx90a") { + stage("Set System Property on Lockhart nodes") { + nodeUtils.showEnv() + setHeartbeat() + } + } + + stage("Tune rocMLIR") { + buildUtils.buildProject('check-rocmlir-build-only ci-performance-scripts', '') + dir('build') { + def tuningLog = "tune_rocmlir_${CHIP}.log" + buildUtils.shStrict """echo "=== Tuning rocMLIR for ${CHIP} ===" | tee ${tuningLog}""" + // Tune gemms with default datatypes, fail if the tuning DB is not created + // (Includes int8xint8->int8 for performance comparisons against CK.) + buildUtils.shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ + --op gemm \ + -c ../mlir/utils/performance/configs/tier1-gemm-configs \ + -o mlir_tuning_${CHIP}.tsv 2>&1 | tee -a ${tuningLog} + [ -f mlir_tuning_${CHIP}.tsv ]""" + // Tune resnet50 and unet configs + buildUtils.shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ + --op conv \ + -c ../mlir/utils/performance/configs/tier1-conv-configs \ + -o mlir_tuning_${CHIP}.tsv 2>&1 | tee -a ${tuningLog}""" + // Tune attention configs + buildUtils.shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ + --op attention \ + -c ../mlir/utils/performance/configs/tier1-attention-configs \ + -o mlir_tuning_${CHIP}.tsv 2>&1 | tee -a ${tuningLog}""" + // Tune gemm_gemm configs + buildUtils.shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ + --op gemm_gemm \ + -c ../mlir/utils/performance/configs/tier1-gemmgemm-configs \ + -o mlir_tuning_${CHIP}.tsv 2>&1 | tee -a ${tuningLog}""" + // Quick tuning + buildUtils.shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ + --op gemm --tuning-space quick \ + -c ../mlir/utils/performance/configs/tier1-gemm-configs \ + -o mlir_quick_tuning_${CHIP}.tsv 2>&1 | tee -a ${tuningLog} + [ -f mlir_quick_tuning_${CHIP}.tsv ]""" + buildUtils.shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ + --op conv --tuning-space quick \ + -c ../mlir/utils/performance/configs/tier1-conv-configs \ + -o mlir_quick_tuning_${CHIP}.tsv 2>&1 | tee -a ${tuningLog}""" + buildUtils.shStrict """python3 ./bin/tuningRunner.py --abort-on-error \ + --op attention --tuning-space quick \ + -c ../mlir/utils/performance/configs/tier1-attention-configs \ + -o mlir_quick_tuning_${CHIP}.tsv 2>&1 | tee -a ${tuningLog}""" + buildUtils.shStrict """echo "=== Tuning rocMLIR for ${CHIP} completed ===" | tee -a ${tuningLog}""" + // Check for errors in the tuning log + script { + def tuneLog = readFile(tuningLog).split('\n') + // Find errors that are not part of a warning line + def errors = tuneLog.findAll { it =~ /(?i)error/ && !(it =~ /(?i)\bWARNING\b.*error/) } + + if (errors) { + currentBuild.result = 'FAILURE' + echo "Detected ${errors.size()} error(s) in tuning log:" + errors.each { echo "ERROR LINE: ${it}" } + error("Tuning failed: Detected errors in tuning log") + } else { + echo "No errors found in tuning log" + } + } + } + } + + stage("Tune Fusion") { + dir('build') { + // Tune resnet50 + buildUtils.shStrict """python3 ./bin/tuningRunner.py --abort-on-error --op fusion --test-dir ../mlir/test/fusion/resnet50-e2e/ -o tuning_fusion_${CHIP}.tsv""" + + // Tune bert + buildUtils.shStrict """python3 ./bin/tuningRunner.py --abort-on-error --op fusion --test-dir ../mlir/test/xmir/bert-torch-tosa-e2e/ -o tuning_fusion_${CHIP}.tsv""" + } + sh 'rm -f build/CMakeCache.txt' + } + + stage("Stash Databases") { + // Save user database for nightly jobs + dir ('build') { + stash name: "MLIR-PerfDB-${params.canXdlops ? CHIP : 'vanilla'}", includes: "*.tsv" + } + } + } + } + } + // Post block in scripted Jenkins works as try {} catch {}. + catch (e) { + throw e + } + finally { + // Publish per-arch tuning DBs as soon as this branch finishes so artifacts published without waiting for other parallel branches + try { + archiveArtifacts artifacts: "build/*.tsv", + allowEmptyArchive: true, onlyIfSuccessful: false + } catch (Exception archiveErr) { + echo "[CI] archiveArtifacts of tuning DBs failed: ${archiveErr.message}" + } + // Workspace cleanup is centralized in withHealthyNode so the per-row log survives. + } + } + ) +} + + +// runBenchmarkMatrixRow: Matrix-row body for the "Benchmark and Report Performance" stage. +// Extracted verbatim from Jenkinsfile to keep the main pipeline body +// within JVM's 64KB CPS bytecode limit. +def runBenchmarkMatrixRow(String CHIP) { + // Prepare node + nodeUtils.withHealthyNode( + getLabelFromChip(CHIP), + { + nodeUtils.checkNodeHealth() + }, + { + stage("SCM Checkout") { + try { + scmUtils.robustScmCheckout() + } catch (e) { + error "[SCM] Checkout failed on ${env.NODE_NAME}: ${e}" + } + } + try { + String args = '' + def img = null + stage("Prepare Docker environment") { + // Fill in the docker args from the node + nodeUtils.dockerArgs() + + args = nodeUtils.DOCKER_ARGS_BY_NODE[env.NODE_NAME] + // Check these args + echo "Running ${CHIP} on ${env.NODE_NAME} with: ${args}" + nodeUtils.explicitDockerLogin() + img = nodeUtils.pullDockerImage(nodeUtils.dockerImage()) + } + // Spin up ONE container and stay in it for all substages + img.inside(args) { + // The only way the env variables worked with all other changes + withEnv([ + "HOME=${env.WORKSPACE}", + "PATH=/opt/rocm/llvm/bin:${env.PATH}" + ]) { + stage("Copy tuning database") { + echo "chip is ${CHIP}" + echo "Container environment:" + nodeUtils.showEnv() + copyArtifacts filter: 'build/perfDB/**',\ + optional: true,\ + flatten: true,\ + projectName: "/MLIR/mlir-weekly",\ + selector: lastSuccessful(),\ + target: 'build' + sh 'ls build' + sh 'cat build/tuning-date' + } + + stage("Build MLIR") { + // Clean up build settings to disable static library and allow ROCm testing + buildUtils.buildProject( + 'check-rocmlir-build-only ci-performance-scripts hipblaslt-benchmark-driver', + '-DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ ' + + '-DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang ' + + '-DROCMLIR_ENABLE_BENCHMARKS=hipblaslt' + ) + } + + stage("Copy earlier performance results") { + copyArtifacts filter: 'build/*.csv,build/perf-run-date',\ + optional: true,\ + flatten: true,\ + projectName: "/${JOB_NAME}",\ + selector: lastSuccessful(),\ + target: 'build/oldData' + } + + stage("Test MLIR vs MIOpen/hipBLASLt") { + dir('build') { + def convInput = "${WORKSPACE}/mlir/utils/performance/configs/tier1-conv-configs" + def convToUse = "${WORKSPACE}/build/tier1-conv-configs" + def gemmInput = "${WORKSPACE}/mlir/utils/performance/configs/tier1-gemm-configs" + def gemmToUse = "${WORKSPACE}/build/tier1-gemm-configs" + script { + if (params.nightly) { + def runIndex = ((env.BUILD_NUMBER as int) - 1) % 5 + 1 + splitConfigFile(convInput, convToUse, runIndex) + splitConfigFile(gemmInput, gemmToUse, runIndex) + } + } + sh 'date --utc +%Y-%m-%d > perf-run-date' + sh 'ls -l /dev/kfd' + sh 'ls -l /dev/dri' + // Run MLIR vs MIOpen perf benchmarks. + buildUtils.shStrict """python3 ./bin/perfRunner.py --op=conv --batch-all \ + --configs-file=${convToUse} \ + --tuning-db=${WORKSPACE}/build/mlir_tuning_${CHIP}.tsv \ + --quick-tuning-db=${WORKSPACE}/build/mlir_quick_tuning_${CHIP}.tsv""" + // Run MLIR vs hipBLASLt perf benchmarks + buildUtils.shStrict """python3 ./bin/perfRunner.py --op=gemm --batch-all \ + --configs-file=${gemmToUse} \ + --tuning-db=${WORKSPACE}/build/mlir_tuning_${CHIP}.tsv \ + --quick-tuning-db=${WORKSPACE}/build/mlir_quick_tuning_${CHIP}.tsv""" + } + } + + stage("Test Fusion") { + dir('build') { + // Run fusion resnet50 perf benchmarks + buildUtils.shStrict """python3 ./bin/perfRunner.py --op=fusion --test-dir=${WORKSPACE}/mlir/test/fusion/resnet50-e2e/ --tuning-db=${WORKSPACE}/build/tuning_fusion_${CHIP}.tsv""" + // Run bert perf benchmarks + buildUtils.shStrict """python3 ./bin/perfRunner.py --op fusion --test-dir=${WORKSPACE}/mlir/test/xmir/bert-torch-tosa-e2e/ --tuning-db=${WORKSPACE}/build/tuning_fusion_${CHIP}.tsv""" + } + } + + if (isNotGfx11x(CHIP)) { + stage("Test Attention") { + dir('build') { + def attnInput = "${WORKSPACE}/mlir/utils/performance/configs/tier1-attention-configs" + def attnToUse = "${WORKSPACE}/build/tier1-attention-configs" + script { + if (params.nightly) { + def runIndex = ((env.BUILD_NUMBER as int) - 1) % 5 + 1 + splitConfigFile(attnInput, attnToUse, runIndex) + } + } + // Run attention benchmarks + buildUtils.shStrict """python3 ./bin/perfRunner.py --op=attention -b \ + --configs-file=${attnToUse} \ + --tuning-db=${WORKSPACE}/build/mlir_tuning_${CHIP}.tsv""" + } + } + + stage("Test Gemm+Gemm") { + dir('build') { + def gemmGemmInput = "${WORKSPACE}/mlir/utils/performance/configs/tier1-gemmgemm-configs" + def gemmGemmToUse = "${WORKSPACE}/build/tier1-gemmgemm-configs" + script { + if (params.nightly) { + def runIndex = ((env.BUILD_NUMBER as int) - 1) % 5 + 1 + splitConfigFile(gemmGemmInput, gemmGemmToUse, runIndex) + } + } + buildUtils.shStrict """python3 ./bin/perfRunner.py --op=gemm_gemm -b \ + --configs-file=${gemmGemmToUse} \ + --tuning-db=${WORKSPACE}/build/mlir_tuning_${CHIP}.tsv""" + } + } + } + + if (params.checkCK && supportsCKBenchmark(CHIP)) { + stage("Test MLIR vs CK") { + catchError (buildResult: null) { // This is an optional stage + def ckInstallDir = "${WORKSPACE}/composable_kernel/build/CKInstallDir" + dir('composable_kernel') { + sh 'rm -rf composable_kernel' + buildUtils.getAndBuildCK(''' + -DGPU_TARGETS=${CHIP} + -DCMAKE_PREFIX_PATH="/opt/rocm" + -DCMAKE_INSTALL_PREFIX=''' + ckInstallDir + ''' + -DCMAKE_BUILD_TYPE=Release + -DBUILD_TESTING=OFF + -DBUILD_CK_EXAMPLES=OFF + -DBUILD_CK_TUTORIALS=OFF + -DBUILD_CK_PROFILER=OFF + -DENABLE_CLANG_CPP_CHECKS=OFF + ''' + buildUtils.ckDtypesCmakeOptions(CHIP) + ''' + ''' + buildUtils.ckFp8CmakeOptions(CHIP) + ''' + ''', + 'device_gemm_operations') + buildUtils.installCKGemmOnly(ckInstallDir) + sh 'echo `git rev-parse HEAD`' + } + sh 'rm -f build/CMakeCache.txt' + buildUtils.buildProject("ck-benchmark-driver", + '''-DCMAKE_PREFIX_PATH=''' + ckInstallDir + ''' + -DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ + -DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang + -DROCMLIR_ENABLE_BENCHMARKS=ck''') + + + dir('build') { + def gemmInput = "${WORKSPACE}/mlir/utils/performance/configs/tier1-gemm-configs" + def gemmToUse = "${WORKSPACE}/build/tier1-gemm-configs" + script { + if (params.nightly) { + def runIndex = ((env.BUILD_NUMBER as int) - 1) % 5 + 1 + splitConfigFile(gemmInput, gemmToUse, runIndex) + } + } + buildUtils.shStrict """python3 ./bin/perfRunner.py --op=gemm --batch-all \ + --configs-file=${gemmToUse} \ + --tuning-db=${WORKSPACE}/build/mlir_tuning_${CHIP}.tsv --data-type f32 f16 i8_i8 --external-gemm-library CK""" + def ckChip = nodeUtils.get_gpu_architecture() + sh "python3 ./bin/createPerformanceReports.py ${ckChip} CK" + } + } + } + } + + stage("Create performance reports") { + dir('build') { + sh 'ls -l' + def reportChip = nodeUtils.get_gpu_architecture() + echo "Detected GPU chip for reports: ${reportChip} (CHIP matrix value: ${CHIP})" + sh "python3 ./bin/createPerformanceReports.py ${reportChip} MIOpen" + sh "python3 ./bin/createPerformanceReports.py ${reportChip} hipBLASLt" + sh "python3 ./bin/createFusionPerformanceReports.py ${reportChip}" + sh "python3 ./bin/perfRegressionReport.py ${reportChip}" + sh "python3 ./bin/perfRegressionReport.py ${reportChip} ./oldData/${reportChip}_mlir_vs_hipblaslt_perf.csv ./${reportChip}_mlir_vs_hipblaslt_perf.csv" + sh 'mkdir -p reports && cp ./*.html reports' + } + reportUtils.postProcessPerfRes(nodeUtils.get_gpu_architecture()) + } + } + } + } + // Post block in scripted Jenkins works as try {} catch {}. + // Workspace cleanup is centralized in withHealthyNode so the per-row log survives. + catch (e) { + throw e + } + } + ) +} + + +// runMIGraphXMatrixRow: Matrix-row body for the "MIGraphX" stage. +// Extracted verbatim from Jenkinsfile to keep the main pipeline body +// within JVM's 64KB CPS bytecode limit. +def runMIGraphXMatrixRow(String CODEPATH) { + // Prepare node + nodeUtils.withHealthyNode( + getLabelFromCodepath(CODEPATH), + { + nodeUtils.checkNodeHealth() + }, + { + stage("SCM Checkout") { + try { + scmUtils.robustScmCheckout() + } catch (e) { + error "[SCM] Checkout failed on ${env.NODE_NAME}: ${e}" + } + } + try { + String args = '' + def img = null + stage("Prepare Docker environment") { + // Fill in the docker args from the node + nodeUtils.dockerArgs() + + args = nodeUtils.DOCKER_ARGS_BY_NODE[env.NODE_NAME] + // Check these args + echo "Running ${CODEPATH} on ${env.NODE_NAME} with: ${args}" + + // Explicit docker login since this repo is private + nodeUtils.explicitDockerLogin() + img = nodeUtils.pullDockerImage(nodeUtils.dockerImageCIMIGraphX()) + } + // Spin up ONE container and stay in it for all substages + img.inside(args) { + // The only way the env variables worked with all other changes + withEnv([ + "HOME=${env.WORKSPACE}", + "PYTHONPATH=${env.WORKSPACE}/MIGraphX/build/lib:${env.PYTHONPATH}" + ]) { + stage("Install MIGraphX Dependencies") { + echo "codepath is ${CODEPATH}" + echo "Container environment:" + nodeUtils.showEnv() + // Package and install current checkout of rocMLIR as MIGraphX dependency (builds the fat lib, can OOM). + buildUtils.shStrict 'cget -p ${WORKSPACE}/MIGraphXDeps install ${WORKSPACE} -DBUILD_FAT_LIBROCKCOMPILER=On -DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ -DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang' + } + + stage("Build and Verify MIGraphX with MLIR") { + def gpu_arch = nodeUtils.get_gpu_architecture() + sh 'rm -rf MIGraphX' + dir('MIGraphX') { + buildUtils.getAndBuildMIGraphX(""" + -DCMAKE_PREFIX_PATH='${WORKSPACE}/MIGraphXDeps;/MIGraphXDeps;/opt/rocm' + -DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ + -DGPU_TARGETS="${gpu_arch}" + """) + } + } + + stage("Verify MIGraphX with MLIR") { + // f32 attention is unsupported on RDNA (no f32 WMMA), and --int8 + // does not quantize attention ops leaving them in f32. Exclude + // attention from MLIR ops on RDNA for int8 + def mlirOps = 'convolution,fused,dot,attention' + def mlirOpsInt8 = (CODEPATH == 'gfx103x' || CODEPATH == 'gfx120x') + ? 'convolution,fused,dot' + : 'convolution,fused,dot,attention' + + dir('MIGraphX/build') { + timeout(time: 120, activity: true, unit: 'MINUTES') { + // run test_verify for accuracy and run MLIR related unit-tests + withEnv(["MIGRAPHX_MLIR_USE_SPECIFIC_OPS=${mlirOps}", 'MIGRAPHX_ENABLE_MLIR_INPUT_FUSION=1', 'MIGRAPHX_ENABLE_MLIR_REDUCE_FUSION=1', 'MIGRAPHX_MLIR_ENABLE_SPLITK=1', 'MIGRAPHX_ENABLE_EXTRA_MLIR=1', 'MIGRAPHX_DISABLE_LAYERNORM_FUSION=1', 'MIGRAPHX_ENABLE_SPLIT_REDUCE=1']) { + buildUtils.shStrict 'make -j$(nproc) test_verify test_gpu_mlir test_gpu_fuse_mlir' + // Verify ResNet50, Bert, Gpt2 with fp16 + buildUtils.shStrict './bin/migraphx-driver verify --gpu --onnx /MIGraphXDeps/resnet50-v1-7.onnx --fp16' + buildUtils.shStrict './bin/migraphx-driver verify --gpu --onnx /MIGraphXDeps/bert_base_cased_1.onnx --fill1 input_ids --input-dim @input_ids 1 384 --fp16' + buildUtils.shStrict './bin/migraphx-driver verify --gpu --onnx /MIGraphXDeps/distilgpt2_1.onnx --fill1 input_ids --input-dim @input_ids 1 384 --fp16' + } + // int8 runs: exclude attention on RDNA to avoid f32 WMMA failure + withEnv(["MIGRAPHX_MLIR_USE_SPECIFIC_OPS=${mlirOpsInt8}", 'MIGRAPHX_ENABLE_MLIR_INPUT_FUSION=1', 'MIGRAPHX_ENABLE_MLIR_REDUCE_FUSION=1', 'MIGRAPHX_MLIR_ENABLE_SPLITK=1', 'MIGRAPHX_ENABLE_EXTRA_MLIR=1', 'MIGRAPHX_DISABLE_LAYERNORM_FUSION=1', 'MIGRAPHX_ENABLE_SPLIT_REDUCE=1']) { + buildUtils.shStrict './bin/migraphx-driver verify --gpu --onnx /MIGraphXDeps/resnet50-v1-7.onnx --int8' + buildUtils.shStrict './bin/migraphx-driver verify --gpu --onnx /MIGraphXDeps/bert_base_cased_1.onnx --fill1 input_ids --input-dim @input_ids 1 384 --int8' + buildUtils.shStrict './bin/migraphx-driver verify --gpu --onnx /MIGraphXDeps/distilgpt2_1.onnx --fill1 input_ids --input-dim @input_ids 1 384 --int8' + } + } + } + //Accuracy_checker will compare outputs from MIGraphX and onnx runtime + dir('MIGraphX/tools/accuracy') { + withEnv(["MIGRAPHX_MLIR_USE_SPECIFIC_OPS=${mlirOpsInt8}", 'MIGRAPHX_ENABLE_MLIR_INPUT_FUSION=1', 'MIGRAPHX_ENABLE_MLIR_REDUCE_FUSION=1', 'MIGRAPHX_MLIR_ENABLE_SPLITK=1', 'MIGRAPHX_ENABLE_EXTRA_MLIR=1', 'MIGRAPHX_DISABLE_LAYERNORM_FUSION=1', 'MIGRAPHX_ENABLE_SPLIT_REDUCE=1']) { + buildUtils.shStrict 'python3 accuracy_checker.py --onnx /MIGraphXDeps/resnet50-v1-7.onnx' + buildUtils.shStrict 'python3 accuracy_checker.py --fill1 --onnx /MIGraphXDeps/bert_base_cased_1.onnx --input-dim input_ids:1,384' + buildUtils.shStrict 'python3 accuracy_checker.py --fill1 --onnx /MIGraphXDeps/distilgpt2_1.onnx --input-dim input_ids:1,384' + } + } + } + } + } + } + // Post block in scripted Jenkins works as try {} catch {}. + // Workspace cleanup is centralized in withHealthyNode so the per-row log survives. + catch (e) { + throw e + } + } + ) +} + + +// runCodeCoverageMatrixRow: Matrix-row body for the "Code coverage" stage. +// Extracted verbatim from Jenkinsfile to keep the main pipeline body +// within JVM's 64KB CPS bytecode limit. +def runCodeCoverageMatrixRow(String CODEPATH) { + // Prepare node + nodeUtils.withHealthyNode( + getLabelFromCodepath(CODEPATH), + { + nodeUtils.checkNodeHealth() + }, + { + stage("SCM Checkout") { + try { + scmUtils.robustScmCheckout() + } catch (e) { + error "[SCM] Checkout failed on ${env.NODE_NAME}: ${e}" + } + } + try { + String args = '' + def img = null + stage("Prepare Docker environment") { + // Fill in the docker args from the node + nodeUtils.dockerArgs() + + args = nodeUtils.DOCKER_ARGS_BY_NODE[env.NODE_NAME] + // Check these args + echo "Running ${CODEPATH} on ${env.NODE_NAME} with: ${args}" + nodeUtils.explicitDockerLogin() + img = nodeUtils.pullDockerImage(nodeUtils.dockerImage()) + } + // Spin up ONE container and stay in it for all substages + img.inside(args) { + // The only way the env variables worked with all other changes + withEnv([ + "HOME=${env.WORKSPACE}", + "PYTHONPATH=${env.WORKSPACE}/MIGraphX/build/lib:${env.PYTHONPATH}", + // Note the %m to avoid issues with threads and dynamic libraries. + "LLVM_PROFILE_FILE=${env.WORKSPACE}/build/%m-%p.profraw", + "LLVM_PROFDATA=/opt/rocm/llvm/bin/llvm-profdata", + "LLVM_COV=/opt/rocm/llvm/bin/llvm-cov" + ]) { + stage ("body") { + echo "Container environment:" + nodeUtils.showEnv() + // Build with profiling on, and just code-generation tests. + try { + // Wall-clock timeout (no `activity:`): the coverage stage has long + // silent phases (lit buffers progress; `llvm-profdata merge` on + // ~125 GB of *.profraw produces no output for several minutes), + // which makes activity-based timeouts fire spuriously and the + // Codecov upload never run. 180 min covers ~60 min tests + + // profdata merge + three llvm-cov calls + upload with margin. + timeout(time: 180, unit: 'MINUTES') { + sh 'rm -f build/CMakeCache.txt' + sh 'rm -f build/*.profraw' + buildUtils.buildProject('check-rocmlir-build-only', + '-DBUILD_FAT_LIBROCKCOMPILER=ON -DCMAKE_BUILD_TYPE=debug -DLLVM_BUILD_INSTRUMENTED_COVERAGE=ON') + dir ('build') { + // Run tests. + testUtils.collectCoverageData("${LLVM_PROFDATA}", "${LLVM_COV}", "${CODEPATH}") + // Interpolate ${CODEPATH} via Groovy at the call site (bare CODEPATH is empty inside nested closures); shell-side vars are escaped with \$. + withCredentials([string(credentialsId: params.codecovCredentialsId ?: 'codecov-token-rocmlir', + variable: 'CODECOV_TOKEN')]) { + def uploadStatus = sh(script: """ + curl -Os https://uploader.codecov.io/latest/linux/codecov && chmod +x ./codecov + proxy_opt="" + if [ -n "\${http_proxy}" ]; then + proxy_opt="-U \${http_proxy}" + fi + set +e + ./codecov -t "\${CODECOV_TOKEN}" --flags "${CODEPATH}" -f ./coverage_${CODEPATH}.lcov \${proxy_opt} -Z + codecov_exit=\$? + set -e + echo "Codecov upload exit code: \${codecov_exit}" + exit \${codecov_exit} + """, returnStatus: true) + if (uploadStatus != 0) { + // Yellow stage but green build: don't block PR merge on Codecov hiccups. + catchError(buildResult: 'SUCCESS', stageResult: 'UNSTABLE', + message: "Codecov upload failed (exit ${uploadStatus})") { + error("Codecov upload failed (exit code ${uploadStatus}). Check that credential '${params.codecovCredentialsId ?: 'codecov-token-rocmlir'}' contains a valid token from codecov.io and that the repo is linked.") + } + } + } + // Best-effort HTML coverage report: runs AFTER the Codecov upload so a slow/timed-out llvm-cov show cannot prevent the LCOV upload, and is bounded by its own timeout so it cannot block archiveArtifacts. Gated by the runCoverageHtml parameter for builds that only need the Codecov upload. + if (params.runCoverageHtml) { + catchError(buildResult: 'SUCCESS', stageResult: 'UNSTABLE', message: 'Skipped HTML coverage report (llvm-cov show was slow or failed)') { + timeout(time: 45, unit: 'MINUTES') { + testUtils.produceCoverageHtml("${LLVM_COV}", "${CODEPATH}") + } + } + } + } + } + } catch (Exception e) { + // Yellow stage but green build, same rationale as above. + catchError(buildResult: 'SUCCESS', stageResult: 'UNSTABLE', + message: 'Code coverage stage had an error or timeout') { + error("Code coverage stage had an error or timeout: ${e}") + } + } finally { + // Always archive whatever was produced, even on UNSTABLE / timeout / exception. + archiveArtifacts artifacts: 'build/coverage*.report, build/coverage*.lcov, build/coverage*.html', allowEmptyArchive: true + } + } + } + } + } + // Post block in scripted Jenkins works as try {} catch {}. + // Workspace cleanup is centralized in withHealthyNode so the per-row log survives. + catch (e) { + throw e + } + } + ) +} + +return this diff --git a/mlir/utils/jenkins/helpers/nodeUtils.groovy b/mlir/utils/jenkins/helpers/nodeUtils.groovy new file mode 100644 index 000000000000..63409a50f8d6 --- /dev/null +++ b/mlir/utils/jenkins/helpers/nodeUtils.groovy @@ -0,0 +1,365 @@ +// Node lifecycle helpers: GPU reset, health checks, Docker discovery, +// docker image names, and the withHealthyNode() retry harness. +// Loaded by Jenkinsfile's Bootstrap stage; consumed as nodeUtils.(). +// ON CHANGING THESE, ALSO CHANGE Jenkinsfile.downstream + +import groovy.transform.Field +import java.util.concurrent.ConcurrentHashMap + +// ConcurrentHashMap helps when we need to write variables in parallel +// one instance for the whole run +@Field +ConcurrentHashMap DOCKER_ARGS_BY_NODE = new ConcurrentHashMap<>() + +// Characters from the end of the per-row log scanned to classify a transient failure; +// the decisive cause sits at the end. +@Field +final int FAILURE_LOG_TAIL_CHARS = 1000000 + +// Cross-helper handle, populated by Jenkinsfile's Bootstrap stage: +// nodeUtils.scmUtils = scmUtils +// Used by withHealthyNode() to invoke scmUtils.gitHealthCheck(). +@Field +def scmUtils + +def resetGPUs() { + // Abort this if runs longer than 10 minutes + timeout(time: 10, unit: 'MINUTES') { + // Run the reset, but don't fail the build if anything is wrong + def rc = sh( + script: ''' + reset_all_gpus() { + echo "Scanning GPUs..." + GPU_IDS=$(rocm-smi | awk '/^[0-9]+[[:space:]]+[0-9]+[[:space:]]+0x/ { print $1 }') + if [ -z "$GPU_IDS" ]; then + echo "WARNING: No GPUs found to reset." + return 0 + fi + for id in $GPU_IDS; do + echo "Resetting GPU ID: $id" + if ! rocm-smi --gpureset -d $id; then + echo "WARNING: Unable to reset GPU $id" + fi + sleep 2 + done + return 0 + } + reset_all_gpus + ''', + returnStatus: true + ) + if (rc != 0) { + echo "WARNING: reset_all_gpus exited with code ${rc}, but continuing anyway" + } + } +} + +def advancedNodeCheck(Map params) { + script { + echo "Jenkins-side PATH = '${env.PATH}'" + } + boolean doCleanWs = params.doCleanWs + boolean doGPUcheck = params.doGPUcheck + + if (doCleanWs) { + timeout(time: 15, unit: 'MINUTES', activity: true) { + cleanWs() + } + } + + resetGPUs() + + timeout(time: 5, unit: 'MINUTES', activity: true) { sh 'docker version' } + + ['ls -la /dev/kfd', 'ls -la /dev/dri'].each { cmd -> + timeout(time: 5, unit: 'MINUTES', activity: true) { sh cmd } + } + + String nodeSpecMessage = "\nNode specification:\n" + timeout(time: 5, unit: 'MINUTES', activity: true) { + nodeSpecMessage += "\nOS info:\n" + sh(script: 'sudo dkms status', returnStdout: true).trim() + '\n' + } + echo nodeSpecMessage + + if (env.NODE_LABELS && !env.NODE_LABELS.contains('build-only')) { + timeout(time: 5, unit: 'MINUTES', activity: true) { sh 'rocminfo' } + timeout(time: 5, unit: 'MINUTES', activity: true) { sh 'rocm-smi' } + timeout(time: 5, unit: 'MINUTES', activity: true) { sh 'cat /opt/rocm/.info/version' } + if (doGPUcheck) { + timeout(time: 5, unit: 'MINUTES', activity: true) { + def n = sh(script: "lspci | grep -e 'controller' -e 'accelerators' | grep 'AMD/ATI' | wc -l", + returnStdout: true).trim().toInteger() + if (n == 0) { + error "No GPUs detected on ${env.NODE_NAME}" + } + echo "Number of GPUs on ${env.NODE_NAME}: ${n}" + } + } + } else { + echo 'Skipping GPU checks…' + } +} + +def checkNodeHealth(Map opts = [:]) { + advancedNodeCheck( + doCleanWs: opts.get('doCleanWs', true), + doGPUcheck: opts.get('doGPUcheck', true) + ) +} + +Map dockerArgs() { + echo "Getting Docker args from ${env.NODE_NAME}..." + def run = { cmd -> sh(script: cmd, returnStdout: true).trim() } + // discover devices + String renderFlags = run("ls -1 /dev/dri/renderD* 2>/dev/null || true") + .split() + .collect { "--device=${it}" } + .join(' ') + // /dev/kfd appears only on GPU-enabled nodes + boolean haveKfd = sh(script: '[ -e /dev/kfd ]', returnStatus: true) == 0 + String kfdFlg = haveKfd ? '--device=/dev/kfd' : '' + + // Get the GIDs of the render and video groups + String renderGid = run("getent group render | cut -d':' -f3") + String videoGid = run("getent group video | cut -d':' -f3") + + String args = """ + ${kfdFlg} \ + ${renderFlags} \ + --group-add ${renderGid} --group-add ${videoGid} + """.trim().replaceAll(/\s+/, ' ') + + DOCKER_ARGS_BY_NODE[env.NODE_NAME] = args + echo "Received Docker args for ${env.NODE_NAME}: ${args}" + return DOCKER_ARGS_BY_NODE // ConcurrentHashMap +} + +void showEnv() { + echo "$env.NODE_NAME" + sh 'cat /etc/os-release' + sh 'ulimit -a' + // Ignore rocm-smi failures in ixt-sjc2-05 + sh '/opt/rocm/bin/rocm-smi || true' + sh '/opt/rocm/bin/rocm_agent_enumerator' + sh 'id' + sh 'printenv' +} + +String dockerImage() { + // If this is being changed please change Dockerfile.migraphx-ci's base image as well + return 'rocm/mlir:rocm7.2-latest' +} + +String dockerImageCIMIGraphX() { + return 'rocm/mlir-migraphx-ci:rocm7.2-latest' +} + +def retryDockerOperation(String description, Closure operation) { + int attempt = 0 + def result = null + retry(10) { + attempt += 1 + try { + result = operation() + } catch (err) { + echo "[Docker retry] ${description} failed on attempt ${attempt}/10 on ${env.NODE_NAME}: ${err}" + if (attempt < 10) { + echo "[Docker retry] Waiting 5 seconds before retrying ${description}" + sleep(time: 5, unit: 'SECONDS') + } + throw err + } + } + return result +} + +// For when the docker image is in a private repo +void explicitDockerLogin() { + withCredentials([usernamePassword(credentialsId: 'DOCKER_HUB_CREDS', + usernameVariable: 'D_USER', + passwordVariable: 'D_PASS')]) { + retryDockerOperation('docker login to DockerHub') { + sh ''' + set +x + printf "%s\n" "$D_PASS" | docker login -u "$D_USER" --password-stdin + ''' + } + } +} + +def pullDockerImage(String imageName) { + def img = docker.image(imageName) + retryDockerOperation("docker pull ${imageName}") { + img?.pull() + } + return img +} + +// Get the base GPU chip name as reported by the runtime (e.g. gfx1200, gfx942). +def get_gpu_architecture() { + try { + def result = sh(script: 'rocminfo', returnStdout: true).trim() + def arch_pattern = /Name:\s+amdgcn-amd-amdhsa--(gfx[0-9a-z]+)/ + def matches = (result =~ arch_pattern) + if (matches) { + return matches[0][1] + } + return 'N/A' + } catch (Exception e) { + echo "Error getting GPU architecture name: ${e}" + return 'N/A' + } +} + +// Run the body on a node that passes the supplied healthChecks() block +// The health check is retried on fresh executors; the body is not retried. +// This function also retries the main 'body' if it fails due to a recoverable node-related issue (e.g., agent disconnect). +// Genuine logic/test failures that must never be retried/re-kicked. Shared veto for the +// per-server classifier below and ciLogic's whole-job classifier (invoked via its nodeUtils +// handle as nodeUtils.realTestFailureSignals()). +List realTestFailureSignals() { + return [ + 'failed tests (', // lit + 'error: no match found', // FileCheck + 'filecheck error', + '*** summary of failures ***', // conv/perf sweeps + 'failing configurations', // attention sweeps + 'tuning failed: detected errors', + 'invalid mlir created', // MIGraphX + ] +} + +// Group-1 transients that can be retried on a fresh node in-pipeline (per matrix row), as opposed +// to whole-job transients like "no healthy node found" (handled by the post-block re-kick). `text` +// is the thrown exception plus the row's console tail. Case-insensitive. Deliberately excludes +// "no healthy node found"/"[withHealthyNode] transient" (nothing to retry on), "InterruptedException" +// (failFast collateral), and all genuine test-failure markers. +boolean isPerServerTransient(String text) { + if (!text) return false + String t = text.toLowerCase() + + // GPU lost/hung on this node; these surface in the test stdout, not in the thrown exception. + // Pre-veto: a dead GPU also makes lit report spurious test failures, so it wins over realSignals. + def gpuSignals = [ + 'hiperror_t.hiperrornodevice', + 'unable to reset gpu', + 'unsupported hip gpu architecture: n/a', + 'no performance report found for n/a', + 'gpu hang', + 'hw exception by gpu', + ] + if (gpuSignals.any { t.contains(it) }) return true + + // Veto: genuine test failures are never per-server retried (mirror the whole-job classifier). + if (realTestFailureSignals().any { t.contains(it) }) return false + if (t =~ /no performance report found for gfx/) return false + + // Node/agent died mid-run, or docker/OOM on this node; these surface in the exception. + def nodeSignals = [ + 'seems to be removed or offline', + 'agentofflineexception', + 'issue with creating launcher for agent', + 'closedchannelexception', + 'requestabortedexception', + 'broken pipe', + 'script returned exit code -1', + 'script returned exit code -2', + 'failed to run image', + 'outofmemoryerror', + 'ninja exited with error code 137', + 'maximum checkout retry attempts reached', + 'error cloning remote repo', + 'error fetching remote repo', + ] + if (nodeSignals.any { t.contains(it) }) return true + return scmUtils.isRetriableScmCheckoutError(t) +} + +def withHealthyNode(String baseLabel, Closure healthChecks, Closure body, int maxAttempts = 3) { + def blacklist = [] // nodes and pods that already failed the check + int attempt = 0 + boolean done = false + + while (!done && attempt < maxAttempts) { + attempt += 1 + + // Build a dynamic label that excludes everything that failed before + def expr = new StringBuilder(baseLabel) + blacklist.each { expr.append(' && !').append(it) } + + echo "[withHealthyNode] attempt #${attempt}: looking for '${expr}'" + node(expr.toString()) { + // Retry ONLY the health-check. We don't want to retry the actual stages + try { + stage("Health checks on ${env.NODE_NAME}") { + echo 'Cleaning up old Docker images...' + def pruneStatus = sh(script: 'docker image prune -af --filter "until=720h"', returnStatus: true) + if (pruneStatus != 0) { + echo "[withHealthyNode] WARNING: Docker image prune failed with exit code ${pruneStatus}. Continuing health check." + } + healthChecks() + scmUtils.gitHealthCheck() + } + } catch (Exception err) { + echo "[withHealthyNode] ❌ ${env.NODE_NAME} rejected: ${err}" + blacklist << env.NODE_NAME + // return exits the node {} block here, not the whole function. Some groovy magic + return + } + stage("Node selected") { + // Health-check passed. Do real work + echo "[withHealthyNode] βœ… using ${env.NODE_NAME}" + } + // Per-row console log: shStrict mirrors output here so we can classify transient + // failures (e.g. GPU hang) that only appear in stdout, not in the thrown exception. + String rowLog = "${env.WORKSPACE}/.rekick-row.log" + try { + withEnv(["REKICK_ROW_LOG=${rowLog}"]) { + body() + } + // If body succeeds, we're done with the loop + done = true + } catch (Exception err) { + String rowText = '' + try { + if (fileExists(rowLog)) { + rowText = readFile(rowLog) + if (rowText.length() > FAILURE_LOG_TAIL_CHARS) { + rowText = rowText.substring(rowText.length() - FAILURE_LOG_TAIL_CHARS) + } + } + } catch (Exception ignored) { } + + if (isPerServerTransient("${err}\n${rowText}")) { + // Group-1 transient on this node: blacklist it and retry the same arch on a + // fresh node. The while loop continues (done still false); if attempts run out + // this becomes "no healthy node found", which the post-block re-kicks whole-job. + echo "[withHealthyNode] Per-server transient on ${env.NODE_NAME}. Blacklisting the node and retrying.." + echo "[withHealthyNode] Error was: ${err}" + blacklist << env.NODE_NAME + return + } + // Real failure (or a whole-job transient like no-healthy-node): fail immediately. + echo "[withHealthyNode] Execution failed with a non-recoverable error on ${env.NODE_NAME}" + echo "[withHealthyNode] Error was: ${err}" + throw err + } finally { + // Clean here (moved out of the matrix bodies) so the per-row log above survives + // until it has been classified. Never let cleanup mask the body's exception. + try { + cleanWs() + } catch (Exception cleanErr) { + echo "[withHealthyNode] cleanWs failed: ${cleanErr}" + } + } + } + } + + if (!done) { + // In-stage breadcrumb: the post block reads the log before the final "error" line is printed. + echo "[withHealthyNode] TRANSIENT: no healthy node for '${baseLabel}' after ${maxAttempts} attempts" + error "No healthy node found for '${baseLabel}' after ${maxAttempts} attempts" + } +} + +return this diff --git a/mlir/utils/jenkins/helpers/reportUtils.groovy b/mlir/utils/jenkins/helpers/reportUtils.groovy new file mode 100644 index 000000000000..9c4a6997a4e8 --- /dev/null +++ b/mlir/utils/jenkins/helpers/reportUtils.groovy @@ -0,0 +1,52 @@ +// Reporting helpers: HTML/perf-plot publication and perfDB archival. +// Loaded by Jenkinsfile's Bootstrap stage; consumed as reportUtils.(). +// ON CHANGING THESE, ALSO CHANGE Jenkinsfile.downstream + +void postProcessPerfRes(String chip) { + publishHTML (target: [ + allowMissing: false, + alwaysLinkToLastBuild: false, + keepAll: true, + reportDir: 'build/reports', + reportFiles: "${chip}_MLIR_Performance_Changes.html,${chip}_MLIR_vs_MIOpen.html,${chip}_MLIR_Performance_Changes_Gemm.html,${chip}_MLIR_vs_hipBLASLt.html,${chip}_MLIR_vs_CK.html,${chip}_conv_fusion.html,${chip}_gemm_fusion.html", + reportName: "Performance report for ${chip}" + ]) + + if (fileExists("build/${chip}_mlir_vs_miopen_perf_for_plot.csv")) { + plot csvFileName: "${chip}_plot-nightly-perf-results-000001.csv",\ + csvSeries: [[file: "build/${chip}_mlir_vs_miopen_perf_for_plot.csv", displayTableFlag: false]],\ + title: "Test performance summary ${chip}, Conv",\ + yaxis: 'TFlops',\ + style: 'line',\ + group: 'Performance plots' + } + if (fileExists("build/${chip}_mlir_vs_hipblaslt_perf_for_plot.csv")) { + plot csvFileName: "${chip}_plot-nightly-perf-results-gemm-000001.csv",\ + csvSeries: [[file: "build/${chip}_mlir_vs_hipblaslt_perf_for_plot.csv", displayTableFlag: false]],\ + title: "Test performance summary ${chip}, GEMM",\ + yaxis: 'TFlops',\ + style: 'line',\ + group: 'Performance plots' + } + // Save results for future comparison + archiveArtifacts artifacts: 'build/*_mlir_*.csv,build/perf-run-date', allowEmptyArchive: true, onlyIfSuccessful: true +} + +void archivePerfDB() { + // Note: add additional architectures here + dir ('build/perfDB') { + def architectures = params.canXdlops ? ['gfx908', 'gfx90a', 'gfx942', 'gfx950', 'gfx1100', 'gfx1201'] : ['vanilla'] + for (arch in architectures) { + try { + unstash name: "MLIR-PerfDB-${arch}" + } catch (Exception e) { + echo "No stash found for MLIR-PerfDB-${arch}, skipping." + } + } + sh 'date --utc +%Y-%m-%d >tuning-date' + } + archiveArtifacts artifacts: 'build/perfDB/**',\ + onlyIfSuccessful: true +} + +return this diff --git a/mlir/utils/jenkins/helpers/scmUtils.groovy b/mlir/utils/jenkins/helpers/scmUtils.groovy new file mode 100644 index 000000000000..5f056c855da0 --- /dev/null +++ b/mlir/utils/jenkins/helpers/scmUtils.groovy @@ -0,0 +1,239 @@ +// SCM helpers: git health check and a robust checkout with fallback to a deep clone. +// Loaded by Jenkinsfile's Bootstrap stage; consumed as scmUtils.(). +// ON CHANGING THESE, ALSO CHANGE Jenkinsfile.downstream + +import com.cloudbees.groovy.cps.NonCPS +import groovy.transform.Field +import hudson.plugins.git.extensions.impl.CheckoutOption +import hudson.plugins.git.extensions.impl.CloneOption + +// Jenkins Git plugin defaults to 10 minutes per command. Use 2h for fetches and +// checkouts that can exceed that on a slow network. +@Field +final int GIT_SCM_TIMEOUT_MINUTES = 120 + +// Lightweight Git probe: verifies auth + network + ref exists +void gitHealthCheck() { + // Check if git installed + sh "git --version" + + // Check if git commands are healthy + String repo = scm?.userRemoteConfigs?.getAt(0)?.url + String cred = scm?.userRemoteConfigs?.getAt(0)?.credentialsId + String ref = env.CHANGE_ID ? "refs/pull/${env.CHANGE_ID}/head" + : env.BRANCH_NAME ? "refs/heads/${env.BRANCH_NAME}" + : "HEAD" + + if (!repo || !cred) { + error "[healthcheck] SCM not configured (repo='${repo}', cred='${cred}')" + } + echo "[healthcheck] Probing git: repo=${repo}, ref=${ref}" + + timeout(time: 2, unit: 'MINUTES') { + withCredentials([usernamePassword(credentialsId: cred, + usernameVariable: 'GIT_USER', + passwordVariable: 'GIT_PASS')]) { + withEnv(["REPO=${repo}", "REF=${ref}"]) { + sh ''' + set -eu + ASK="$(mktemp)"; trap 'rm -f "$ASK"' EXIT + printf '#!/bin/sh\nprintf %s "$GIT_PASS"\n' > "$ASK" + chmod +x "$ASK" + GIT_ASKPASS="$ASK" \ + git -c credential.username="$GIT_USER" \ + ls-remote --exit-code "$REPO" "$REF" >/dev/null + ''' + } + } + } + echo "[healthcheck] Git OK" +} + +@NonCPS +Map scmWithGitTimeout(Object baseScm) { + List extensions = [] + boolean hasCloneOption = false + boolean hasCheckoutOption = false + + (baseScm.extensions ?: []).each { ext -> + if (ext instanceof CloneOption) { + extensions << [ + $class: 'CloneOption', + depth: ext.depth ?: 0, + shallow: ext.shallow ?: false, + noTags: ext.noTags ?: false, + reference: ext.reference ?: '', + honorRefspec: ext.honorRefspec ?: false, + timeout: gitTimeoutAtLeast(ext.timeout) + ] + hasCloneOption = true + } else if (ext instanceof CheckoutOption) { + extensions << [$class: 'CheckoutOption', timeout: gitTimeoutAtLeast(ext.timeout)] + hasCheckoutOption = true + } else { + extensions << ext + } + } + + if (!hasCloneOption) { + extensions << [$class: 'CloneOption', timeout: GIT_SCM_TIMEOUT_MINUTES] + } + if (!hasCheckoutOption) { + extensions << [$class: 'CheckoutOption', timeout: GIT_SCM_TIMEOUT_MINUTES] + } + + Map checkoutScm = [ + $class: 'GitSCM', + branches: baseScm.branches, + doGenerateSubmoduleConfigurations: baseScm.doGenerateSubmoduleConfigurations ?: false, + extensions: extensions, + submoduleCfg: baseScm.submoduleCfg ?: [], + userRemoteConfigs: baseScm.userRemoteConfigs + ] + if (baseScm.gitTool) { + checkoutScm.gitTool = baseScm.gitTool + } + if (baseScm.browser) { + checkoutScm.browser = baseScm.browser + } + return checkoutScm +} + +@NonCPS +int gitTimeoutAtLeast(Integer timeout) { + int currentTimeout = timeout ?: 0 + return Math.max(currentTimeout, GIT_SCM_TIMEOUT_MINUTES) +} + +String scmCheckoutRetryContext(Object err) { + String msg = "${err}".toLowerCase() + try { + def logLines = currentBuild?.rawBuild?.getLog(500) ?: [] + msg = msg + '\n' + logLines.join('\n').toLowerCase() + } catch (ignored) { + // Fall back to the exception text; retry classification should not mask the checkout failure. + } + return msg +} + +boolean isRetriableScmCheckoutError(String msg) { + return [ + "connection reset by peer", + "curl 18", + "transfer closed with outstanding read data remaining", + "bytes of body are still expected", + "unexpected disconnect while reading sideband packet", + "bad pack header", + "git-remote-https died of signal 15", + "early eof", + "invalid index-pack output" + ].any { msg.contains(it) } +} + +// Retry checkout without shallow clone if GitSCM chokes on a specific SHA +void robustScmCheckout() { + int maxAttempts = 2 + for (int attempt = 1; attempt <= maxAttempts; attempt++) { + try { + // This inner 'try' handles the "reference is not a tree" fallback + try { + echo "[SCM] Attempting checkout (${attempt}/${maxAttempts})..." + checkout(scmWithGitTimeout(scm)) + echo "[SCM] Checkout successful" + // If checkout succeeds, exit the function immediately + return + } catch (err) { + def msg = "${err}".toLowerCase() + if (!msg.contains("reference is not a tree") && !msg.contains("could not checkout")) { + // If it's not a known transient error, re-throw it to be caught by the outer block + throw err + } + + // This is the fallback logic for the "reference is not a tree" error + echo "[SCM] Default checkout failed: ${err}. Retrying ONCE with robust deep clone" + String repo = scm?.userRemoteConfigs?.getAt(0)?.url + String cred = scm?.userRemoteConfigs?.getAt(0)?.credentialsId + String ref = env.CHANGE_ID ? "refs/pull/${env.CHANGE_ID}/head" + : env.BRANCH_NAME ? "refs/heads/${env.BRANCH_NAME}" + : "HEAD" + + def deepScm = [ + $class: 'GitSCM', + userRemoteConfigs: [[url: repo, credentialsId: cred, refspec: "+${ref}:${ref}"]], + branches: [[name: ref]], + doGenerateSubmoduleConfigurations: false, + extensions: [ + [$class: 'CloneOption', depth: 0, shallow: false, noTags: false, honorRefspec: true, timeout: GIT_SCM_TIMEOUT_MINUTES], + [$class: 'CheckoutOption', timeout: GIT_SCM_TIMEOUT_MINUTES] + ] + ] + checkout(deepScm) + echo "[SCM] Deep clone checkout successful." + // If the deep clone succeeds, exit the function + return + } + } catch (err) { + // This outer 'catch' block is specifically for retrying network errors + def msg = scmCheckoutRetryContext(err) + if (isRetriableScmCheckoutError(msg) && attempt < maxAttempts) { + echo "[SCM] Attempt ${attempt}/${maxAttempts} failed due to a transient git fetch error." + echo "[SCM] Waiting 2 minutes before retrying..." + sleep(time: 2, unit: 'MINUTES') + // The loop will now continue to the next attempt. + } else { + // This is either not a network error, or it was the final attempt. Fail the build + echo "[SCM] Unrecoverable SCM error after ${attempt} attempt(s)." + throw err + } + } + } +} + +Map externalGitScm(String url, String branch) { + return [ + $class: 'GitSCM', + branches: [[name: "*/${branch}"]], + doGenerateSubmoduleConfigurations: false, + extensions: [ + [ + $class: 'CloneOption', + depth: 0, + shallow: false, + noTags: false, + reference: '', + honorRefspec: false, + timeout: GIT_SCM_TIMEOUT_MINUTES + ], + [$class: 'CheckoutOption', timeout: GIT_SCM_TIMEOUT_MINUTES] + ], + submoduleCfg: [], + userRemoteConfigs: [[url: url]] + ] +} + +void robustExternalCheckout(String url, String branch) { + int maxAttempts = 2 + for (int attempt = 1; attempt <= maxAttempts; attempt++) { + try { + // Discard a partial pack before retrying the clone. + deleteDir() + checkout( + changelog: false, + poll: false, + scm: externalGitScm(url, branch) + ) + return + } catch (err) { + String context = scmCheckoutRetryContext(err) + if (attempt == maxAttempts || !isRetriableScmCheckoutError(context)) { + throw err + } + + echo "[SCM] External checkout attempt ${attempt}/${maxAttempts} failed due to a transient git fetch error." + echo "[SCM] Waiting 2 minutes before retrying..." + sleep(time: 2, unit: 'MINUTES') + } + } +} + +return this diff --git a/mlir/utils/jenkins/helpers/testUtils.groovy b/mlir/utils/jenkins/helpers/testUtils.groovy new file mode 100644 index 000000000000..93f6efba1208 --- /dev/null +++ b/mlir/utils/jenkins/helpers/testUtils.groovy @@ -0,0 +1,154 @@ +// Test helpers: pre-merge static checks, fixed/random E2E test suites, +// parameter sweeps, lit worker sizing, and coverage collection. +// Loaded by Jenkinsfile's Bootstrap stage; consumed as testUtils.(). +// ON CHANGING THESE, ALSO CHANGE Jenkinsfile.downstream + +import groovy.transform.Field + +// Cross-helper handles, populated by Jenkinsfile's Bootstrap stage: +// testUtils.nodeUtils = nodeUtils +// testUtils.buildUtils = buildUtils +@Field def nodeUtils +@Field def buildUtils + +void preMergeCheck(String codepath) { + // Only do static check on mfma codepath during PR CI + if ( (params.nightly == false) && (codepath == "mfma") ) { + echo "Performing Static Test (preMergeCheck)" + sh ''' + if [ ! -f ./build/compile_commands.json ]; then + echo "No compile commands, bailing." + exit 1 + fi + if [ ! -f ./compile_commands.json ]; then + ln -s build/compile_commands.json compile_commands.json + fi + ''' + def targetBranch = env.CHANGE_TARGET + if (!targetBranch) { + targetBranch = "develop" + } + if (params.ignoreExternalLinting == true) { + sh "python3 ./mlir/utils/jenkins/static-checks/premerge-checks.py --base-commit=origin/${targetBranch} --ignore-external" + } + else { + sh "python3 ./mlir/utils/jenkins/static-checks/premerge-checks.py --base-commit=origin/${targetBranch}" + } + } else { + echo "Static Test step skipped" + } +} + +void preMergeCheckPackage(String codepath) { + // Only do static check on mfma codepath during PR CI + if ( (params.nightly == false) && (codepath == "mfma") ) { + echo "Checking if the fat library target list is accurate" + dir('build') { + sh '../mlir/utils/jenkins/static-checks/get_fat_library_deps_list.pl > ./librockcompiler_deps.cmake.new' + } + sh 'diff -up mlir/tools/rocmlir-lib/librockcompiler_deps.cmake ./build/librockcompiler_deps.cmake.new' + } else { + echo "Skipping fat library target list check" + } +} + +int setLitWorkerCount() { + int limit_lit_workers = 8 + def gpu_arch = nodeUtils.get_gpu_architecture() + if (gpu_arch.contains('gfx908') || gpu_arch.contains('gfx90a')) { + limit_lit_workers = 20 + } else if (gpu_arch.contains('gfx942')) { + limit_lit_workers = 64 + } + return limit_lit_workers +} + +void build_fixedE2ETests(String codepath) { + // Limit the number of lit workers for gfx908, gfx90a to (8, 30) on CI as a workaround for issue #1845 and #1841 + int limit_lit_workers = setLitWorkerCount() + buildUtils.buildProject("check-mlir-build-only check-rocmlir-build-only${params.nightly ? ' hipblaslt-benchmark-driver' : ''}", """ + -DROCMLIR_DRIVER_PR_E2E_TEST_ENABLED=${params.nightly ? '0' : '1'} + -DROCMLIR_DRIVER_E2E_TEST_ENABLED=${params.nightly ? '1' : '0'} + -DROCK_E2E_TEST_ENABLED=${params.nightly ? '1' : '0'} + -DROCMLIR_DRIVER_TEST_GPU_VALIDATION=1 + -DROCMLIR_ENABLE_BENCHMARKS=${params.nightly ? 'hipblaslt' : ''} + -DLLVM_LIT_ARGS='-v --time-tests --timeout=3600 --max-failures=1 -j ${limit_lit_workers}' + -DCMAKE_EXPORT_COMPILE_COMMANDS=1 + """) +} + +void check_randomE2ETests(String codepath) { + // Limit the number of lit workers for gfx908, gfx90a to (8, 30) on CI as a workaround for issue #1845 and #1841 + int limit_lit_workers = setLitWorkerCount() + // Configure and build the E2E deps without running the tests, then run the GPU tests via + // shStrict so their stdout is mirrored to the per-row log (withHealthyNode classifies GPU + // hangs there and retries only this node). Running check-rocmlir directly through cmakeBuild + // would bypass shStrict and force a whole-job re-kick instead. + buildUtils.buildProject('check-rocmlir-build-only', """ + -DROCMLIR_DRIVER_PR_E2E_TEST_ENABLED=0 + -DROCMLIR_DRIVER_E2E_TEST_ENABLED=1 + -DROCK_E2E_TEST_ENABLED=1 + -DROCMLIR_DRIVER_RANDOM_DATA_SEED=1 + -DROCMLIR_DRIVER_TEST_GPU_VALIDATION=0 + -DLLVM_LIT_ARGS='-v --time-tests --timeout=3600 --max-failures=1 -j ${limit_lit_workers}' + -DCMAKE_EXPORT_COMPILE_COMMANDS=1 + """) + timeout(time: 60, activity: true, unit: 'MINUTES') { + buildUtils.shStrict 'cd build; ninja check-rocmlir' + } +} + +void parameterSweep(String CONFIG, String sweepType = "default") { + int limit_lit_workers = setLitWorkerCount() + timeout(time: 300, activity: true, unit: 'MINUTES') { + dir('build') { + if (sweepType == "attention" || sweepType == "gemm_gemm") { + String accelCodepath = "auto" + if (CONFIG == "mfma" || CONFIG == "gfx950") { + accelCodepath = "mfma" + } else if (CONFIG == "gfx103x" || CONFIG == "gfx110x" || CONFIG == "gfx120x") { + accelCodepath = "wmma" + } + buildUtils.shStrict """python3 ./bin/attentionSweeps.py --op ${sweepType} -j ${limit_lit_workers} --codepath ${accelCodepath} --log-failures --debug-fails""" + } else { + buildUtils.shStrict """python3 ./bin/parameterSweeps.py -j ${limit_lit_workers} ${CONFIG} --log-failures""" + } + } + } +} + +void collectCoverageData(String profdata, String cov, String cpath) { + // Runs `ninja check-rocmlir` (GPU E2E), so use shStrict to mirror output to the per-row log. + buildUtils.shStrict """ + rm -f *.profraw + # Arbitrarily 150 GB; we typically see 125 GB of *.profraw. + if [ `df --output=avail -k . | tail -n 1` -lt 153600000 ]; then + echo Not enough free disk space for profiling. + exit 1 + fi + ninja check-rocmlir + # Profile processing. + ${profdata} merge -sparse ./*.profraw -o ./coverage.profdata + rm -f build/*.profraw + ${cov} report --object ./bin/rocmlir-opt --object ./bin/rocmlir-driver \ + --object ./bin/rocmlir-gen --instr-profile ./coverage.profdata \ + --ignore-filename-regex='external/llvm-project|mlir/test/|mlir/unittests/' > ./coverage_${cpath}.report + cat ./coverage_${cpath}.report + ${cov} export --object ./bin/rocmlir-opt --object ./bin/rocmlir-driver \ + --object ./bin/rocmlir-gen --instr-profile ./coverage.profdata \ + --ignore-filename-regex='external/llvm-project|mlir/test/|mlir/unittests/' --format=lcov \ + --compilation-dir ${WORKSPACE} > ./coverage_${cpath}.lcov + """ +} + +// Produce the HTML coverage report +void produceCoverageHtml(String cov, String cpath) { + sh """ + ${cov} show --object ./bin/rocmlir-opt --object ./bin/rocmlir-driver \ + --object ./bin/rocmlir-gen --instr-profile ./coverage.profdata \ + --ignore-filename-regex='external/llvm-project|mlir/test/|mlir/unittests/' -Xdemangler=llvm-cxxfilt \ + --format=html > ./coverage_${cpath}.html + """ +} + +return this