Migration doc fixes by jayhawk-commits · Pull Request #4 · ROCm/rocm-systems

jayhawk-commits · 2025-04-07T22:21:04Z

No description provided.

The Omnitrace program is being renamed. Full name: "ROCm Systems Profiler" Package name: "rocprofiler-systems" Binary / Library names: "rocprof-sys-*" --------- Co-authored-by: Xuan Chen <xuchen@amd.com> Signed-off-by: David Galiffi <David.Galiffi@amd.com>

Fix for ROCm/legacy-rocm-build#3997. When compiling a C program that includes rdc/rdc.h, multiple assertion errors are thrown without this header included. Change-Id: Ie5b5c1a1a17c8207cf9b1be23b31193e260d5c1a Co-authored-by: harkgill-amd <harkgill@amd.com>

* Add rocDecode API support * Update CI * CI update: Sanitizers run on ubuntu 22.04 --------- Co-authored-by: Jonathan R. Madsen <jonathanrmadsen@gmail.com>

Update LICENSE Co-authored-by: srawat <120587655+SwRaw@users.noreply.github.com>

#4) SWDEV-521313 Streamline Version Management: Retain Only version-rocm File

* Add rocDecode API support * Update CI * CI update: Sanitizers run on ubuntu 22.04 --------- Co-authored-by: Jonathan R. Madsen <jonathanrmadsen@gmail.com> [ROCm/rocprofiler-register commit: de3a874]

Update LICENSE Co-authored-by: srawat <120587655+SwRaw@users.noreply.github.com> [ROCm/rocprofiler commit: 8c62af8]

The Omnitrace program is being renamed. Full name: "ROCm Systems Profiler" Package name: "rocprofiler-systems" Binary / Library names: "rocprof-sys-*" --------- Co-authored-by: Xuan Chen <xuchen@amd.com> Signed-off-by: David Galiffi <David.Galiffi@amd.com> [ROCm/rocprofiler-systems commit: d07bf50]

Closes #4

Fix for ROCm/legacy-rocm-build#3997. When compiling a C program that includes rdc/rdc.h, multiple assertion errors are thrown without this header included. Change-Id: Ie5b5c1a1a17c8207cf9b1be23b31193e260d5c1a Co-authored-by: harkgill-amd <harkgill@amd.com> [ROCm/rdc commit: 83f36f1]

1.Remove clock functions from some tests that don't need them. 2.In some memory pool tests and coherency tests, timer-based kernel delay isn't reliable, use pinned host based notification instead. 3.Add CHECK_PCIE_ATOMICS_SUPPORT before some tests. 4.catch/unit/memory/hipMemoryAllocateCoherent.cc is removed as it is useless and originally excluded in building. 5.Some tests can still pass even if clock rate =0, thus they will be kept as is. 6.Some logic and format improvement in some tests. Change-Id: I6b3c6bf54c61cffd45cd6f17c75998f751b75725

#4) SWDEV-521313 Streamline Version Management: Retain Only version-rocm File [ROCm/rocm-core commit: cf86775]

1.Remove clock functions from some tests that don't need them. 2.In some memory pool tests and coherency tests, timer-based kernel delay isn't reliable, use pinned host based notification instead. 3.Add CHECK_PCIE_ATOMICS_SUPPORT before some tests. 4.catch/unit/memory/hipMemoryAllocateCoherent.cc is removed as it is useless and originally excluded in building. 5.Some tests can still pass even if clock rate =0, thus they will be kept as is. 6.Some logic and format improvement in some tests. Change-Id: I6b3c6bf54c61cffd45cd6f17c75998f751b75725 [ROCm/hip-tests commit: ec8ff45]

1. Get the bad pages threshold of a processor. 2. Verify the checksum of RAS EEPROM Signed-off-by: Meng Li <li.meng@amd.com>

1. Get the bad pages threshold of a processor. 2. Verify the checksum of RAS EEPROM Signed-off-by: Meng Li <li.meng@amd.com> [ROCm/amdsmi commit: dc400d9]

hyd-7c

hyd-7c [ROCm/rocshmem commit: d054fc6]

