Wykres commitów

4332 Commity

Autor SHA1 Wiadomość Data
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 67e8b8bcd2 Merge "Start the lifetime of the texture reference" into amd-master-next
[ROCm/hip commit: 816a56a70b]
2020-03-18 18:07:24 -04:00
Jiabao Xie c12977513d Merge "Cleaned include statements and deleted unnecessary operator overloads" into amd-master-next
[ROCm/hip commit: ae6dd6c2fc]
2020-03-18 14:37:19 -04:00
Vladislav Sytchenko a7ff065adc Start the lifetime of the texture reference
reinterpret_cast<> doesn't create an object, so the texref is actually unitiliazed. This may lead to garbage data in some of its struct members.

Initialize it by performing a placement new. The constructer should set all of its members to default values. There's no way currently to extract the channel type, so use single channel char for now.

Change-Id: I41b305a75bb3f30130324de785099f55b3e130c7


[ROCm/hip commit: 292d008a64]
2020-03-18 12:30:11 -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 ac7f37d9c3 Merge "[hip-vdi]Fix for TF build failure [SWDEV-225827]" into amd-master-next
[ROCm/hip commit: 0a35286988]
2020-03-18 11:52:46 -04:00
Jatin Chaudhary a9bccfb6b2 Merge "Enable saxpy Test" into amd-master-next
[ROCm/hip commit: b68b2884ba]
2020-03-18 06:32:05 -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 ac1cc9bacd Enable simpleTexture2DLayered test for VDI
Change-Id: I420f68824c6825152ac50e5a129b11b6ad88deb9


[ROCm/hip commit: 1b288456ae]
2020-03-17 17:50:46 -04:00
Vladislav Sytchenko 520fe971a9 Enable hipNormalizedFloatValueTex test for VDI
Change-Id: Iac8631312f34821d919b69b3fcb92b9387ba31dd


[ROCm/hip commit: a605458a0c]
2020-03-17 17:50:30 -04:00
Tao Sang bc3232490f Merge "Fix failure to get global variables" into amd-master-next
[ROCm/hip commit: 2827eb110e]
2020-03-17 17:08:14 -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
Vladislav Sytchenko ac4378a16a Rework hipNormalizedFloatValueTex test
This is currently so buggy that it causes a runtime crash on Nvidia platfrom...

Disable the new version for hcc and vdi, header fixes are required for it to pass.

Currently tex1D<char, hipTextureType1D, hipReadModeNormalizedFloat> returns a char, when the actual sampled pixel value is a float, so the hi 3 bytes get lost.

Change-Id: I8222a4d8d1d8b101eb43f3f8dfbe4818f885f8ea


[ROCm/hip commit: cc134f7c58]
2020-03-17 17:04:17 -04:00
Saleel Kudchadker 2d79f035fe Merge "libhiprtc.so needs to be installed in DEST dir" into amd-master-next
[ROCm/hip commit: 9343eee07c]
2020-03-17 16:17:55 -04:00
Tao Sang acff9e29dd Fix failure to get global variables
Implement _ihipGetGlobalVar() and ihipGetGlobalVar() to
get global variables.

Change-Id: I442ab6712e12306c3316f114f5dc42f6daefaad9


[ROCm/hip commit: d432dbfe20]
2020-03-17 16:14:16 -04:00
Vladislav Sytchenko 526bfd939f Merge "Enable simpleTexture3D test for VDI" into amd-master-next
[ROCm/hip commit: c94160aafb]
2020-03-17 11:41:05 -04:00
Jatin e8f737c277 Enable saxpy Test
Change-Id: Iadb5f631e1cebaf016b1835510771b3b7fac0a55


[ROCm/hip commit: 2dac197548]
2020-03-17 01:15:20 -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 d48472b772 Enable simpleTexture3D test for VDI
Change-Id: Ida65de6e1cae64ac97b624eaab1340877a6fba73


