Commit Graph

1194 Commits

Author SHA1 Message Date
Vladislav Sytchenko cfc508eef0 Add initial entry points for mipmapped array API
Change-Id: Icd59cc7323ddcb6773da6105260415a1e6f4cdcb


[ROCm/clr commit: 06bb183558]
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko df2f07ce18 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/clr commit: 428b56e411]
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko b7f6419f6a Set textureObject to nullptr
This avoids dangling pointers for newly initiazlied textures

Change-Id: Ia444b91fe17fd756ed583ec595ae1febbdfbd034


[ROCm/clr commit: 3ab8ff87cc]
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko c0a3cb04fe Correct typos in texture function declarations
Change-Id: I492995e984eda2e8a5e806c5d4c9c78da09ac483


[ROCm/clr commit: dc015d27c1]
2020-03-26 12:43:17 -04:00
Vladislav Sytchenko fd286a205b 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/clr commit: 78f7954765]
2020-03-18 18:15:17 -04:00
Vladislav Sytchenko 42f9702627 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/clr commit: 53dd6b7a66]
2020-03-18 18:14:53 -04:00
Vladislav Sytchenko 39c78589e2 Correct the declaration of hipBindTexture2D()
The texture reference needs to be passed as a constant pointer.

Change-Id: Idde461f0f328ac87ce677b6bab3203161b514cbf


[ROCm/clr commit: 7a81e55de7]
2020-03-18 18:08:23 -04:00
Vladislav Sytchenko 77867d5700 Correct the declaration of hipBindTextureToArray()
The texture reference needs to be passed as a constant pointer.

Change-Id: Iff171626536071fb2020cfff7132ec930577b1b9


[ROCm/clr commit: 7765792a42]
2020-03-18 18:08:13 -04:00
Vladislav Sytchenko 900f2a76c0 Correct the declaration of hipBindTexture()
The texture reference needs to be passed as a constant pointer.

Change-Id: I36ca0bddaba30becfc2ce70dd9e5b7db66c57f27


[ROCm/clr commit: 600b1006c8]
2020-03-18 18:08:01 -04:00
Vladislav Sytchenko 415070fbfc Add missing mipmap API entries
Introduce hipFreeMipmappedArray(), hipMallocMipmappedArray() and hipGetMipmappedArrayLevel() APIs.

Change-Id: I878228c79fa1c54536c17d6baf45f83d51d2b1c7


[ROCm/clr commit: 6ff0d15a62]
2020-03-18 18:07:45 -04:00
Vladislav Sytchenko 762d6f3212 Don't hardcode the texture read mode
The readmode needs to be inferred from the template arguments.

Change-Id: I067037035e2492a24eac47e16d4015f879be0ea7


[ROCm/clr commit: b4b05e2ab0]
2020-03-18 18:07:33 -04:00
Vladislav Sytchenko ebd1dec272 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/clr commit: 8667e21b25]
2020-03-18 12:24:40 -04:00
Vladislav Sytchenko 8b211978a7 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/clr commit: a994da8af6]
2020-03-18 12:24:40 -04:00
Sarbojit Sarkar 7cac7ff3b5 [hip-vdi]Fix for TF build failure [SWDEV-225827]
Change-Id: I8478779bef92bad8353b8d066b28c220bb59b98d


[ROCm/clr commit: 957caff8b4]
2020-03-17 22:52:01 -04:00
Vladislav Sytchenko 8d204f649e 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/clr commit: 3d7945faae]
2020-03-17 17:04:37 -04:00
Sameer Sahasrabuddhe f12c0f1dbb 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/clr commit: d48738856c]
2020-03-16 04:00:39 -04:00
Vladislav Sytchenko 98761c5721 Fix typo in device __shfl_xor function
Change-Id: I8bcdd53ced00c596a0af013a0c34e37aa67c93ae


[ROCm/clr commit: caea3f9b32]
2020-03-10 13:23:08 -04:00
Vladislav Sytchenko 952502e2ed Add hipDrvMemcpy3D.
This is the equivalent of cuMemcpy3D.

Change-Id: Ib2e06dbd6f5093c931cdfd36c87617f32acffc2d


[ROCm/clr commit: cd76eb7486]
2020-03-09 16:11:25 -04:00
Lad, Aditya eca6c4f001 Merge branch 'master' into amd-master-next
Conflicts:
	CMakeLists.txt
	tests/src/texture/simpleTexture2DLayered.cpp
	tests/src/texture/simpleTexture3D.cpp

