If a user passes a ptr that was allocated by hipMalloc(), the offset is guaranteed to be 0 and NULL may be passed as the offset parameter. We shouldn't return an error in this case.
Change-Id: I4a8d645121e5a17d5e2861a0629356a3599de9ee
Also make D2H and H2D keep track of the chain of events
when we need to use a different HostQueue.
Change-Id: I1c5da6ea6104b37ad7aac00f0eb8ea9371e6ba1c
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
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
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
hipEventRecord should always create a new marker so it can track work going on at the time the API is called.
Change-Id: I10ce98044be894fbacab8798441ec3d3f2753b93
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
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
SWDEV-125823 - Fix the build issue due to API interface change.
- `hipOccupancyMaxActiveBlocksPerMultiprocessor` interface is revised
and the runtime needs updating.
Affected files ...
... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#60 edit
Change-Id: Ia7901b0dbbfd37977ce4adf2ae1a821aba0ac044
Vladislav Sytchenko
Saleel Kudchadker
Christophe Paquot
Payam Ghafari
Karthik Jayaprakash
Evgeny Shcherbakov
Tao Sang
Laurent Morichetti
Mark Searles