Review comments - generate hiprtc lib everytime when HIP_PLATFORM is hcc
Changes for hip-clang
Removing pre processor directive to simplify
Change-Id: Id38ab368362b58ee0458baeb8051fea709ae6bba
[ROCm/hip commit: 54447268cc]
Make __gnu_h2f_ieee and __gnu_f2h_ieee visible so that hipTestHalf
test can succeed in Clang compiler + Hcc RT.
Change-Id: I5f7d5db19e559b3b66356f0170a8dbc1e5505f3e
[ROCm/hip commit: ceb942375d]
For hipLaunchKernelGGL(), hmod->kernargs is empty, thus we need
insert hmod->kernargs[name_str] which is empty.
Change-Id: I95f818d0525da84452e66c5778f0648a643843c7
[ROCm/hip commit: 30da92e2af]
Fix the following issues:
1.Ignore hidden arguments of kernel functions.
2.Look up both origial function name and function name with .kd postfix
when argments are retrived from module.
3.Addition, fix compiling issue of LaunchKernel test app.
Change-Id: I9400943f2f02433cb4409b19c0cac3626c2bc454
[ROCm/hip commit: 9a005eacf3]
* Device texture functions should not normalize the sampled pixel. This is already done by HW.
* Add support to use h/w capability for normalized float data convertion for driver API's
Co-authored-by: ansurya <50609411+ansurya@users.noreply.github.com>
[ROCm/hip commit: ef514eef71]
* Fix bug in LaunchKernel test
Instead of passing the address of the gpu buffer, pass the address
of the pointer that holds the address of the gpu buffer
* Fix hipLaunchKernel's kernarg buffer construction.
The hipLaunchKernel implementation should rely on ihipModuleLaunchKernel
to construct the kernarg buffer correctly based on kernel metadata.
* Fix a bug in get_functions where the Kernel_descriptor wasn't constructed with the correct kernarg layout information.
* Fix a bug in kernarg layout parsing dealing with kernel without any arg
* teach ihipModuleLaunchKernel to handle kernel without any arg
* Add a more interesting test
[ROCm/hip commit: bff8e15e13]
* Add missing texturePitchAlignment member to the hipDeviceProp_t struct.
* Add missing hipDeviceAttributeTexturePitchAlignment enumerator to the hipDeviceAttribute_t enum.
* Initialize texturePitchAlignment to 256. This works for gfx9+, but is technically overaligned in most cases for pre-gfx9.
* Add the texturePitchAlignment property to the NVCC path.
[ROCm/hip commit: f72a669487]
The current implementation skips this procedure for a given device
object when a global symbol is found in the cache. This is incorrect:
- There could be other undefined globals that have not been previously
encountered further down the list
- If a symbol is found in the cache, it doesn't need to be pinned again
but it still need to be defined for the current executable
Added special case for the printf buffer symbol (already pinned by HCC)
The bug was exposed by running printf on different GPUs.
[ROCm/hip commit: 6613a37b3b]
=> New ROCr calculates pitch as per HSA specification and addrlib is used to check whether HW can support that configuration. Hence few texture tests are failing with HSA_EXT_STATUS_ERROR_IMAGE_PITCH_UNSUPPORTED.
=> Determine pitch for linear images and always pass rowpitch to HSA API's.
[ROCm/hip commit: df20c17f12]
Fixes SWDEV-218626 and SWDEV-218629
Changes:
- Revert "`static inline` in a header, just like excess sugar in a diet, causes bloat (#1692)"
This reverts commit b7803eb0c6.
- Revert "Fix rocFFT build failure (#1777)"
This reverts commit 9c8c3df294.
[ROCm/hip commit: 00bd5d1cec]
* Detect when an explicit printf buffer flush is required
in a device/stream synchronization function.
* hip_module.cpp: add missing hc_am.hpp header
[ROCm/hip commit: f4555c835a]
This PR is a follow-up on PR# #1698 and it makes two more APIs (hipLaunchCooperativeKernel/hipLaunchCooperativeKernelMultiDevice) inline so that they can work correctly with lazy binding.
[ROCm/hip commit: 6968aeb841]
SWDEV-151670: Issue with 3D texture with 4 components
SWDEV-151671: Issue with 2D layered texture with 4 components
Fixed memcpy when memory is allocated with driver API's.
Github issues: #1755
Fixed 3D default case when array type is not set during memory allocation.
[ROCm/hip commit: 391e99cd12]
Reverting #1673, #1697 and #1707.
Support for hipMemcpyWithStream and memcpy optimizations, will be brought in again once issues seen with these are resolved independently.
[ROCm/hip commit: e53fc316f1]
[Background] it was found that if lazy linking used for a library that calls hipExtLaunchMultiKernelMultiDevice API then this API can get the wrong program_state object for looking up device kernels leading to a "No device code available" error in this API.
To fix this issue, the API was refactored to be inline and get and pass the correct program_state to an internal hip API to request a multi-device kernel launch.
[ROCm/hip commit: 68cc787781]
Handled the HCC version check appropriately as few of the directed tests (SWDEV-212161) were failing when hcc was bumped to 3.0.
[ROCm/hip commit: 6b06911ef1]
* hipMemset et al can use HSA API directly for synchronous cases
* lock and flush stream in hipMemset, hold lock until complete
* move hipMemset async check to front of conditional
* use hsa_amd_memory_fill for additional sync memset cases
code cleanup/review for all memset calls
* Fix inversion of execution mutating value.
* ihipMemsetSync fall back to kernel if HSA memset fails
* Never fallback, never surrender.
* Allow NULL stream.
* Optimise memset kernel. Remove deadwood.
* Update hip_memory.cpp
* Clean up stream logic in sync memset
* Revert "Clean up stream logic in sync memset"
This reverts commit 6117dedf673367f44cc704192573a117a3d92477.
[ROCm/hip commit: e31e0ca12e]
Jatin Chaudhary
Tao Sang
Christophe Paquot
Maneesh Gupta
vsytch
saleelk
Siu Chi Chan
Sarbojit2019
Evgeny Mankov
satyanveshd
ansurya
Jatin Chaudhary
Rahul Garg
Aryan Salmanpour
Alex Voicu
Jeff Daily
Rahul Garg