Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
331 commits
Select commit Hold shift + click to select a range
4c66d60
[Comgr][Cache] Enably the cache by default
jmmartinez Mar 10, 2025
3f3b249
[Comgr] Do not compute opencl-c.h hash (#1269)
Mar 20, 2025
20fbc4b
[Comgr][Cache] Fix broken test: spirv-translator-cached.cl (#1270)
Mar 20, 2025
9d5047e
[Comgr][Cache] Late code reviews cherry-picked from staging to mainli…
Mar 20, 2025
85a5451
[Comgr][Cache] Enable the cache by default (#1272)
Mar 20, 2025
3680aa6
[Cache][SPIRV] Fix flacky test... again
jmmartinez Mar 21, 2025
ca47f24
[Cache][SPIRV] Fix flacky test... again (#1318)
lamb-j Mar 25, 2025
842526e
[Comgr] Fix disassem-instruction memory corruption (#1336)
lamb-j Mar 27, 2025
b8906e0
[Comgr] Add new Action to compile SPIR-V to Relocatable
lamb-j Mar 27, 2025
84c7a38
[Comgr][Cache] Fix broken test: spirv-translator-cached.cl
jmmartinez Mar 12, 2025
a661302
[Cache][SPIRV] Fix flacky test... again
jmmartinez Mar 21, 2025
a893c39
[Comgr][Merge] Two dependent changes to fix spirv test (#1317) (#1365)
Mar 27, 2025
39a37df
[AMDGPU] Auto generated check lines for two tests (#1126)
shiltian Mar 13, 2025
4a86a97
[AMDGPU] Auto generated check lines for two tests (#1126) (#1368)
Mar 27, 2025
ec4fa6b
[AMDGPU] Change SGPR layout to striped caller/callee saved (#127353)
shiltian Mar 8, 2025
bb5598a
[AMDGPU] Change SGPR layout to striped caller/callee saved (#127353) …
Mar 27, 2025
e6c5ed5
[Attributor] Do not optimize away externally_initialized loads. (#128…
Pierre-vh Mar 3, 2025
3dbccb7
[Attributor] Do not optimize away externally_initialized loads. (#128…
Mar 27, 2025
715b0e7
[Hipcc] Mark ~HipBinBase as virutal
jmmartinez Mar 19, 2025
92fc035
[Comgr] Fix disassem-instruction memory corruption (#1336) (#1354)
Mar 27, 2025
7aac4cb
[Comgr] Add new Action to compile SPIR-V to Relocatable (#1362)
Mar 27, 2025
c27d4e8
Add Env var control on enabling device-to-device memory access
hidekisaito Mar 24, 2025
0475427
[Comgr][V3] Increment Version number to 3.0
lamb-j Dec 19, 2024
0f51431
Add test for warning about using amdflang-new
mjklemm Apr 2, 2025
0f9a104
Emit warning when amdflang-new is invoked
mjklemm Apr 2, 2025
cdead4f
Add Env var control on enabling device-to-device memory access (#1384)
Apr 3, 2025
28b85e3
[Comgr][V3] Increment Version number to 3.0 (#1423)
Apr 3, 2025
aaed638
Combined ASAN device malloc patches to eliminate false reports
b-sumner Mar 31, 2025
9b71db5
Combined ASAN device malloc patches to eliminate false reports (#1522)
Apr 7, 2025
af6fb56
Bring tgamma patches 1397 and 1429 over to mainline
b-sumner Apr 1, 2025
3de1e5d
Bring tgamma patches 1397 and 1429 over to mainline (#1523)
Apr 7, 2025
fa50204
[Flang] Emit warning when AMD Next-gen Fortran Compiler is invoked as…
Apr 7, 2025
c13f8d5
[clang][Driver] Default to DIOp-based DIExpressions in SPIRV (#1278)
epilk Mar 26, 2025
499cece
[clang][Driver] Default to DIOp-based DIExpressions in SPIRV (#1278) …
Apr 7, 2025
5ba6cb8
[CUDA] Increment VTable index for device thunks (#124989)
gandhi56 Feb 20, 2025
847c570
[CUDA] Increment VTable index for device thunks (#124989) (#1528)
Apr 7, 2025
c3c6266
Cherry-pick missing part of [Clang][Sema] Revert changes to operator=…
yxsamliu Mar 27, 2025
5900c9f
Cherry-pick missing part of [Clang][Sema] Revert changes to operator=…
Apr 7, 2025
f795529
Update amd/comgr/test-lit/spirv-to-reloc-debuginfo.hip
lamb-j Mar 27, 2025
538b36b
Update amd/comgr/test-lit/spirv-to-reloc-debuginfo.hip (#1532)
Apr 8, 2025
38e98c3
[clang][AMDGPU] Enable module splitting by default (#128509)
Pierre-vh Mar 24, 2025
05183b9
[clang][AMDGPU] Enable module splitting by default (#128509) (#1535)
Apr 8, 2025
8594151
AMDGPU: Handle gfx950 XDL-write-VGPR-VALU-Mem-Exp wait state change (…
VigneshwarJ Feb 11, 2025
09566fa
AMDGPU: Handle gfx950 XDL-write-VGPR-VALU-Mem-Exp wait state change (…
Apr 8, 2025
c9a8ed5
AMDGPU: Handle gfx950 XDL Write-VGPR-VALU-WAW wait state change (#126…
VigneshwarJ Feb 11, 2025
5e7fd37
AMDGPU: Handle gfx950 XDL Write-VGPR-VALU-WAW wait state change (#126…
Apr 8, 2025
9f7960a
Reapply "[AMDGPU][GlobalISel] Fix load/store of pointer vectors, buff…
krzysz00 Oct 9, 2024
0d40750
Reapply "[AMDGPU][GlobalISel] Fix load/store of pointer vectors, buff…
Apr 8, 2025
03ed033
Reapply "[AMDGPU][GlobalISel] Fix load/store of pointer vectors, buff…
krzysz00 Nov 1, 2024
c95a090
Reapply "[AMDGPU][GlobalISel] Fix load/store of pointer vectors, buff…
Apr 8, 2025
b5bb324
Revert "[NFC][AMDGPU] Pre-commit clang and llvm tests for dynamic all…
easyonaadit Dec 18, 2024
8595f09
Revert "[NFC][AMDGPU] Pre-commit clang and llvm tests for dynamic all…
Apr 8, 2025
f8f2573
(cherry-pick) [AMDGPU] Prevent m0 from being used as v_readlane/v_rea…
Pierre-vh Apr 9, 2025
499e10d
[Offload] [OMPT] Fixed missing trace records in ompt_flush_trace.
dhruvachak Apr 3, 2025
b6216c4
[Offload] [OMPT] Fixed missing trace records in ompt_flush_trace. (#1…
Apr 9, 2025
4c11670
[Comgr] Re-add SPIR-V translator options dropped in rebase
lamb-j Apr 9, 2025
3cc45f3
[Comgr] Re-add SPIR-V translator options dropped in rebase (#1614)
Apr 15, 2025
62f6fff
[Comgr] Add Image Support to ISA metadata
lamb-j Apr 15, 2025
2199fbd
[Clang][AMDGPU] Expose buffer load lds as a clang builtin (#132048)
jmmartinez Apr 3, 2025
3b7c371
[AMDGPU][SelectionDAG] Use COPY instead of S_MOV_B32 to assign values…
jmmartinez Mar 25, 2025
4f5b614
[AMDGPU] Use a target feature to enable __builtin_amdgcn_global_load_…
jmmartinez Apr 2, 2025
af6adb6
[Clang][AMDGPU] Add __builtin_amdgcn_cvt_off_f32_i4 (#133741)
jmmartinez Apr 2, 2025
9c59fa1
Collection of LDS builtin changes (#1720)
Apr 17, 2025
3600741
[AMDGPU] S_SET_GPR_IDX_ON can be passed an immediate index (#125086)
JonChesterfield Jan 30, 2025
60bf26c
[amdgpu][nfc] Post-commit feedback on c39fba209
JonChesterfield Jan 30, 2025
70c627c
fixes the error in rocfft and hipfft
Apr 21, 2025
9a98832
(cherry-pick) Fixes and Improvements for -flto-partitions option (#1416)
Pierre-vh Apr 10, 2025
3a5421c
(cherry-pick) Fixes and Improvements for -flto-partitions option (#14…
Apr 29, 2025
521620c
SWDEV-520417 - Cherry-pick fixes from Shore to amd-mainline (#1470)
kzhuravl Apr 11, 2025
bdf914f
SWDEV-520417 - Cherry-pick fixes from Shore to amd-mainline (#1470) (…
Apr 29, 2025
66e34ba
[SLP] Sort PHIs by ExtractElements when relevant (#131229)
jrbyrnes Mar 17, 2025
a097554
[SLP] Sort PHIs by ExtractElements when relevant (#131229) (#1900)
Apr 29, 2025
0cb9ca8
[Clang][LLVM] Port ZCFS from staging (#1617)
searlmc1 Apr 11, 2025
8c1e1a0
[Clang][LLVM] Port ZCFS from staging (#1617) (#1901)
Apr 29, 2025
3066e7c
[LiveDebugValues] Disable spill tracking on AMDGPU (#561)
epilk Feb 20, 2025
9d3e440
[LiveDebugValues] Disable spill tracking on AMDGPU (#561) (#1902)
Apr 29, 2025
c414d7c
AMDGPU: Fix the double rounding issue in v2f64 -> v2f16 conversion (#…
changpeng Apr 17, 2025
d0bfeed
AMDGPU: Fix the double rounding issue in v2f64 -> v2f16 conversion (#…
Apr 29, 2025
6c17104
[OMPD] Remove deprecated/unused module that is causing error (#127434…
vigbalu Mar 6, 2025
c4ffb6f
[OMPD] Remove deprecated/unused module that is causing error (#127434…
May 1, 2025
2154555
XFAIL w-option.f90
mjklemm May 2, 2025
669d67d
[Flang][SWDEV-530408] XFAIL the problematic test (#1948)
May 2, 2025
861d8cd
[AMDGPU] Legalize vectorization of i8 types
gandhi56 Mar 6, 2025
c941ee4
[AMDGPU] Legalize vectorization of i8 types (#2013)
May 8, 2025
dd5dd3a
[Comgr] Fix broken SPIRV LIT test
lamb-j May 7, 2025
e97f023
[Comgr] Fix broken SPIRV LIT test (#2063)
May 8, 2025
dfbbc6a
[AMDGPU] Correctly merge noalias scopes during lowering of LDS data. …
srpande Apr 28, 2025
efd85c7
[AMDGPU] Correctly merge noalias scopes during lowering of LDS data. …
May 20, 2025
37411f5
Reapply "[clang][HIP] Make some math not not work with AMDGCN SPIR-V …
AlexVlx Feb 28, 2025
160ae96
Reapply "[clang][HIP] Make some math not not work with AMDGCN SPIR-V …
May 20, 2025
8cb07ce
[StructurizeCFG] Stop setting DebugLocs in flow blocks (#139088)
epilk May 9, 2025
3a4a864
[StructurizeCFG] Stop setting DebugLocs in flow blocks (#139088) (#2281)
May 21, 2025
2b7634d
[clang][CodeGen] Fix crash on non-natural type in CheckAtomicAlignmen…
May 22, 2025
9fd5af2
[clang][CodeGen] Fix crash on non-natural type in CheckAtomicAlignmen…
May 22, 2025
1c6d7b8
[Hipcc] Mark ~HipBinBase as virutal (#1373)
May 27, 2025
53e78fa
[SWDEV-534361] Fix error "-gz" option not recognized.
ampandey-1995 May 22, 2025
5ac13ad
Support sanitizer flag for Flang Driver.
ampandey-1995 Apr 29, 2025
1afc1c3
Fixes SWDEV-534361 - [Compiler Stg] [ASAN] Hipblaslt build failed
May 27, 2025
5ac52c0
Extend llvm objdump fatbin (#140286)
david-salinas May 23, 2025
7d7d22c
Extend llvm objdump fatbin (#140286) (#2427)
May 30, 2025
cade41d
Add Offoading to llvm-readobj and llvm-objcopy
david-salinas Aug 28, 2024
74010fc
[AMDGPU][clang] provide device implementation for __builtin_logb and …
choikwa May 19, 2025
7101fb3
[AMDGPU] Make the iterative schedulers selectable via amdgpu-sched-st…
jrbyrnes Apr 10, 2025
b5f132a
[AMDGPU] Teach iterative schedulers about IGLP (#134953)
jrbyrnes Apr 11, 2025
bdfe571
[ScheduleDAG] Allow disabling the SchedModel / Itineraries during Sch…
jrbyrnes May 5, 2025
6f3b1b8
[Compiler CI/CD] fix psdb mainline trigger
skganesan008 Jun 3, 2025
1c23824
[Compiler CI/CD] fix psdb amd-mainline trigger (#2469)
skganesan008 Jun 3, 2025
d23d92d
[OpenMP] Fix atomic compare handling with overloaded operators (#141142)
jhuber6 May 29, 2025
cf61b59
Get the type of the used array from its elements.
Jan 29, 2025
1015967
Temporarily unconditionally emit used / compiler.used array elements …
searlmc1 Jan 24, 2025
6a62f34
You shall now pass!
Jan 29, 2025
b4f6f2b
[AMDGPU] Handle lowering addrspace casts from LDS to FLAT address in …
skc7 Feb 19, 2025
0c5187d
[AMDGPU] Lower LDS in functions without sanitize_address in amdgpu-sw…
skc7 Apr 9, 2025
c22262c
[Flang][MLIR][OpenMP] Fix Target Data if (present(...)) causing LLVM-…
agozillon Jan 30, 2025
4917540
[AMDGPU][clang] provide device implementation for __builtin_logb and …
ronlieb Jun 6, 2025
f5dcf1e
Add Offoading to llvm-readobj and llvm-objcopy (#2455)
ronlieb Jun 6, 2025
a0b7748
Disable flang warning, its erroneously firing on all uses
ronlieb Jun 8, 2025
5b0c3d9
MachineCopyPropagation: Do not remove copies preserved by regmask (#1…
shiltian Jun 9, 2025
d9acca4
[OpenMP] Fix atomic compare handling with overloaded operators (#1411…
ronlieb Jun 9, 2025
7ec8b71
Disable flang warning, its erroneously firing on all uses (#2546)
ronlieb Jun 9, 2025
236ab35
[Flang][MLIR][OpenMP] Fix Target Data if (present(...)) causing LLVM-…
ronlieb Jun 9, 2025
3ab962c
[COMGR] Add VFS support for device library linking (#2472)
chinmaydd Jun 9, 2025
8dc9657
[flang][OpenMP] Use function symbol on DECLARE TARGET (#134107)
kparzysz Apr 2, 2025
4da8cee
rocm 70 : amdflang-legacy -> amdflang-classic
ronlieb Feb 24, 2025
3a48adc
rocm 70 : amdflang-legacy -> amdflang-classic (#2589)
ronlieb Jun 12, 2025
df93f70
[flang][OpenMP] Use function symbol on DECLARE TARGET (#134107) (#2584)
ronlieb Jun 12, 2025
ce5378f
[offload] Remove bad assert in StaticLoopChunker::Distribute (#132705)
macurtis-amd Mar 28, 2025
d317f02
[Flang][OpenMP] Add PointerAssociateScalar to Cray Pointer used in th…
Thirumalai-Shaktivel Mar 29, 2025
49a9ae4
Cherry-pick collection of spir-v and hipstdpar related patches. (#2603)
searlmc1 Jun 12, 2025
064591f
[Flang][OpenMP] Add PointerAssociateScalar to Cray Pointer used in th…
ronlieb Jun 12, 2025
ae71285
[offload] Remove bad assert in StaticLoopChunker::Distribute (#132705…
ronlieb Jun 12, 2025
1a2e2ae
[Flang][MLIR] - Handle the mapping of subroutine arguments when they …
bhandarkar-pranav May 13, 2025
6a5af62
[MLIR][OpenMP] Use correct DebugLoc in target construct callbacks. (#…
abidh Feb 5, 2025
0edb80c
[flang][OMPIRbuilder] Set debug loc on terminator created by splitBB.…
abidh Feb 5, 2025
d6a61ad
[Flang][OpenMP] Update semantics checks for 'teams' nesting (#126922)…
estewart08 Jun 13, 2025
c8272c5
[XRay] Fix argument parsing with offloading (#140748) (#141043)
tcgu-amd Jun 13, 2025
2b6562d
[Flang][MLIR] - Handle the mapping of subroutine arguments when they …
bhandarkar-pranav May 13, 2025
37dca93
[Flang] Generate math ops for non-precise calls to acosh, asin, asinh…
jsjodin Feb 13, 2025
4296542
Revert "[Flang][OpenMP] Add PointerAssociateScalar to Cray Pointer us…
ronlieb Jun 14, 2025
4c4dca8
Revert "[Flang][OpenMP] Add PointerAssociateScalar to Cray Pointer us…
ronlieb Jun 14, 2025
cc42135
[Flang][MLIR] - Handle the mapping of subroutine arguments when they …
ronlieb Jun 14, 2025
ddde5bd
Amd/dev/rlieberm/handle mapping (#2607)
ronlieb Jun 14, 2025
a06c2c8
[AMDGPU] Fix scale opsel flags for scaled MFMA operations (#140183) (…
bcahoon Jun 16, 2025
b09f4db
[AMDGPU] Bugfix for scaled MFMA parsing FP literals (#142493)
VigneshwarJ Jun 4, 2025
faf1611
[AMDGPU] Relax lds dma waitcnt with no aliasing pair (#131842) (#2463)
bcahoon Jun 17, 2025
7755b9f
[flang][driver] Introduce FCC_OVERRIDE_OPTIONS. (#140556)
abidh Jun 2, 2025
4b1b232
[AMDGPU] Bugfix for scaled MFMA parsing FP literals (#142493) (#2660)
VigneshwarJ Jun 17, 2025
f3128a9
[AMDGPU] Fix SIFoldOperandsImpl::canUseImmWithOpSel() for VOP3 packed…
bcahoon Jun 17, 2025
1e4eb9f
[libomptarget] [ompt] Added envars controlling buffer flush.
dhruvachak May 1, 2025
b3dac00
[libomptarget] [ompt] Added envars controlling buffer flush. (#2693)
ronlieb Jun 18, 2025
37783b6
[flang][driver] Introduce FCC_OVERRIDE_OPTIONS. (#140556) (#2671)
ronlieb Jun 18, 2025
34e6e84
[Flang] Generate math ops for non-precise calls to acosh, asin, asinh…
estewart08 Jun 18, 2025
4f49035
[IndVarsSimplify] sinkUnusedInvariants is skipping instructions while…
srpande Jun 18, 2025
2da56cf
AMDGPU: Make v2f32 -> v2f16 legal when target supports v_cvt_pk_f16_f…
changpeng Jun 18, 2025
d2408e7
[openmp] - add support for openmp-config.cmake
estewart08 May 12, 2025
46702f1
[AMDGPU] Don't unify divergent exit nodes with `musttail` calls (#126…
shiltian Feb 10, 2025
4c5731b
[Comgr] Add new Action to compile SPIR-V to Relocatable
lamb-j Jun 18, 2025
36f33cb
[openmp] - add support for openmp-config.cmake (#2706)
ronlieb Jun 19, 2025
b701d28
[OffloadBundler] Rework the ctor of `OffloadTargetInfo` to support AM…
lamb-j Jun 19, 2025
0203927
[AMDGPU] Don't unify divergent exit nodes with `musttail` calls (#126…
CatherineMoore Jun 19, 2025
d0c8ed3
Fix SPIR-V invalid casts for @__hip_cuid_ to ptr in @llvm.used (#2461)
jmmartinez Jun 19, 2025
d428028
[OMPIRBuilder][debug] Don't drop debug info for loop constructs. (#14…
abidh Jun 17, 2025
1dc2de7
Reland: [LoopUnroll] Make use of MaxTripCount for loops with pragma u…
doru1004 May 21, 2025
0ca677b
[OMPIRBuilder][debug] Don't drop debug info for loop constructs. (#14…
ronlieb Jun 19, 2025
0d8974d
[AMDGPU] Teach iterative schedulers about IGLP (#134953) (#2464)
bcahoon Jun 20, 2025
66bacea
[HIP] Remove dots in HIP runtime path (#143792)
scchan Jun 19, 2025
cb30313
[Clang][OpenMP] Process debug info for reduction - fix SWDEV-535727
Kewen12 Jun 16, 2025
67a3e1c
[HIPCC] Stop including perl scripts in hipcc packages (#2474)
Jun 20, 2025
75e7bb2
suppress intermittent mlir dbg test for autopromo to pass
ronlieb Jun 20, 2025
f8c83de
[Clang][OpenMP] Process debug info for reduction - fix SWDEV-535727 (…
ronlieb Jun 20, 2025
1f3ce4a
suppress intermittent mlir dbg test for autopromo to pass (#2751)
ronlieb Jun 21, 2025
64e29ec
[HIP] Remove dots in HIP runtime path (#143792) (#2749)
scchan Jun 23, 2025
e56c3d4
[SWDEV-520916] amdgpu-sw-lower-lds amd-mainline PRs (#2516)
skc7 Jun 24, 2025
4640ce4
Port #140158 whilst upstream review continues (#2675)
searlmc1 Jun 24, 2025
42894d3
[Reland][InstCombine] Iterative replacement in PtrReplacer (#145410)
gandhi56 Jun 24, 2025
7bd7e24
[Comgr] Fix SPIRV tests
lamb-j Jun 24, 2025
f0c1f2b
SWDEV-527443 - Port NV4x Intrinsic changes for ROCm 7.0 (#2801)
chinmaydd Jun 24, 2025
4c8ef2e
Fix unused variable warning. NFC. (#2788)
shiltian Jun 25, 2025
21033ba
Port ZCFS improvements (#2817)
searlmc1 Jun 25, 2025
d7ddf53
[openmp] - Ensure openmp-config.cmake honors proper install suffix (#…
estewart08 Jun 25, 2025
4cd1404
Fix all compile warnings (#2832)
shiltian Jun 26, 2025
a05f8bb
StructurizeCFG: Use poison instead of undef (#130459)
shiltian Jun 23, 2025
1116f71
[NFC][AMDGPU] Add a test that can cause backend crash (#145777)
shiltian Jun 26, 2025
66e9c65
[NFC] Update check lines for `llvm/test/CodeGen/AMDGPU/undef-handling…
shiltian Jun 26, 2025
3b74517
[OpenMP] Remove declaration and usage of __AMDGCN_WAVEFRONT_SIZE (#14…
CatherineMoore Jun 26, 2025
503b0dd
[AMDGPU] Fix to prevent sinking of PERMLANE_SWAP instruction (#144423…
VigneshwarJ Jun 26, 2025
fbbfa1f
Reland: [LoopUnroll] Make use of MaxTripCount for loops with pragma u…
ronlieb Jun 26, 2025
63fa22f
[AMDGPU] Max. WG size-induced occupancy limits max. waves/EU (#2711)
lucas-rami Jun 26, 2025
1aa5652
StructurizeCFG: Use poison instead of undef (#130459) (#2793)
shiltian Jun 26, 2025
bbc64d3
Revert "Revert "[RemoveDIs] Enable direct-to-bitcode writing by defau…
epilk Jun 27, 2025
3987e44
[openmp] - Rework openmp-config for portability (#2857)
estewart08 Jun 27, 2025
2ba33cb
[clang] fix matching of nested template template parameters (#130447)…
changpeng Jul 2, 2025
55f615e
[InstCombine] Iterative replacement in PtrReplacer (#145410) (#2803)
gandhi56 Jul 2, 2025
a74ba40
[openmp] - Add openmp_INCLUDE_DIR to openmp-config.cmake (#2972)
estewart08 Jul 5, 2025
6b86029
[Comgr] Remove file reorg backward compatibility macro (#2989)
lamb-j Jul 8, 2025
f6b8388
[AMDGPU][MC] Fix disassembly for v_permlane16_swap_b32 for GFX950 (#1…
Jul 11, 2025
a3d0e5f
[AMDGPU] Visit all PHIs in each call to optimizeLiveType (#3030)
kzhuravl Jul 11, 2025
c1a4a7a
disable comgr lit due to build-infra
ronlieb Jul 12, 2025
03b6e54
disable comgr lit due to build-infra (#3076)
ronlieb Jul 12, 2025
511f469
[AMDGPU] Move kernarg preload logic to separate pass (#130434) (#2992)
bcahoon Jul 15, 2025
50d2f40
[openmp] - Changes to support building openmp/offload in LLVM_ENABLE_…
estewart08 Apr 15, 2025
6c63bea
add sentinel value to flang/EnableFlangBuild to disable classic
ronlieb Feb 22, 2025
ff58d3a
[amdllvm] - Remove symlink for flang-classic
estewart08 Jul 10, 2025
a05f5ad
Amd/dev/estewart/cherry pick openmp offload runtime build support (#3…
estewart08 Jul 16, 2025
d45f1ae
AMDGPU: Fix runtime unrolling when cascaded GEPs present (#147700)
macurtis-amd Jul 10, 2025
fd89fa2
[runtimes] - Use extra_cmake_args for rocr/rocm-device-libs external …
estewart08 Jul 16, 2025
1a2974a
AMDGPU: Fix runtime unrolling when cascaded GEPs present (#147700) (#…
ronlieb Jul 17, 2025
579872b
Revert "AMDGPU: Fix runtime unrolling when cascaded GEPs present (#14…
ronlieb Jul 17, 2025
e13573e
[InstCombine]PtrReplacer: Correctly handle select with unavailable op…
Pierre-vh Jul 17, 2025
5e2093d
Revert "AMDGPU: Fix runtime unrolling when cascaded GEPs present (#14…
ronlieb Jul 17, 2025
1b038f2
Revert "[StructurizeCFG] Refactor insertConditions. NFC. (#115476)" (…
changpeng Jul 18, 2025
3269740
[runtimes] - Use cmake semicolon generator expression for rocr-runtim…
estewart08 Jul 22, 2025
9364925
[NFC][AMDGPU] Auto generate check lines for `llvm/test/CodeGen/AMDGPU…
shiltian May 1, 2025
734c629
[AMDGPU][Attributor] Rework update of `AAAMDWavesPerEU` (#123995)
shiltian May 17, 2025
9a093fe
[SLPVectorizer] Use accurate cost for external users of resize shuffl…
bcahoon Jul 24, 2025
e53b1a5
[Comgr] Turn LIT tests back on
lamb-j Jul 25, 2025
42375bd
SWDEV-522811 - fix compress/decompress in LLVM Offloading API
Jul 25, 2025
d852555
[AMDGPU] Allow unaligned VGPR for ds_read_b96_tr_b6 (#125169)
Jan 31, 2025
dd188aa
[AMDGPU] Allow unaligned VGPR for ds_read_b96_tr_b6 (#125169) (#3329)
Jul 30, 2025
b600e2d
[Comgr] Fix VGPR values for gfx90a, gfx942, gfx950, gfx9-4-generic
lamb-j Jul 30, 2025
3dce363
[AMDGPU] Fix op_sel settings for v_cvt_scale32_* and v_cvt_sr_* (#151…
changpeng Aug 1, 2025
cdc1ea8
[Comgr] Handle amdgcnspirv when used as an ISA name
lamb-j Jun 19, 2025
7d26f23
[Comgr] Support 4-field-triple for spirv
lamb-j Aug 4, 2025
7efab10
[AMDGPU] Ensure non-reserved CSR spilled regs are live-in (#146427)
macurtis-amd Aug 1, 2025
9c796bd
[HIP] Claim `--offload-compress` for `-M` (#133456)
yxsamliu Apr 2, 2025
3a0e77d
[CUDA][HIP] capture possible ODR-used var (#136645)
yxsamliu Apr 23, 2025
7f61756
[AMDGPU][Attributor] Rework update of `AAAMDWavesPerEU` (#123995) (#3…
bcahoon Aug 5, 2025
1ab8080
[AMDGPU] SelDAG: fix lowering of undefined workitem intrinsics (#1260…
choikwa Aug 5, 2025
60f9156
[CUDA][HIP] capture possible ODR-used var (#136645) (#3443)
yxsamliu Aug 8, 2025
7e0e42b
[HIP] Claim `--offload-compress` for `-M` (#133456) (#3442)
yxsamliu Aug 8, 2025
d785555
[HIP] compressed bundle format defaults to v3
yxsamliu Aug 7, 2025
eeb5d84
[HIP] compressed bundle format defaults to v3 (#3503)
yxsamliu Aug 8, 2025
185ddcf
Cherry-picking 71d6762309a7db67770bdbd39572ef04e6a1ea59
zGoldthorpe Aug 14, 2025
94b2eae
Reapply "[CUDA][HIP] Add a __device__ version of std::__glibcxx_asser…
jmmartinez Aug 14, 2025
2b07f34
[AMDGPU] Ensure non-reserved CSR spilled regs are live-in (#146427) (…
ronlieb Aug 18, 2025
1b5ca05
[AMDGPU] Hoist readlane/readfirstlane through unary/binary operands (…
VigneshwarJ Aug 18, 2025
12fb44f
SWDEV-465041 - Enable queue write index programming (#2404) (#3749)
kzhuravl Aug 27, 2025
9461df3
[NFC][OffloadBundle] Fix compile warnings (#3700)
shiltian Aug 27, 2025
c7019de
[AMDGPU][Attributor] Remove final update of waves-per-eu after the at…
shiltian Aug 28, 2025
7757f13
[NFC] Fix compile warnings in `llvm/unittests/Object/OffloadingBundle…
shiltian Aug 28, 2025
2c69a9e
[InstCombine] Added pattern for recognising the construction of packe…
zGoldthorpe Sep 3, 2025
cc250b8
[compiler-rt]: fix CodeQL errors (#3798)
ampandey-AMD Sep 5, 2025
a76bc2d
[ASan] Intercept rocr api's 'hsa_amd_vmem_address_reserve_align' and …
ampandey-AMD Sep 5, 2025
9227599
[NFC][Clang] Fix typo in feature predicate insertion
lamb-j Sep 5, 2025
b2c1136
[AMDGPU] Fix a crash by skipping DBG instrs at start of sched region …
JanekvO Sep 11, 2025
320bbe6
[AMDGPU] Make ds/global load intrinsics IntrArgMemOnly (#152792) (#3748)
choikwa Sep 11, 2025
a96e394
[AMDGPU] Restrict to VGPR only for mfma scale operands (#158117) (#4011)
changpeng Sep 12, 2025
afe89d2
[Comgr] Fix memory leak in name expression API
lamb-j Sep 15, 2025
1e6a516
[hipcc] Remove PERL scripts, add GitHub repo link
lamb-j Sep 15, 2025
f9f614c
Hostexec: Enable building with cmake4
HereThereBeDragons Sep 16, 2025
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
The table of contents is too big for display.
Diff view
Diff view
  •  
  •  
  •  
91 changes: 56 additions & 35 deletions .github/workflows/PSDB-amd-mainline.yml
Original file line number Diff line number Diff line change
@@ -1,5 +1,4 @@
# This workflow is used to invoke the PSDB jenkins job for ROCm Compiler CI. The python script can be used to invoke any jenkins job but input params needs to be configured properly
name: Compiler CI test
name: Compiler CI PSDB trigger on amd-mainline branch

# Controls when the workflow will run
on:
Expand All @@ -10,46 +9,64 @@ on:

# A workflow run is made up of one or more jobs that can run sequentially or in parallel, below is a single job called invoke jenkins jobs
jobs:
# This workflow contains a single job called "build"
# This workflow contains a single job called "invoke_jenkins_PSDB"
invoke_jenkins_PSDB:
# The type of runner that the job will run on. For github hosted runner use (${{ 'ubuntu-latest' }}) or self-hosted for sel-hosted runner.
#runs-on: ubuntu-latest
if: github.event.pull_request.draft == false
runs-on: self-hosted
container:
image: compute-artifactory.amd.com:5000/rocm-base-images/ghemu-action-ubuntu-24.04:2024101101
runs-on:
group: compiler-generic-runners
env:
svc_acc_org_secret: ${{'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'}}
svc_acc_org_secret: ${{secrets.CI_GITHUB_TOKEN}}
input_sha: ${{ github.event.pull_request.head.sha != '' && github.event.pull_request.head.sha || github.sha }}
input_pr_num: ${{ github.event.pull_request.number != '' && github.event.pull_request.number || 0 }}
input_pr_url: ${{ github.event.pull_request.html_url != '' && github.event.pull_request.html_url || '' }}
input_pr_title: ${{ github.event.pull_request.title != '' && github.event.pull_request.title || '' }}
# set the pipeline name here based on branch name
pipeline_name: ${{ 'compiler-psdb-amd-mainline' }}
JENKINS_URL: ${{'https://compiler-ci.amd.com/'}}
pipeline_name: ${{secrets.CI_JENKINS_MAINLINE_JOB_NAME}}
JENKINS_URL: ${{secrets.CI_JENKINS_URL}}
CONTAINER_IMAGE: ${{ secrets.JENKINS_TRIGGER_DOCKER_IMAGE }}

# Steps represent a sequence of tasks that will be executed as part of the job
steps:
# Checks-out your repository under $GITHUB_WORKSPACE, so your job can access it
- name: Check out rocm_ci_infra private repo
uses: actions/checkout@main
with:
#ref: ci-utils-dev-siva
#fetch-depth: 2
repository: AMD-Lightning-Internal/ci-utils
token: ${{ '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' }}
#token: ${{'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'}}
steps:

# Checks-out your repository under $GITHUB_WORKSPACE, so your job can access it
- name: Set environment variable for container image
run: |
echo "CONTAINER_IMAGE=${{ secrets.JENKINS_TRIGGER_DOCKER_IMAGE }}" >> $GITHUB_ENV
echo "CONTAINER_NAME=my_container_${{ github.run_id }}" >> $GITHUB_ENV


- name: Pull container image
run: docker pull "${{env.CONTAINER_IMAGE}}"


- name: Run container
run: |
docker run -d --name "${{env.CONTAINER_NAME}}" $CONTAINER_IMAGE sleep infinity
#docker exec "${{env.CONTAINER_NAME}}" /bin/bash -c "git clone ${{secrets.CI_UTILS_REPO}} ."
docker exec "${{env.CONTAINER_NAME}}" /bin/bash -c "echo 'Running commands inside the container'"

- name: Escape pull request title
run: |
import json
import os
import shlex
with open('${{ github.event_path }}') as fh:
event = json.load(fh)
escaped = event['pull_request']['title']
with open(os.environ['GITHUB_ENV'], 'a') as fh:
print(f'PR_TITLE={escaped}', file=fh)
shell: python3 {0}

- name: Run Jenkins Cancel Script
env:
JENKINS_URL: ${{ 'https://compiler-ci.amd.com/' }}
JENKINS_USER: ${{ 'z1_cciauto' }}
JENKINS_API_TOKEN: ${{ '11bdb3dcd61f1a00f3999c8e3a0d6da9a7' }}
JENKINS_JOB_NAME: ${{ 'compiler-psdb-amd-mainline' }}
JENKINS_URL: ${{secrets.CI_JENKINS_URL}}
JENKINS_USER: ${{secrets.CI_JENKINS_USER}}
JENKINS_API_TOKEN: ${{secrets.CI_JENKINS_TOKEN}}
JENKINS_JOB_NAME: ${{secrets.CI_JENKINS_JOB_NAME}}
PR_NUMBER: ${{ github.event.pull_request.number }}
COMMIT_HASH: ${{ github.event.after }}
run: |
python3 cancel_previous_build.py
docker exec -e JENKINS_JOB_NAME=${{secrets.CI_JENKINS_JOB_NAME}} -e PR_NUMBER=${{ github.event.pull_request.number }} -e COMMIT_HASH=${{ github.event.after }} -e JENKINS_URL=${{secrets.CI_JENKINS_URL}} -e JENKINS_USER=${{secrets.CI_JENKINS_USER}} -e JENKINS_API_TOKEN=${{secrets.CI_JENKINS_TOKEN}} "${{env.CONTAINER_NAME}}" /bin/bash -c "PYTHONHTTPSVERIFY=0 python3 cancel_previous_build.py"


# Runs a set of commands using the runners shell
Expand All @@ -65,21 +82,25 @@ jobs:
echo "GITHUB_REF_NAME is: $GITHUB_REF_NAME"
echo "github.event.pull_request.id is: ${{github.event.pull_request.id}}"
echo "github.event.pull_request.html_url is: ${{github.event.pull_request.html_url}}"
echo "github.event.pull_request.number is: ${{github.event.pull_request.number}}"
echo "github.event.pull_request.number is: ${{github.event.pull_request.number}}"
echo "github.event.pull_request.url is: ${{github.event.pull_request.url}}"
echo "github.event.pull_request.issue_url is: ${{github.event.pull_request.issue_url}}"
echo "github.event.pull_request.comments_url is: ${{github.event.pull_request.comments_url}}"
echo "github.event.pull_request.statuses_url is: ${{github.event.pull_request.statuses_url}}"
echo "github.event.pull_request.head.sha is: ${{github.event.pull_request.head.sha}}"
echo "github.event.pull_request.base.ref is: ${{github.event.pull_request.base.ref}}"
echo "github.event.pull_request.merge_commit_sha is: ${{github.event.pull_request.merge_commit_sha}}"
echo "github.event.pull_request is: ${{github.event.pull_request}}"
pip3 show python-jenkins || echo "python-jenkins is not installed"
#sudo -H pip3 install --upgrade python-jenkins
# pipeline name shuould be unique to the workfow yml for a given repository
#curl -L -X POST -H "Accept: application/vnd.github+json" -H "Authorization: Bearer ${{'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'}}" ${{github.event.pull_request.comments_url}} -d '{"body":"Github action triggered jenkins job for compute-psdb-staging-smi-libs-ghemu "}'


- name: Trigger Jenkins Pipeline
if: steps.check_changes.outcome != 'failure'
run: |
echo "running jenkins_api.py with input sha - $input_sha for pull request - $input_pr_url"
python3 jenkins_api.py -s $JENKINS_URL -jn $pipeline_name -ghr $GITHUB_REPOSITORY -ghsha $input_sha -ghprn $input_pr_num -ghpru "$input_pr_url" -ghprt "$input_pr_title" -ghpat $svc_acc_org_secret
echo "--Running jenkins_api.py with input sha - $input_sha for pull request - $input_pr_url"
docker exec -e GITHUB_REPOSITORY="$GITHUB_REPOSITORY" -e svc_acc_org_secret="$svc_acc_org_secret" -e input_sha="$input_sha" -e input_pr_url="$input_pr_url" -e pipeline_name="$pipeline_name" \
-e input_pr_num="$input_pr_num" -e PR_TITLE="$PR_TITLE" -e JENKINS_URL="$JENKINS_URL" -e GITHUB_PAT="$svc_acc_org_secret" "${{env.CONTAINER_NAME}}" \
/bin/bash -c 'echo \"PR NUM: "$input_pr_num"\" && PYTHONHTTPSVERIFY=0 python3 jenkins_api.py -s \"${JENKINS_URL}\" -jn "$pipeline_name" -ghr "$GITHUB_REPOSITORY" -ghsha "$input_sha" -ghprn "$input_pr_num" -ghpru "$input_pr_url" -ghprt "$PR_TITLE" -ghpat="$svc_acc_org_secret"'

- name: Stop and remove container
if: always()
run: |
docker stop "${{env.CONTAINER_NAME}}"
docker rm "${{env.CONTAINER_NAME}}"
85 changes: 0 additions & 85 deletions .github/workflows/PSDB-amd-staging.yml

This file was deleted.

33 changes: 10 additions & 23 deletions amd/comgr/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -38,7 +38,6 @@ endif(EXISTS "${CMAKE_SOURCE_DIR}/../../.git")

include(GNUInstallDirs)
include(CMakePackageConfigHelpers)
option(FILE_REORG_BACKWARD_COMPATIBILITY "Enable File Reorg with backward compatibility" OFF)
# Optionally, build Compiler Support with ccache.
set(ROCM_CCACHE_BUILD OFF CACHE BOOL "Set to ON for a ccache enabled build")
if (ROCM_CCACHE_BUILD)
Expand Down Expand Up @@ -71,6 +70,10 @@ option(COMGR_BUILD_SHARED_LIBS "Build the shared library"
${build_shared_libs_default})

set(SOURCES
src/comgr-cache.cpp
src/comgr-cache-command.cpp
src/comgr-cache-bundler-command.cpp
src/comgr-clang-command.cpp
src/comgr-compiler.cpp
src/comgr.cpp
src/comgr-device-libs.cpp
Expand All @@ -81,6 +84,7 @@ set(SOURCES
src/comgr-metadata.cpp
src/comgr-objdump.cpp
src/comgr-signal.cpp
src/comgr-spirv-command.cpp
src/comgr-symbol.cpp
src/comgr-symbolizer.cpp
src/time-stat/time-stat.cpp)
Expand Down Expand Up @@ -178,9 +182,12 @@ message("")
option(COMGR_DISABLE_SPIRV "To disable SPIRV in Comgr" OFF)

if (NOT COMGR_DISABLE_SPIRV)
CHECK_INCLUDE_FILE_CXX(LLVMSPIRVLib/LLVMSPIRVLib.h HAVE_LLVMSPIRVLIB_H)
if (NOT HAVE_LLVMSPIRVLIB_H)
# TODO: Explore switching this to CHECK_INCLUDE_FILE_CXX() macro
if (NOT EXISTS "${LLVM_INCLUDE_DIRS}/LLVMSPIRVLib/LLVMSPIRVLib.h")
message("-- LLVMSPIRVLib/LLVMSPIRVLib.h not found")
set(COMGR_DISABLE_SPIRV ON)
else()
message("-- LLVMSPIRVLib/LLVMSPIRVLib.h found")
endif()
endif()

Expand Down Expand Up @@ -319,26 +326,6 @@ install(FILES
COMPONENT amd-comgr
DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/${amd_comgr_NAME})

#File reorg Backward compatibility function
if(FILE_REORG_BACKWARD_COMPATIBILITY)
# To enable/disable #error in wrapper header files
if(NOT DEFINED ROCM_HEADER_WRAPPER_WERROR)
if(DEFINED ENV{ROCM_HEADER_WRAPPER_WERROR})
set(ROCM_HEADER_WRAPPER_WERROR "$ENV{ROCM_HEADER_WRAPPER_WERROR}"
CACHE STRING "Header wrapper warnings as errors.")
else()
set(ROCM_HEADER_WRAPPER_WERROR "OFF" CACHE STRING "Header wrapper warnings as errors.")
endif()
endif()
if(ROCM_HEADER_WRAPPER_WERROR)
set(deprecated_error 1)
else()
set(deprecated_error 0)
endif()

include(comgr-backward-compat.cmake)
endif()

if(ENABLE_ASAN_PACKAGING)
install(FILES
"LICENSE.txt"
Expand Down
36 changes: 35 additions & 1 deletion amd/comgr/README.md
Original file line number Diff line number Diff line change
Expand Up @@ -125,7 +125,27 @@ These include:
certain runtime headers. If this is not set, it has a default value of
"${ROCM_PATH}/llvm".

Comgr also supports some environment variables to aid in debugging. These
Comgr utilizes a cache to preserve the results of compilations between executions.
The cache's status (enabled/disabled), storage location for its results,
and eviction policy can be manipulated through specific environment variables.
If an issue arises during cache initialization, the execution will proceed with
the cache turned off.

By default, the cache is enabled.

* `AMD_COMGR_CACHE`: When unset or set to a value different than "0", the cache is enabled.
Disabled when set to "0".
* `AMD_COMGR_CACHE_DIR`: If assigned a non-empty value, that value is used as
the path for cache storage. If the variable is unset or set to an empty string `""`,
it is directed to "$XDG_CACHE_HOME/comgr" (which defaults to
"$USER/.cache/comgr" on Linux, and "%LOCALAPPDATA%\cache\comgr"
on Microsoft Windows).
* `AMD_COMGR_CACHE_POLICY`: If assigned a value, the string is interpreted and
applied to the cache pruning policy. The cache is pruned only upon program
termination. The string format aligns with [Clang's ThinLTO cache pruning policy](https://clang.llvm.org/docs/ThinLTO.html#cache-pruning).
The default policy is set as: "prune_interval=1h:prune_expiration=0h:cache_size=75%:cache_size_bytes=30g:cache_size_files=0".

Comgr supports some environment variables to aid in debugging. These
include:

* `AMD_COMGR_SAVE_TEMPS`: If this is set, and is not "0", Comgr does not delete
Expand All @@ -143,6 +163,20 @@ include:
* `AMD_COMGR_TIME_STATISTICS`: If this is set, and is not "0", logs will
include additional Comgr-specific timing information for compilation actions.

Comgr implements support for an in-memory, virtual filesystem (VFS) for storing
temporaries generated during intermediate compilation steps. This is aimed at
improving performance by reducing on-disk file I/O. Currently, VFS is only supported
for the device library link step, but we aim to progressively add support for
more actions.

By default, VFS is turned on.

* `AMD_COMGR_USE_VFS`: When set to "0", VFS support is turned off.
* Users may use the API `amd_comgr_action_info_set_vfs` to disable VFS for individual actions
without having to modify system-wide environment variables.
* If `AMD_COMGR_SAVE_TEMPS` is set and not "0", VFS support is turned off irrespective
of `AMD_COMGR_USE_VFS` or the use of `amd_comgr_action_info_set_vfs`.

Versioning
----------

Expand Down
4 changes: 2 additions & 2 deletions amd/comgr/VERSION.txt
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
#COMGR_VERSION_MAJOR
2
3
#COMGR_VERSION_MINOR
9
0
14 changes: 14 additions & 0 deletions amd/comgr/cmake/DeviceLibs.cmake
Original file line number Diff line number Diff line change
Expand Up @@ -59,6 +59,7 @@ foreach(AMDGCN_LIB_TARGET ${AMD_DEVICE_LIBS_TARGETS})
add_dependencies(amd_comgr ${AMDGCN_LIB_TARGET}_header)

list(APPEND TARGETS_INCLUDES "#include \"${header}\"")
list(APPEND TARGETS_HEADERS "${INC_DIR}/${header}")
endforeach()

list(JOIN TARGETS_INCLUDES "\n" TARGETS_INCLUDES)
Expand Down Expand Up @@ -110,4 +111,17 @@ list(APPEND TARGETS_DEFS "#undef AMD_DEVICE_LIBS_FUNCTION")
list(JOIN TARGETS_DEFS "\n" TARGETS_DEFS)
file(GENERATE OUTPUT ${GEN_LIBRARY_DEFS_INC_FILE} CONTENT "${TARGETS_DEFS}")

# compute the sha256 of the device libraries to detect changes and pass them to comgr (used by the cache)
find_package(Python3 REQUIRED Interpreter)
set(DEVICE_LIBS_ID_SCRIPT "${CMAKE_CURRENT_SOURCE_DIR}/cmake/device-libs-id.py")
set(DEVICE_LIBS_ID_HEADER ${INC_DIR}/libraries_sha.inc)
add_custom_command(OUTPUT ${DEVICE_LIBS_ID_HEADER}
COMMAND ${Python3_EXECUTABLE} ${DEVICE_LIBS_ID_SCRIPT} --varname DEVICE_LIBS_ID --output ${DEVICE_LIBS_ID_HEADER} ${TARGETS_HEADERS}
DEPENDS ${DEVICE_LIBS_ID_SCRIPT} ${TARGETS_HEADERS}
COMMENT "Generating ${INC_DIR}/libraries_sha.inc"
)
set_property(DIRECTORY APPEND PROPERTY ADDITIONAL_MAKE_CLEAN_FILES ${INC_DIR}/libraries_sha.inc)
add_custom_target(libraries_sha_header DEPENDS ${INC_DIR}/libraries_sha.inc)
add_dependencies(amd_comgr libraries_sha_header)

include_directories(${INC_DIR})
Loading