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

5092 Коммитов

Автор SHA1 Сообщение Дата
Vladislav Sytchenko 1ce5ee07d2 Add entry points for hipTexObject*() API
Even though the runtime and driver texture object API is one to one, the structs used by these APIs are not. See hipResourceDesc vs HIP_RESOURCE_DESC differences.

These differences are not trivial and most likely won't be able to handled by hipify, so we need new API entry points.

Change-Id: Id4bcb1ad0ae15378dbdb5a2ed07e5ea30f320082


[ROCm/hip commit: aea688b79c]
2020-04-01 14:51:51 -04:00
Saleel Kudchadker 1bd55da10a Merge "Cleanup stream from hip:Event class." into amd-master-next
[ROCm/hip commit: 34a3ed9c1b]
2020-03-31 20:14:48 -04:00
Vladislav Sytchenko deacfcd063 (SWDEV-229354)
This patch is a workaround to support user pitch for hipMemcpy{2D/3D}.

Historically OpenCL didn't support pitch with clEnqueueFillBuffer(), so neither did we in VDI. Adding it now will be slightly nontrivial, since the fill kernel and runtime in many places will need to be modified.

As a temporary workaround for cases when pitch > width, we can just enqueue a fill for each row separately. This implementation is slow, but it satisfies the correctness criteria.

Change-Id: Idfeca349288b51d6ff84a7cf001fb63c6a66818a


