PAL optimized the logic for the barriers, which caused failures with CP DMA on Navi4x.
Change barrier's code to match the most recent PAL optimizations.

Change-Id: I55eeab20f51eb8e920bcbb4b55fbe3c7f77fd3fa

Change-Id: If1bcca45825c9899462bb95ed6f637f5af806cc8

This no longer does anything.

Change-Id: I0643198a46a534a76454a5b461d010ed1776a89a

Added call to hipDeviceSynchronize in __hipUnregisterFatBinary
to ensure that all HSA async signal handlers complete before removeFatBinary

Change-Id: I756fecca1c2a5eae092613d8079de266399e5685

… elementSize should be 1 as width is in bytes while capturing hipMemset2DAsync.

Change-Id: I8f9122a30cba0a07c097dfd7609432090caab142

This reverts commit 5f68a45.

Reason for revert: due to windows staging branch using Opencl-icd-loader master

Change-Id: I9cca7564a21de1733665a34da6f0322aa3b886e7

…function to be called at exit"

This reverts commit 5e294f8.

Change-Id: Ib9cb1cc0c3903bfba56c9a5d05ae8afe96be583a

Under ROCr physical allocations don't have initial VA and require extra
flag in ROCclr. Add an option to have a mempool of physical allocations.

Change-Id: I4d062fe0dd8113d4eaf6e8b51749ed56d8701d1e

Change-Id: I33d1359d5e4c871f63350d8300f726e039664d86

Change-Id: I7b58177c41dc0c6c59813977cb90e65a6cb3be72

Change-Id: Ica32017ef7b00326dfb6d1f604e126d40ad5b786

Change-Id: Ib479c744b90125b74d99cbf18b7f4b8cf765bf1c

- With https://gerrit-git.amd.com/c/lightning/ec/llvm-project/+/1002628 applied, at -O0 Kernel::dynamicParallelism() returns true but virtual queue isn't created
- This causes segfault inside VirtualGPU::submitKernelInternal() when getVQVirtualAddress() is called

Change-Id: Ia7af042adad2329e870c142caaac3e8fa886f8b8

- Create a vector to allow multiple TS to be stored in Command.
- This would mean we dont wait for entire batch in Accumulate command
to finish when we exhaust signals.
- Reduce the number of signals created at init to 64. This min value
may still need to be tuned but the KFD allows max of 4094 interrupt
signals per device.
- Store kernel names whenever they are available and not just when
profiling. If we dynamically enable profiling like for Torch, a crash
can happen if hipGraphInstantiate wasnt included in Torch profile scope
beacuse we previously entered kernel names only when profiler is
attached.

Change-Id: I34e7881a25bbc763f82fdeb3408a8ea58e1ec006

…ead is being destroied by app/test's unload libamdhip64.so call.

Change-Id: I8d4a8d8b6801d9f6eb745c45adf831597def0cb5

alwaysResident setting doesn't require per queue residency tracking.
Thus, the logic can be skipped to avoid the lock of queues.

Change-Id: Ib5cff5b79d3ecb8c2f2eb2565cf069f9a69438b0

The new logic has a lock for PAL call and doesn't require the lock for queues.

Change-Id: I61b67c3c4abd2ede44809de1d6beed756766032e

… during capture

=> hipDeviceSynchronize is not allowed during capture.
=> hipEventSynchronize during capture should return hipErrorCapturedEvent error
=> hipEventQuery during capture should return hipErrorCapturedEvent error
hipStreamSynchronize, hipEventSynchronize, hipStreamWaitEvent, hipStreamQuery
For Side Stream(Stream that is not currently under capture):
=> If current thread is capturing in relaxed mode, calls are allowed
=> If any stream in current/concurrent thread is capturing in global mode, calls are not allowed
=> If any stream in current thread is capturing in ThreadLocal mode, calls are not allowed
For Stream that is currently under capture
=> calls are not allowed
=> Any call that is not allowed during capture invalidates the capture sequence
=> It is invalid to call synchronous APIs during capture. Synchronous APIs,
such as hipMemcpy(), enqueue work to the legacy stream and synchronize it before returning.

Change-Id: I201c6e63e1a5d93fd416a3b520264c0fdbe31237

Change-Id: Id58f982edd4f17d675f7a0f61a9b4dea0baebd9b

Change-Id: Icbe67024297c92bf59139b6a2ccd2ba3674f60b1

Add cltrace compile definition for CL_TARGET_OPENCL_VERSION to OpenCL 2.2

Change-Id: Ie868ab0a6e86951afc6d07da58be942c3b736d15

When large bar is enabled, persistent memory leads to overallocation for 32 bit architecture.

Change-Id: Iae39359d8128588de02e42d77fe58e868b8e71fd

