- 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
[ROCm/clr commit: d69cc35750]
Since we made the members public, we can optimize some operations which
do not require redundant conversions to half_raw types.
Change-Id: I31555ef18e695d8e24b89f0418187fa4e932a38a
[ROCm/clr commit: 6a655a77e7]
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
[ROCm/clr commit: 152f343124]
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
[ROCm/clr commit: 819e537dc5]
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
[ROCm/clr commit: 49349f168c]
- 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
[ROCm/clr commit: d84c5ae3af]
This should be used in place of dlsym or GetProcAddress (linux and windows respectively)
Change-Id: I5501b538e03892e8e5a2282678d848fcaf21d911
[ROCm/clr commit: 0479cdb3dd]
- 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
[ROCm/clr commit: 0f1fdc7794]
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
[ROCm/clr commit: c5ab5680b4]
Sameer Sahasrabuddhe
Satyanvesh Dittakavi
Rahul Manocha
Jaydeep Patel
Sourabh Betigeri
Jatin Chaudhary
Julia Jiang
Marko Arandjelovic
sdashmiz
Ioannis Assiouras
Fabian Ritter
Saleel Kudchadker
Anusha GodavarthySurya
kjayapra-amd
Rahul Manocha
Jonathan R. Madsen
jiabaxie