Commit Graph

1570 Commits

Author SHA1 Message Date
Michael LIAO 5e5f4e23d5 [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


[ROCm/hip commit: b72196613a]
2020-03-31 00:26:53 -04:00
Maneesh Gupta 0ee021d444 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

[ROCm/hip commit: cbc3d1713f]
2020-03-28 17:28:33 +05:30
Sameer Sahasrabuddhe 3e7d08f5ee 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


[ROCm/hip commit: 9a0c5d0653]
2020-03-28 00:18:21 -04:00
Siu Chi Chan e58a0d06f7 don't expose symbols from code_object_bundle (#1971)
Change-Id: I56479485aad42c3d517fe6d9055be1cd846eeb00

[ROCm/hip commit: 43abf84f54]
2020-03-27 14:09:07 +05:30
Vladislav Sytchenko c1b7b2276e Add initial entry points for mipmapped array API
Change-Id: Icd59cc7323ddcb6773da6105260415a1e6f4cdcb


[ROCm/hip commit: e0187ba405]
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko 1ff602f312 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


[ROCm/hip commit: 2028b6eb29]
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko ab94b954f9 Set textureObject to nullptr
This avoids dangling pointers for newly initiazlied textures

Change-Id: Ia444b91fe17fd756ed583ec595ae1febbdfbd034


[ROCm/hip commit: ced0582a52]
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko 462dcfe244 Correct typos in texture function declarations
Change-Id: I492995e984eda2e8a5e806c5d4c9c78da09ac483


[ROCm/hip commit: b09fe1280e]
2020-03-26 12:43:17 -04:00
Sarbojit2019 59cf17fb32 Fix for __usad issue (#1972)
Fixes #1930

[ROCm/hip commit: 5024f9057a]
2020-03-26 17:09:44 +05:30
Benjamin Sherman d9d1017168 Add const qualifiers to HIP_vector_type unary arithmetic operators (#1965)
Resolves issue #1960

[ROCm/hip commit: 3d38135ae2]
2020-03-26 17:09:00 +05:30
Joseph Greathouse cb69c6037c Fix cooperative launch APIs to set hipGetLastError (#1935)
* Fix cooperative launch APIs to set hipGetLastError

Previously, the cooperative launch APIs did not properly log their
errors in the global hipGetLastError variable before returning back
to the user. As such, the APIs would leave hipSuccess in the
last error, which would break some use cases.

This fixes that problem by making a trampoline function that does
the HIP_INIT_API and ihipLogStatus.

* Add missing flag to the log of multi-GPU launch

[ROCm/hip commit: f61b79d9a3]
2020-03-25 14:39:24 -07:00
Nick Curtis 15ea4194c4 Update hip_runtime_api.h (#1966)
Correct URL for deprecated api list

[ROCm/hip commit: b4c69a2e4a]
2020-03-23 10:16:24 -07:00
Vladislav Sytchenko 2847cdc6f0 Add support for creating typed buffers
What Cuda refers to "linear texture memory" is the OpenCL equivalent of CL_MEM_OBJECT_IMAGE1D_BUFFER. For these types of allocations we should create a typed buffer instead of an image.

Currently there is no check in the texture fetch functions as to what kind of SRD is written into the texture object, so any kind of incorrect programming will cause the TA to hang. Fortunately for us, every one writes correct code :)

Change-Id: I80dab85a992f2c0754ebf303d40ac6b5e045c7c1


[ROCm/hip commit: 4829a7c215]
2020-03-18 18:15:17 -04:00
Vladislav Sytchenko 357404e25f Rework the texture C++ API
Currently the texture C++ API is forwarded to the ihip*Impl() calls, which are not even a part of Cuda. These should be forwarded to their respective Cuda C APIs instead.

This change also fixes a bug with hipUnbindTexture() creating a dangling pointer.

Change-Id: Ifafc9d106855a11bec84a18ea214b3d89e39990d


[ROCm/hip commit: 5429b40afe]
2020-03-18 18:14:53 -04:00
Vladislav Sytchenko b61a7da1ec Correct the declaration of hipBindTexture2D()
The texture reference needs to be passed as a constant pointer.

Change-Id: Idde461f0f328ac87ce677b6bab3203161b514cbf


[ROCm/hip commit: 3e460ab514]
2020-03-18 18:08:23 -04:00
Vladislav Sytchenko 4723ceced8 Correct the declaration of hipBindTextureToArray()
The texture reference needs to be passed as a constant pointer.

Change-Id: Iff171626536071fb2020cfff7132ec930577b1b9


[ROCm/hip commit: 2d77399747]
2020-03-18 18:08:13 -04:00
Vladislav Sytchenko efca40a693 Correct the declaration of hipBindTexture()
The texture reference needs to be passed as a constant pointer.

Change-Id: I36ca0bddaba30becfc2ce70dd9e5b7db66c57f27


[ROCm/hip commit: 7190fa518e]
2020-03-18 18:08:01 -04:00
Vladislav Sytchenko 42064e3222 Add missing mipmap API entries
Introduce hipFreeMipmappedArray(), hipMallocMipmappedArray() and hipGetMipmappedArrayLevel() APIs.

Change-Id: I878228c79fa1c54536c17d6baf45f83d51d2b1c7


[ROCm/hip commit: 551bcc6293]
2020-03-18 18:07:45 -04:00
Vladislav Sytchenko 07f864f128 Don't hardcode the texture read mode
The readmode needs to be inferred from the template arguments.

Change-Id: I067037035e2492a24eac47e16d4015f879be0ea7


[ROCm/hip commit: 99e744ab4a]
2020-03-18 18:07:33 -04:00
Vladislav Sytchenko 34bf0bd816 Add constraints to texture indirect functions
Similar to the previous patch, this change adds type constraints to texture indirect functions. Since we don't have to deduce the return type for these, we simply just have to check if the user provided a valid channel type.

Change-Id: Ia094bd6126e01df2ea90902c9aa59cb6cfe85773


[ROCm/hip commit: 117f0ab102]
2020-03-18 12:24:40 -04:00
Vladislav Sytchenko e120f9164e Add constraints to texture fetch functions
When sampling a pixel the hw always returns a float4. The type in the texture reference controls the bitcast that we perform before returning the sampled pixel. Creating a texture with an unsupported will lead to potential UB.

This change makes it so that it's only possible to use textures with a type that makes sense. Using something like texture<int, hipTextureType1D, hipReadModeNormalizedFloat> will now lead to a compilation error with a message "Invalid channel type!".

Change-Id: I7fde44cb1d4b9737e0c48c28cb59c018c59ccaa2


[ROCm/hip commit: ef2415edc7]
2020-03-18 12:24:40 -04:00
Yaxun (Sam) Liu e05e7d063e Workaround for libc++ include path for HIP-Clang (#1917)
HIP-Clang cuda_wrapper headers require clang include path before standard C++ include path.
However libc++ include path requires to be before clang include path.
To workaround this, we pass -isystem with the parent directory of clang include
path instead of the clang include path itself.

[ROCm/hip commit: 08d9759eba]
2020-03-18 11:20:21 +05:30
Sarbojit Sarkar 4135ec890a [hip-vdi]Fix for TF build failure [SWDEV-225827]
Change-Id: I8478779bef92bad8353b8d066b28c220bb59b98d


[ROCm/hip commit: 82926666c4]
2020-03-17 22:52:01 -04:00
Vladislav Sytchenko 31bd87b9eb Rework device texture headers
This change addresses three things.

First the available APIs are brought up to par with Cuda (missing ones are added and incorrect ones removed).

Second the size of hip/hcc_detail/texture_functions.h. Using some template magic we can bring down the code size down from ~11k lines to only ~900 lines in total.

Third this change fixes some bugs in the declaration of the texture fetch funcitons. Currently the return type for textures with read mode set to hipReadModeNormalizedFloat is not float. This causes pixel data to be lost during the bitcast when the texture pixel element size is less than the size of float.

The new headers will only be enabled for VDI to avoid breaking HCC.

Change-Id: I77cb29293fb79e55681be094c37702a48d80b64c


[ROCm/hip commit: a0751402d8]
2020-03-17 17:04:37 -04:00
Jatin Chaudhary 1bc892c6f2 Adding Half Abs APIs (#1902)
[ROCm/hip commit: 16a6a94fbf]
2020-03-17 14:13:19 +05:30
Sameer Sahasrabuddhe 7ab3583c60 enable HCC printf when using hip-clang (#1947)
This allows printf to work with hip-clang and HCC runtime. See comments under #1919 for a reported bug and feature request.

[ROCm/hip commit: 899c878703]
2020-03-17 14:03:27 +05:30
Joseph Greathouse 504ba0a4c9 Fix compiler warning on NVCC path (#1942)
GCC emits a warning about using static functions like
hipCUDAErrorTohipError inside this function, because it has an
inline directive, but it's not static. Adding static to this function
to silence warnings (and prevent potential problems in the future).

[ROCm/hip commit: f7e85649f4]
2020-03-17 14:02:59 +05:30
Joseph Greathouse 122c2f9034 Fix occupancy calculations API on NVCC (#1941)
NVCC warned if you tried to use hipOccupancyMaxActiveBlocksPerMultiprocessor
because when passing in a device function pointer, "const void* func" was
insufficient to describe it accurately. Adding a C++ templated class type
definition for this function.

[ROCm/hip commit: 4128d68ed7]
2020-03-17 14:02:48 +05:30
Sarbojit2019 0310e4d7f8 Fix __sad signature match with Cuda (#1936)
Fix for issue #1930

[ROCm/hip commit: 320742e8a0]
2020-03-17 14:02:00 +05:30
Aryan Salmanpour 03654845a4 [HIP] add cooperative kernel launch APIs on NVCC (#1929)
[ROCm/hip commit: 015895a265]
2020-03-17 14:01:11 +05:30
Maneesh Gupta 8feab1161e Annotate __constant__ (#1901)
[ROCm/hip commit: eee5cc8621]
2020-03-17 13:59:44 +05:30
mhbliao e0da34b5b1 [hip] Improve the portability of the header for vector type support. (#1873)
- Need to check the availability of `__has_attribute` builtin macro
  instead of compiler versions. That's more reliable and portable among
  various compilers.
- Provides a very basic support of vectors for unknown compilers.

[ROCm/hip commit: 774035d869]
2020-03-17 13:59:24 +05:30
Sameer Sahasrabuddhe be216b9ab6 SWDEV-204784: separate printf declaration for vdi/clang
There are now two implementations of printf in HIP:

1. The implemenation for HCC is controlled by the HC_FEATURE_PRINTF
   macro, and it works only with the HCC compiler used in combination
   with the HCC runtime.

2. The implementation for hip-clang requires the VDI runtime, and is
   always enabled with that combination.

Change-Id: Ibaeda7900ffe2ce602ca0094aafed0f1147ac2b6


[ROCm/hip commit: 64cd527335]
2020-03-16 04:00:39 -04:00
Evgeny Mankov 6039400c04 Merge pull request #1908 from asalmanp/prop_mulit_coop
[HIP] add hip specific properties for cooperative kernel multi device

[ROCm/hip commit: 70f5646f8a]
2020-03-12 19:12:11 +03:00
Alex Voicu 8cecdaa704 Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP into feature_robust_constant
[ROCm/hip commit: 1c5f526e6b]
2020-03-12 14:20:26 +00:00
Maneesh Gupta d44d5f8cdd Expose support for non-returning atomic FADD (#1909)
Change-Id: If5359488324477315a9bd4f308a75f606c065b39

[ROCm/hip commit: 0726abf424]
2020-03-11 14:33:15 +05:30
Vladislav Sytchenko c5a842de46 Fix typo in device __shfl_xor function
Change-Id: I8bcdd53ced00c596a0af013a0c34e37aa67c93ae


[ROCm/hip commit: 4ca9cda372]
2020-03-10 13:23:08 -04:00
Nick Curtis 98b7cb62aa Fix incorrect shfl_xor for Windows
copy/paste error, need __shfl_xor w/ lane_mask

[ROCm/hip commit: 09edc7e49c]
2020-03-10 12:04:05 -05:00
Vladislav Sytchenko 88ae7337b3 Add hipDrvMemcpy3D.
This is the equivalent of cuMemcpy3D.

Change-Id: Ib2e06dbd6f5093c931cdfd36c87617f32acffc2d


[ROCm/hip commit: ecd7c99b49]
2020-03-09 16:11:25 -04:00
Sameer Sahasrabuddhe 07abb2d633 separate printf declaration for vdi/clang
There are now two implementations of printf in HIP:

1. The implemenation for HCC is controlled by the HC_FEATURE_PRINTF
   macro, and it works only with the HCC compiler used in combination
   with the HCC runtime.

2. The implementation for hip-clang requires the VDI runtime, and is
   always enabled with that combination.


[ROCm/hip commit: 09130b3b92]
2020-03-09 09:40:05 +05:30
Lad, Aditya 044a1924ba Merge branch 'master' into amd-master-next
Conflicts:
	CMakeLists.txt
	tests/src/texture/simpleTexture2DLayered.cpp
	tests/src/texture/simpleTexture3D.cpp

Change-Id: I4aa4754d391b5f37ddf15fa0bcfc84d9da020119


[ROCm/hip commit: d80edf9541]
2020-03-06 14:10:44 -05:00
Aryan Salmanpour feb28352cb move new enums to the end to maintain compatibility
[ROCm/hip commit: 7e45c54ea6]
2020-03-06 11:38:44 -05:00
Maneesh Gupta 08ed6ab780 Expose support for non-returning atomic FADD
Change-Id: If5359488324477315a9bd4f308a75f606c065b39


[ROCm/hip commit: 4a40010ac6]
2020-03-05 10:30:52 +05:30
agodavar 76ee85ff82 Fix hipExtLaunchMultiKernelMultiDevice compilation issue
Fix compilation error on hip-hcc+clang , hip-vdi+clang
Enabled hipExtLaunchMultiKernelMultiDevice test on hip-vdi path
hipExtLaunchMultiKernelMultiDevice common declaration for all paths

Change-Id: I76031840614fce8e12a8e845548fa43a389a741a


[ROCm/hip commit: 6a5d04209c]
2020-03-04 15:38:14 -05:00
Aryan Salmanpour df81136734 [HIP] add hip specific properties for cooperative kernel multi device
[ROCm/hip commit: 03797ae986]
2020-03-03 13:25:36 -05:00
Alex Voicu abbb594f7c Annotate __constant__
[ROCm/hip commit: 27480ff5a2]
2020-02-28 22:54:00 +02:00
saleelk 3c66b171e1 Fix HIPRTC headers to export C style symbols (#1879)
[ROCm/hip commit: 3e1f41c165]
2020-02-28 16:47:29 +05:30
Rahul Garg c34c9a4b4d Remove deprecated HIP markers (#1876)
[ROCm/hip commit: 6c5fa32815]
2020-02-28 16:47:15 +05:30
Rahul Garg 5229ffff99 Add hipDrvOccupancyMaxActiveBlocksPerMultiprocessor[WithFlags] (#1854)
Equivalent to cuOccupancyMaxActiveBlocksPerMultiprocessor[WithFlags].

[ROCm/hip commit: edc97f3073]
2020-02-28 16:46:55 +05:30
Nick Curtis 2715d1b036 fix long shuffle implementations for windows (#1895)
Fixes for SWDEV-223694

[ROCm/hip commit: b7dd073d93]
2020-02-26 15:53:56 +05:30