Commit Graph

64 Commits

Author SHA1 Message Date
Vladislav Sytchenko 600b1006c8 Correct the declaration of hipBindTexture()
The texture reference needs to be passed as a constant pointer.

Change-Id: I36ca0bddaba30becfc2ce70dd9e5b7db66c57f27
2020-03-18 18:08:01 -04:00
Vladislav Sytchenko da5f852cdd Merge "Start the lifetime of the texture reference" into amd-master-next 2020-03-18 18:07:24 -04:00
Jiabao Xie bbfba3bae1 Merge "Cleaned include statements and deleted unnecessary operator overloads" into amd-master-next 2020-03-18 14:37:19 -04:00
Vladislav Sytchenko b4a0008b36 Start the lifetime of the texture reference
reinterpret_cast<> doesn't create an object, so the texref is actually unitiliazed. This may lead to garbage data in some of its struct members.

Initialize it by performing a placement new. The constructer should set all of its members to default values. There's no way currently to extract the channel type, so use single channel char for now.

Change-Id: I41b305a75bb3f30130324de785099f55b3e130c7
2020-03-18 12:30:11 -04:00
Sarbojit Sarkar 957caff8b4 [hip-vdi]Fix for TF build failure [SWDEV-225827]
Change-Id: I8478779bef92bad8353b8d066b28c220bb59b98d
2020-03-17 22:52:01 -04:00
Jiabao Xie f81175c3bc Cleaned include statements and deleted unnecessary operator overloads
Change-Id: I5eb322c05082e8f27584446af2f6a19243dc63c9
2020-03-12 14:13:01 -04:00
Saleel Kudchadker 7d86f10c3c Merge "Fix P4 linux make builds" into amd-master-next 2020-03-11 16:17:18 -04:00
Jiabao Xie 214c07a5b8 Merge "Structs expanded to print struct information than address. Moved to hip_formatting.hpp. Reformatted to follow Google standards" into amd-master-next 2020-03-10 15:56:49 -04:00
Jiabao Xie 10d90ab2ad Structs expanded to print struct information than address. Moved to hip_formatting.hpp. Reformatted to follow Google standards
Change-Id: I08695058c11db51b9f3cbe1deb4af944ebf9e64d
2020-03-10 13:12:28 -04:00
Saleel Kudchadker 7bba326ed1 Fix P4 linux make builds
Change-Id: I6cc4e69a914389b53bc3b52535eb6faf24897a09
2020-03-10 00:02:10 -07:00
Saleel Kudchadker 8d05d06dce Merge "Change HIPRTC Version to 9.0" into amd-master-next 2020-03-09 22:55:27 -04:00
Vladislav Sytchenko cd76eb7486 Add hipDrvMemcpy3D.
This is the equivalent of cuMemcpy3D.

Change-Id: Ib2e06dbd6f5093c931cdfd36c87617f32acffc2d
2020-03-09 16:11:25 -04:00
Saleel Kudchadker f34d5bdfb1 Change HIPRTC Version to 9.0
Change-Id: I7e88df61248f0fe6e33c2315805f9e49a493cf29
2020-03-09 11:28:55 -07:00
Saleel Kudchadker 514c7c8ef8 Fix HIPRTC headers to export C style symbols
Change-Id: I3e0d2b19ace4a9096e3e46bd22f420483da51a8a
2020-03-09 14:18:46 -04:00
Vladislav Sytchenko a15e895cfd Fix Windows build.
extern "C" on Windows implies nothrow. We shouldn't be throwing exceptions either way.

Change-Id: If0ed1f7ec194bf7f65b7cea1a5c250e768a8f190
2020-03-09 11:46:55 -04:00
Vladislav Sytchenko 59fb85f94d Add hipMemcpy3DParams -> HIP_MEMCPY3D conversion
Change-Id: I66dd1ae722b8c0d7ddb2562c958b14854e6b86bc
2020-03-06 14:33:52 -05:00
Vladislav Sytchenko 7ffa6eb963 Fix dangling pointer after hipUnbindTexture() call
Change-Id: Ic4b476c62ebfae31e94dd139b20b6aaaa52bb866
2020-03-06 14:10:56 -05:00
Vladislav Sytchenko e746584d24 Add missing coversion cases.
Handle converting signed int cases of hipResourceViewFormat to number of channels.

