~45% to 50% of Performance drop on rocBLAS_int8 test
Use the last command in the queue for a wait.
Add extra print information about processed commands.
Add an option to disable file location printing.
Change-Id: I4187883e1a90e571fde3128af98368108fda8785
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
When we're aligning rowPitch to imagePitchAlignment, rowPitch is in pixels,
but imagePitchAlignment_ is bytes, so we end up overaligning the pitch.
Convert imagePitchAlignment_ to pixels before doing any logic.
Change-Id: Ia5ab9d54bed150fe974e86b060dbadc196165b29
Make __gnu_h2f_ieee and __gnu_f2h_ieee visible so that hipTestHalf
test can succeed in Clang compiler + Hcc RT.
Change-Id: I5f7d5db19e559b3b66356f0170a8dbc1e5505f3e
hip_threadfence_system passes locally with this change. This also fixes
hipHostMalloc() failures when hipHostMallocMapped flag is used.
Change-Id: Id412efe502accc7c6e7676b52c05ccb9d8fbbe67
Remove a workaround to CS_PARTIAL_FLUSH added in CL#1495187,
since PAL is no longer uses CS_PARTIAL_FLUSH.
Change-Id: I03edc7595459e19aad33b2b0901f0ebe4754d310
[hipclang-vdi-rocm][perf]~45% to 50% of Performance drop on
rocBLAS_int8 test
- Enable AMD_OPT_FLUSH optimization by default to match HCC
- Disable CPU writes to GPU memory on boards with large bar,
because it requires HDP flush tracking.
- Enable L2 cache on kernel arguments, because L2 will be
invalidated on memory reuse .
Change-Id: I124cf250bdd4d19c523ce542c163813828f8fbdc
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.
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
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.
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
For hipLaunchKernelGGL(), hmod->kernargs is empty, thus we need
insert hmod->kernargs[name_str] which is empty.
Change-Id: I95f818d0525da84452e66c5778f0648a643843c7
Vladislav Sytchenko
German Andryeyev
Chauncey Hui
Tao Sang
Karthik Jayaprakash
Jason Tang
Christophe Paquot
Payam
Jatin Chaudhary
agodavar
Sarbojit2019
eshcherb
Evgeny Shcherbakov
Alex Voicu
Yaxun (Sam) Liu
Rahul Garg
Maneesh Gupta