Commit Graph

5436 Commits

Author SHA1 Message Date
Rahul Garg 017a27214c Fix 2D and 3D memset (#1987) 2020-04-06 15:35:59 +05:30
ansurya ff7aaedca2 Initial support for bfloat16 (#1980) 2020-04-06 15:35:43 +05:30
Yaxun (Sam) Liu 6f439f0232 Fix ambiguity of fma for _Float16 for libc++ (#1976)
libc++ defines fma as template function for auto promotion of mixed-type
arguments. libc++ does not handle _Float16 as _Float16 is not a supported
type by C++ standard. As such, it is unlikely we can commit our fix for
_Float16 to libc++ trunk.

Therefore we handle _Float16 with a template specialization of
__numeric_type in HIP headers.

Change-Id: If01960a657ebf1a7a67463cdcf66fab7458dff3c
2020-04-06 15:35:18 +05:30
Sameer Sahasrabuddhe 01d4117789 SWDEV-227201: Introduce tests for printf on hostcall
Tests that check POSIX specifiers with a single thread:
 - hipPrintfSpecifiers.cpp     : all conversion specifiers
 - hipPrintfFlags.cpp          : common flags that modify conversions
 - hipPrintfAltForms.cpp       : alternate forms ('#')
 - hipPrintfStar.cpp           : additional arguments ('*')
 - hipPrintfWidthPrecision.cpp : floating point details

Tests that check functionality on top of hostcall
 - hipPrintfBasic.cpp       : divergent calls, series of calls, return value, etc
 - hipPrintfManyWaves.cpp   : many waves printing together
 - hipPrintfManyDevices.cpp : many waves on many devices

Change-Id: I35e069f4c542f896999239996dc89eda0faad7b8
2020-04-06 00:49:34 -04:00
eshcherb e524d1d54a adding HIP_PROF version macro (#2000) 2020-04-04 13:10:28 -07:00
Christophe Paquot b820c66c55 Default HostMalloc to uncached memory
Change-Id: I72e19c7f7820a77fd5afc09f09cfea9acd0b8e84
2020-04-03 19:19:33 -04:00
Saleel Kudchadker 74dc32537e Merge "Wake up commandQueue before returning" into amd-master-next 2020-04-03 18:27:03 -04:00
Michael LIAO 6fe3edc5a8 [vdi] Add hipFreeHost
Change-Id: I8a5b7ff3f0ab4f5674efd6723c18808ad6ef33f5
2020-04-03 16:34:28 -04:00
German Andryeyev 9e4aeb0f67 Merge "SWDEV-184709 - support hipLaunchCooperativeKernel()" into amd-master-next 2020-04-03 16:23:05 -04:00
Vladislav Sytchenko c9084d0ad2 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 5efb3f26c0 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
Michael Hong Bin Liao 16ac35c4d5 Merge "Fix size type in __hipRegisterVar" into amd-master-next 2020-04-03 15:12:21 -04:00
Saleel Kudchadker e2def55164 Wake up commandQueue before returning
Change-Id: I87eb5a22c81a9cb807474a960b5987d5fb6c2b86
2020-04-03 10:23:36 -07:00
Aaron En Ye Shi 0bb217245c Merge "Fix path for hip-clang when using hipcc (#1961)" into amd-master-next 2020-04-03 12:28:50 -04:00
Michael LIAO d904f30f9e Fix size type in __hipRegisterVar
Change-Id: I6b667600ae8f133583b768ab963318882b84179f
2020-04-03 10:51:58 -04:00
Michael Hong Bin Liao 66aca97d6f Merge "[hip] Clean up unnecessary casting." into amd-master-next 2020-04-03 10:50:40 -04:00
German Andryeyev 2e948e4034 SWDEV-184709 - support hipLaunchCooperativeKernel()
- Enable cooperative tests for single and multiple devices

Change-Id: I54b6713f578b6b5e670f117b17469c0091028c99
2020-04-02 12:55:05 -04:00
Michael LIAO e3795436b2 [hip] Clean up unnecessary casting.
Change-Id: I64b08aaef5c67ffb49330c9c605611f1fbd3f5a2
2020-04-02 12:46:15 -04:00
Paul Fultz II 58f04fb774 Fix path for hip-clang when using hipcc (#1961)
* Fix path for hip-clang when using hipcc

* Fix typo

* Update regex

Change-Id: I31bbee2e70d58b89191f970f5c6ae7e1c8b40900
2020-04-02 12:09:31 -04:00
Paul Fultz II 631f82ab82 Add missing flags for hip::device target on hip-clang (#1230)
This adds the missing compilation flags to hip::device so it can compile with hip-clang compiler.

Change-Id: Ie2b30ea606bfca385a0e84ae03ee0a8d828ad16a
2020-04-02 12:09:03 -04:00
Saleel Kudchadker 5a357a795a Merge "OpenCL2.2 Header changes" into amd-master-next 2020-04-02 02:46:49 -04:00
Vladislav Sytchenko a09fadecf2 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
eshcherb 1e5ea70588 adding hipApiString (#1989) 2020-04-01 12:10:33 +05:30
ansurya 68f1b814b1 tex1Dfetch behaviour for different address mode and filter mode (#1772)
Fixes github issue: #1754

- When ResourceDesc::resType is hipResourceTypeLinear ignore address mode and filter mode.
- When textureDesc::normalizedCoords is set to zero, AddressModeWrap and AddressModeMirror won't be supported and will be switched to AddressModeClamp.
2020-04-01 12:10:17 +05:30
Saleel Kudchadker 99a024fc14 Merge "Cleanup stream from hip:Event class." into amd-master-next 2020-03-31 20:14:48 -04:00
Vladislav Sytchenko 6e0722a5d0 (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
Reshabh Sharma 36f24a40e5 Output file name should not change flags picked for compiler (#1938)
Fixes SWDEV-207362,

The output file name should not contribute to picking up the right flags for the compiler. This fix solves issues when the output has conflicting extensions which confuses hipcc to treat them as the source files and add the required flags for them.

PS: Output file refers to the file followed by -o

Change-Id: I1095966c11143ad73e81fabc35b4e9de5d3afada
Example: hipcc test.o -o test.hip will add the flags for .hip compilation ignoring the fact that it is an output file
2020-03-31 16:13:43 -04:00
Saleel Kudchadker 65722f8cca Cleanup stream from hip:Event class.
Change-Id: I98de07d33bb7fea8f5e2d32b288c15f10ce58902
2020-03-31 11:22:00 -07:00
Michael Hong Bin Liao 9420f56c66 Merge "[hipcc] Remove the previous workaround." into amd-master-next 2020-03-31 10:05:56 -04:00
Michael LIAO a14695d4eb [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 47718cbf16 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 7356a74d35 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 9d213070b0 Merge "Update Enable/Disable peers to match cuda behaviour" into amd-master-next 2020-03-30 13:46:42 -04:00
agodavar 96d119ac77 Update Enable/Disable peers to match cuda behaviour
Change-Id: I67194ccf77a0019368579ff7d95b7790fcf228f3
2020-03-30 12:49:16 -04:00
Tao Sang 44a145b2a3 Merge "enable HCC printf when using hip-clang" into amd-master-next 2020-03-30 11:37:23 -04:00
Michael LIAO 64bbce8e73 [hipcc] Remove the previous workaround.
Change-Id: Ie1c2aa21689bc343c10fd4b699546cf500ced5b2
2020-03-30 10:18:47 -04:00
Michael LIAO 55d869df99 [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
Sarbojit2019 b55775a49e Fix for segfault seen in hipMemcpyDtoD (#1964)
* Fixes SWDEV-227444.
2020-03-28 17:29:49 +05:30
Maneesh Gupta 062938a14c Remove address_space(1) typecast and use __ockl_atomic_add_noret_f32 (#1956)
* Remove address_space(1) typecast for ockl_global_atomic_add_f32
* use __ockl_atomic_add_noret_f32
2020-03-28 17:28:33 +05:30
Sameer Sahasrabuddhe c014c733c6 enable HCC printf when using hip-clang
This is cherry-picked from PR#1947 that was committed to the
github repo. It allows printf to work with hip-clang and HCC
runtime.

Change-Id: I754753250ea1e694cf3441722e2d4c9d25fa75bc
2020-03-28 00:18:21 -04:00
Saleel Kudchadker a5c287554e Merge "Check event status before notify" into amd-master-next 2020-03-27 20:19:40 -04:00
Vladislav Sytchenko 86c969fcea (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 08729d0f43 (SWDEV-228794)
Add hipMallocHost()

Change-Id: Ia3c7c5ca94b39fe30f3a51d1b60782d3472259ff
2020-03-27 15:57:48 -04:00
Vladislav Sytchenko 111897dbea (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
Tao Sang b2afe3c250 Merge "Support hipFuncGetAttributes with hip-clang+Hcc RT" into amd-master-next 2020-03-27 14:36:41 -04:00
Saleel Kudchadker fa7dea42b5 Check event status before notify
Change-Id: I68f6bbbf236e49b859be2d5afbe0c8282fe15dd3
2020-03-27 11:32:46 -07:00
satyanveshd 4554bbd9f9 [dtests] Added few Negative tests (#1735) 2020-03-27 14:10:12 +05:30
Siu Chi Chan 6ab1e864b6 don't expose symbols from code_object_bundle (#1971)
Change-Id: I56479485aad42c3d517fe6d9055be1cd846eeb00
2020-03-27 14:09:07 +05:30
Sarbojit2019 4a68ab5a8c Fix few memory leaks in HIP (#1969) 2020-03-27 14:08:30 +05:30
Jatin Chaudhary 2449e0de8b Packaging changes to bring them to sync with generated ones (#1918) 2020-03-27 14:08:12 +05:30