* 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/clr commit: 6da0c82b78]
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/clr commit: 8fc7cad90f]
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/clr commit: 9f5a4148ce]
* Use hipExtLaunchKernelGGL in dispatchlatency sample
* Let it run on NVCC path too
* Refactoring
* Add test_kernel source
* Remove ResultDB
* Remove error checks
[ROCm/clr commit: 8b8e4a66b4]
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/clr commit: 37bb58b556]
=> 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/clr commit: 688b6e6af0]
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/clr commit: ec3fbfd60c]
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/clr commit: 55525fa466]
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 47725baf8c.
- Revert "Fix rocFFT build failure (#1777)"
This reverts commit f7bac3e057.
[ROCm/clr commit: c091635f5a]
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/clr commit: e5c5338146]
* 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/clr commit: 26b50e1e1b]
SWDEV-213000 - Backout change to notifyCmdQueue as the real issue has been identified
Affected files ...
... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#53 edit
[ROCm/clr commit: a941903848]
HIP_VERSION_MAJOR, HIP_VERSION_MINOR, HIP_VERSION_PATCH and HIP_VERSION pre-processor macros are now defined in hipVersion.h instead of being set by hipcc.
[ROCm/clr commit: 4fd77f8336]
Laurent Morichetti
satyanveshd
vsytch
Evgeny Mankov
aakanksha555
mshivama
paulfreddy
Siu Chi Chan
Alexey Chernov
foreman
Rahul Garg
ansurya
kpyzhov
Aakanksha Patil
Jatin Chaudhary
Yaxun (Sam) Liu
Maneesh Gupta
Rick Ho
Sarbojit2019