…views, optimize flat batch→SDMA 1. Replace last_graph_signal_ with GraphSignalPool::GetLastAcquired() query via gpu_.command()->graphSignalPool(). Eliminates dangling pointer risk (Copilot #1) — no retain/release needed since the pool outlives all graph operations and the CallbackCommand marker handles the boundary. 2. Add ResetLastAcquired() to GraphSignalPool. Called after SyncPlan HW event pre-allocation to prevent GetLastAcquired() from returning a pre-allocated (not yet dispatched) signal, which caused circular dependencies across parallel streams. 3. Fix data race in GraphSignalPool::Acquire/AcquireIrq (Copilot #3): replace signals_.size() reads with atomic capacity_/irq_capacity_ fields, updated after Allocate/GrowAndAcquire under lock. 4. Add null checks on Allocate/AllocateIrq/Acquire return values in EnqueueSegmentedGraph with hipErrorOutOfMemory early returns (Copilot #4). 5. Clear cmd->HwEvent() on ActiveSignal pool allocation failure (Copilot #5). 6. Flat batch → SDMA barrier elimination: - dispatchCurrentBatch accepts attach_signal; look-ahead checks if next node is non-captured hipGraphNodeTypeMemcpy. - attach_signal=true places a completion signal on the last flat batch packet so the SDMA engine waits via WaitingSignal, not a barrier. - hasPendingDispatch_ set to false when last packet has a signal. - Reorder setFenceDirty in dispatchAqlPacketBatchFlat to match single dispatch path (set dirty on reserve, clear if system scope). - Move attach_signal fixup before per-packet log loop so logs show the correct completion_signal. Made-with: Cursor

…#6076) ## Motivation Some of the comments were not correct ## Technical Details Will be more clear for the user ## JIRA ID   ## Test Plan CTest should pass ## Test Result Test project /home/rajy/work/rocm-systems/projects/rocdecode/build Start 1: video_decodeRaw-HEVC 1/6 Test #1: video_decodeRaw-HEVC ............. Passed 2.67 sec Start 2: video_decodeRaw-AVC 2/6 Test #2: video_decodeRaw-AVC .............. Passed 1.68 sec Start 3: video_decodeRaw-AV1 3/6 Test #3: video_decodeRaw-AV1 .............. Passed 1.77 sec Start 4: video_decodeRaw-VP9 4/6 Test #4: video_decodeRaw-VP9 .............. Passed 1.82 sec Start 5: rocdec_Decode-HEVC 5/6 Test #5: rocdec_Decode-HEVC ............... Passed 0.72 sec Start 6: rocDecode_Negative_API_Tests 6/6 Test #6: rocDecode_Negative_API_Tests ..... Passed 0.53 sec 100% tests passed, 0 tests failed out of 6 ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

…xing is difficult to read (#5453) ## Motivation The `[join_prof] Insufficient number of kernel calls` warning printed the kernels as one comma separated string. Since these kernels are extremely long and mangled C++ names, and we are not truncating them, the message was an extremely hard to read wall of text where individual kernels couldn't be visually separated. ## Technical Details In `src/rocprof_compute_analyze/analysis_base.py::detect_missing_counters`, we can format the kernel list as a numbered newline separated block instead of a single joined string showing the same data with no truncation. No behavioral changes. ## JIRA ID AIPROFCOMP-294 ## Test Plan Run in the rocprof-compute dev container with a Mistral-7B inference workload that triggers the warning: ```bash rocprof-compute profile -n mistral_iter --iteration-multiplexing -- \ python3 /root/torch_mistral_workload/torch.py rocprof-compute analyze -p workloads/mistral_iter/MI300X_A1 ``` Captured the analyze output before and after: ## Test Result Warning now lists each kernel on its own line and is fully readable. **Before** (develop): `WARNING Reading intermediate results_*.csv files is deprecated and will be removed in a future release. WARNING [join_prof] Insufficient number of kernel calls for kernels: __amd_rocclr_fillBufferAligned, __amd_rocclr_copyBuffer, void at::native::vectorized_elementwise_kernel<4, at::native::compare_scalar_kernel<long>(at::TensorIteratorBase&, at::native::(anonymous namespace)::OpType, long)::{lambda(long)#1}, std::array<char*, 2ul> >(int, at::native::compare_scalar_kernel<long>(at::TensorIteratorBase&, at::native::(anonymous namespace)::OpType, long)::{lambda(long)#1}, std::array<char*, 2ul>), void at::native::vectorized_elementwise_kernel<4, at::native::(anonymous namespace)::masked_fill_kernel(at::TensorIterator&, c10::Scalar const&)::{lambda()#1}::operator()() const::{lambda()#4}::operator()() const::{lambda(long, bool)#1}, std::array<char*, 3ul> >(int, at::native::(anonymous namespace)::masked_fill_kernel(at::TensorIterator&, c10::Scalar const&)::{lambda()#1}::operator()() const::{lambda()#4}::operator()() const::{lambda(long, bool)#1}, std::array<char*, 3ul>), Cijk_Ailk_Bljk_S_B_Bias_HA_S_SAV_UserArgs_MT16x16x128_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA4_GRVWB4_GSUAMB_GLS0_ISA942_IU1_K1_LDSTI0_LBSPPA256_LBSPPB512_LBSPPM0_LPA16_LPB8_LPM0_LRVW4_LWPMn1_MIAV0_MIWT1_1_MO40_NTn1_NTA0_NTB0_NTC0_NTD4_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK0_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_4, void at::native::(anonymous namespace)::CatArrayBatchedCopy<at::native::(anonymous namespace)::OpaqueType<4u>, unsigned int, 3, 64, 64>(at::native::(anonymous namespace)::OpaqueType<4u>*, at::native::(anonymous namespace)::CatArrInputTensorMetadata<at::native::(anonymous namespace)::OpaqueType<4u>, unsigned int, 64, 64>, at::native::(anonymous namespace)::TensorSizeStride<unsigned int, 4u>, int, unsigned int), void rocprim::ROCPRIM_400200_NS::detail::trampoline_kernel<rocprim::ROCPRIM_400200_NS::detail::wrapped_scan_config<rocprim::ROCPRIM_400200_NS::default_config, long>, (rocprim::ROCPRIM_400200_NS::detail::target_arch)942, rocprim::ROCPRIM_400200_NS::detail::scan_impl<(rocprim::ROCPRIM_400200_NS::detail::lookback_scan_determinism)0, false, false, rocprim::ROCPRIM_400200_NS::default_config, long const*, long*, long, std::plus<long>, long>(void*, unsigned long&, long const*, long*, long, unsigned long, std::plus<long>, ihipStream_t*, bool)::{lambda(auto:1, auto:2)#1}::operator()<std::integral_constant<bool, false>, std::integral_constant<bool, true> >(std::integral_constant<bool, false>, std::integral_constant<bool, true>) const::{lambda(auto:1)#2}, rocprim::ROCPRIM_400200_NS::detail::default_config_selector>(rocprim::ROCPRIM_400200_NS::detail::scan_impl<(rocprim::ROCPRIM_400200_NS::detail::lookback_scan_determinism)0, false, false, rocprim::ROCPRIM_400200_NS::default_config, long const*, long*, long, std::plus<long>, long>(void*, unsigned long&, long const*, long*, long, unsigned long, std::plus<long>, ihipStream_t*, bool)::{lambda(auto:1, auto:2)#1}::operator()<std::integral_constant<bool, false>, std::integral_constant<bool, true> >(std::integral_constant<bool, false>, std::integral_constant<bool, true>) const::{lambda(auto:1)#2}), void at::native::reduce_kernel<512, 1, at::native::ReduceOp<bool, at::native::func_wrapper_t<bool, at::native::and_kernel_cuda(at::TensorIterator&)::{lambda()#1}::operator()() const::{lambda()#12}::operator()() const::{lambda(bool, bool)#1}>, unsigned int, bool, 4, 4> >(at::native::ReduceOp<bool, at::native::func_wrapper_t<bool, at::native::and_kernel_cuda(at::TensorIterator&)::{lambda()#1}::operator()() const::{lambda()#12}::operator()() const::{lambda(bool, bool)#1}>, unsigned int, bool, 4, 4>) to collect all counters using iteration multiplexing. These kernels do not have enough dispatches to fill all 13 counter sets. Use kernel filtering (-k) to profile only the kernels of interest, e.g.: rocprof-compute profile -k <your_kernel> --iteration-multiplexing -- <app>. Alternatively, remove --iteration-multiplexing to use application replay instead. INFO Created /root/torch_mistral_workload/workloads/mistral_iter/MI300X_A1/pmc_perf.csv` **After** (this PR): <img width="933" height="864" alt="image" src="https://github.com/user-attachments/assets/89567418-c6d5-438f-9424-0cf771ea83c8" /> ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

…error diagnostics in parser (#7010) ## Motivation This PR enhances rocJPEG logging by adding structured, offset-annotated debug logging for each major JPEG marker (SOF0, DQT, DHT, SOS, EOI), improves clarity of error messages for unsupported formats, and explicitly rejects progressive JPEGs (SOF2). sample output: ``` ROCJPEG_LOG_LEVEL=4 ./jpegdecode -i ~/install/share/rocjpeg/images/mug_420.jpg Reading images from disk, please wait! Using GPU device 0: AMD Radeon RX 7900 XT[gfx1100] on PCI bus 03:00.0 [3, Info] rocjpeg_api.cpp:123: 2274381572050 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] rocJpegCreate( 0, 0, 0x7fff0596b718 ): entry ... [3, Info] rocjpeg_decoder.cpp:81: 2274381572066 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] InitializeDecoder( ): entry ... [3, Info] rocjpeg_decoder.cpp:50: 2274381572069 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] InitHIP( 0 ): entry ... [3, Info] rocjpeg_decoder.cpp:50: 2274381592272 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] InitHIP( 0 ): exit (20203 us) ... [3, Info] rocjpeg_vaapi_decoder.cpp:416: 2274381592356 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] InitializeDecoder(): gpu_uuids_to_render_nodes_map_: {e0b64ad979bec1fd: 128} [3, Info] rocjpeg_vaapi_decoder.cpp:417: 2274381592362 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] InitializeDecoder(): Selected GPU UUID: e0b64ad979bec1fd [3, Info] rocjpeg_vaapi_decoder.cpp:475: 2274381592366 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] InitVAAPI(): Opening DRM node: /dev/dri/renderD128 [3, Info] rocjpeg_vaapi_decoder.cpp:500: 2274381593624 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] InitVAAPI(): VA-API version 1.22 [3, Info] rocjpeg_vaapi_decoder.cpp:502: 2274381593634 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] InitVAAPI(): VA-API vendor: Mesa Gallium driver 26.0.0-devel for AMD Radeon RX 7900 XT (radeonsi, navi31, , DRM 3.64, 6.17.0-23-generic) [3, Info] rocjpeg_vaapi_decoder.cpp:504: 2274381593638 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] InitVAAPI(): Trying to open /home/aryan/install/lib/rocm_sysdeps/lib/radeonsi_drv_video.so [3, Info] rocjpeg_decoder.cpp:81: 2274381593728 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] InitializeDecoder( ): exit (21662 us) ... [3, Info] rocjpeg_api.cpp:123: 2274381593732 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] rocJpegCreate( 0, 0, 0x7fff0596b718 ): exit (21682 us) ... [3, Info] rocjpeg_api.cpp:40: 2274381593735 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] rocJpegStreamCreate( 0x7fff0596b710 ): entry ... [3, Info] rocjpeg_api.cpp:40: 2274381593738 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] rocJpegStreamCreate( 0x7fff0596b710 ): exit (3 us) ... Input file name: /home/aryan/install/share/rocjpeg/images/mug_420.jpg [3, Info] rocjpeg_api.cpp:75: 2274381594731 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] rocJpegStreamParse( 0x70b1f41c4010, 2340067, 0x55f2ebfec840 ): entry ... [3, Info] rocjpeg_parser.cpp:48: 2274381594734 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] ParseJpegStream( 0x70b1f41c4010, 2340067 ): entry ... [4, Debug] rocjpeg_parser.cpp:408: 2274381594759 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] ParseDQT(): *** Marker: DQT (xFFDB) *** Define a Quantization Table. OFFSET: 0x00001296 Table length = 67 ---- Precision=8 bits Destination ID=0 (Luminance) DQT, Row #0: 1 1 1 1 1 1 1 1 DQT, Row #1: 1 1 1 1 1 1 1 1 DQT, Row #2: 1 1 1 1 1 1 1 1 DQT, Row #3: 1 1 1 1 1 1 1 1 DQT, Row #4: 1 1 1 1 1 1 1 1 DQT, Row #5: 1 1 1 1 1 1 1 1 DQT, Row #6: 1 1 1 1 1 1 1 1 DQT, Row #7: 1 1 1 1 1 1 1 1 Approx quality factor = 100.00 (scaling=2.99 variance=6.13) [4, Debug] rocjpeg_parser.cpp:408: 2274381594772 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] ParseDQT(): *** Marker: DQT (xFFDB) *** Define a Quantization Table. OFFSET: 0x000012DB Table length = 67 ---- Precision=8 bits Destination ID=1 (Chrominance) DQT, Row #0: 1 1 1 1 1 1 1 1 DQT, Row #1: 1 1 1 1 1 1 1 1 DQT, Row #2: 1 1 1 1 1 1 1 1 DQT, Row #3: 1 1 1 1 1 1 1 1 DQT, Row #4: 1 1 1 1 1 1 1 1 DQT, Row #5: 1 1 1 1 1 1 1 1 DQT, Row #6: 1 1 1 1 1 1 1 1 DQT, Row #7: 1 1 1 1 1 1 1 1 Approx quality factor = 100.00 (scaling=1.54 variance=1.58) [4, Debug] rocjpeg_parser.cpp:297: 2274381594790 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] ParseSOF(): *** Marker: SOF0 (Baseline DCT) (xFFC0) *** OFFSET: 0x00001320 Frame header length = 17 Precision = 8 Number of Lines = 2160 Samples per Line = 3840 Image Size = 3840 x 2160 Chroma subsampling: 4:2:0 Raw Image Orientation = Landscape Number of Img components = 3 Component[1]: ID=0x01, Samp Fac=0x22 (Subsamp 1 x 1), Quant Tbl Sel=0x00 (Lum: Y) Component[2]: ID=0x02, Samp Fac=0x11 (Subsamp 2 x 2), Quant Tbl Sel=0x01 (Chrom: Cb) Component[3]: ID=0x03, Samp Fac=0x11 (Subsamp 2 x 2), Quant Tbl Sel=0x01 (Chrom: Cr) [4, Debug] rocjpeg_parser.cpp:515: 2274381594802 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] ParseDHT(): *** Marker: DHT (Define Huffman Table) (xFFC4) *** OFFSET: 0x00001333 Huffman table length = 31 ---- Destination ID = 0 Class = 0 (DC / Lossless Table) Codes of length 01 bits (000 total): Codes of length 02 bits (002 total): 03 04 Codes of length 03 bits (002 total): 01 02 Codes of length 04 bits (003 total): 00 05 06 Codes of length 05 bits (001 total): 07 Codes of length 06 bits (001 total): 08 Codes of length 07 bits (001 total): 09 Codes of length 08 bits (001 total): 0A Codes of length 09 bits (001 total): 0B Codes of length 0A bits (000 total): Codes of length 11 bits (000 total): Codes of length 12 bits (000 total): Codes of length 13 bits (000 total): Codes of length 14 bits (000 total): Codes of length 15 bits (000 total): Codes of length 16 bits (000 total): Total number of codes: 012 [4, Debug] rocjpeg_parser.cpp:515: 2274381594816 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] ParseDHT(): *** Marker: DHT (Define Huffman Table) (xFFC4) *** OFFSET: 0x00001354 Huffman table length = 91 ---- Destination ID = 0 Class = 1 (AC Table) Codes of length 01 bits (000 total): Codes of length 02 bits (002 total): 01 02 Codes of length 03 bits (001 total): 11 Codes of length 04 bits (002 total): 03 21 Codes of length 05 bits (004 total): 00 04 12 31 Codes of length 06 bits (003 total): 05 41 51 Codes of length 07 bits (005 total): 06 13 22 61 71 Codes of length 08 bits (005 total): 32 81 91 A1 D1 Codes of length 09 bits (005 total): 07 52 B1 C1 F0 Codes of length 0A bits (005 total): 08 14 23 42 E1 Codes of length 0B bits (005 total): 62 72 92 A2 D2 Codes of length 0C bits (007 total): 09 15 33 53 82 E2 F1 Codes of length 0D bits (002 total): 24 B2 Codes of length 0E bits (001 total): C2 Codes of length 0F bits (002 total): 16 43 Codes of length 10 bits (023 total): 93 63 D3 0A 17 25 73 A3 34 54 83 18 19 B3 26 27 44 F2 35 64 C3 E3 46 Total number of codes: 072 [4, Debug] rocjpeg_parser.cpp:515: 2274381594829 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] ParseDHT(): *** Marker: DHT (Define Huffman Table) (xFFC4) *** OFFSET: 0x000013B1 Huffman table length = 29 ---- Destination ID = 1 Class = 0 (DC / Lossless Table) Codes of length 01 bits (001 total): 00 Codes of length 02 bits (000 total): Codes of length 03 bits (003 total): 01 02 03 Codes of length 04 bits (001 total): 04 Codes of length 05 bits (001 total): 05 Codes of length 06 bits (001 total): 06 Codes of length 07 bits (001 total): 07 Codes of length 08 bits (001 total): 08 Codes of length 09 bits (001 total): 09 Codes of length 0A bits (000 total): Codes of length 11 bits (000 total): Codes of length 12 bits (000 total): Codes of length 13 bits (000 total): Codes of length 14 bits (000 total): Codes of length 15 bits (000 total): Codes of length 16 bits (000 total): Total number of codes: 010 [4, Debug] rocjpeg_parser.cpp:515: 2274381594842 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] ParseDHT(): *** Marker: DHT (Define Huffman Table) (xFFC4) *** OFFSET: 0x000013D0 Huffman table length = 83 ---- Destination ID = 1 Class = 1 (AC Table) Codes of length 01 bits (000 total): Codes of length 02 bits (001 total): 01 Codes of length 03 bits (001 total): 11 Codes of length 04 bits (004 total): 00 02 21 61 Codes of length 05 bits (007 total): 03 31 41 51 71 91 F0 Codes of length 06 bits (006 total): 12 22 52 81 A1 D1 Codes of length 07 bits (004 total): 04 62 B1 E1 Codes of length 08 bits (005 total): 13 32 82 A2 C1 Codes of length 09 bits (004 total): 05 42 92 F1 Codes of length 0A bits (000 total): Codes of length 11 bits (004 total): 06 14 23 72 Codes of length 0C bits (004 total): 07 33 53 B2 Codes of length 0D bits (005 total): 15 43 C2 D2 E2 Codes of length 0E bits (001 total): 16 Codes of length 0F bits (001 total): 63 Codes of length 10 bits (017 total): 93 24 34 54 08 73 83 17 18 44 A3 35 F2 25 74 B3 D3 Total number of codes: 064 [4, Debug] rocjpeg_parser.cpp:606: 2274381594856 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] ParseSOS(): *** Marker: SOS (Start of Scan) (xFFDA) *** OFFSET: 0x00001425 Scan header length = 12 Number of img components = 3 Component[1]: selector=0x01, table=0(DC),0(AC) Component[2]: selector=0x02, table=1(DC),1(AC) Component[3]: selector=0x03, table=1(DC),1(AC) Spectral selection = 0 .. 63 Successive approximation = 0x00 [4, Debug] rocjpeg_parser.cpp:669: 2274381595378 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] ParseEOI(): *** Marker: EOI (End of Image) (xFFD9) *** OFFSET: 0x0023B4E1 [3, Info] rocjpeg_parser.cpp:48: 2274381595383 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] ParseJpegStream( 0x70b1f41c4010, 2340067 ): exit (649 us) ... [3, Info] rocjpeg_api.cpp:75: 2274381595386 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] rocJpegStreamParse( 0x70b1f41c4010, 2340067, 0x55f2ebfec840 ): exit (655 us) ... [3, Info] rocjpeg_api.cpp:188: 2274381595393 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] rocJpegGetImageInfo( 0x55f2ebf6dde0, 0x55f2ebfec840, 0x7fff0596b82a, 0x7fff0596b728, 0x7fff0596b810, 0x7fff0596b800 ): entry ... [3, Info] rocjpeg_decoder.cpp:332: 2274381595397 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] GetImageInfo( 0x55f2ebfec840, 0x7fff0596b82a, 0x7fff0596b728, 0x7fff0596b810, 0x7fff0596b800 ): entry ... [3, Info] rocjpeg_decoder.cpp:332: 2274381595400 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] GetImageInfo( 0x55f2ebfec840, 0x7fff0596b82a, 0x7fff0596b728, 0x7fff0596b810, 0x7fff0596b800 ): exit (3 us) ... [3, Info] rocjpeg_api.cpp:188: 2274381595403 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] rocJpegGetImageInfo( 0x55f2ebf6dde0, 0x55f2ebfec840, 0x7fff0596b82a, 0x7fff0596b728, 0x7fff0596b810, 0x7fff0596b800 ): exit (10 us) ... Input image resolution: 3840x2160 Chroma subsampling: YUV 4:2:0 Decoding started, please wait! ... [3, Info] rocjpeg_api.cpp:222: 2274381595488 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] rocJpegDecode( 0x55f2ebf6dde0, 0x55f2ebfec840, 0x7fff0596b6c0, 0x7fff0596b6e0 ): entry ... [3, Info] rocjpeg_decoder.cpp:120: 2274381595492 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] Decode( 0x55f2ebfec840, 0x7fff0596b6c0, 0x7fff0596b6e0 ): entry ... [3, Info] rocjpeg_decoder.cpp:120: 2274381607931 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] Decode( 0x55f2ebfec840, 0x7fff0596b6c0, 0x7fff0596b6e0 ): exit (12439 us) ... [3, Info] rocjpeg_api.cpp:222: 2274381607937 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] rocJpegDecode( 0x55f2ebf6dde0, 0x55f2ebfec840, 0x7fff0596b6c0, 0x7fff0596b6e0 ): exit (12449 us) ... Average processing time per image (ms): 12.4542 Average images per sec: 80.2941 [3, Info] rocjpeg_api.cpp:154: 2274381608091 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] rocJpegDestroy( 0x55f2ebf6dde0 ): entry ... [3, Info] rocjpeg_api.cpp:154: 2274381610464 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] rocJpegDestroy( 0x55f2ebf6dde0 ): exit (2373 us) ... [3, Info] rocjpeg_api.cpp:99: 2274381610471 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] rocJpegStreamDestroy( 0x55f2ebfec840 ): entry ... [3, Info] rocjpeg_api.cpp:99: 2274381610474 us: [pid:1740825 tid:1740825 hashid:0xbc0a7] rocJpegStreamDestroy( 0x55f2ebfec840 ): exit (3 us) ... Decoding completed! ``` ## Technical Details  ## JIRA ID   ## Test Plan  ## Test Result  ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

…is full ## Summary The env-driven plugin loop in ncclRmaInitPluginLibs() reserves NCCL_RMA_NUM_INTERNAL_PLUGINS trailing slots in pluginLibs[], but the constant was set to 1 while the post-loop code unconditionally appends up to three internal entries: the GIN RMA plugin, the NET-plugin RMA fallback, and the internal ncclRmaIbProxy. With a full NCCL_RMA_PLUGIN list (15 user entries), the NET and IB additions wrote past pluginLibs[16]. Bump NCCL_RMA_NUM_INTERNAL_PLUGINS to 3 so the env-loop guard reserves the slots the post-loop code actually uses. The other consumers of the macro (rma.cc:137-139, 216, 228) already treat it as "trailing internal slot count" and remain consistent. ## Tracking [internal link] plugin/rma: fix pluginLibs[] out-of-bounds write when env list is full The env-driven plugin loop in ncclRmaInitPluginLibs() reserves NCCL_RMA_NUM_INTERNAL_PLUGINS trailing slots in pluginLibs[], but the constant was set to 1 while the post-loop code unconditionally appends up to three internal entries: the GIN RMA plugin, the NET-plugin RMA fallback, and the internal ncclRmaIbProxy. With a full NCCL_RMA_PLUGIN list (15 user entries), the NET and IB additions wrote past pluginLibs[16]. Bump NCCL_RMA_NUM_INTERNAL_PLUGINS to 3 so the env-loop guard reserves the slots the post-loop code actually uses. The other consumers of the macro (rma.cc:137-139, 216, 228) already treat it as "trailing internal slot count" and remain consistent. Reported by Coverity (OVERRUN, MK 8d2e9b1dcf8a602dad9f4c493ba66f22, run nccl_v2.30u1_coverity_2026_5_13_633091, defect #4 of 8). Signed-off-by: Fancy Fan <ffan@nvidia.com> Signed-off-by: Xiaofan Li <xiaofanl@nvidia.com> Mirrored-from: d9cc756c3dd59205f96b15922191c046a4559d77

## Motivation The unit tests fails in an asan build currently, since `asan.so` is not set in the LD_PRELOAD. This PR fixes that issue. ## Technical Details Check for an ASAN environment first, and if enabled, finds the `asan.so` in the environment and adds it to the LD_PRELOAD path ## JIRA ID ROCM-26593 ## Test Plan Run ctests on regular build and ASAN ## Test Result On a non ASAN build. ``` lakshmi@kapu:~/work/rocdecode-test$ cmake $ROCM_PATH/share/rocdecode/test -- The C compiler identification is Clang 23.0.0 -- The CXX compiler identification is Clang 23.0.0 -- Detecting C compiler ABI info -- Detecting C compiler ABI info - done -- Check for working C compiler: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/llvm/bin/amdclang - skipped -- Detecting C compile features -- Detecting C compile features - done -- Detecting CXX compiler ABI info -- Detecting CXX compiler ABI info - done -- Check for working CXX compiler: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/llvm/bin/amdclang++ - skipped -- Detecting CXX compile features -- Detecting CXX compile features - done -- rocdecode-test: rocdecode found with find_package(rocdecode QUIET) -- rocdecode_INCLUDE_DIR -- /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/include -- rocdecode_LIB_DIR -- /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib -- rocdecode_LIBRARY -- /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/librocdecode.so -- rocdecode_FOUND -- 1 -- rocdecode_VERSION -- 1.8.0 -- rocdecode_VERSION_MAJOR -- 1 -- rocdecode_VERSION_MINOR -- 8 -- rocdecode_VERSION_PATCH -- 0 -- rocdecode-test: rocdecode-host found with find_package(rocdecode-host QUIET) -- rocdecode-host_INCLUDE_DIR -- /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/include -- rocdecode-host_LIB_DIR -- /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib -- rocdecode-host_LIBRARY -- /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/librocdecode-host.so -- rocdecode-host_FOUND -- 1 -- rocdecode-host_VERSION -- 1.0.0 -- rocdecode-host_VERSION_MAJOR -- 1 -- rocdecode-host_VERSION_MINOR -- 0 -- rocdecode-host_VERSION_PATCH -- 0 -- Found PkgConfig: /usr/bin/pkg-config (found version "0.29.2") -- Checking for module 'libavcodec' -- Found libavcodec, version 58.134.100 -- Checking for module 'libavformat' -- Found libavformat, version 58.76.100 -- Checking for module 'libavutil' -- Found libavutil, version 56.70.100 -- Using FFMPEG -- Libraries:/usr/lib/x86_64-linux-gnu/libavcodec.so;/usr/lib/x86_64-linux-gnu/libavformat.so;/usr/lib/x86_64-linux-gnu/libavutil.so Includes:/usr/include/x86_64-linux-gnu -- Configuring done -- Generating done -- Build files have been written to: /home/lakshmi/work/rocdecode-test (rock_0408) lakshmi@kapu:~/work/rocdecode-test$ ctest -VV UpdateCTestConfiguration from :/home/lakshmi/work/rocdecode-test/DartConfiguration.tcl Parse Config file:/home/lakshmi/work/rocdecode-test/DartConfiguration.tcl UpdateCTestConfiguration from :/home/lakshmi/work/rocdecode-test/DartConfiguration.tcl Parse Config file:/home/lakshmi/work/rocdecode-test/DartConfiguration.tcl Test project /home/lakshmi/work/rocdecode-test Constructing a list of tests Done constructing a list of tests Updating test list for fixtures Added 0 tests to meet fixture requirements Checking test dependency graph... Checking test dependency graph end test 1 Start 1: video_decodeRaw-HEVC 1: Test command: /usr/local/bin/ctest "--build-and-test" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/samples/videoDecodeRaw" "/home/lakshmi/work/rocdecode-test/videoDecodeRaw-HEVC" "--build-generator" "Unix Makefiles" "--test-command" "videodecoderaw" "-i" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/video/AMD_driving_virtual_20-H265.265" 1: Working Directory: /home/lakshmi/work/rocdecode-test 1: Environment variables: 1: LD_PRELOAD=/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib/librocm_sysdeps_va.so.2:/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib/librocm_sysdeps_va-drm.so.2 1: LIBVA_DRIVERS_PATH=/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib 1: Test timeout computed to be: 1500 1: Internal cmake changing into directory: /home/lakshmi/work/rocdecode-test/videoDecodeRaw-HEVC 1: ======== CMake output ====== 1: The C compiler identification is Clang 23.0.0 1: The CXX compiler identification is Clang 23.0.0 1: Detecting C compiler ABI info 1: Detecting C compiler ABI info - done 1: Check for working C compiler: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/llvm/bin/amdclang - skipped 1: Detecting C compile features 1: Detecting C compile features - done 1: Detecting CXX compiler ABI info 1: Detecting CXX compiler ABI info - done 1: Check for working CXX compiler: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/llvm/bin/amdclang++ - skipped 1: Detecting CXX compile features 1: Detecting CXX compile features - done 1: HIP: Using hipcc from relative path: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/bin/hipcc 1: Performing Test HIP_CLANG_SUPPORTS_PARALLEL_JOBS 1: Performing Test HIP_CLANG_SUPPORTS_PARALLEL_JOBS - Success 1: Configuring done 1: Generating done 1: Build files have been written to: /home/lakshmi/work/rocdecode-test/videoDecodeRaw-HEVC 1: ======== End CMake output ====== 1: Change Dir: /home/lakshmi/work/rocdecode-test/videoDecodeRaw-HEVC 1: 1: Run Clean Command:/usr/bin/gmake -f Makefile clean 1: 1: Run Build Command(s):/usr/bin/gmake -f Makefile && [ 33%] Building CXX object CMakeFiles/videodecoderaw.dir/videodecoderaw.cpp.o 1: [ 66%] Building CXX object CMakeFiles/videodecoderaw.dir/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/utils/rocvideodecode/roc_video_dec.cpp.o 1: [100%] Linking CXX executable videodecoderaw 1: [100%] Built target videodecoderaw 1: 1: Running test command: "/home/lakshmi/work/rocdecode-test/videoDecodeRaw-HEVC/videodecoderaw" "-i" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/video/AMD_driving_virtual_20-H265.265" 1: info: Input file: AMD_driving_virtual_20-H265.265 1: info: Using built-in bitstream reader 1: info: Using GPU device 0 - AMD Radeon RX 7900 XT[gfx1100] on PCI bus 03:00.0 1: info: decoding started, please wait! 1: Input Video Information 1: Codec : H.265/HEVC 1: Frame rate : 30000/1001 = 29.97 fps 1: Sequence : Progressive 1: Coded size : [2048, 1024] 1: Display area : [0, 0, 2048, 1024] 1: Chroma : YUV 420 1: Bit depth : 8 1: Video Decoding Params: 1: Num Surfaces : 7 1: Crop : [0, 0, 2048, 1024] 1: Resize : 2048x1024 1: 1: info: Total pictures decoded: 601 1: info: Total frames output/displayed: 601 1: info: avg decoding time per picture: 1.02047 ms 1: info: avg decode FPS: 979.937 1: info: avg output/display time per frame: 1.02047 ms 1: info: avg output/display FPS: 979.937 1: 1/6 Test #1: video_decodeRaw-HEVC ............. Passed 3.91 sec test 2 Start 2: video_decodeRaw-AVC 2: Test command: /usr/local/bin/ctest "--build-and-test" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/samples/videoDecodeRaw" "/home/lakshmi/work/rocdecode-test/videoDecodeRaw-AVC" "--build-generator" "Unix Makefiles" "--test-command" "videodecoderaw" "-i" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/video/AMD_driving_virtual_20-H264.264" 2: Working Directory: /home/lakshmi/work/rocdecode-test 2: Environment variables: 2: LD_PRELOAD=/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib/librocm_sysdeps_va.so.2:/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib/librocm_sysdeps_va-drm.so.2 2: LIBVA_DRIVERS_PATH=/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib 2: Test timeout computed to be: 1500 2: Internal cmake changing into directory: /home/lakshmi/work/rocdecode-test/videoDecodeRaw-AVC 2: ======== CMake output ====== 2: The C compiler identification is Clang 23.0.0 2: The CXX compiler identification is Clang 23.0.0 2: Detecting C compiler ABI info 2: Detecting C compiler ABI info - done 2: Check for working C compiler: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/llvm/bin/amdclang - skipped 2: Detecting C compile features 2: Detecting C compile features - done 2: Detecting CXX compiler ABI info 2: Detecting CXX compiler ABI info - done 2: Check for working CXX compiler: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/llvm/bin/amdclang++ - skipped 2: Detecting CXX compile features 2: Detecting CXX compile features - done 2: HIP: Using hipcc from relative path: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/bin/hipcc 2: Performing Test HIP_CLANG_SUPPORTS_PARALLEL_JOBS 2: Performing Test HIP_CLANG_SUPPORTS_PARALLEL_JOBS - Success 2: Configuring done 2: Generating done 2: Build files have been written to: /home/lakshmi/work/rocdecode-test/videoDecodeRaw-AVC 2: ======== End CMake output ====== 2: Change Dir: /home/lakshmi/work/rocdecode-test/videoDecodeRaw-AVC 2: 2: Run Clean Command:/usr/bin/gmake -f Makefile clean 2: 2: Run Build Command(s):/usr/bin/gmake -f Makefile && [ 33%] Building CXX object CMakeFiles/videodecoderaw.dir/videodecoderaw.cpp.o 2: [ 66%] Building CXX object CMakeFiles/videodecoderaw.dir/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/utils/rocvideodecode/roc_video_dec.cpp.o 2: [100%] Linking CXX executable videodecoderaw 2: [100%] Built target videodecoderaw 2: 2: Running test command: "/home/lakshmi/work/rocdecode-test/videoDecodeRaw-AVC/videodecoderaw" "-i" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/video/AMD_driving_virtual_20-H264.264" 2: info: Input file: AMD_driving_virtual_20-H264.264 2: info: Using built-in bitstream reader 2: info: Using GPU device 0 - AMD Radeon RX 7900 XT[gfx1100] on PCI bus 03:00.0 2: info: decoding started, please wait! 2: Input Video Information 2: Codec : AVC/H.264 2: Frame rate : 60000/2002 = 29.97 fps 2: Sequence : Progressive 2: Coded size : [2048, 1024] 2: Display area : [0, 0, 2048, 1024] 2: Chroma : YUV 420 2: Bit depth : 8 2: Video Decoding Params: 2: Num Surfaces : 7 2: Crop : [0, 0, 2048, 1024] 2: Resize : 2048x1024 2: 2: info: Total pictures decoded: 601 2: info: Total frames output/displayed: 601 2: info: avg decoding time per picture: 0.990579 ms 2: info: avg decode FPS: 1009.51 2: info: avg output/display time per frame: 0.990579 ms 2: info: avg output/display FPS: 1009.51 2: 2/6 Test #2: video_decodeRaw-AVC .............. Passed 3.86 sec test 3 Start 3: video_decodeRaw-AV1 3: Test command: /usr/local/bin/ctest "--build-and-test" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/samples/videoDecodeRaw" "/home/lakshmi/work/rocdecode-test/videoDecodeRaw-AV1" "--build-generator" "Unix Makefiles" "--test-command" "videodecoderaw" "-i" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/video/AMD_driving_virtual_20-AV1.ivf" 3: Working Directory: /home/lakshmi/work/rocdecode-test 3: Environment variables: 3: LD_PRELOAD=/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib/librocm_sysdeps_va.so.2:/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib/librocm_sysdeps_va-drm.so.2 3: LIBVA_DRIVERS_PATH=/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib 3: Test timeout computed to be: 1500 3: Internal cmake changing into directory: /home/lakshmi/work/rocdecode-test/videoDecodeRaw-AV1 3: ======== CMake output ====== 3: The C compiler identification is Clang 23.0.0 3: The CXX compiler identification is Clang 23.0.0 3: Detecting C compiler ABI info 3: Detecting C compiler ABI info - done 3: Check for working C compiler: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/llvm/bin/amdclang - skipped 3: Detecting C compile features 3: Detecting C compile features - done 3: Detecting CXX compiler ABI info 3: Detecting CXX compiler ABI info - done 3: Check for working CXX compiler: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/llvm/bin/amdclang++ - skipped 3: Detecting CXX compile features 3: Detecting CXX compile features - done 3: HIP: Using hipcc from relative path: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/bin/hipcc 3: Performing Test HIP_CLANG_SUPPORTS_PARALLEL_JOBS 3: Performing Test HIP_CLANG_SUPPORTS_PARALLEL_JOBS - Success 3: Configuring done 3: Generating done 3: Build files have been written to: /home/lakshmi/work/rocdecode-test/videoDecodeRaw-AV1 3: ======== End CMake output ====== 3: Change Dir: /home/lakshmi/work/rocdecode-test/videoDecodeRaw-AV1 3: 3: Run Clean Command:/usr/bin/gmake -f Makefile clean 3: 3: Run Build Command(s):/usr/bin/gmake -f Makefile && [ 33%] Building CXX object CMakeFiles/videodecoderaw.dir/videodecoderaw.cpp.o 3: [ 66%] Building CXX object CMakeFiles/videodecoderaw.dir/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/utils/rocvideodecode/roc_video_dec.cpp.o 3: [100%] Linking CXX executable videodecoderaw 3: [100%] Built target videodecoderaw 3: 3: Running test command: "/home/lakshmi/work/rocdecode-test/videoDecodeRaw-AV1/videodecoderaw" "-i" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/video/AMD_driving_virtual_20-AV1.ivf" 3: info: Input file: AMD_driving_virtual_20-AV1.ivf 3: info: Using built-in bitstream reader 3: info: Using GPU device 0 - AMD Radeon RX 7900 XT[gfx1100] on PCI bus 03:00.0 3: info: decoding started, please wait! 3: Input Video Information 3: Codec : AV1 3: Sequence : Progressive 3: Coded size : [2048, 1024] 3: Display area : [0, 0, 2048, 1024] 3: Chroma : YUV 420 3: Bit depth : 8 3: Video Decoding Params: 3: Num Surfaces : 12 3: Crop : [0, 0, 2048, 1024] 3: Resize : 2048x1024 3: 3: info: Total pictures decoded: 602 3: info: Total frames output/displayed: 601 3: info: avg decoding time per picture: 1.02181 ms 3: info: avg decode FPS: 978.654 3: info: avg output/display time per frame: 1.02351 ms 3: info: avg output/display FPS: 977.028 3: 3/6 Test #3: video_decodeRaw-AV1 .............. Passed 3.87 sec test 4 Start 4: video_decodeRaw-VP9 4: Test command: /usr/local/bin/ctest "--build-and-test" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/samples/videoDecodeRaw" "/home/lakshmi/work/rocdecode-test/videoDecodeRaw-VP9" "--build-generator" "Unix Makefiles" "--test-command" "videodecoderaw" "-i" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/video/AMD_driving_virtual_20-VP9.ivf" 4: Working Directory: /home/lakshmi/work/rocdecode-test 4: Environment variables: 4: LD_PRELOAD=/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib/librocm_sysdeps_va.so.2:/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib/librocm_sysdeps_va-drm.so.2 4: LIBVA_DRIVERS_PATH=/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib 4: Test timeout computed to be: 1500 4: Internal cmake changing into directory: /home/lakshmi/work/rocdecode-test/videoDecodeRaw-VP9 4: ======== CMake output ====== 4: The C compiler identification is Clang 23.0.0 4: The CXX compiler identification is Clang 23.0.0 4: Detecting C compiler ABI info 4: Detecting C compiler ABI info - done 4: Check for working C compiler: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/llvm/bin/amdclang - skipped 4: Detecting C compile features 4: Detecting C compile features - done 4: Detecting CXX compiler ABI info 4: Detecting CXX compiler ABI info - done 4: Check for working CXX compiler: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/llvm/bin/amdclang++ - skipped 4: Detecting CXX compile features 4: Detecting CXX compile features - done 4: HIP: Using hipcc from relative path: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/bin/hipcc 4: Performing Test HIP_CLANG_SUPPORTS_PARALLEL_JOBS 4: Performing Test HIP_CLANG_SUPPORTS_PARALLEL_JOBS - Success 4: Configuring done 4: Generating done 4: Build files have been written to: /home/lakshmi/work/rocdecode-test/videoDecodeRaw-VP9 4: ======== End CMake output ====== 4: Change Dir: /home/lakshmi/work/rocdecode-test/videoDecodeRaw-VP9 4: 4: Run Clean Command:/usr/bin/gmake -f Makefile clean 4: 4: Run Build Command(s):/usr/bin/gmake -f Makefile && [ 33%] Building CXX object CMakeFiles/videodecoderaw.dir/videodecoderaw.cpp.o 4: [ 66%] Building CXX object CMakeFiles/videodecoderaw.dir/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/utils/rocvideodecode/roc_video_dec.cpp.o 4: [100%] Linking CXX executable videodecoderaw 4: [100%] Built target videodecoderaw 4: 4: Running test command: "/home/lakshmi/work/rocdecode-test/videoDecodeRaw-VP9/videodecoderaw" "-i" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/video/AMD_driving_virtual_20-VP9.ivf" 4: info: Input file: AMD_driving_virtual_20-VP9.ivf 4: info: Using built-in bitstream reader 4: info: Using GPU device 0 - AMD Radeon RX 7900 XT[gfx1100] on PCI bus 03:00.0 4: info: decoding started, please wait! 4: Input Video Information 4: Codec : VP9 4: Sequence : Progressive 4: Coded size : [2048, 1024] 4: Display area : [0, 0, 2048, 1024] 4: Chroma : YUV 420 4: Bit depth : 8 4: Video Decoding Params: 4: Num Surfaces : 12 4: Crop : [0, 0, 2048, 1024] 4: Resize : 2048x1024 4: 4: info: Total pictures decoded: 601 4: info: Total frames output/displayed: 601 4: info: avg decoding time per picture: 1.11495 ms 4: info: avg decode FPS: 896.899 4: info: avg output/display time per frame: 1.11495 ms 4: info: avg output/display FPS: 896.899 4: 4/6 Test #4: video_decodeRaw-VP9 .............. Passed 3.91 sec test 5 Start 5: rocdec_Decode-HEVC 5: Test command: /usr/local/bin/ctest "--build-and-test" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/samples/rocdecDecode" "/home/lakshmi/work/rocdecode-test/rocdecDecode" "--build-generator" "Unix Makefiles" "--test-command" "rocdecdecode" "-i" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/frames" 5: Working Directory: /home/lakshmi/work/rocdecode-test 5: Environment variables: 5: LD_PRELOAD=/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib/librocm_sysdeps_va.so.2:/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib/librocm_sysdeps_va-drm.so.2 5: LIBVA_DRIVERS_PATH=/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib 5: Test timeout computed to be: 1500 5: Internal cmake changing into directory: /home/lakshmi/work/rocdecode-test/rocdecDecode 5: ======== CMake output ====== 5: The C compiler identification is Clang 23.0.0 5: The CXX compiler identification is Clang 23.0.0 5: Detecting C compiler ABI info 5: Detecting C compiler ABI info - done 5: Check for working C compiler: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/llvm/bin/amdclang - skipped 5: Detecting C compile features 5: Detecting C compile features - done 5: Detecting CXX compiler ABI info 5: Detecting CXX compiler ABI info - done 5: Check for working CXX compiler: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/llvm/bin/amdclang++ - skipped 5: Detecting CXX compile features 5: Detecting CXX compile features - done 5: HIP: Using hipcc from relative path: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/bin/hipcc 5: Performing Test HIP_CLANG_SUPPORTS_PARALLEL_JOBS 5: Performing Test HIP_CLANG_SUPPORTS_PARALLEL_JOBS - Success 5: Found PkgConfig: /usr/bin/pkg-config (found version "0.29.2") 5: Checking for module 'libavcodec' 5: Found libavcodec, version 58.134.100 5: Checking for module 'libavformat' 5: Found libavformat, version 58.76.100 5: Checking for module 'libavutil' 5: Found libavutil, version 56.70.100 5: -- Using FFMPEG -- 5: Libraries:/usr/lib/x86_64-linux-gnu/libavcodec.so;/usr/lib/x86_64-linux-gnu/libavformat.so;/usr/lib/x86_64-linux-gnu/libavutil.so 5: Includes:/usr/include/x86_64-linux-gnu 5: Configuring done 5: Generating done 5: Build files have been written to: /home/lakshmi/work/rocdecode-test/rocdecDecode 5: ======== End CMake output ====== 5: Change Dir: /home/lakshmi/work/rocdecode-test/rocdecDecode 5: 5: Run Clean Command:/usr/bin/gmake -f Makefile clean 5: 5: Run Build Command(s):/usr/bin/gmake -f Makefile && [ 50%] Building CXX object CMakeFiles/rocdecdecode.dir/rocdecdecode.cpp.o 5: [100%] Linking CXX executable rocdecdecode 5: [100%] Built target rocdecdecode 5: 5: Running test command: "/home/lakshmi/work/rocdecode-test/rocdecDecode/rocdecdecode" "-i" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/frames" 5: Read 53 frames from disk. 5: Input Video Information 5: Codec : 4 5: Frame rate : 30000/1001 = 29.97 fps 5: Sequence : Progressive 5: Coded size : [2048, 1024] 5: Display area : [0, 0, 2048, 1024] 5: Bit depth : 8 5: Total decoded frames in iteration 0: 53 5: Decoding time: 58858 microseconds 5: Success. 5: 5: 5: 5/6 Test #5: rocdec_Decode-HEVC ............... Passed 2.32 sec test 6 Start 6: rocDecode_Negative_API_Tests 6: Test command: /usr/local/bin/ctest "--build-and-test" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rocdecode/test/rocDecodeNegativeApiTests" "/home/lakshmi/work/rocdecode-test/rocdecodenegativetest" "--build-generator" "Unix Makefiles" "--test-command" "rocdecodenegativetest" 6: Working Directory: /home/lakshmi/work/rocdecode-test 6: Environment variables: 6: LD_PRELOAD=/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib/librocm_sysdeps_va.so.2:/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib/librocm_sysdeps_va-drm.so.2 6: LIBVA_DRIVERS_PATH=/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/rocm_sysdeps/lib 6: Test timeout computed to be: 1500 6: Internal cmake changing into directory: /home/lakshmi/work/rocdecode-test/rocdecodenegativetest 6: ======== CMake output ====== 6: The C compiler identification is Clang 23.0.0 6: The CXX compiler identification is Clang 23.0.0 6: Detecting C compiler ABI info 6: Detecting C compiler ABI info - done 6: Check for working C compiler: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/llvm/bin/amdclang - skipped 6: Detecting C compile features 6: Detecting C compile features - done 6: Detecting CXX compiler ABI info 6: Detecting CXX compiler ABI info - done 6: Check for working CXX compiler: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/llvm/bin/amdclang++ - skipped 6: Detecting CXX compile features 6: Detecting CXX compile features - done 6: HIP: Using hipcc from relative path: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/bin/hipcc 6: Performing Test HIP_CLANG_SUPPORTS_PARALLEL_JOBS 6: Performing Test HIP_CLANG_SUPPORTS_PARALLEL_JOBS - Success 6: Configuring done 6: Generating done 6: Build files have been written to: /home/lakshmi/work/rocdecode-test/rocdecodenegativetest 6: ======== End CMake output ====== 6: Change Dir: /home/lakshmi/work/rocdecode-test/rocdecodenegativetest 6: 6: Run Clean Command:/usr/bin/gmake -f Makefile clean 6: 6: Run Build Command(s):/usr/bin/gmake -f Makefile && [ 33%] Building CXX object CMakeFiles/rocdecodenegativetest.dir/rocdecodenegativetest.cpp.o 6: [ 66%] Building CXX object CMakeFiles/rocdecodenegativetest.dir/rocdecode_api_negative_tests.cpp.o 6: [100%] Linking CXX executable rocdecodenegativetest 6: [100%] Built target rocdecodenegativetest 6: 6: Running test command: "/home/lakshmi/work/rocdecode-test/rocdecodenegativetest/rocdecodenegativetest" 6: info: Executing negative test cases for the rocDecCreateDecoder API 6: [0, Critical] rocdecode_api.cpp:37: 849626476281 us: [pid:1526419 tid:1526419 hashid:0x91e74] rocDecCreateDecoder(): Null pointer 6: [0, Critical] roc_decoder.cpp:51: 849626476299 us: [pid:1526419 tid:1526419 hashid:0x91e74] InitializeDecoder(): Invalid number of decode surfaces. 6: [0, Critical] vaapi_videodecoder.cpp:1028: 849626508117 us: [pid:1526419 tid:1526419 hashid:0x91e74] InitHIP(): ERROR: the requested device_id is not found! 6: [0, Critical] vaapi_videodecoder.cpp:686: 849626508131 us: [pid:1526419 tid:1526419 hashid:0x91e74] GetVaContext(): Failed to initialize the HIP. 6: [0, Critical] vaapi_videodecoder.cpp:824: 849626508137 us: [pid:1526419 tid:1526419 hashid:0x91e74] CheckDecCapForCodecType(): Failed to initialize. 6: [0, Critical] rocdecode_api.cpp:90: 849626508143 us: [pid:1526419 tid:1526419 hashid:0x91e74] rocDecGetDecoderCaps(): Error: Failed to obtain decoder capabilities from driver. 6: [0, Critical] vaapi_videodecoder.cpp:130: 849626508148 us: [pid:1526419 tid:1526419 hashid:0x91e74] InitializeDecoder(): The codec config combination is not supported. 6: [0, Critical] roc_decoder.cpp:61: 849626508153 us: [pid:1526419 tid:1526419 hashid:0x91e74] InitializeDecoder(): Failed to initialize the VAAPI Video decoder. 6: [0, Critical] vaapi_videodecoder.cpp:130: 849626509822 us: [pid:1526419 tid:1526419 hashid:0x91e74] InitializeDecoder(): The codec config combination is not supported. 6: [0, Critical] roc_decoder.cpp:61: 849626509829 us: [pid:1526419 tid:1526419 hashid:0x91e74] InitializeDecoder(): Failed to initialize the VAAPI Video decoder. 6: [0, Critical] vaapi_videodecoder.cpp:130: 849626509835 us: [pid:1526419 tid:1526419 hashid:0x91e74] InitializeDecoder(): The codec config combination is not supported. 6: [0, Critical] roc_decoder.cpp:61: 849626509838 us: [pid:1526419 tid:1526419 hashid:0x91e74] InitializeDecoder(): Failed to initialize the VAAPI Video decoder. 6: [0, Critical] vaapi_videodecoder.cpp:130: 849626509849 us: [pid:1526419 tid:1526419 hashid:0x91e74] InitializeDecoder(): The codec config combination is not supported. 6: [0, Critical] roc_decoder.cpp:61: 849626509852 us: [pid:1526419 tid:1526419 hashid:0x91e74] InitializeDecoder(): Failed to initialize the VAAPI Video decoder. 6: [0, Critical] vaapi_videodecoder.cpp:894: 849626509858 us: [pid:1526419 tid:1526419 hashid:0x91e74] CheckDecCapForCodecType(): VAAPI failure: vaCreateConfig(va_contexts_[va_ctx_id].va_display, va_contexts_[va_ctx_id].va_profile, VAEntrypointVLD, &va_config_attrib, 1, &va_contexts_[va_ctx_id].va_config_id) failed with 'status: 13: the requested VAEntryPoint is not supported' at /home/lakshmi/work/rocm-systems/projects/rocdecode/src/rocdecode/vaapi/vaapi_videodecoder.cpp:894 6: [0, Critical] rocdecode_api.cpp:90: 849626509865 us: [pid:1526419 tid:1526419 hashid:0x91e74] rocDecGetDecoderCaps(): Error: Failed to obtain decoder capabilities from driver. 6: [0, Critical] vaapi_videodecoder.cpp:130: 849626509868 us: [pid:1526419 tid:1526419 hashid:0x91e74] InitializeDecoder(): The codec config combination is not supported. 6: [0, Critical] roc_decoder.cpp:61: 849626509870 us: [pid:1526419 tid:1526419 hashid:0x91e74] InitializeDecoder(): Failed to initialize the VAAPI Video decoder. 6: info: Executing negative test cases for the rocDecDestroyDecoder API 6: info: Executing negative test cases for the rocDecGetDecoderCaps API 6: info: Executing negative test cases for the rocDecDecodeFrame API 6: info: Executing negative test cases for the rocDecGetDecodeStatus API 6: info: Executing negative test cases for the rocDecReconfigureDecoder API 6: info: Executing negative test cases for the rocDecGetVideoFrame API 6: info: Executing negative test cases for the rocDecGetErrorName API 6: info: Executing negative test cases for the rocDecCreateVideoParser API 6: [0, Critical] roc_decoder.cpp:124: 849626510695 us: [pid:1526419 tid:1526419 hashid:0x91e74] GetVideoFrame(): Invalid parameters. 6: [0, Critical] rocparser_api.cpp:44: 849626510705 us: [pid:1526419 tid:1526419 hashid:0x91e74] rocDecCreateVideoParser(): Error: The current version of rocDecode officially supports only the H.265 (HEVC), H.264 (AVC), AV1 and VP9 codecs. 6: info: Executing negative test cases for the rocDecParseVideoData API 6: info: Executing negative test cases for the rocDecDestroyVideoParser API 6: info: Executing negative test cases for the rocDecCreateBitstreamReader API 6: info: Executing negative test cases for the rocDecGetBitstreamCodecType API 6: info: Executing negative test cases for the rocDecGetBitstreamBitDepth API 6: [0, Critical] roc_bs_reader_api.cpp:31: 849626511165 us: [pid:1526419 tid:1526419 hashid:0x91e74] rocDecCreateBitstreamReader(): Null pointer 6: [0, Critical] roc_bs_reader_api.cpp:31: 849626511169 us: [pid:1526419 tid:1526419 hashid:0x91e74] rocDecCreateBitstreamReader(): Null pointer 6: [0, Critical] roc_bs_reader_api.cpp:52: 849626511174 us: [pid:1526419 tid:1526419 hashid:0x91e74] rocDecGetBitstreamCodecType(): Null pointer 6: [0, Critical] roc_bs_reader_api.cpp:52: 849626511177 us: [pid:1526419 tid:1526419 hashid:0x91e74] rocDecGetBitstreamCodecType(): Null pointer 6: [0, Critical] roc_bs_reader_api.cpp:74: 849626511182 us: [pid:1526419 tid:1526419 hashid:0x91e74] rocDecGetBitstreamBitDepth(): Null pointer 6: [0, Critical] roc_bs_reader_api.cpp:74: 849626511185 us: [pid:1526419 tid:1526419 hashid:0x91e74] rocDecGetBitstreamBitDepth(): Null pointer 6: [0, Critical] roc_bs_reader_api.cpp:96: 849626511189 us: [pid:1526419 tid:1526419 hashid:0x91e74] rocDecGetBitstreamPicData(): Null pointer 6: [0, Critical] roc_bs_reader_api.cpinfo: Executing negative test cases for the rocDecGetBitstreamPicData API 6: info: Executing negative test cases for the rocDecDestroyBitstreamReader API 6: Test Passed! 6: p:96: 849626511193 us: [pid:1526419 tid:1526419 hashid:0x91e74] rocDecGetBitstreamPicData(): Null pointer 6: [0, Critical] roc_bs_reader_api.cpp:118: 849626511197 us: [pid:1526419 tid:1526419 hashid:0x91e74] rocDecDestroyBitstreamReader(): Null pointer 6: 6/6 Test #6: rocDecode_Negative_API_Tests ..... Passed 1.99 sec 100% tests passed, 0 tests failed out of 6 Total Test time (real) = 19.86 sec ``` ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

## Motivation On certain test systems we do not have a shared filesystem, the deadlock_analysis script harness would depend on a shared filesystem to collect traces and artifacts. Remove that limitation. ## Technical Details Use MPI stdout redirection to coalesce the trace artifacts to the headnode in a headnode local output directory, perform analysis on the headnode. Prior filesystem based mode remains available and default. ## JIRA ID ## Test Plan Intentionally added a deadlock to a functional test, compiled with Debug. Run mpiexec_deadlock_attach script and observe artifacts. ## Test Result ``` ROCSHMEM_DEBUG_LEVEL=trace:stats ~/ompi/bin/mpirun -n 2 tests/functional_tests/rocshmem_functional_tests -n 1 -a 26 & (sleep 2; /home/abouteil/rocm-systems/projects/rocshmem/scripts/debug/mpiexec_deadlock_analysis.sh rocshmem_functional_tests --stats --color --no-shared-fs ) ### Creating Test: Blocking WG level Puts B=ipc PE=2 W=1 Z=64 ### Deploying deadlock analysis... Executable: rocshmem_functional_tests Output directory: ./rocshmem_deadlock_analysis_20260530_013437 Filesystem mode: no shared fs (--output-filename) Stats: yes Color: --color Found 4 process(es) matching 'rocshmem_functional_tests': 717472 717473 751554 751555 Output directory: /tmp/rocshmem_deadlock_20260530_013437 Per-process timeout: 120s Analyzing PID 717472 (PE 0) ... OK Analyzing PID 717473 (PE 1) ... OK (no GPU threads) Analyzing PID 751554 (PE 0) ... OK Analyzing PID 751555 (PE 1) ... OK (no GPU threads) === Attach and Analyze Summary === Processes analyzed: 4 / 4 PE 0 PID 751554 groups=1 \033[1;31m1\033[0m wf(s) in rocSHMEM; \033[1;36mWaiting for remote PE memory update (barrier/sync or user rocshmem_wait_until). Check if the remote PE is alive and making progress. (at: #0 0x0000792efa5af3c8 in rocshmem::Context::wait_until<long> (this=0x792eed407000, ivars=0x792efac04008, cmp=0, val=1) at src/context_tmpl_device.hpp:224)\033[0m PE 1 PID 751555 groups=? (no GPU threads found) === BEGIN FILE: pe0_pid_717472.txt === [New LWP 717493] [New LWP 717490] [New LWP 717488] [New LWP 717485] [New LWP 717481] [New LWP 717480] [Thread debugging using libthread_db enabled] Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1". 0x000077a09e06695f in ?? () from /opt/rocm-7.2.1/lib/libhsa-runtime64.so.1 === rocSHMEM Deadlock Analysis === Process(es): 717472 Total GPU wavefronts analyzed: 1 Unique backtrace groups: 1 --- Group 1 (1 wavefront(s)) --- WFs: (0,0,0)/0 Backtrace: #0 0x0000776005fafce0 in rocshmem::uncached_load<long volatile> (src=0x776006604008) at src/assembly.hpp:125 #1 rocshmem::Context::test<long> (this=0x775ff8e07000, ivars=0x776006604008, cmp=0, val=1) at src/context_tmpl_device.hpp:432 #2 rocshmem::Context::wait_until<long> (this=0x775ff8e07000, ivars=0x776006604008, cmp=0, val=1) at src/context_tmpl_device.hpp:224 #3 rocshmem::IPCContext::internal_direct_barrier (this=0x775ff8e07000, pe=0, PE_start=0, stride=1, n_pes=2, pSync=0x776006604000) at build/src/ipc/context_ipc_device_coll.cpp:44 #4 0x0000776005fb7c24 in rocshmem::IPCContext::internal_sync_wg (this=0x775ff8e07000, pe=0, PE_start=0, stride=1, PE_size=2, pSync=0x776006604000) at build/src/ipc/context_ipc_device_coll.cpp:115 #5 0x000077600603d018 in rocshmem::IPCContext::sync_all_wg (this=0x775ff8e07000) at build/src/ipc/context_ipc_device_coll.cpp:168 #6 0x000077600603d558 in rocshmem::IPCContext::barrier_all_wg (this=0x775ff8e07000) at build/src/ipc/context_ipc_device_coll.cpp:187 #7 0x000077600603dedc in rocshmem::Context::barrier_all_wg (this=0x775ff8e07000) at build/src/context_device.cpp:151 #8 0x000077600603e33c in rocshmem::rocshmem_barrier_all_wg () at build/src/rocshmem_gpu.cpp:885 #9 0x00007760060437cc in WorkGroupPrimitiveTest (loop=1, skip=10, start_time=0x775ff8e31000, end_time=0x775ff8e32000, source=0x776013c00480 "abcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqr"..., dest=0x776013d00500 "1", size=1, type=WGPutTestType, ctx_type=4) at /home/abouteil/rocm-systems/projects/rocshmem/tests/functional_tests/workgroup_primitives.cpp:72 [HINT] Waiting for remote PE memory update (barrier/sync or user rocshmem_wait_until). Check if the remote PE is alive and making progress. (at: #2 rocshmem::Context::wait_until<long> (this=0x775ff8e07000, ivars=0x776006604008, cmp=0, val=1) at src/context_tmpl_device.hpp:224) ... === END FILE: pe1_pid_751555.txt === Reconstructing per-PE output files from mpiexec output... Reconstructed output from 1 rank file(s). Running cross-rank coalescing pass... === Cross-Rank Deadlock Analysis === Directory: ./rocshmem_deadlock_analysis_20260530_013437 Ranks with analysis files: 2 (0, 1) Unique backtrace patterns: 2 --- Pattern 1 [UNIQUE: rank 0] --- Ranks: 0 Total wavefronts: 1 Backtrace: #0 0x0000776005fafce0 in rocshmem::uncached_load<long volatile> (src=0x776006604008) at src/assembly.hpp:125 #1 rocshmem::Context::test<long> (this=0x775ff8e07000, ivars=0x776006604008, cmp=0, val=1) at src/context_tmpl_device.hpp:432 #2 rocshmem::Context::wait_until<long> (this=0x775ff8e07000, ivars=0x776006604008, cmp=0, val=1) at src/context_tmpl_device.hpp:224 #3 rocshmem::IPCContext::internal_direct_barrier (this=0x775ff8e07000, pe=0, PE_start=0, stride=1, n_pes=2, pSync=0x776006604000) at build/src/ipc/context_ipc_device_coll.cpp:44 #4 0x0000776005fb7c24 in rocshmem::IPCContext::internal_sync_wg (this=0x775ff8e07000, pe=0, PE_start=0, stride=1, PE_size=2, pSync=0x776006604000) at build/src/ipc/context_ipc_device_coll.cpp:115 #5 0x000077600603d018 in rocshmem::IPCContext::sync_all_wg (this=0x775ff8e07000) at build/src/ipc/context_ipc_device_coll.cpp:168 #6 0x000077600603d558 in rocshmem::IPCContext::barrier_all_wg (this=0x775ff8e07000) at build/src/ipc/context_ipc_device_coll.cpp:187 #7 0x000077600603dedc in rocshmem::Context::barrier_all_wg (this=0x775ff8e07000) at build/src/context_device.cpp:151 #8 0x000077600603e33c in rocshmem::rocshmem_barrier_all_wg () at build/src/rocshmem_gpu.cpp:885 #9 0x00007760060437cc in WorkGroupPrimitiveTest (loop=1, skip=10, start_time=0x775ff8e31000, end_time=0x775ff8e32000, source=0x776013c00480 "abcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqr"..., dest=0x776013d00500 "1", size=1, type=WGPutTestType, ctx_type=4) at /home/abouteil/rocm-systems/projects/rocshmem/tests/functional_tests/workgroup_primitives.cpp:72 [HINT] Waiting for remote PE memory update (barrier/sync or user rocshmem_wait_until). Check if the remote PE is alive and making progress. (at: #2 rocshmem::Context::wait_until<long> (this=0x775ff8e07000, ivars=0x776006604008, cmp=0, val=1) at src/context_tmpl_device.hpp:224) === Per-Rank rocSHMEM Wavefront Count === Rank 0: 2 wf(s) inside rocSHMEM Rank 1: 0 wf(s) inside rocSHMEM === Summary === 0 pattern(s) common to all 2 rank(s) 1 pattern(s) present in exactly one rank Most common hint: Waiting for remote PE memory update (barrier/sync or user rocshmem_wait_until). Check if the remote PE is alive and making progress. (at: #2 rocshmem::Context::wait_until<long> (this=0x775ff8e07000, ivars=0x776006604008, cmp=0, val=1) at src/context_tmpl_device.hpp:224) ```

… progress tracker (#6616) ## Motivation When the program deadlocks it can be hard to know how much work was performed thus far, for example did it deadlock after only a couple of communication calls or after a long time, is there only 1 call of 1 type? etc. ## Technical Details builds with PROFILE=ON can produce call statistics (slowing down execution btw), this can be triggered from the gdb attachement to get the stat dump. ## JIRA ID   ## Test Plan modified a functional test to intentionally deadlock. build with Debug and PROFILE=ON. ``` ROCSHMEM_DEBUG_LEVEL=trace:stats ~/ompi/bin/mpirun -n 2 tests/functional_tests/rocshmem_functional_tests -n 1 -a 26 & (sleep 2; ../scripts/debug/attach_deadlock_analysis.sh rocshmem_functional_tests --stats ) ``` Output file ``` === rocSHMEM Deadlock Analysis === Process(es): 845512 Total GPU wavefronts analyzed: 1 Unique backtrace groups: 1 --- Group 1 (1 wavefront(s)) --- WFs: (0,0,0)/0 Backtrace: #0 0x00007ae55f5afcf0 in rocshmem::uncached_load<long volatile> (src=0x7ae55fc04008) at src/assembly.hpp:125 #1 rocshmem::Context::test<long> (this=0x7ae552407000, ivars=0x7ae55fc04008, cmp=0, val=1) at src/context_tmpl_device.hpp:432 #2 rocshmem::Context::wait_until<long> (this=0x7ae552407000, ivars=0x7ae55fc04008, cmp=0, val=1) at src/context_tmpl_device.hpp:224 #3 rocshmem::IPCContext::internal_direct_barrier (this=0x7ae552407000, pe=0, PE_start=0, stride=1, n_pes=2, pSync=0x7ae55fc04000) at build/src/ipc/context_ipc_device_coll.cpp:44 #4 0x00007ae55f5b7c24 in rocshmem::IPCContext::internal_sync_wg (this=0x7ae552407000, pe=0, PE_start=0, stride=1, PE_size=2, pSync=0x7ae55fc04000) at build/src/ipc/context_ipc_device_coll.cpp:115 #5 0x00007ae55f63d018 in rocshmem::IPCContext::sync_all_wg (this=0x7ae552407000) at build/src/ipc/context_ipc_device_coll.cpp:168 #6 0x00007ae55f63d558 in rocshmem::IPCContext::barrier_all_wg (this=0x7ae552407000) at build/src/ipc/context_ipc_device_coll.cpp:187 #7 0x00007ae55f63dedc in rocshmem::Context::barrier_all_wg (this=0x7ae552407000) at build/src/context_device.cpp:151 #8 0x00007ae55f63e33c in rocshmem::rocshmem_barrier_all_wg () at build/src/rocshmem_gpu.cpp:885 #9 0x00007ae55f6437cc in WorkGroupPrimitiveTest (loop=1, skip=10, start_time=0x7ae552431000, end_time=0x7ae552432000, source=0x7ae56d200480 "abcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqrstuvwxyzabcdefghijklmnopqr"..., dest=0x7ae56d300500 "1", size=1, type=WGPutTestType, ctx_type=4) at /home/abouteil/rocm-systems/projects/rocshmem/tests/functional_tests/workgroup_primitives.cpp:72 [HINT] Waiting for remote PE memory update (barrier/sync or user rocshmem_wait_until). Check if the remote PE is alive and making progress. (at: #2 rocshmem::Context::wait_until<long> (this=0x7ae552407000, ivars=0x7ae55fc04008, cmp=0, val=1) at src/context_tmpl_device.hpp:224) === rocSHMEM Stats === PE 0 (counters reflect activity since process start) Device: quiet 1 put_wg 704 Host: (all zero) === Summary === 1 wavefront(s) inside rocSHMEM 0 wavefront(s) in user code / other ``` Note the == rocshmem Stats == section contains the number of put and quiet calls. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Signed-off-by: Aurelien Bouteiller <aurelien.bouteiller@amd.com> Signed-off-by: Aurelien Bouteiller <abouteil@amd.com> Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>

…ssembly - comgr/: ROCm/llvm-project amd/comgr source (commit b72b062dfe0), needed to disassemble gfx1250 new instruction encodings (e.g. v_writelane_b32 VOP3 0xd7610000) that the container LLVM22 comgr rejects as illegal .long. - build_comgr_gfx1250.sh: one-click build against an external LLVM23 build (LLVM_BUILD_DIR), with zstd cmake prefix so LLVM export target zstd::libzstd_shared resolves. - README: document comgr build usage and root cause ROCm#4 (comgr LLVM version mismatch causing truncated code.json).

jayhawk-commits and others added 3 commits April 7, 2025 18:13

Initial migration document draft

4c3f2ee

Update migration.md

6c1861a

Merge branch 'develop' into joseph/migrationDoc

276ddb3

jayhawk-commits merged commit 9c836cb into develop Apr 7, 2025

jayhawk-commits deleted the joseph/migrationDoc branch April 7, 2025 22:22

systems-assistant Bot pushed a commit that referenced this pull request Jul 22, 2025

Add rocDecode API support (#4)

de3a874

* Add rocDecode API support * Update CI * CI update: Sanitizers run on ubuntu 22.04 --------- Co-authored-by: Jonathan R. Madsen <jonathanrmadsen@gmail.com>

systems-assistant Bot pushed a commit that referenced this pull request Jul 22, 2025

Update copyright date to 2025 in License.md (#4)

8c62af8

Update LICENSE Co-authored-by: srawat <120587655+SwRaw@users.noreply.github.com>

systems-assistant Bot pushed a commit that referenced this pull request Aug 5, 2025

SWDEV-521313 - Streamline Version Management: Retain Only version-roc… (

cf86775

#4) SWDEV-521313 Streamline Version Management: Retain Only version-rocm File

jayhawk-commits pushed a commit that referenced this pull request Aug 5, 2025

Update copyright date to 2025 in License.md (#4)

c9a4af3

Update LICENSE Co-authored-by: srawat <120587655+SwRaw@users.noreply.github.com> [ROCm/rocprofiler commit: 8c62af8]

systems-assistant Bot pushed a commit that referenced this pull request Aug 7, 2025

Fix static deinitialization fiasco in destructor

231eae5

Closes #4

This was referenced Aug 7, 2025

Fix static deinitialization fiasco in destructor #126

Closed

[Issue]: hsa_init()/hsa_shut_down() and PAPI rocp_sdk inconsistent behavior #137

Open

jayhawk-commits pushed a commit that referenced this pull request Aug 11, 2025

SWDEV-521313 - Streamline Version Management: Retain Only version-roc… (

e12743f

#4) SWDEV-521313 Streamline Version Management: Retain Only version-rocm File [ROCm/rocm-core commit: cf86775]

jmellorcrummey mentioned this pull request Sep 25, 2025

[Issue]: PC sampling on MI300A eventually dies with an HSA error #1141

Open

ammallya pushed a commit that referenced this pull request Nov 17, 2025

[SWDEV-230863] add two new interfaces for background health check (#4)

dc400d9

1. Get the bad pages threshold of a processor. 2. Verify the checksum of RAS EEPROM Signed-off-by: Meng Li <li.meng@amd.com>

oliveiradan mentioned this pull request Dec 18, 2025

Add support for newer GPU metrics versions #1961

Closed

habajpai-amd mentioned this pull request Dec 18, 2025

fix: resolve crash when profiling TensorFlow GPU application #2381

Merged

1 task done

ammallya pushed a commit that referenced this pull request Jan 21, 2026

Merge pull request #4 from BKP/hyd-7c_07-11-24

d054fc6

hyd-7c

ammallya pushed a commit that referenced this pull request Jan 21, 2026

Merge pull request #4 from BKP/hyd-7c_07-11-24

fc7f3a2

hyd-7c [ROCm/rocshmem commit: d054fc6]

systems-assistant Bot mentioned this pull request Jan 22, 2026

[Issue]: RCCL is compleatly broken in most configurations #2779

Open

jeffqjiangNew mentioned this pull request Mar 10, 2026

[rocDecode] Add core function entry and exit logs. #3851

Merged

1 task

rrawther mentioned this pull request May 13, 2026

Users/rrawther/rocdecode: Update api comments for rocdecDecode sample #6076

Merged

1 task

vlaindic mentioned this pull request May 18, 2026

PC sampling API refactoring #4548

Closed

1 task

andrew-aladjev mentioned this pull request Jun 24, 2026

[Issue]: amd-smi set -L sclk throws TypeError: '<' not supported between instances of 'int' and 'str' #7583

Closed

atgutier mentioned this pull request Jul 13, 2026

feat(rocjitsu): route AMDKFD_IOC_DBG_TRAP into a debug-session dispatcher #8393

Merged

1 task done

oliveiradan mentioned this pull request Aug 25, 2026

test(amdsmi): [AILITOOLS-296] Amdsmi refactor C tests #10081

Closed

bwelton mentioned this pull request Aug 27, 2026

fix(rocprofiler-sdk): KFD dispatch-log signal-less re-review fixes #10809

Merged

vlaindic mentioned this pull request Sep 8, 2026

feat(rocprofiler-sdk): Add SPM aqlprofile support for gfx12/gfx1250 #9646

Merged

1 task done

dlamd1dai mentioned this pull request Sep 22, 2026

feat(rccl/gin): honour ncclGinFenceLevel on the Anvil SDMA backend #11021

Merged

1 task done

SakaSitharammurthy mentioned this pull request Oct 1, 2026

feat(amdsmi): add CPU apb_status, xgmi_link_width, df_pstate_range metrics #7521

Open

This was referenced Oct 14, 2026

[RCCL] gfx1250 AllToAll: LL128 V2, abortable waits, async device errors #13465

Draft

[RCCL] gfx1250 AllToAll: one-shot push for out-of-place 16 MiB-1 GiB #13475

Draft

shwetagkhatri mentioned this pull request Oct 23, 2026

rocr: serve swap copies with a shader blit where the DMA engine cannot #14190

Draft

gerardpamd mentioned this pull request Oct 30, 2026

fix(rccl): honour RCCL_DIRECT_REDUCE_SCATTER_THRESHOLD in rcclUseReduceScatterDirect #14316

Merged

Navigation Menu

Repository navigation

Migration doc fixes - #4

Migration doc fixes#4
jayhawk-commits merged 3 commits into
developfrom
joseph/migrationDoc

jayhawk-commits commented Apr 7, 2025

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Conversation

jayhawk-commits commented Apr 7, 2025

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant