نمودار کامیت

6604 کامیت‌ها

مولف SHA1 پیام تاریخ
Christophe Paquot f037cb7285 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/clr commit: 9f81498bb5]
2020-03-20 13:07:34 -07:00
Saleel Kudchadker adc574b2ed Sync streams when freeing or destroying mem
Change-Id: I6932f225a8b932bb2adbd5e37880f7e604496809


[ROCm/clr commit: 68df8efe90]
2020-03-20 10:53:23 -07:00
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
Yaxun (Sam) Liu 95dc897815 Let hipcc treat .c program as C program (#1945)
nvcc treats .c program as C program and .cpp program as C++ program.

Currently hipcc treats .c and .cpp programs as HIP programs.

It is desirable to let hipcc behave like nvcc.

Currently it is not feasible to let hipcc treat .cpp programs as C++ program
since there are too many HIP applications use .cpp as extension for HIP programs.
However we should be able let hipcc treat .c program C program since there
are few applications use .c as extensioin for HIP programs.

[ROCm/clr commit: 31d9ea89f1]
2020-03-19 16:16:31 +05:30
Jeff Daily 1a434dccd3 fix hipStreamAddCallback, block future work on stream (#1934)
[ROCm/clr commit: 1444f850ac]
2020-03-19 16:16:04 +05:30
nelsonc-amd 006f7e0b06 packaging: change Provides to use dash instead of underscore (#1913)
Addresses an installation problem for several HIP packages.
Packages builds have been tested locally, validation in progress.

[ROCm/clr commit: d90a0c05c0]
2020-03-19 16:15:09 +05:30
Sarbojit Sarkar aa49455725 Enable VDI Occupancy test
Change-Id: Ia89155ca909b9b7346584d19962fa0a94033bb17


[ROCm/clr commit: 6cf86b20c6]
2020-03-19 10:43:40 +05:30
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
Evgeny Mankov 0844a41baa Merge pull request #1955 from emankov/hipify-clang
[HIPIFY][perl] Fix missed sync with hipify-clang

[ROCm/clr commit: 01bcc0015c]
2020-03-18 14:30:11 +03:00
Evgeny Mankov 677f9a3eb7 [HIPIFY][perl] Fix missed sync with hipify-clang
[ROCm/clr commit: 758b3864b7]
2020-03-18 14:24:02 +03:00
Evgeny Mankov 02294f492f Merge pull request #1954 from emankov/hipify-clang
[HIPIFY][SPARSE] sync with hipSPARSE's "added csrsm2" (#83)

[ROCm/clr commit: 6ce54a6e25]
2020-03-18 13:46:05 +03:00
Evgeny Mankov a82c69da48 [HIPIFY][SPARSE] sync with hipSPARSE's "added csrsm2" (#83)
+ https://github.com/ROCmSoftwarePlatform/hipSPARSE/pull/83
+ Update doc and hipify-perl accordingly


[ROCm/clr commit: 126ebadcc6]
2020-03-18 13:44:50 +03:00
Jatin Chaudhary b77ba57b22 Merge "Enable saxpy Test" into amd-master-next
[ROCm/clr commit: d32802a176]
2020-03-18 06:32:05 -04:00
Evgeny Mankov a983f3cdde Merge pull request #1953 from emankov/hipify-clang
[HIPIFY][SPARSE] sync with hipSPARSE

[ROCm/clr commit: aeac6239ff]
2020-03-18 13:09:21 +03:00
Evgeny Mankov be4ba44d12 [HIPIFY][SPARSE] sync with hipSPARSE
+ Update doc and hipify-perl accordingly


[ROCm/clr commit: baa0c68ab4]
2020-03-18 13:07:46 +03:00
Nagy-Egri Máté Ferenc f3c3070404 Remove debug message (#1950)
[ROCm/clr commit: b519d1d681]
2020-03-18 11:20:53 +05:30
Aryan Salmanpour 1c8dc380bc [HIP] use markers to sync cooperative and normal queues (#1948)
[ROCm/clr commit: 66735bff13]
2020-03-18 11:20:43 +05:30
Yaxun (Sam) Liu 5fffd89ae8 Workaround for libc++ include path for HIP-Clang (#1917)
HIP-Clang cuda_wrapper headers require clang include path before standard C++ include path.
However libc++ include path requires to be before clang include path.
To workaround this, we pass -isystem with the parent directory of clang include
path instead of the clang include path itself.

[ROCm/clr commit: 0307ae4c9a]
2020-03-18 11:20:21 +05:30
jglaser 76759ae09f Implement accurate max block size in hipFuncGetAttributes() (#1676)
This PR takes ensures that the maxThreadsPerBlock returned by hipFuncGetAttributes is both a multiple of the warp size and that the register usage of the maximum block does not exceed the number of available registers.

Fixes #1662

[ROCm/clr commit: ea28d64297]
2020-03-18 11:20:06 +05:30
Paul Fultz II 85fdd5afc9 Add missing flags for hip::device target on hip-clang (#1230)
This adds the missing compilation flags to hip::device so it can compile with hip-clang compiler.

[ROCm/clr commit: 9a47358324]
2020-03-18 11:18:58 +05:30
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
Evgeny Mankov a521966508 Merge pull request #1951 from emankov/hipify-clang
[HIPIFY][doc] Update README.md: LLVM 10.0.0-rc4 is supported

[ROCm/clr commit: dddc145893]
2020-03-17 18:57:13 +03:00
Evgeny Mankov 0d11749db2 [HIPIFY][doc] Update README.md: LLVM 10.0.0-rc4 is supported
[ROCm/clr commit: 7f51cdb9da]
2020-03-17 18:51:46 +03: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 Chaudhary 6d39b6fbb8 Adding Half Abs APIs (#1902)
[ROCm/clr commit: 8a3ce17ba6]
2020-03-17 14:13:19 +05:30