Commit graph

123 Commits

Autor SHA1 Nachricht Datum
Christophe Paquot 4d9c540f4f Remove a map lookup whenever we were getting the default stream
Change-Id: I64b6d1deea41d81e94a58a83de287e78923656b3
2020-04-09 12:44:21 -07:00
Evgeny 1e4e4cb229 enabling hipPrintString (to master-next)
Change-Id: I28859f3dbe5b867a858ca1d76c93e6fab6a68d1f
2020-04-09 09:57:27 -05:00
Saleel Kudchadker f99db12c82 Merge "Revert "Wake up commandQueue before returning"" into amd-master-next 2020-04-06 18:52:29 -04:00
Saleel Kudchadker 71f850f9ec Revert "Wake up commandQueue before returning"
This reverts commit bc8d6ac97c.

Reason for revert: German advised againt this change.

Change-Id: Ia1b1b9db60c965b2d9c006bd7d20012a9d7697e1
2020-04-06 16:46:50 -05:00
Payam 8c26e0b3b6 updated LOG_LEVEL prints to print pid and tid
Change-Id: I8a9212b26bb7e312408a222823efcfd00344094b
2020-04-06 16:58:25 -04:00
German Andryeyev 5fe91ccb1b SWDEV-184710
Support hipLaunchCooperativeKernelMultiDevice()

- Add validation logic for MGPU launches to pass a cuda test

Change-Id: Iccca7fde43493fc3bc6685512d39202271ae3e92
2020-04-06 16:38:27 -04:00
German Andryeyev 382d5ce77f Merge "(SWDEV-228488)" into amd-master-next 2020-04-06 16:30:37 -04:00
German Andryeyev da1444bfc8 SWDEV-184710
Support hipLaunchCooperativeKernelMultiDevice()

- Add hipCooperativeLaunchMultiDeviceNoPreSync and
hipCooperativeLaunchMultiDeviceNoPostSync support to pass a cuda test

Change-Id: If518f11ef2636a2235e5df9e77f879d8ced68102
2020-04-06 15:29:03 -04:00
Vladislav Sytchenko 1bd640b659 (SWDEV-228488)
These fixes address regressions caused by http://gerrit-git.amd.com/c/compute/ec/hip/+/337601

Currently we're converting a 1D offset into a 3D offset, which doesn't make much sense once you consider the fact that this offset is relative to a different origin than our current 3D offset.

I traced through our blit kernels in VDI - the copy buffer rect path is able to handle immediate offsets in the 3D buffer via the amd::BufferRect::start_ parameter.

Instead of adjusting the offset, simply adjust the start of the region.

Change-Id: Ic8797a2c8ac0ad106f246f61ff06ca1ca03d3058
2020-04-06 14:17:11 -04:00
Michael LIAO 7e051d8a96 [vdi] Fix -Wsign-compare warning. NFC.
- TeamCity build failed as `-Werror` is turned on.

Change-Id: Icd2cbd45f60e3c296894e8e73685e1d177f125a8
2020-04-06 12:16:07 -04:00
Christophe Paquot fa5a9b3810 Default HostMalloc to uncached memory
Change-Id: I72e19c7f7820a77fd5afc09f09cfea9acd0b8e84
2020-04-03 19:19:33 -04:00
Saleel Kudchadker f1a9a4f22a Merge "Wake up commandQueue before returning" into amd-master-next 2020-04-03 18:27:03 -04:00
Michael LIAO 9e619430f4 [vdi] Add hipFreeHost
Change-Id: I8a5b7ff3f0ab4f5674efd6723c18808ad6ef33f5
2020-04-03 16:34:28 -04:00
German Andryeyev 6f19e77f69 Merge "SWDEV-184709 - support hipLaunchCooperativeKernel()" into amd-master-next 2020-04-03 16:23:05 -04:00
Vladislav Sytchenko 5f14ae1161 Take into an account the number of channels...
when querying the element size of an array.

Change-Id: Id57d3374b14d80a59230ec8286704f2fbabb0fae
2020-04-03 15:43:18 -04:00
German Andryeyev 7820018037 SWDEV-184709 - support hipLaunchCooperativeKernel()
- Add validation checks for cooperative launch to pass Cuda test

