The current implementation of surd2D{read/write} directly addresses into
the image buffer via the hipArray::data ptr. This is incorrect to do
since we don't know the layout of the image. Also with VDI we won't have
access to the underlying image buffer.
Disable the surface api untill the device functions are switched to
using __ockl_image_{load/store}().
Change-Id: I19a33680176812d5aad3660e9045812061a1c443
[ROCm/clr commit: 8c8d963c65]
Make __gnu_h2f_ieee and __gnu_f2h_ieee visible so that hipTestHalf
test can succeed in Clang compiler + Hcc RT.
Change-Id: I5f7d5db19e559b3b66356f0170a8dbc1e5505f3e
[ROCm/clr commit: 314766b4c2]
The existing one can have issues on certain systems, therefore this limits use of direct memcpy via largeBAR to sizes where it is unequivocally better.
Also addresses SWDEV-220030 and SWDEV-222237.
[ROCm/clr commit: 730f23829e]
And also don't optimize the case where start==stop event to compute
elapsed time since the command can be a NDRange one.
HIP directed test will need to be fixed for that.
Change-Id: I64fadd6ab8ab1a490e7a2b7165a591df5a5cf3a2
[ROCm/clr commit: 1f5ae789bb]
This is a quick workaround to match HCC behavior for performance since inlining usually
results in more optimization opportunities therefore better performance.
We will fine tuning inline threashold later.
[ROCm/clr commit: f4120043ed]
Temporarily comment out Hcc-specific template functions
hipExtLaunchKernelGGL and hipOccupancyMaxPotentialBlockSize for CLang
compiler so that all test cases under hip/samples can be built
successfully for Clang + Hip/Hcc runtime.
Change-Id: Iafc761257be4a7b34eafa6759a01f369570cd6ce
[ROCm/clr commit: 15b91e9f1e]
For hipLaunchKernelGGL(), hmod->kernargs is empty, thus we need
insert hmod->kernargs[name_str] which is empty.
Change-Id: I95f818d0525da84452e66c5778f0648a643843c7
[ROCm/clr commit: 05785772e3]
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/clr commit: eb8c4d1ce2]
* Use deque instead of vector for code readers so that the iterators and references will be stable
* Fix compile error
* Assign the iterator
* Add multithreaded test
* Make threads a multiple of hardware concurrency
* Output on failure
* Add setDevice to try and initialize the context on cuda
* Create context for cuda
* Set context on each thread
* Reduce threads on cuda
* Skip test on cuda
* Try to initialize the primary context on cuda
* Push ctx to the stack as current
* Revert "Push ctx to the stack as current"
This reverts commit c1b707ff63.
* Revert "Try to initialize the primary context on cuda"
This reverts commit d9cbd5a143.
* updated test for nvidia path
* Add c++11 option for nvcc
Co-authored-by: satyanveshd <53337087+satyanveshd@users.noreply.github.com>
[ROCm/clr commit: e085417370]
Some PyTorch unit tests have regression. Disabling cov3 to allow more
time to debug and unblock PyTorch
Change-Id: Iba7f425ef3499c20c42ec45d9152b5d27ce97d03
[ROCm/clr commit: 13052f27c2]
hipEventRecord should always create a new marker so it can track work going on at the time the API is called.
Change-Id: I10ce98044be894fbacab8798441ec3d3f2753b93
[ROCm/clr commit: eb132ccef8]
Those APIs should look at the device associated with the stream first.
If that stream is null then get the current device ID.
Change-Id: Iedde1d1644818ba64f128b988f0bd9674f5b8ad6
[ROCm/clr commit: 8f5a70a150]
Vladislav Sytchenko
Tao Sang
Karthik Jayaprakash
Christophe Paquot
Jatin Chaudhary
agodavar
Payam
Sarbojit2019
eshcherb
Evgeny Shcherbakov
Alex Voicu
Yaxun (Sam) Liu
Rahul Garg
Maneesh Gupta
Nick Curtis
Siu Chi Chan
Evgeny Mankov