Change-Id: Ica8ae6f644edfaa0d4803d0b8e90e320479118e2
2020-03-06 14:06:21 -05:00
Vladislav Sytchenko 801c70279f Correct logic in ihipMemcpyAtoA()
HIP assumes that image width is in bytes, but VDI assumes that image width in pixels. Need to perform byte -> pixel coversion before doing anything.

Change-Id: Ia9fd1f46d05db3fbe8049add10b4d7e5118a2b9a
2020-03-05 15:03:18 -05:00
Vladislav Sytchenko aa490c96a1 Fix typos in ihipMemcpy3D()
Change-Id: I8720f113642f00bb013cf46284e9b13cc932bf4a
2020-03-05 14:11:55 -05:00
agodavar f360ff6479 Enable tests that are passing on hip-vdi
Change-Id: I7de965f7e8bb7e4d0fa61cd584f0cd118c1e212e
2020-03-05 16:08:51 +05:30
German Andryeyev a083dc8547 Merge "SWDEV-193956 - [hipclang-vdi-rocm][perf]" into amd-master-next 2020-03-04 18:04:52 -05:00
Michael LIAO d2d0225c2e Fix -Werror=switch.
Change-Id: Ib0850a3f79308efeeecd18bb03ecd82bd355ef85
2020-03-04 17:52:39 -05:00
Saleel Kudchadker f4caaeb308 Merge "Add -Wno-ignored-attributes flag to mute warnings" into amd-master-next 2020-03-03 20:39:16 -05:00
Vladislav Sytchenko 14810044f7 Fix Windows build.
Calm down warning C4715: not all control paths return a value.

Change-Id: I99d939b84770a499bb6a0edcc3bc0bf12961a711
2020-03-03 16:14:12 -05:00
German Andryeyev b93d997fb7 SWDEV-193956 - [hipclang-vdi-rocm][perf]
~45% to 50% of Performance drop on rocBLAS_int8 test

Enable cudaSetDeviceFlags() api call. Use active wait by default
for all devices.

Change-Id: Ifc2ebe3dd9b0aa3fdbfbc9cb5c2cd8b3b726124f
2020-03-02 18:35:56 -05:00
Vladislav Sytchenko d28b77bf23 Fix hipMemcpy3d (partially)
Incoming changes from upstream split the struct hipMemcpy3DParms into two separate ones - hipMemcpy3DParms and HIP_MEMCPY3D, which are cudaMemcpy3DParms and CUDA_MEMCPY3D equivalents respectively.