Change-Id: Ie296f0c3f113909d9a357879db3b2a833ab314c5
2020-04-03 15:18:21 -04:00
Saleel Kudchadker bc8d6ac97c Wake up commandQueue before returning
Change-Id: I87eb5a22c81a9cb807474a960b5987d5fb6c2b86
2020-04-03 10:23:36 -07:00
Michael LIAO 00bb2ce04a Fix size type in __hipRegisterVar
Change-Id: I6b667600ae8f133583b768ab963318882b84179f
2020-04-03 10:51:58 -04:00
Michael LIAO 79d8f7e47e [hip] Clean up unnecessary casting.
Change-Id: I64b08aaef5c67ffb49330c9c605611f1fbd3f5a2
2020-04-02 12:46:15 -04:00
Saleel Kudchadker 68d013d030 Merge "OpenCL2.2 Header changes" into amd-master-next 2020-04-02 02:46:49 -04:00
Vladislav Sytchenko aea688b79c Add entry points for hipTexObject*() API
Even though the runtime and driver texture object API is one to one, the structs used by these APIs are not. See hipResourceDesc vs HIP_RESOURCE_DESC differences.

These differences are not trivial and most likely won't be able to handled by hipify, so we need new API entry points.

Change-Id: Id4bcb1ad0ae15378dbdb5a2ed07e5ea30f320082
2020-04-01 14:51:51 -04:00
Saleel Kudchadker 34a3ed9c1b Merge "Cleanup stream from hip:Event class." into amd-master-next 2020-03-31 20:14:48 -04:00
Vladislav Sytchenko 77223a8eca (SWDEV-229354)
This patch is a workaround to support user pitch for hipMemcpy{2D/3D}.

Historically OpenCL didn't support pitch with clEnqueueFillBuffer(), so neither did we in VDI. Adding it now will be slightly nontrivial, since the fill kernel and runtime in many places will need to be modified.

As a temporary workaround for cases when pitch > width, we can just enqueue a fill for each row separately. This implementation is slow, but it satisfies the correctness criteria.

Change-Id: Idfeca349288b51d6ff84a7cf001fb63c6a66818a
2020-03-31 18:12:56 -04:00
Saleel Kudchadker e4ea1f60bb Cleanup stream from hip:Event class.
Change-Id: I98de07d33bb7fea8f5e2d32b288c15f10ce58902
2020-03-31 11:22:00 -07:00
Michael LIAO b72196613a [vdi] Fix hipGetSymbol{Address|Size}
- Use symbol value as the qeury key. Compared to the symbol name, the
  symbol value is more robust as developers may use unqualified or
  qualified identifiers. It also removes the mangling and/or demangling
  requirement for the runtime API.

Change-Id: I9d4259f3842612c7cc98551269fc2092d8b5c19e
2020-03-31 00:26:53 -04:00
Christophe Paquot 94a7ef6ed1 Do not retry to allocate when OOM. Shouldn't be needed since we idle on Free.
SWDEV-229214

