* 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]
Don't allow `HIP_PATH` to be propagated to `hipconfig`,
when run by CMake to detect the package version, as it
leads to the wrong version is detected: when there's
already HIP of some different version installed in the
system and `HIP_PATH` points to its location, `hipconfig`
tends to return the version of the installed HIP, rather
than the value defined for the distribution. The compiled
results report wrong version and spoils the rest of the
stack in this case.
[ROCm/hip commit: 8fc262ef23]
* Use hipExtLaunchKernelGGL in dispatchlatency sample
* Let it run on NVCC path too
* Refactoring
* Add test_kernel source
* Remove ResultDB
* Remove error checks
[ROCm/hip commit: 2eb7ae4f0c]
SWDEV-215533 - [HIP][Windows]Output mismatch with tex3D()
When calculating the pitch for hipMallocPitch() align it to the device's image pitch alignment. If we don't, PAL backend will overalign the pitch when creating an image buffer.
ReviewBoardURL = http://ocltc.amd.com/reviews/r/18496/diff/
Affected files ...
... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#96 edit
[ROCm/hip commit: d18baabc70]
=> 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]
SWDEV-205925 - Update HIP texture APIs for issue in hipTexRefSetAddress in HIP/PAL on Windows
-Update HIP texture function parameter to use hipChannelFormatDesc& desc
http://ocltc.amd.com/reviews/r/18456/
Affected files ...
... //depot/stg/opencl/drivers/opencl/api/hip/hip_texture.cpp#33 edit
[ROCm/hip commit: d8c1215113]
The randomly generated offset+width may exceeds 32, which causes
a left shift operation with 32-offset-width. As an unsigned number
that is greater than 32 and causes undefined behavior. When the
test is compiled without -mavx it is still OK. However when
the test is compiled with -mavx, the undefined behavior causes
wrong results and test failure.
This patch adjusts width so that offset+width<=32 always.
[ROCm/hip commit: 02848e0b80]
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]
SWDEV-215533 - [HIP][Windows]Output mismatch with tex3D()
1. Implement hipMemcpy3DAsync().
2. Add logic in hipMemcpy3D() to determine based on src/dst parameters if the user intended on calling cudaMemcpy3D() or cuMemcpy3D().
ReviewBoardURL = http://ocltc.amd.com/reviews/r/18444/diff/
Affected files ...
... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.def.in#38 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.map.in#36 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#95 edit
[ROCm/hip commit: 3cbd083c8d]
* 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]
Laurent Morichetti
Payam
gandryey
Maneesh Gupta
Evgeny Mankov
jujiang-del
satyanveshd
vsytch
aakanksha555
mshivama
paulfreddy
Siu Chi Chan
Alexey Chernov
foreman
Rahul Garg
ansurya
kpyzhov
Aakanksha Patil
Jatin Chaudhary
Yaxun (Sam) Liu