[ROCm/hip commit: bfcce529fa]
2020-03-13 18:34:39 -04:00
Saleel Kudchadker 7bed9a90b7 libhiprtc.so needs to be installed in DEST dir
Change-Id: I3ae8b4f18a05a707b87fa73873c1721928ebe904


[ROCm/hip commit: 9ecd0e79b6]
2020-03-12 12:23:36 -07:00
Jiabao Xie 1e76a4a5be Cleaned include statements and deleted unnecessary operator overloads
Change-Id: I5eb322c05082e8f27584446af2f6a19243dc63c9


[ROCm/hip commit: e7abc7a553]
2020-03-12 14:13:01 -04:00
Saleel Kudchadker 27288c9358 Merge "Fix P4 linux make builds" into amd-master-next
[ROCm/hip commit: ffcbd7e633]
2020-03-11 16:17:18 -04:00
Jiabao Xie 0c637f136f Merge "Structs expanded to print struct information than address. Moved to hip_formatting.hpp. Reformatted to follow Google standards" into amd-master-next
[ROCm/hip commit: edce444e33]
2020-03-10 15:56:49 -04:00
Yaxun Liu 87054b217a Merge "Let hipcc not pass -mllvm option to HIP-Clang on Windows" into amd-master-next
[ROCm/hip commit: 027d6fb928]
2020-03-10 14:12:06 -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
Jiabao Xie 608d85a005 Structs expanded to print struct information than address. Moved to hip_formatting.hpp. Reformatted to follow Google standards
Change-Id: I08695058c11db51b9f3cbe1deb4af944ebf9e64d


[ROCm/hip commit: 604befb472]
2020-03-10 13:12:28 -04:00
Yaxun (Sam) Liu 9683880363 Let hipcc not pass -mllvm option to HIP-Clang on Windows
Currently there is a clang bug on Windows causing duplicate -mllvm options in clang -cc1.

Tempoarily disable -mllvm options for HIP-Clang on Windows until the bug is fixed.

Change-Id: I3a4393ba7745989398dc6c6001722837dad18704


[ROCm/hip commit: 92af5e4375]
2020-03-10 13:07:26 -04:00
Saleel Kudchadker 411708c6c1 Fix P4 linux make builds
Change-Id: I6cc4e69a914389b53bc3b52535eb6faf24897a09


[ROCm/hip commit: aa04b964cf]
2020-03-10 00:02:10 -07:00
Saleel Kudchadker 2b4b2b2dc9 Merge "Change HIPRTC Version to 9.0" into amd-master-next
[ROCm/hip commit: bcd23335fb]
2020-03-09 22:55:27 -04:00
Vladislav Sytchenko 3082077cc9 Merge "Add hipDrvMemcpy3D." into amd-master-next
[ROCm/hip commit: 7fe7edd77d]
2020-03-09 18:13:16 -04:00
Julia Jiang a4255f8089 Merge "SWDEV-225337-Fix test failure in hipNormalizedFloatValueTex kernel" into amd-master-next
[ROCm/hip commit: 1fa3bf5123]
2020-03-09 16:58:10 -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
jujiang d745491b17 SWDEV-225337-Fix test failure in hipNormalizedFloatValueTex kernel
Change-Id: I3a4c3b58578703993640a8d28242ec1a0ed5ff60


[ROCm/hip commit: 47cc94b1b4]
2020-03-09 15:21:29 -04:00
Sameer Sahasrabuddhe dd7cad9ed4 enable hostcall tests
Change-Id: Ic6efe71c868defc53be214acd7e0907109ec1410


[ROCm/hip commit: c7d7640d87]
2020-03-09 14:41:37 -04:00
Saleel Kudchadker 0585e40955 Change HIPRTC Version to 9.0
Change-Id: I7e88df61248f0fe6e33c2315805f9e49a493cf29


