提交線圖

5070 次程式碼提交

作者 SHA1 備註 日期
Vladislav Sytchenko 2b538b09e6 Add support for formating hipExtent objects
Change-Id: Iea54a510e81a856c0c450305b3e5a7179ee48295
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko 06bb183558 Add initial entry points for mipmapped array API
Change-Id: Icd59cc7323ddcb6773da6105260415a1e6f4cdcb
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko 6cfbe19160 Replace hip::TextureObject with __hip_texture
This avoids the use of extra casts when obtaining a texture object handle.

Change-Id: I42df22bdad0ab9ac6c33cb8b282dee65fe7cfd6e
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko 428b56e411 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
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko 23211f05d6 Enable initial sRGB support
Instead of using the sampler field force_degamma to perform sRGB->linear conversion during pixel sampling, we use an appropriate image format instead. The overhead of this is having to create an image view when creating a texture object from an array.

Change-Id: I1ca368c312c1fd4b6f784a3a1b35b5eeb28070ff
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko af1c5a0015 Handle offsets for dptr <-> image copies
Change-Id: I7a4a56ee07a26a741d2aac35502446d248f720ad
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko 1cacb5c9b7 Correctly format hipResourceDesc objects
The struct consists of a union - only the active object should be read.

Change-Id: I1c40965b61518acd91a2dcbae92a015ac9be346a
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko 3ab8ff87cc Set textureObject to nullptr
This avoids dangling pointers for newly initiazlied textures

Change-Id: Ia444b91fe17fd756ed583ec595ae1febbdfbd034
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko d662abaf55 Modify formatting for textureReferences
We don't program the numChannels and format members (these are HCC specific), so printing these will only display garbage.

Change-Id: I83dc8be9a3cae2659c64f4594d07c05330d2dd14
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko dc015d27c1 Correct typos in texture function declarations
Change-Id: I492995e984eda2e8a5e806c5d4c9c78da09ac483
2020-03-26 12:43:17 -04:00
Vladislav Sytchenko 9823ccf5b0 Allow creating texture from unaligned user ptr
All we have to do is align the ptr to HW requirments an if it's not zero, then return the offset to the user.

We currently don't have anywhere to store this offset, so hipGetTextureAlignmentOffset() will still always return 0.

Change-Id: If31998127d99a2a3222a026d88249519d6102505
2020-03-26 12:43:17 -04:00
Payam Ghafari 857990e3f6 Merge "updated cmake to create libamdhip64 static file as well" into amd-master-next 2020-03-26 01:10:12 -04:00
jujiang 01272807ea SWDEV-226708-hipMemcpyPeerAsync failed-->Update test app to create stream on proper device
Change-Id: Ia9f51435ffb5a8fbbff39e363acecabb338bf34c
2020-03-25 17:22:56 -04:00
Payam b3d705f7b3 updated cmake to create libamdhip64 static file as well
Change-Id: I2054b9501cefa232abbf398524ab62450ab6805d
2020-03-25 16:37:57 -04:00
Vladislav Sytchenko 684cdd4a18 Disable failing tests for VDI
The hipExtModuleLaunchKernel and hipModuleLoadDataMultiThreaded tests keeps randomly failing on Jenkins.

Change-Id: I87e5d54fb7429c14ff1dcecb20e03a7816670fae
2020-03-23 17:09:57 -04:00
Christophe Paquot cf09b9ca1e Merge "Enable VDI Occupancy test" into amd-master-next 2020-03-23 13:38:57 -04:00
Saleel Kudchadker 436e5bb664 Merge "Sync streams when freeing or destroying mem" into amd-master-next 2020-03-21 13:33:41 -04:00
Christophe Paquot 9f81498bb5 hipStreamAddCallback test seg faults
Change-Id: If419d2fad490d0ed50eb1315af809fc1deda1ce3
SWDEV-227875: Add a lock in streams to lock when the callback is call so we make sure things aren't moving forward in the stream
2020-03-20 13:07:34 -07:00
Saleel Kudchadker 68df8efe90 Sync streams when freeing or destroying mem
Change-Id: I6932f225a8b932bb2adbd5e37880f7e604496809
2020-03-20 10:53:23 -07:00
Christophe Paquot 446c8685e2 Merge "hipStreamAddCallback test seg faults" into amd-master-next 2020-03-20 12:05:37 -04:00
kjayapra-amd 334a1f8770 SWDEV-216213 - Lookup module functions from PlatformState::functions_.
Change-Id: I91dfe327f2ebdcf4c9b39ddd14d60aa0ce2fa9f4
2020-03-20 11:52:28 -04:00
Christophe Paquot 31df9b358d hipStreamAddCallback test seg faults
Change-Id: I1f107fc8a5c586cd571f0280ed8716c5f89d25b7
SWDEV-227875: Need to add a dummy marker in case the stream is empty.
2020-03-19 11:11:59 -07:00
Sarbojit Sarkar 6cf86b20c6 Enable VDI Occupancy test
Change-Id: Ia89155ca909b9b7346584d19962fa0a94033bb17
2020-03-19 10:43:40 +05:30
Vladislav Sytchenko 78f7954765 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
2020-03-18 18:15:17 -04:00
Vladislav Sytchenko ea7170d33c Program texture flags in a better way
Not sure what I was thinking when initially implementing this...