Note that HIP_MEMCPY3D is missing half the members of CUDA_MEMCPY3D (this should be fixed in PR#1887). Work around this by using a substitute _HIP_MEMCPY3D struct for now.

Change-Id: Ic15134e6deb260189b662b3804d2309a9b8473e9
2020-03-01 13:52:05 -05:00
Saleel Kudchadker d292b4eee9 Add -Wno-ignored-attributes flag to mute warnings
Change-Id: Ib9ad75554c0ddfe025e7d500f057830f7ad2ef5b
2020-02-27 10:52:13 -08:00
Christophe Paquot e9af4c8794 Blocking and default streams' sync:
Add hip::syncStreams(dev) to sync blocking streams on a given device.
hip::syncStreams(void) should only sync streams on the current device.

Change-Id: Ib6b0735215fa0ed12c646ebd029e9763ee3712ce
2020-02-26 08:54:00 -08:00
Payam Ghafari 630e28fcf1 Merge "removed excutable permission from source files" into amd-master-next 2020-02-25 17:44:43 -05:00
Payam eb4bd38c27 removed excutable permission from source files
Change-Id: Iae8639a96c55a098e28de41c5a3f38a07acbe25c
2020-02-25 16:16:47 -05:00
Vladislav Sytchenko 21ec6d005e Bump c++ version to 14.
This is to match our p4 build.

Change-Id: I084b07968ced98cca216146a0228cc36e9f56ab3
2020-02-25 14:41:55 -05:00
Vladislav Sytchenko fd76d220a5 hipBindTexture() should handle nullptr offset.
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
2020-02-25 13:50:34 -05:00
Saleel Kudchadker 1ccaea7ca8 Use the context variant of getNullStream
Do not create a new queue to call finish in hipFree if none was
created earlier elsewhere.
Change-Id: I87bb191e6b186ddbe607ab29d11e3ae5bc2ac8e6
2020-02-25 00:13:43 -08:00
Christophe Paquot 2bdfc73649 Fixed a few multithreaded potential issues
Also make D2H and H2D keep track of the chain of events
when we need to use a different HostQueue.

Change-Id: I1c5da6ea6104b37ad7aac00f0eb8ea9371e6ba1c
2020-02-24 20:14:10 -08:00
Payam Ghafari 45b1691e0a Merge "updated lib elf path and clean ups" into amd-master-next 2020-02-24 17:35:36 -05:00
Vladislav Sytchenko f71817a342 HIP-VDI texture rework
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
2020-02-24 15:23:45 -05:00
Payam 3beac02cbd updated lib elf path and clean ups
Change-Id: Id0b3c295fa8353a5da8517204bf53dab9887defb
2020-02-24 15:09:34 -05:00
Payam Ghafari 9518af6083 Merge "export hip::host and hip::device" into amd-master-next 2020-02-24 13:29:01 -05:00
Payam e5c8903925 export hip::host and hip::device
Change-Id: If1427a180f91d3f8bae203d956f21cd69345c060
2020-02-21 19:58:42 -05:00
Vladislav Sytchenko 16ffc6fe3e Update hipGetErroName() to match hipError_t
Change-Id: I8f7fe0cca01ddec5d6333ba6e876128276323be9
2020-02-21 18:31:27 -05:00
Vladislav Sytchenko d934d731a0 Fix Windows build
MSVC unlike gcc doesn't add colons for you.

Change-Id: I06d81a9a9b346065d0452fe7117ab82144a06f74
2020-02-21 14:37:41 -05:00
Vladislav Sytchenko 8be685e7b9 Report the HW requirments for pitch alignment
Change-Id: Iaaa9d597dff57cfad5d07d931f881aba1a5f98f1
2020-02-21 11:09:47 -05:00
Vladislav Sytchenko 8c8d963c65 Disable hip{Create/Destroy}SurfaceObject
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
2020-02-21 11:09:28 -05:00
Karthik Jayaprakash 4b2ff6ec91 SWDEV-223674 - Return hipErrorNoBinaryForGpu in case particular binary is not found in clang offload bundler.
Change-Id: Iaa08fcdc8ecb719edd9f81e4a1456ea642f362f4
2020-02-19 20:01:36 -05:00
Christophe Paquot 13bf30569e SWDEV-223262
hipMemcpyWithStream is supposed to be synchronous.

Change-Id: Ie44e37ecc9246e26a6b315c01e88a279f9e42fd7
2020-02-19 14:08:12 -08:00
Evgeny Shcherbakov 9c4e11f217 Merge "adding 'hipHccModuleLaunchKernel' and 'hipExtModuleLaunchKernel'" into amd-master-next 2020-02-18 18:10:22 -05:00
Christophe Paquot b4ad4262cc Introducing hip::Device which wraps around amd::Context and deviceId
Change-Id: Ie35a6edb65c001b35eb9f5d2af26e765dc41c00e
2020-02-18 17:18:56 -05:00
Karthik Jayaprakash e066cc6f60 SWDEV-223394 - Pass module info from hipModuleGetTexRef to internal Platformstate:: functions.
Change-Id: I7d1ba3f940f595c3fca74a57fa20f484c52d4741
2020-02-18 11:23:03 -05:00
Christophe Paquot 1f5ae789bb Don't create a marker for start event in hipModuleLaunchKernel
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
2020-02-17 14:16:31 -08:00