[ROCm/hip commit: c540a55d21]
2020-03-09 11:28:55 -07:00
Saleel Kudchadker 043b44934a Fix HIPRTC headers to export C style symbols
Change-Id: I3e0d2b19ace4a9096e3e46bd22f420483da51a8a


[ROCm/hip commit: 08c6d941ca]
2020-03-09 14:18:46 -04:00
Vladislav Sytchenko 7affd35b9b Fix Windows build.
extern "C" on Windows implies nothrow. We shouldn't be throwing exceptions either way.

Change-Id: If0ed1f7ec194bf7f65b7cea1a5c250e768a8f190


[ROCm/hip commit: 38004b6ef0]
2020-03-09 11:46:55 -04:00
Saleel Kudchadker cfcf0d025a Merge "Merge branch 'master' into amd-master-next" into amd-master-next
[ROCm/hip commit: a30f9ab2db]
2020-03-06 16:29:16 -05:00
Vladislav Sytchenko 2f8a37880d Add hipMemcpy3DParams -> HIP_MEMCPY3D conversion
Change-Id: I66dd1ae722b8c0d7ddb2562c958b14854e6b86bc


[ROCm/hip commit: 2e9d177304]
2020-03-06 14:33:52 -05:00
Vladislav Sytchenko 42d9576cbe Fix dangling pointer after hipUnbindTexture() call
Change-Id: Ic4b476c62ebfae31e94dd139b20b6aaaa52bb866


[ROCm/hip commit: dbb8f96a8e]
2020-03-06 14:10:56 -05: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
Vladislav Sytchenko 8bf13d8d20 Add missing coversion cases.
Handle converting signed int cases of hipResourceViewFormat to number of channels.

Change-Id: Ica8ae6f644edfaa0d4803d0b8e90e320479118e2


[ROCm/hip commit: 90feaf046f]
2020-03-06 14:06:21 -05:00
Tao Sang 9862e33448 Merge "Fix hipExtLaunchMultiKernelMultiDevice compilation issue" into amd-master-next
[ROCm/hip commit: 4b312caca4]
2020-03-06 12:32:55 -05:00
Evgeny Mankov c48bf1b4b0 [HIPIFY][doc] Update README.md: LLVM 10.0.0-rc3 is supported
+ Add -DLLVM_TEMPORARILY_ALLOW_OLD_TOOLCHAIN=ON for LLVM 10.0.0 or newer
+ Supported versions update


[ROCm/hip commit: dd5f3fd282]
2020-03-06 18:17:05 +03:00
jujiang 127e56e482 SWDEV-225338-simpletexture3D failed-Fix bug in test app
Change-Id: I106d1951368e0c7a8ccd6aa462c22f81e8332cdc


[ROCm/hip commit: c5aa54cbd2]
2020-03-05 16:50:27 -05:00
Vladislav Sytchenko 9a62b2ffe7 Correct logic in ihipMemcpyAtoA()
HIP assumes that image width is in bytes, but VDI assumes that image width in pixels. Need to perform byte -> pixel coversion before doing anything.

Change-Id: Ia9fd1f46d05db3fbe8049add10b4d7e5118a2b9a


[ROCm/hip commit: 3fc50731ae]
2020-03-05 15:03:18 -05:00
Vladislav Sytchenko 2f2a2843dc Fix typos in ihipMemcpy3D()
Change-Id: I8720f113642f00bb013cf46284e9b13cc932bf4a


[ROCm/hip commit: f28e2ebb42]
2020-03-05 14:11:55 -05:00
agodavar de0b92948e Enable tests that are passing on hip-vdi
Change-Id: I7de965f7e8bb7e4d0fa61cd584f0cd118c1e212e


[ROCm/hip commit: 65ffaf82d8]
2020-03-05 16:08:51 +05:30
German Andryeyev c503c44b67 Merge "SWDEV-193956 - [hipclang-vdi-rocm][perf]" into amd-master-next
[ROCm/hip commit: 33da3f5c50]
2020-03-04 18:04:52 -05:00