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]
SWDEV-2 - Change OpenCL version number from 3085 to 3086.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#2833 edit
[ROCm/clr commit: 914c2eb8b6]
SWDEV-197836 - Drop the use of llvm header files in opencl runtime
- COv2 doesn't report HostCall argument properly. Make a workaround for it.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/devkernel.cpp#34 edit
[ROCm/clr commit: 6ace9abcd6]
SWDEV-2 - Change OpenCL version number from 3084 to 3085.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#2832 edit
[ROCm/clr commit: 0edb23bf71]
* 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-2 - Change OpenCL version number from 3083 to 3084.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#2831 edit
[ROCm/clr commit: 7bbcdc4186]
SWDEV-2 - Change OpenCL version number from 3082 to 3083.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#2830 edit
[ROCm/clr commit: 2e299c52b0]
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]
SWDEV-2 - Change OpenCL version number from 3081 to 3082.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#2829 edit
[ROCm/clr commit: 357a13b91c]
SWDEV-2 - Change OpenCL version number from 3080 to 3081.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#2828 edit
[ROCm/clr commit: 1fe7807642]
=> 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-79445 - OCL generic changes and code clean-up
Make the conversion from amd::Coord3D to size_t* be explicit.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/platform/object.hpp#21 edit
[ROCm/clr commit: 65b47104c7]
SWDEV-2 - Change OpenCL version number from 3079 to 3080.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#2827 edit
[ROCm/clr commit: 43ad5275cb]
SWDEV-79445 - OCL generic changes and code clean-up
Alllow amd::Coord3D to decay into size_t*. This allows creating an amd::BufferRect obect without the need of explicitly passing size_t[3] arguments.
ReviewBoardURL = http://ocltc.amd.com/reviews/r/18473/diff/
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/platform/object.hpp#20 edit
[ROCm/clr commit: a9b467ba7b]
SWDEV-2 - Change OpenCL version number from 3078 to 3079.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#2826 edit
[ROCm/clr commit: 7ceef75528]
SWDEV-219121 - [Navi][ROCm] Add performance counter support
This is initial implementation for Navi 10 performance counter in OpenCL runtime.
Tests:
1. http://ocltc.amd.com:8111/viewModification.html?modId=130609&personal=true&tab=vcsModificationBuilds
2. ./ocltst -m oclruntime.so -t OCLPerfCounters
Before this code change, the segmenation fault happens inside OpenCL runtime. After this code change, the error happened inside HSA. Error message is generated in function hsa_ven_amd_aqlprofile_start. Inside this function, the C++ try block create exception. The exception error in HSA is "GFXIP is not supported(gfx1010)". We need HSA to add support for Navi.
ReviewBoards: http://ocltc.amd.com/reviews/r/18463/
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/roccounters.cpp#5 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/roccounters.hpp#4 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#94 edit
... //depot/stg/opencl/drivers/opencl/tests/ocltst/module/runtime/OCLPerfCounters.cpp#50 edit
[ROCm/clr commit: 68107b160f]
SWDEV-2 - Change OpenCL version number from 3077 to 3078.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#2825 edit
[ROCm/clr commit: e3e2727c47]
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]
SWDEV-2 - Change OpenCL version number from 3076 to 3077.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#2824 edit
[ROCm/clr commit: c583dfddf8]
SWDEV-2 - Change OpenCL version number from 3075 to 3076.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#2823 edit
[ROCm/clr commit: b7dabf5c3b]
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-2 - Change OpenCL version number from 3074 to 3075.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#2822 edit
[ROCm/clr commit: bc91a60139]
SWDEV-218016 - There is a bug in the OCL/DX interop
- Initialize event query value to FALSE and correct the loop condition to make sure query waits for the end.
http://ocltc.amd.com/reviews/r/18447/
Affected files ...
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_d3d10.cpp#18 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_d3d11.cpp#28 edit
[ROCm/clr commit: 0a27603178]
paulfreddy
Siu Chi Chan
Alexey Chernov
foreman
Evgeny Mankov
Rahul Garg
ansurya
kpyzhov
Aakanksha Patil
Jatin Chaudhary
Yaxun (Sam) Liu
Maneesh Gupta