Change-Id: Ib82f0f5a86683c08823dd4b59c98259d27151822
2020-03-18 18:15:09 -04:00
Vladislav Sytchenko 8083935855 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
2020-03-18 18:15:01 -04:00
Vladislav Sytchenko 53dd6b7a66 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
2020-03-18 18:14:53 -04:00
Vladislav Sytchenko 42b149fe3c 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
2020-03-18 18:14:45 -04:00
Vladislav Sytchenko 9731b61a60 Correct the definition of ...
hipBindTextureToMipmappedArray()

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

Change-Id: I6d31204c7f2325a5bc1e8b6e089fd9f8d21d1d78
2020-03-18 18:14:36 -04:00
Vladislav Sytchenko 7a81e55de7 Correct the declaration of hipBindTexture2D()
The texture reference needs to be passed as a constant pointer.

Change-Id: Idde461f0f328ac87ce677b6bab3203161b514cbf
2020-03-18 18:08:23 -04:00
Vladislav Sytchenko 7765792a42 Correct the declaration of hipBindTextureToArray()
The texture reference needs to be passed as a constant pointer.

Change-Id: Iff171626536071fb2020cfff7132ec930577b1b9
2020-03-18 18:08:13 -04:00
Vladislav Sytchenko 600b1006c8 Correct the declaration of hipBindTexture()
The texture reference needs to be passed as a constant pointer.

Change-Id: I36ca0bddaba30becfc2ce70dd9e5b7db66c57f27
2020-03-18 18:08:01 -04:00
Vladislav Sytchenko 6ff0d15a62 Add missing mipmap API entries
Introduce hipFreeMipmappedArray(), hipMallocMipmappedArray() and hipGetMipmappedArrayLevel() APIs.

Change-Id: I878228c79fa1c54536c17d6baf45f83d51d2b1c7
2020-03-18 18:07:45 -04:00
Vladislav Sytchenko b4b05e2ab0 Don't hardcode the texture read mode
The readmode needs to be inferred from the template arguments.

Change-Id: I067037035e2492a24eac47e16d4015f879be0ea7
2020-03-18 18:07:33 -04:00
Vladislav Sytchenko da5f852cdd Merge "Start the lifetime of the texture reference" into amd-master-next 2020-03-18 18:07:24 -04:00
Jiabao Xie bbfba3bae1 Merge "Cleaned include statements and deleted unnecessary operator overloads" into amd-master-next 2020-03-18 14:37:19 -04:00
Vladislav Sytchenko b4a0008b36 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
2020-03-18 12:30:11 -04:00
Vladislav Sytchenko 8667e21b25 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
2020-03-18 12:24:40 -04:00
Vladislav Sytchenko a994da8af6 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
2020-03-18 12:24:40 -04:00
Sarbojit Sarkar 4cd02f20cf Merge "[hip-vdi]Fix for TF build failure [SWDEV-225827]" into amd-master-next 2020-03-18 11:52:46 -04:00
Jatin Chaudhary d32802a176 Merge "Enable saxpy Test" into amd-master-next 2020-03-18 06:32:05 -04:00
Sarbojit Sarkar 957caff8b4 [hip-vdi]Fix for TF build failure [SWDEV-225827]
Change-Id: I8478779bef92bad8353b8d066b28c220bb59b98d
2020-03-17 22:52:01 -04:00
Vladislav Sytchenko 3150e1ef8d Enable simpleTexture2DLayered test for VDI
Change-Id: I420f68824c6825152ac50e5a129b11b6ad88deb9
2020-03-17 17:50:46 -04:00
Vladislav Sytchenko 8d01fad7ef Enable hipNormalizedFloatValueTex test for VDI
Change-Id: Iac8631312f34821d919b69b3fcb92b9387ba31dd
2020-03-17 17:50:30 -04:00
Tao Sang b0defd2676 Merge "Fix failure to get global variables" into amd-master-next 2020-03-17 17:08:14 -04:00
Vladislav Sytchenko 3d7945faae 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
2020-03-17 17:04:37 -04:00
Vladislav Sytchenko 3abd89925a 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
2020-03-17 17:04:17 -04:00
Saleel Kudchadker 36a07434ea Merge "libhiprtc.so needs to be installed in DEST dir" into amd-master-next 2020-03-17 16:17:55 -04:00
Tao Sang 8e9e6a44a4 Fix failure to get global variables
Implement _ihipGetGlobalVar() and ihipGetGlobalVar() to
get global variables.

Change-Id: I442ab6712e12306c3316f114f5dc42f6daefaad9
2020-03-17 16:14:16 -04:00
Vladislav Sytchenko 7da159bac2 Merge "Enable simpleTexture3D test for VDI" into amd-master-next 2020-03-17 11:41:05 -04:00