Commit grafiek

5050 Commits

Auteur SHA1 Bericht Datum
Christophe Paquot ee23705de2 Merge "hipStreamAddCallback test seg faults" into amd-master-next
[ROCm/clr commit: 446c8685e2]
2020-03-20 12:05:37 -04:00
kjayapra-amd 216aba5512 SWDEV-216213 - Lookup module functions from PlatformState::functions_.
Change-Id: I91dfe327f2ebdcf4c9b39ddd14d60aa0ce2fa9f4


[ROCm/clr commit: 334a1f8770]
2020-03-20 11:52:28 -04:00
Christophe Paquot 9130e160d8 hipStreamAddCallback test seg faults
Change-Id: I1f107fc8a5c586cd571f0280ed8716c5f89d25b7
SWDEV-227875: Need to add a dummy marker in case the stream is empty.


[ROCm/clr commit: 31df9b358d]
2020-03-19 11:11:59 -07: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 fadfaf9c80 Program texture flags in a better way
Not sure what I was thinking when initially implementing this...

Change-Id: Ib82f0f5a86683c08823dd4b59c98259d27151822


[ROCm/clr commit: ea7170d33c]
2020-03-18 18:15:09 -04:00
Vladislav Sytchenko b2e593fc62 Purge the use of ihip*impl() texture APIs
These are artifacts left from HIP-HCC and now are not needed by HIP-VDI.

Change-Id: Ib25a1081fe6146c8a89659395151e9d5bdaf7519


[ROCm/clr commit: 8083935855]
2020-03-18 18:15:01 -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 5dc14caf87 Correctly infer the texture read mode
Currently we extract the read mode from the ihip*impl() calls, which is not correct. We should be getting it from the texture itself directly.

Change-Id: Idf6449fefa395a887138a252e8ea937a6897e600


[ROCm/clr commit: 42b149fe3c]
2020-03-18 18:14:45 -04:00
Vladislav Sytchenko 6f93cfe58e Correct the definition of ...
hipBindTextureToMipmappedArray()

The texture reference needs to be passed as a constant pointer.

Change-Id: I6d31204c7f2325a5bc1e8b6e089fd9f8d21d1d78


[ROCm/clr commit: 9731b61a60]
2020-03-18 18:14:36 -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 b255d351b0 Merge "Start the lifetime of the texture reference" into amd-master-next
[ROCm/clr commit: da5f852cdd]
2020-03-18 18:07:24 -04:00
Jiabao Xie 7c5db60141 Merge "Cleaned include statements and deleted unnecessary operator overloads" into amd-master-next
[ROCm/clr commit: bbfba3bae1]
2020-03-18 14:37:19 -04:00
Vladislav Sytchenko 84ba26979a 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/clr commit: b4a0008b36]
2020-03-18 12:30:11 -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 bfc7d1c663 Merge "[hip-vdi]Fix for TF build failure [SWDEV-225827]" into amd-master-next
[ROCm/clr commit: 4cd02f20cf]
2020-03-18 11:52:46 -04:00
Jatin Chaudhary b77ba57b22 Merge "Enable saxpy Test" into amd-master-next
[ROCm/clr commit: d32802a176]
2020-03-18 06:32:05 -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 84a6c0001d Enable simpleTexture2DLayered test for VDI
Change-Id: I420f68824c6825152ac50e5a129b11b6ad88deb9


[ROCm/clr commit: 3150e1ef8d]
2020-03-17 17:50:46 -04:00
Vladislav Sytchenko 76b3ece656 Enable hipNormalizedFloatValueTex test for VDI
Change-Id: Iac8631312f34821d919b69b3fcb92b9387ba31dd


[ROCm/clr commit: 8d01fad7ef]
2020-03-17 17:50:30 -04:00
Tao Sang 7cfb4bac80 Merge "Fix failure to get global variables" into amd-master-next
[ROCm/clr commit: b0defd2676]
2020-03-17 17:08:14 -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
Vladislav Sytchenko 364377eda0 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/clr commit: 3abd89925a]
2020-03-17 17:04:17 -04:00
Saleel Kudchadker cbaa4b75b2 Merge "libhiprtc.so needs to be installed in DEST dir" into amd-master-next
[ROCm/clr commit: 36a07434ea]
2020-03-17 16:17:55 -04:00
Tao Sang 7a171f28c7 Fix failure to get global variables
Implement _ihipGetGlobalVar() and ihipGetGlobalVar() to
get global variables.