Change-Id: I183006f409388e3c7981f2569649d01d6378be46
2020-03-30 12:49:48 -07:00
kjayapra-amd b081912962 SWDEV-216213 - Use different static & dynamic module maps for faster lookup.
Change-Id: Ia605e76a411ad5be04046b9d61f1ac111d49bb4a
2020-03-30 14:28:07 -04:00
Anusha Godavarthy Surya caa083357a Merge "Update Enable/Disable peers to match cuda behaviour" into amd-master-next 2020-03-30 13:46:42 -04:00
agodavar bdb3a4b393 Update Enable/Disable peers to match cuda behaviour
Change-Id: I67194ccf77a0019368579ff7d95b7790fcf228f3
2020-03-30 12:49:16 -04:00
Michael LIAO 3a690e960f [vdi] Fix calculation of MaxWaves
- Consider the case where `usedVGPRs` is zero.
- This fixes [SWDEV-228537](http://ontrack-internal.amd.com/browse/SWDEV-228537)

Change-Id: I8675311f5fe24fb59c5d45bada122afefb55b128
2020-03-30 09:10:16 -04:00
Saleel Kudchadker 5a8add03a5 Merge "Check event status before notify" into amd-master-next 2020-03-27 20:19:40 -04:00
Vladislav Sytchenko 6ed73f50f7 (SWDEV-228794)
Adjust the origin of the copy if the user passes a pointer that wasn't allocated by the runtime.

Change-Id: I0aeb20195ed730857a461a53f537626ec2573fd1
2020-03-27 16:33:16 -04:00
Vladislav Sytchenko fd7a8f0367 (SWDEV-228794)
Add hipMallocHost()

Change-Id: Ia3c7c5ca94b39fe30f3a51d1b60782d3472259ff
2020-03-27 15:57:48 -04:00
Vladislav Sytchenko a91b82f00e (SWDEV-228782)
The only requirment from hipMallocPitch() is that the returned pitch is aligned to the HW image pitch alignment. There is no restriction on the size of the allocation, since the memory might not be used for images.

Change-Id: I97438e5fe4012ca4721b14b85f514dbac803c17c
2020-03-27 15:52:17 -04:00
Saleel Kudchadker 178fc17846 Check event status before notify
Change-Id: I68f6bbbf236e49b859be2d5afbe0c8282fe15dd3
2020-03-27 11:32:46 -07:00
Saleel Kudchadker a5fb06f4db OpenCL2.2 Header changes
Change-Id: I3285a8ffc7e34333feb68e925ebd460831c8e3d4
2020-03-26 17:35:09 -07:00
Vladislav Sytchenko 2a2b9e47a0 Add support for formating hipExtent objects
Change-Id: Iea54a510e81a856c0c450305b3e5a7179ee48295
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko 8ddeeb4551 Enable initial sRGB support
Instead of using the sampler field force_degamma to perform sRGB->linear conversion during pixel sampling, we use an appropriate image format instead. The overhead of this is having to create an image view when creating a texture object from an array.

Change-Id: I1ca368c312c1fd4b6f784a3a1b35b5eeb28070ff
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko e0187ba405 Add initial entry points for mipmapped array API
Change-Id: Icd59cc7323ddcb6773da6105260415a1e6f4cdcb
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko faf3b83594 Handle offsets for dptr <-> image copies
Change-Id: I7a4a56ee07a26a741d2aac35502446d248f720ad
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko c1475f948e Replace hip::TextureObject with __hip_texture
This avoids the use of extra casts when obtaining a texture object handle.

Change-Id: I42df22bdad0ab9ac6c33cb8b282dee65fe7cfd6e
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko fbe30090a2 Correctly format hipResourceDesc objects
The struct consists of a union - only the active object should be read.

Change-Id: I1c40965b61518acd91a2dcbae92a015ac9be346a
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko 2028b6eb29 Headers need to export C symbols for texture API
This also adds declarations of all the missing texture APIs.

hipTexRefSet*() functions need to take a textureReference as a ptr for type erasure to work. Runtime has been modified to accomodate this.

This change only applies to VDI.

Change-Id: Icf43cc5bd44dfc2c39084b7fe56d5a793bf7319f
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko 39558dc1b9 Modify formatting for textureReferences
We don't program the numChannels and format members (these are HCC specific), so printing these will only display garbage.

Change-Id: I83dc8be9a3cae2659c64f4594d07c05330d2dd14
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko ea701777d8 Allow creating texture from unaligned user ptr
All we have to do is align the ptr to HW requirments an if it's not zero, then return the offset to the user.

We currently don't have anywhere to store this offset, so hipGetTextureAlignmentOffset() will still always return 0.

Change-Id: If31998127d99a2a3222a026d88249519d6102505
2020-03-26 12:43:17 -04:00
Payam 76703acb1d updated cmake to create libamdhip64 static file as well
Change-Id: I2054b9501cefa232abbf398524ab62450ab6805d
2020-03-25 16:37:57 -04:00
Saleel Kudchadker 91069c0de4 Merge "Sync streams when freeing or destroying mem" into amd-master-next 2020-03-21 13:33:41 -04:00
Christophe Paquot 653277bd3f hipStreamAddCallback test seg faults
Change-Id: If419d2fad490d0ed50eb1315af809fc1deda1ce3
SWDEV-227875: Add a lock in streams to lock when the callback is call so we make sure things aren't moving forward in the stream
2020-03-20 13:07:34 -07:00
Saleel Kudchadker 8b39e0b74e Sync streams when freeing or destroying mem
Change-Id: I6932f225a8b932bb2adbd5e37880f7e604496809
2020-03-20 10:53:23 -07:00
Christophe Paquot d98879a250 Merge "hipStreamAddCallback test seg faults" into amd-master-next 2020-03-20 12:05:37 -04:00