The current texture implementation is based off the one for HIP-HCC. There's a lot of problems with it - only creating images from buffers, hard coding logic and ignoring user parameters. This leads to a whole lot of UB even with simple examples (as seen with RedShift's code).
This CL is aimed to bring the HIP-VDI texture implementation closer to what is described by Cuda.
hipMemcpyAtoA() - image to image copy.
hipMemcpyHtoA()/hipMemcpyDtoA() - buffer to image copy.
hipMemcpyAtoH()/hipMemcpyAtoD() - image to buffer copy.
hipArrayCreate()/hipArray3DCreate()/hipMallocArray()/hipMalloc3DArray() - creates 1D/2D/3D/1D Array/2D Array images.
hipCreateTextureObject() - creates sampler, (optional) creates 1D/2D image from buffer, (optional) creates image views.
hipBindTexture() - creates 1D image from buffer (should create a typed buffer, however this is not compatible with HIP-HCC).
hipBindTexture2D() - creates 2D image form buffer.
hipBindTextureToArray() - creates image view.
hipTexRefSetAddress() - creates 1D image from buffer (should create a typed buffer, however this is not compatible with HIP-HCC).
hipTexRefSetAddress2D() - creates 2D image from buffer.
hipTexRefSetArray() - creates image view.
There are still a lot of TODOs in the code, here's a few important ones:
1. VDI doesn't support a lot of sampler flags.
2. VDI doesn't support device to image 2D/3D copy.
3. Mipmaps implementation is incomplete.
4. Image view implementation is incomplete.
Change-Id: Ia374ee27aa14f76451fee7667495036f4419a487
[ROCm/clr commit: f71817a342]
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]
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]
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]
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]
The issues of the following functions have been fixed.
hipModuleLoad: Make Hip-Vdi runtime able to read code object module
generated by Hcc compiler.
hipLaunchKernel: Use introspect method to find function if it cannot
be found from platform state instance.
Change-Id: Id740e5a96614ec6a0b6c704f8f74600bfdc4983e
[ROCm/clr commit: 62ef029288]
Vladislav Sytchenko
Payam Ghafari
Tao Sang
Karthik Jayaprakash
Christophe Paquot
Jatin Chaudhary
agodavar
Evgeny Shcherbakov
Aaron Enye Shi
Laurent Morichetti
Mark Searles
Jatin Chaudhary