Change-Id: I442ab6712e12306c3316f114f5dc42f6daefaad9


[ROCm/clr commit: 8e9e6a44a4]
2020-03-17 16:14:16 -04:00
Vladislav Sytchenko 0c7735ad28 Merge "Enable simpleTexture3D test for VDI" into amd-master-next
[ROCm/clr commit: 7da159bac2]
2020-03-17 11:41:05 -04:00
Jatin b6f99e9332 Enable saxpy Test
Change-Id: Iadb5f631e1cebaf016b1835510771b3b7fac0a55


[ROCm/clr commit: fa10ee3301]
2020-03-17 01:15:20 -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 303e04fa55 Enable simpleTexture3D test for VDI
Change-Id: Ida65de6e1cae64ac97b624eaab1340877a6fba73


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


[ROCm/clr commit: f663c93cd6]
2020-03-12 12:23:36 -07:00
Jiabao Xie 3b776dfe6a Cleaned include statements and deleted unnecessary operator overloads
Change-Id: I5eb322c05082e8f27584446af2f6a19243dc63c9


[ROCm/clr commit: f81175c3bc]
2020-03-12 14:13:01 -04:00
Saleel Kudchadker 80c688d9f5 Merge "Fix P4 linux make builds" into amd-master-next
[ROCm/clr commit: 7d86f10c3c]
2020-03-11 16:17:18 -04:00
Jiabao Xie 9bfc5ba7a9 Merge "Structs expanded to print struct information than address. Moved to hip_formatting.hpp. Reformatted to follow Google standards" into amd-master-next
[ROCm/clr commit: 214c07a5b8]
2020-03-10 15:56:49 -04:00
Yaxun Liu 4042bf151b Merge "Let hipcc not pass -mllvm option to HIP-Clang on Windows" into amd-master-next
[ROCm/clr commit: 3fbd8e0dc2]
2020-03-10 14:12:06 -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
Jiabao Xie 3f843860d0 Structs expanded to print struct information than address. Moved to hip_formatting.hpp. Reformatted to follow Google standards
Change-Id: I08695058c11db51b9f3cbe1deb4af944ebf9e64d


[ROCm/clr commit: 10d90ab2ad]
2020-03-10 13:12:28 -04:00
Yaxun (Sam) Liu 23a8dfc8f3 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/clr commit: e796a1ed78]
2020-03-10 13:07:26 -04:00
Saleel Kudchadker 049340f08e Fix P4 linux make builds
Change-Id: I6cc4e69a914389b53bc3b52535eb6faf24897a09


[ROCm/clr commit: 7bba326ed1]
2020-03-10 00:02:10 -07:00
Saleel Kudchadker 82a836eb22 Merge "Change HIPRTC Version to 9.0" into amd-master-next
[ROCm/clr commit: 8d05d06dce]
2020-03-09 22:55:27 -04:00
Vladislav Sytchenko 426e3ecb06 Merge "Add hipDrvMemcpy3D." into amd-master-next
[ROCm/clr commit: 305e4eedcf]
2020-03-09 18:13:16 -04:00
Julia Jiang cf35a2a1de Merge "SWDEV-225337-Fix test failure in hipNormalizedFloatValueTex kernel" into amd-master-next
[ROCm/clr commit: 9b436779e4]
2020-03-09 16:58:10 -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
jujiang 0279f1d89e SWDEV-225337-Fix test failure in hipNormalizedFloatValueTex kernel
Change-Id: I3a4c3b58578703993640a8d28242ec1a0ed5ff60


[ROCm/clr commit: 09d442d7c3]
2020-03-09 15:21:29 -04:00
Sameer Sahasrabuddhe 7973b6d45d enable hostcall tests
Change-Id: Ic6efe71c868defc53be214acd7e0907109ec1410


[ROCm/clr commit: 5fc32755b7]
2020-03-09 14:41:37 -04:00
Saleel Kudchadker 1834fb4c5e Change HIPRTC Version to 9.0
Change-Id: I7e88df61248f0fe6e33c2315805f9e49a493cf29


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


[ROCm/clr commit: 514c7c8ef8]
2020-03-09 14:18:46 -04:00