Граф коммитов

1194 Коммитов

Автор SHA1 Сообщение Дата
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
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
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
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
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
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
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
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
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
Rahul Garg f1746197c7 Fix hipMemcpy3D (#1798)
Fixes #1790 and #1791. hipMemcpy3D still requires further refactoring for different input and output combinations.

[ROCm/hip commit: 8c5e5e435b]
2020-02-17 19:35:35 +05:30
Tao Sang eb5fed63c1 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/hip commit: b3f445c0f5]
2020-02-16 22:26:47 -05:00
Nick Curtis 6d661ba6ca Implement long / long long shuffles (#1829)
Implement additional data-types for shuffles (long and long long).
Based upon the double implementation.


[ROCm/hip commit: 797a929a65]
2020-02-15 09:51:09 +05:30
ansurya fb53186682 Reduce GPU copying based on arch it runs on (#1751)
Implements SWDEV-213230.

[ROCm/hip commit: 8c6934223b]
2020-02-13 14:21:51 +05:30
Aryan Salmanpour 07c3e4f898 fix build error in nvcc path
[ROCm/hip commit: 959f1b0f0e]
2020-02-11 12:16:51 -05:00
Aryan Salmanpour 5db60dab96 Fix a typo causing a build error
[ROCm/hip commit: 5a29f27455]
2020-02-10 11:44:40 -05:00
Aryan Salmanpour 964a272004 resolve merge conflict
[ROCm/hip commit: 874b201ee2]
2020-02-10 10:30:55 -05:00
Maneesh Gupta 7753b3e827 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 e38db9fb6f.

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

[ROCm/hip commit: dd34ea95d6]
2020-02-10 10:44:49 +05:30
vsytch e5243c1834 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/hip commit: ef514eef71]
2020-02-05 20:56:17 +05:30
Aryan Salmanpour adc9134318 code clean up
[ROCm/hip commit: c8137263d6]
2020-01-31 13:08:25 -05:00
Aryan Salmanpour 78a36b4fb8 [HIP][HIPIFY] Add some missing flags for cooperative launch and occupancy APIs
[ROCm/hip commit: 6e867eacb6]
2020-01-30 15:05:53 -05:00
satyanveshd e38db9fb6f Match Occupancy APIs syntax with CUDA (#1625)
* Match Occupancy APIs syntax with CUDA and fix tests using these APIs


[ROCm/hip commit: fa98798b63]
2020-01-29 13:05:53 -08:00
vsytch 1af35a6044 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/hip commit: f72a669487]
2020-01-27 16:37:00 -08:00
vsytch 95ee059bbb Update the HIP_TRSF_* flags to match their Cuda equivalents. (#1801)
[ROCm/hip commit: 9cfada0f9d]
2020-01-24 11:41:15 -08:00
mshivama c630cb57b2 SWDEV-220503: this_grid().thread_rank() gives incorrect result (#1808)
* fix a minor bug while computing this.grid()::thread_rank()

[ROCm/hip commit: bed8f1c1b8]
2020-01-24 16:23:28 +05:30
kpyzhov dabbde21c6 Don't use accelerated vector element access for hip-clang. (#1796)
[ROCm/hip commit: 566adc4594]
2020-01-15 18:17:08 -08:00
kpyzhov 4c457c2bb5 Add missing constructors for Scalar_accessor class. (#1792)
[ROCm/hip commit: fae85cf6d2]
2020-01-14 11:30:21 -08:00
Evgeny Mankov 3d22f829f5 [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/hip commit: a005a8550d]
2020-01-10 13:47:18 +05:30
Maneesh Gupta 6ebf5ca991 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 b7803eb0c6.
- Revert "Fix rocFFT build failure (#1777)"
   This reverts commit 9c8c3df294.

[ROCm/hip commit: 00bd5d1cec]
2020-01-08 15:11:58 +05:30
ansurya 9c8c3df294 Fix rocFFT build failure (#1777)
Fixes SWDEV-217761

[ROCm/hip commit: 753277422a]
2020-01-07 08:12:37 +05:30
Rahul Garg 327a45b18b Add hipBindTexture2D on NVCC path (#1773)
[ROCm/hip commit: a5d7e7d8d3]
2020-01-06 12:33:50 +05:30
Rahul Garg 86dbb13897 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/hip commit: f3cafd5855]
2020-01-06 12:33:23 +05:30
Evgeny Mankov 6ac4d18753 Merge pull request #1759 from emankov/master
[HIP] Unify hipError_t (Step 2)

[ROCm/hip commit: 0dadb23327]
2019-12-30 19:21:09 +03:00
Sarbojit2019 6d2ec032fe 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/hip commit: aa4aea0754]
2019-12-30 12:44:24 +05:30
Aryan Salmanpour ffea90f865 [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/hip commit: 6968aeb841]
2019-12-30 12:42:17 +05:30