[ROCm/hip commit: 77223a8eca]
2020-03-31 18:12:56 -04:00
Reshabh Sharma 264082869c Output file name should not change flags picked for compiler (#1938)
Fixes SWDEV-207362,

The output file name should not contribute to picking up the right flags for the compiler. This fix solves issues when the output has conflicting extensions which confuses hipcc to treat them as the source files and add the required flags for them.

PS: Output file refers to the file followed by -o

Change-Id: I1095966c11143ad73e81fabc35b4e9de5d3afada
Example: hipcc test.o -o test.hip will add the flags for .hip compilation ignoring the fact that it is an output file


[ROCm/hip commit: bb62f8553f]
2020-03-31 16:13:43 -04:00
Saleel Kudchadker 2f1617e21a Cleanup stream from hip:Event class.
Change-Id: I98de07d33bb7fea8f5e2d32b288c15f10ce58902


[ROCm/hip commit: e4ea1f60bb]
2020-03-31 11:22:00 -07:00
Michael Hong Bin Liao e89abbb63b Merge "[hipcc] Remove the previous workaround." into amd-master-next
[ROCm/hip commit: 5b96f38c1e]
2020-03-31 10:05:56 -04:00
Michael LIAO 5e5f4e23d5 [vdi] Fix hipGetSymbol{Address|Size}
- Use symbol value as the qeury key. Compared to the symbol name, the
  symbol value is more robust as developers may use unqualified or
  qualified identifiers. It also removes the mangling and/or demangling
  requirement for the runtime API.

Change-Id: I9d4259f3842612c7cc98551269fc2092d8b5c19e


[ROCm/hip commit: b72196613a]
2020-03-31 00:26:53 -04:00
Christophe Paquot 24cd888b3a Do not retry to allocate when OOM. Shouldn't be needed since we idle on Free.
SWDEV-229214

Change-Id: I183006f409388e3c7981f2569649d01d6378be46


[ROCm/hip commit: 94a7ef6ed1]
2020-03-30 12:49:48 -07:00
kjayapra-amd ce461331d3 SWDEV-216213 - Use different static & dynamic module maps for faster lookup.
Change-Id: Ia605e76a411ad5be04046b9d61f1ac111d49bb4a


[ROCm/hip commit: b081912962]
2020-03-30 14:28:07 -04:00
Anusha Godavarthy Surya 03a8b95eb4 Merge "Update Enable/Disable peers to match cuda behaviour" into amd-master-next
[ROCm/hip commit: caa083357a]
2020-03-30 13:46:42 -04:00
agodavar d7db609b92 Update Enable/Disable peers to match cuda behaviour
Change-Id: I67194ccf77a0019368579ff7d95b7790fcf228f3


[ROCm/hip commit: bdb3a4b393]
2020-03-30 12:49:16 -04:00
Tao Sang 1254e84b7b Merge "enable HCC printf when using hip-clang" into amd-master-next
[ROCm/hip commit: 1697cfc790]
2020-03-30 11:37:23 -04:00
Michael LIAO 9a2e585762 [hipcc] Remove the previous workaround.
Change-Id: Ie1c2aa21689bc343c10fd4b699546cf500ced5b2


[ROCm/hip commit: f1566f6dc5]
2020-03-30 10:18:47 -04:00
Michael LIAO cbac75a003 [vdi] Fix calculation of MaxWaves
- Consider the case where `usedVGPRs` is zero.
- This fixes [SWDEV-228537](http://ontrack-internal.amd.com/browse/SWDEV-228537)

Change-Id: I8675311f5fe24fb59c5d45bada122afefb55b128


[ROCm/hip commit: 3a690e960f]
2020-03-30 09:10:16 -04:00
Sameer Sahasrabuddhe 3e7d08f5ee enable HCC printf when using hip-clang
This is cherry-picked from PR#1947 that was committed to the
github repo. It allows printf to work with hip-clang and HCC
runtime.

Change-Id: I754753250ea1e694cf3441722e2d4c9d25fa75bc


[ROCm/hip commit: 9a0c5d0653]
2020-03-28 00:18:21 -04:00
Saleel Kudchadker 35e9fa08be Merge "Check event status before notify" into amd-master-next
[ROCm/hip commit: 5a8add03a5]
2020-03-27 20:19:40 -04:00
Vladislav Sytchenko 511f2d03a2 (SWDEV-228794)
Adjust the origin of the copy if the user passes a pointer that wasn't allocated by the runtime.

Change-Id: I0aeb20195ed730857a461a53f537626ec2573fd1


[ROCm/hip commit: 6ed73f50f7]
2020-03-27 16:33:16 -04:00
Vladislav Sytchenko d3b9203359 (SWDEV-228794)
Add hipMallocHost()

Change-Id: Ia3c7c5ca94b39fe30f3a51d1b60782d3472259ff


[ROCm/hip commit: fd7a8f0367]
2020-03-27 15:57:48 -04:00
Vladislav Sytchenko 7a867cd98f (SWDEV-228782)
The only requirment from hipMallocPitch() is that the returned pitch is aligned to the HW image pitch alignment. There is no restriction on the size of the allocation, since the memory might not be used for images.

Change-Id: I97438e5fe4012ca4721b14b85f514dbac803c17c


[ROCm/hip commit: a91b82f00e]
2020-03-27 15:52:17 -04:00
Tao Sang aa94d9c654 Merge "Support hipFuncGetAttributes with hip-clang+Hcc RT" into amd-master-next
[ROCm/hip commit: af7d131749]
2020-03-27 14:36:41 -04:00
Saleel Kudchadker 573d676ef0 Check event status before notify
Change-Id: I68f6bbbf236e49b859be2d5afbe0c8282fe15dd3


[ROCm/hip commit: 178fc17846]
2020-03-27 11:32:46 -07:00
Tao Sang c20e13d1f0 Support hipFuncGetAttributes with hip-clang+Hcc RT
Fix issues of missing kernel function symbols and missing argument list via
using __hipRegister* functions.
Then the following tests can pass,
directed_tests/runtimeApi/module/hipFuncGetAttributes
directed_tests/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice
directed_tests/gcc/LaunchKernel

Change-Id: I52135b61e8283eb4f9f10f77895151e4e55418d9


[ROCm/hip commit: d967be8eeb]
2020-03-26 23:28:35 -04:00
Vladislav Sytchenko ec09381120 Add support for formating hipExtent objects
Change-Id: Iea54a510e81a856c0c450305b3e5a7179ee48295


[ROCm/hip commit: 2a2b9e47a0]
2020-03-26 14:45:20 -04:00
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 e149397f73 Replace hip::TextureObject with __hip_texture
This avoids the use of extra casts when obtaining a texture object handle.

Change-Id: I42df22bdad0ab9ac6c33cb8b282dee65fe7cfd6e


[ROCm/hip commit: c1475f948e]
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 99b2084641 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


[ROCm/hip commit: 8ddeeb4551]
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko b48bf1dc8c Handle offsets for dptr <-> image copies
Change-Id: I7a4a56ee07a26a741d2aac35502446d248f720ad


[ROCm/hip commit: faf3b83594]
2020-03-26 14:45:20 -04:00
Vladislav Sytchenko cbeaeeb249 Correctly format hipResourceDesc objects
The struct consists of a union - only the active object should be read.

Change-Id: I1c40965b61518acd91a2dcbae92a015ac9be346a


[ROCm/hip commit: fbe30090a2]
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 fca3ce3008 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


[ROCm/hip commit: 39558dc1b9]
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 f8d11f2f30 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


[ROCm/hip commit: ea701777d8]
2020-03-26 12:43:17 -04:00
Payam Ghafari 3a9576f1e2 Merge "updated cmake to create libamdhip64 static file as well" into amd-master-next
[ROCm/hip commit: 8e8b372726]
2020-03-26 01:10:12 -04:00
jujiang 910a8a5d1e SWDEV-226708-hipMemcpyPeerAsync failed-->Update test app to create stream on proper device
Change-Id: Ia9f51435ffb5a8fbbff39e363acecabb338bf34c


[ROCm/hip commit: 3774a1b106]
2020-03-25 17:22:56 -04:00
Payam 1eb9cea01e updated cmake to create libamdhip64 static file as well
Change-Id: I2054b9501cefa232abbf398524ab62450ab6805d


[ROCm/hip commit: 76703acb1d]
2020-03-25 16:37:57 -04:00
Vladislav Sytchenko a48ad08f24 Disable failing tests for VDI
The hipExtModuleLaunchKernel and hipModuleLoadDataMultiThreaded tests keeps randomly failing on Jenkins.

Change-Id: I87e5d54fb7429c14ff1dcecb20e03a7816670fae


[ROCm/hip commit: e7d91b8230]
2020-03-23 17:09:57 -04:00
Christophe Paquot 58ea10b083 Merge "Enable VDI Occupancy test" into amd-master-next
[ROCm/hip commit: 3dce22538c]
2020-03-23 13:38:57 -04:00
Saleel Kudchadker 09bd0d2138 Merge "Sync streams when freeing or destroying mem" into amd-master-next
[ROCm/hip commit: 91069c0de4]
2020-03-21 13:33:41 -04:00
Christophe Paquot 3c623309cd 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


[ROCm/hip commit: 653277bd3f]
2020-03-20 13:07:34 -07:00
Saleel Kudchadker 1b4bbfcb53 Sync streams when freeing or destroying mem
Change-Id: I6932f225a8b932bb2adbd5e37880f7e604496809


[ROCm/hip commit: 8b39e0b74e]
2020-03-20 10:53:23 -07:00
Christophe Paquot 7a418b2fce Merge "hipStreamAddCallback test seg faults" into amd-master-next
[ROCm/hip commit: d98879a250]
2020-03-20 12:05:37 -04:00
kjayapra-amd 64677e52e2 SWDEV-216213 - Lookup module functions from PlatformState::functions_.
Change-Id: I91dfe327f2ebdcf4c9b39ddd14d60aa0ce2fa9f4


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


[ROCm/hip commit: 3dfbfc408b]
2020-03-19 11:11:59 -07:00
Sarbojit Sarkar 771fb00972 Enable VDI Occupancy test
Change-Id: Ia89155ca909b9b7346584d19962fa0a94033bb17


[ROCm/hip commit: 3baff10618]
2020-03-19 10:43:40 +05:30
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 a6af2fc07f Program texture flags in a better way
Not sure what I was thinking when initially implementing this...

Change-Id: Ib82f0f5a86683c08823dd4b59c98259d27151822


[ROCm/hip commit: e8fa3b2589]
2020-03-18 18:15:09 -04:00
Vladislav Sytchenko 87cb03c270 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/hip commit: 2bad9e2821]
2020-03-18 18:15:01 -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 68c96cca38 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/hip commit: 7385a032ae]
2020-03-18 18:14:45 -04:00