- Application is passing null for parameter stream in API hipStreamWaitEvent
- When event stream isn't capturing and event is not recorded, causes segfault because we are accessing deviceId() from waitStream

Change-Id: I8b87ffd6f234677f68b66dcb7ef44b2ff04a7c91

…hGenericAqlPacket

Dispatching multiple packets with ring the doorbell once is not supported by the lower layers

Change-Id: I7665a2dcdd4ef9e47dadfe410180fed64c5a4ee0

Change-Id: Iadcdadd734e7aeeb23742e426353defa972d3ad5

If we are using the mask returned by getLastUsedSdmaEngine() then we
need to apply the SDMA Read/Write mask to it before using with HSA
copy_on_engine API.

Change-Id: I6e5dc6c187eeb3c61ee159e9d2a0fa7b4737c06e

This also brings bfloat16 implementation closer to CUDA's.

Change-Id: I23f381141faacd6537923ae9b88ada4d661db496

Change-Id: Ic41b1ad1b64cca0e31986337a83a5146d52a7328

Add kernel arguments optimization into blit path.
Enabled by default on MI300.

Change-Id: I2694a81b90d48ad07d86dfe4c0c64fe187bada8e

…lti devices and chunk decommit gets delayed.

Change-Id: Ia4b0d5fbfa8f198776e52d14de8b22c6942f740d

Refactor PrintfDbg::outputArgument() to remove potential risk.
Fix half vector printf issue on all devices.
Fix FEAT-56794 as well.

Change-Id: Iae39359d2128588def2e43d77fe58e868b8e71ff

Remove offset from blit kernels, since it can be applied in setup.

Change-Id: I06b585068d68a0ee8e125ddf46a36fccb372f30d

Change-Id: I10ad3fbfca887e92cd81f68392fa1acf753cbd2b

- UUID is Ascii string with a maximum of 21 chars which uniquely identifies a GPU
- Convert set UUID in HIP_VISIBLE_DEVICES to device index internally
- Then use existing device index logic for HIP_VISIBLE_DEVICES

Change-Id: I8cab4fe42459f8209b97f909300789e6e687b9ac

…ccess.

Change-Id: Iecda74ca0207c95d3fbed8b4e15c8c4c5895d939

…ad of granularity factor.

Change-Id: I1e0db6e0628c09d26850e5a0339e2a4660442db8

… mode

Change-Id: Ib9739f9917bc6ff69cc76f444d909311922ebc1e

Change-Id: Id9a5a85641882961e4d860a815217c641e6f3387

… hipGetDevicePCIBusId.

Change-Id: I932724cc872d7ae2643ce6ac2924901cb49cd7ad

Force implicit runtime teardown with a global destructor.

Change-Id: Iabe63dedf5b94fefc98668585c45a61607120669

AMD CPUs have had avx512_bf16 support for quite some time now (from
consumer Ryzen 7000 series to enterprise grade CPUs). This
patch should allow users to use the hardware bf16 unit when running the
__host__ variants of the function. This can be enabled via `hipcc ...
-mavx512vl -mavx512bf16`.

Change-Id: I67c377afc95ddfe8d45a048dce078a247d4a1878

Change-Id: I5f8b030eefdb37d3d51da3e135e5aa0f18ad9018

Change-Id: If027dd8cbe5cbe142fff353cb72c16f08e9aea8e

for fnuz nums zero val is 0x00, -0 i.e. 0x80 would be a NaN

Change-Id: Ibdc4fb4b9fb307b5952434f08d45a8ddd6262db8

Signed-off-by: sdashmiz <[email protected]>
Change-Id: Ie60aa598dd73df66cdf02c1d96daf2dfccba7a59

Avoid a deadlock on the host call buffer creation. Since the buffer will be
allocated in the queue thread, then use direct device memory allocation
skipping the global context lock.

Change-Id: I09b55ee03bb42ab5d320c152b52a8c842c5fdcc1

Change-Id: I21ec8eadb189147c579ec65acf68de40d604686b

Change-Id: I517c2557573db258b3e3e353f02f6a56652b0fde

…eters

For all windows allocation on SVM memory tagged with flag ROCCLR_MEM_INTERPROCESS.
hipHostMalloc validation is based on flag. So remove ROCCLR_MEM_INTERPROCESS before check.

Change-Id: I823bbf228d9a4a9acb4abffc01ac6b3f544c6e12

Change-Id: Icfeb83ca51e96be35abb67a94d6e3e1a1ca5a934

…valid."

This reverts commit a9ff2c5.

Change-Id: I26c4b3c74b2861afc17f979492d025b59d4388ab

…device kernel args

Change-Id: I6d41a089a17f55306e7ff402588a1e831b20a7a7

…urn the correct value

Change-Id: Ieced2ee61bba28f2d1df96893a661287b0a5c7b7

