This change adds a new HIP API `hipExtHostAlloc` which preserves
the functionality of `hipHostMalloc`.
Change-Id: I13504c6fc13465ddd7aed329795bb4f2fef1baff
Integration into pytorch pointed out some issues, value narrowing, to
fix this we are now using unions. Also removed check for -munsafe*
compiler flag. The check is now just on builtin detection.
Change-Id: I49364503fa429bd862952f9b29879072afa6d553
- Fixes -0.0 and +0.0 comparison. For atomicMax if the value on
address is -0.0 and on val is +0.0, gfx90a's unsafe atomics will swap
them. This behavior should be consistent with cas loop as well.
- _system variants of atomicMax and atomicMin are resulting in
incorrect output. Updated these to use the similar implementation as
atomicMax and atomicMin.
Change-Id: I20df36ee29ae0434a6b564f2ba71193fe41cfa59
Since we made the members public, we can optimize some operations which
do not require redundant conversions to half_raw types.
Change-Id: I31555ef18e695d8e24b89f0418187fa4e932a38a
In case when the tile size is greater than the number of active threads,
the coalesced group size should be equal to the number of active threads.
Change-Id: I1d41322f2428a07862a590cb5d34b01243383b7c
The warpSize variable is set to the value of the __AMDGCN_WAVEFRONT_SIZE macro,
which is a meaningless default in host code.
The resolution for SWDEV-449015 will introduce diagnostics for uses of this
macro in host code, which includes the current definition of the warpSize
variable. With the __device__ specifier, the definition of the warpSize
variable will not cause these diagnostics.
This change does not stop the variable from being used in host code since clang
intentionally does not diagnose uses of __device__ constexpr variables in host
code.
Change-Id: I0317217affe94fdf2dfd9ad0f134e68f5173245f
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
- enforce incrementing the table versioning number when a table size changes outside of ifdef for ROCPROFILER_REGISTER
- add new HIP_ENFORCE_ABI entries
- update the HipDispatchTable size and bump HIP_RUNTIME_API_TABLE_STEP_VERSION to 1
- re-enable rocprofiler-register
Change-Id: Ie0cc1d8491c5640056e5dd393ea243e4dce4e8a9
- Add API table versioning defines to hip_api_trace.hpp
- Add rocprofiler-register implementation in hip_api_trace.cpp
- created ToolInit implementation
- moved some local functions into anonymous namespace
- renamed some function to be overloads
- made most of the initialization code templated (reduced code duplication)
- enforce ABI
- Updated hipamd/src/CMakeLists.txt
- find rocprofiler-register package (enabled by default but not required)
- Address review comments
- simplify size calculation for dispatch table
- remove setting .size field in GetDispatchTableImpl
- Fix windows build
- Do not enforce ABI on Windows
Change-Id: I08766e8a083528a1236996274bf4522e0e8a1b9f
The following builtins from the CUDA spec are implemented:
- __all_sync, __any_sync, __ballot_sync and __activemask
- __match_any_sync and __match_all_sync
- __shfl_sync, __shfl_up_sync, __shfl_down_sync, and __shfl_xor_sync
The following builtins are NOT implemented, pending support in the compiler:
- __reduce_add_sync, __reduce_min_sync, __reduce_max_sync
- __reduce_and_sync, __reduce_or_sync, __reduce_xor_sync
Change-Id: I07dedbbfe5449f4b5c9b040bed59f5603ccec8c3
Ioannis Assiouras
Jatin Chaudhary
Sameer Sahasrabuddhe
Satyanvesh Dittakavi
Rahul Manocha
Jaydeep Patel
Sourabh Betigeri
Julia Jiang
Marko Arandjelovic
sdashmiz
Fabian Ritter
Saleel Kudchadker
Anusha GodavarthySurya
kjayapra-amd
Rahul Manocha
Jonathan R. Madsen
jiabaxie