Change-Id: I4aa4754d391b5f37ddf15fa0bcfc84d9da020119


[ROCm/clr commit: e7fdb3d796]
2020-03-06 14:10:44 -05:00
agodavar 1c828b7575 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/clr commit: 5a6c605730]
2020-03-04 15:38:14 -05:00
saleelk 21712e6fb8 Fix HIPRTC headers to export C style symbols (#1879)
[ROCm/clr commit: 40a28e767e]
2020-02-28 16:47:29 +05:30
Rahul Garg 7395bf9ba9 Remove deprecated HIP markers (#1876)
[ROCm/clr commit: 6be7537cf9]
2020-02-28 16:47:15 +05:30
Rahul Garg ff421fb2d2 Add hipDrvOccupancyMaxActiveBlocksPerMultiprocessor[WithFlags] (#1854)
Equivalent to cuOccupancyMaxActiveBlocksPerMultiprocessor[WithFlags].

[ROCm/clr commit: 1c794045e0]
2020-02-28 16:46:55 +05:30
Nick Curtis e5d029b8a4 fix long shuffle implementations for windows (#1895)
Fixes for SWDEV-223694

[ROCm/clr commit: a094c87038]
2020-02-26 15:53:56 +05:30
Rahul Garg 6f1d14c96d Fix hipMemcpy3D (#1798)
Fixes #1790 and #1791. hipMemcpy3D still requires further refactoring for different input and output combinations.

[ROCm/clr commit: ec84c16d75]
2020-02-17 19:35:35 +05:30
Tao Sang 0292fd9991 Temporarily comment out Hcc-specific APIs for CLang compiler
Temporarily comment out Hcc-specific template functions
hipExtLaunchKernelGGL and hipOccupancyMaxPotentialBlockSize for CLang
compiler so that all test cases under hip/samples can be built
successfully for Clang + Hip/Hcc runtime.

Change-Id: Iafc761257be4a7b34eafa6759a01f369570cd6ce


[ROCm/clr commit: 15b91e9f1e]
2020-02-16 22:26:47 -05:00
Nick Curtis e9ee5a3719 Implement long / long long shuffles (#1829)
Implement additional data-types for shuffles (long and long long).
Based upon the double implementation.


[ROCm/clr commit: 5b0f34c59b]
2020-02-15 09:51:09 +05:30
ansurya 708af7e6ac Reduce GPU copying based on arch it runs on (#1751)
Implements SWDEV-213230.

[ROCm/clr commit: 888a7f2a90]
2020-02-13 14:21:51 +05:30
Aryan Salmanpour 89cb17284b fix build error in nvcc path
[ROCm/clr commit: 1c013bab3e]
2020-02-11 12:16:51 -05:00
Aryan Salmanpour e7059e518d Fix a typo causing a build error
[ROCm/clr commit: 504412c1e4]
2020-02-10 11:44:40 -05:00
Aryan Salmanpour 44ce519211 resolve merge conflict
[ROCm/clr commit: dfa565d76f]
2020-02-10 10:30:55 -05:00
Maneesh Gupta a21b1f84cb Revert "Match Occupancy APIs syntax with CUDA (#1625)" (#1857)
Reverting this for now till we figure out how to avoid the build
breakage.

This reverts commit 2c6e880e05.

[ROCm/clr commit: d032637934]
2020-02-10 10:45:28 +05:30
Alex Voicu 7f9f87b0d9 (Maybe) Match alignment between Clang and GCC. (#1789)
Should fix #1740 and the related internal bug.

[ROCm/clr commit: 8f96152d3a]
2020-02-10 10:44:49 +05:30
vsytch e2b6794a07 Device texture functions should not normalize the sampled pixel (#1826)
* Device texture functions should not normalize the sampled pixel. This is already done by HW.
* Add support to use h/w capability for normalized float data convertion for driver API's

Co-authored-by: ansurya <50609411+ansurya@users.noreply.github.com>


[ROCm/clr commit: fa363ae721]
2020-02-05 20:56:17 +05:30
Aryan Salmanpour cbefd62db9 code clean up
[ROCm/clr commit: 75572a4175]
2020-01-31 13:08:25 -05:00
Aryan Salmanpour 2733df36b5 [HIP][HIPIFY] Add some missing flags for cooperative launch and occupancy APIs
[ROCm/clr commit: 6db9b782be]
2020-01-30 15:05:53 -05:00
satyanveshd 2c6e880e05 Match Occupancy APIs syntax with CUDA (#1625)
* Match Occupancy APIs syntax with CUDA and fix tests using these APIs


[ROCm/clr commit: 262ad13dd5]
2020-01-29 13:05:53 -08:00
vsytch 4162b55dbe Add missing texturePitchAlignment member to the hipDeviceProp_t struct. (#1802)
* Add missing texturePitchAlignment member to the hipDeviceProp_t struct.

* Add missing hipDeviceAttributeTexturePitchAlignment enumerator to the hipDeviceAttribute_t enum.

* Initialize texturePitchAlignment to 256. This works for gfx9+, but is technically overaligned in most cases for pre-gfx9.

* Add the texturePitchAlignment property to the NVCC path.


[ROCm/clr commit: 6da0c82b78]
2020-01-27 16:37:00 -08:00
vsytch f1747af7cd Update the HIP_TRSF_* flags to match their Cuda equivalents. (#1801)
[ROCm/clr commit: e419a1609f]
2020-01-24 11:41:15 -08:00
mshivama a49cbb1580 SWDEV-220503: this_grid().thread_rank() gives incorrect result (#1808)
* fix a minor bug while computing this.grid()::thread_rank()

[ROCm/clr commit: 0605426049]
2020-01-24 16:23:28 +05:30
kpyzhov 2d617bba3b Don't use accelerated vector element access for hip-clang. (#1796)
[ROCm/clr commit: 1d6a2fb064]
2020-01-15 18:17:08 -08:00
kpyzhov 8bf4210095 Add missing constructors for Scalar_accessor class. (#1792)
[ROCm/clr commit: 0918614615]
2020-01-14 11:30:21 -08:00
Evgeny Mankov d1d680a1d6 [HIP] Unify hipError_t (Step 3): Sync nvcc path (#1778)
* [HIP] Unify hipError_t (Step 3): Sync nvcc path

* [HIP][fix] Add CUDA 10.x support to nvcc path


[ROCm/clr commit: 9c2fcff9ce]
2020-01-10 13:47:18 +05:30
Maneesh Gupta d814f1e292 Revert PRs that break ROCm builds (#1781)
Fixes SWDEV-218626 and SWDEV-218629

Changes:
- Revert "`static inline` in a header, just like excess sugar in a diet, causes bloat (#1692)"
   This reverts commit 47725baf8c.
- Revert "Fix rocFFT build failure (#1777)"
   This reverts commit f7bac3e057.

[ROCm/clr commit: c091635f5a]
2020-01-08 15:11:58 +05:30
ansurya f7bac3e057 Fix rocFFT build failure (#1777)
Fixes SWDEV-217761

[ROCm/clr commit: cbd8f604ec]
2020-01-07 08:12:37 +05:30
Rahul Garg 1e78733f53 Add hipBindTexture2D on NVCC path (#1773)
[ROCm/clr commit: 534376cb41]
2020-01-06 12:33:50 +05:30
Rahul Garg 32139b7465 Fix hipcc warning related to hipVersion (#1767)
* Fix hipcc warning related to hipVersion
* Rename hipVersion.h to hip_version.h
* Remove HIP_VERSION splitting
* Update .gitignore
- Ignore generated include/hip/hip_version.h
- Removed some stale entries
- Added executables from samples/1_Utils/*/ for consistency with bin/ entries.

[ROCm/clr commit: 3fe1695c3a]
2020-01-06 12:33:23 +05:30
Evgeny Mankov 8c58286a9b Merge pull request #1759 from emankov/master
[HIP] Unify hipError_t (Step 2)

[ROCm/clr commit: ed43699a3a]
2019-12-30 19:21:09 +03:00
Sarbojit2019 4d2ee146e1 Change to generate hipVersion.h (#1726)
HIP_VERSION_MAJOR, HIP_VERSION_MINOR, HIP_VERSION_PATCH and HIP_VERSION pre-processor macros are now defined in hipVersion.h instead of being set by hipcc.

[ROCm/clr commit: 4fd77f8336]
2019-12-30 12:44:24 +05:30
Aryan Salmanpour ed2609b880 [hip] refactoring cooperative kernel launch APIs (#1737)
This PR is a follow-up on PR# #1698 and it makes two more APIs (hipLaunchCooperativeKernel/hipLaunchCooperativeKernelMultiDevice) inline so that they can work correctly with lazy binding.

[ROCm/clr commit: 857052be1e]
2019-12-30 12:42:17 +05:30