Change-Id: I9ba34454f7487d6bc0d398b322a147cbac6c6443

Mempool may reuse memory without a wait. Hence, the timestamp has
to be preserved and can't be destroyed.

Change-Id: I6f095f44afa69887a4b7aeb3b329804aedd96f3e

Change-Id: I9e2aefdbcbba153c7f1080d80aab7a345eaf1eb4

- UUID needs to be specified in the format GPU-<body>, <body> encodes UUID as a 16 chars
- Convert set UUID in HIP_VISIBLE_DEVICES to device index internally
- Then use existing device index logic for HIP_VISIBLE_DEVICES

Change-Id: I654f492a49cd4d7a9b7339360ab558165240caa5

…ct bundle processing fail.

Change-Id: Iea553ab0265c08341f915644075ce2b6ed9b3200

The "optimized" version of memcpy is outdated and
was used in win32 only.

Change-Id: I7f2e0e9051e37cec95438266824b5b0025c324c6

Change-Id: Ib6a9b8fc9a8692eb17b39b854cefd92c6b59733f

Change-Id: I9a4a4d9087b5daf15e3ba31e786d34db431212a1

Change-Id: Idcd37628a2167f0bd2db2a83132a1862cbd051b0

…error on memcpy direction change

Change-Id: I2c8f5ea394caeaaa6895003e63cd62a052c491f8

Change-Id: Ica87b4ab511d26e0372502f069afc0e3baaa3256

Lower layer ignores alignment

Change-Id: If16df951ecefddc804a6effe013058afc595d30f

…code on windows

Change-Id: Id80cef5a36db8707276de052cbaf73b6826d222f

Change-Id: Id23882286cb2a0d0472964ffc501ab27b7dc7f00

Signed-off-by: shadi <[email protected]>
Change-Id: Iff3ee7dcbcd24836f227fdc9bd5ff4b554ac914f

Change-Id: Ie3f3c0bcea84368c1b0607fd52b4bc7cae41c512

Make sure graph mempool unmaps VA on release

Change-Id: Id3f1bd8d0115b533ae60aa5ba3676b8bf7e5b961

…a single virtual address.

Change-Id: Ie678e607a64f2e5c35a10b9083185f041c5527ac

The .bat file will not be removed from windows pacakge.
Also used cmake install(PROGRAMS …), which will set the correct permission rather than hard coding the permission

Change-Id: I8b57778b59f70e01de949be2ea353b67eec70d2f

Change-Id: Iedd6290a813d6e43a4350709484f78e05b08adc8

Change-Id: I1f8411eae9ed49632667e244a25f223fed92c720

Change-Id: I7fda94e61121f9d3a30f4ad185b8a97712922f3c

Change-Id: Ief73138faed0af70b90186db5bde6689e0a83f88

…ss the flag to HSA API.

Change-Id: I1bafeaa3096395c729723af958d609bc41e7845c

- Add LOG_TS mask for printing signal times
- Read raw ticks from signals

Change-Id: Ibdd0bf06c790729f6c65083a4784c97a3c3219e0

- Fix possible buffer overflow for long kernel names

Change-Id: I3c51669de7ff242d03f9210ee045b6d5e7ac274a

Change-Id: I820cacd75a36363d1387e2e881c96937281bc265

Change-Id: Id3cb4b58b9efd3aceec4377d4d1d3a053c50333e

Change-Id: I3601373f680aa4bff0075f4b3b9e885e54b4600d

1.Make runtime use comgr to unbundle code objects
2.Support compressed/uncompressed modes
3.Remove HIP_USE_RUNTIME_UNBUNDLER and
  HIPRTC_USE_RUNTIME_UNBUNDLER to simplify logics
4.Add comgr wrapper for
  amd_comgr_action_info_set_bundle_entry_ids()

Change-Id: Ic41b1ad1b64cca1e31986437983a5146d52a7329

The Readback and Avoid HDP Flush memory ordering workaround is
used as a fallback solution only when HDP flush register is invalid

Change-Id: Ic284eba1f95ed22b0270d3abeb904fb902015b1a

Unbundle compressed code objects needs comgr 2.8 or higher

Change-Id: I23942d2038e19b02c3ea5d3c9c1fe5367db87136

…ltiple translation unit.

Change-Id: I92179ad198abbdaf5aec9c3c4ba76eeb6b0cc761

Change-Id: I21abe109ddfabfe7640bf78a96c81a1317d31952

Change-Id: Iabb4071bb77201576bc2c0488a04f4fa188815df

Change-Id: I374ea7c3119b0c61f9846a862c4a448ddb179748

Switch commands creation to the new suballocator to avoid
frequent expensive OS calls

Change-Id: I3597c811820e577c15708bad8b8a41aa53acc400