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

1047 Коммитов

Автор SHA1 Сообщение Дата
Maneesh Gupta 4ee600ed5e Merge pull request #1280 from ROCm-Developer-Tools/fix_dont_break_hcc_just_because
This difference makes absolutely no sense.
2019-08-05 09:51:53 +00:00
Sarbojit2019 3bfff0a23d Enabled gcc for hip host code (#1214)
* Enabled gcc for hip host code

* Adding tests for hip code + (gcc & g++), without kernels

* Excluding nvcc platforms for gcc and g++ tests + Addressing review comments

* minor code clean-up

* Add rocm include path

* Added relative path for library

* Hiding non supported functions for gcc

* Incorporating review comments
2019-08-05 09:51:36 +00:00
Jeff Daily 1eb3dbf065 consolidate thread local storage (#915)
* all thread local access now through single struct

* clean up old commented-out code, more use of GET_TLS()

* fewer calls to GET_TLS by passing tls as a funtion argument

* revert unnecessary change to printf

* fix failing tests due to TLS change

* fix merge conflicts in ihipOccupancyMaxActiveBlocksPerMultiprocessor
2019-08-05 09:51:02 +00:00
Alex Voicu a85b0fe68e This difference makes absolutely no sense. 2019-08-02 12:30:03 +03:00
wkwchau e7447d5809 Added query of hipDeviceAttributeHdpMemFlushCntl and hipDeviceAttribu… (#1238)
* Added query of hipDeviceAttributeHdpMemFlushCntl and hipDeviceAttributeHdpRegFlushCntl

* Added NVCC blocker for the hip*FlushCntl test cases
2019-08-01 16:03:35 +00:00
Maneesh Gupta adc95002c6 Merge pull request #1277 from mangupta/nvcc_devprop
[nvcc] Populate missing fields in hipGetDeviceProperties
2019-08-01 08:59:58 +00:00
wkwchau 4b18b321f7 Added support of hipOccupancyMaxActiveBlocksPerMultiprocessor & hipOc… (#1240)
* Added support of hipOccupancyMaxActiveBlocksPerMultiprocessor & hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags APIs

* Taking into account of SGPR usage to determine the max active blocks in hipOccupancyMaxActiveBlocksPerMultiprocessor()
2019-08-01 08:58:48 +00:00
Maneesh Gupta 3ec381d729 [nvcc] Populate missing fields in hipGetDeviceProperties
Change-Id: Ie90e02674d503e385f144f1ead3d53ff7b49cecc
2019-08-01 13:16:39 +05:30
Yaxun (Sam) Liu d7aea38d74 Temporarily disable maxTexture1D in hipDeviceProp_t for VDI (#1237) 2019-07-19 15:38:32 +00:00
ansurya 8e496c09d9 Add Max Texture 1D,2D,3D device properties (#1226)
* Add Max Texture 1D,2D,3D device properties

* Corrected testcase to use enums defined in hipDeviceAttribute_t

* Added texture 1D,2D and 3D support for NVIDIA path
2019-07-18 03:18:50 +00:00
ansurya 6b6258b84f UChar and UShort textures as Normalized Float (#1213)
* UChar and UShort textures as Normalized Float

* UChar and UShort textures as Normalized Float for all float variants

* Handled uninitilaized texture format value
2019-07-18 03:18:28 +00:00
Evgeny Mankov c7117df91b [HIP][HIPIFY] Split HIP_ARRAY_DESCRIPTOR struct to HIP_ARRAY_DESCRIPTOR and HIP_ARRAY3D_DESCRIPTOR
[Reason] To be compatible with CUDA [#1133]

Update HIP code, hipify-clang, tests and docs

[TODO] Add support of the corresponding functions on nvcc fallback path
2019-07-11 14:58:16 +03:00
mhbliao a16cd0c69a Add missing compound assignment operators. (#1216) 2019-07-11 09:16:15 +00:00
Jatin Chaudhary 3e213d94e1 Adding fix to initialize a dim3 variable from blockIdx. (#1212)
Adding gitignore for ctags
2019-07-11 09:16:03 +00:00
Aryan Salmanpour e87b5a1e8a Fix a compilation error when templated operator<< fails if 'U' type is std::ostream. (#1209) 2019-07-10 03:05:04 +00:00
darkbuck c8b76a6a63 Fix even more typo. (#1204) 2019-07-08 11:00:51 +00:00
searlmc1 a9870f0465 Fix typo ( T in Vec3_cmp = needs to be int ). (#1203)
Typo introduced here:
commit 67abac1365
Author: Alex Voicu <alexandru.voicu@amd.com>
Date:   Mon Jun 24 20:02:09 2019 -0500

    Put 3-wide vector types on a ketogenic diet. (#1180)
2019-07-04 00:44:26 +00:00
kpyzhov b6b749f714 Change the type of hipTextureObject_t to pointer to a struct. (#1199)
This is necessary to allow Runtime to perform required texture buffer handling.
2019-07-04 00:44:06 +00:00
Jannik 2a1b0ba27d Fix call to cuda library (#1197)
* Remove flags parameter from hipOccupancyMaxPotentialBlockSize

This commit makes the hipOccupancyMaxPotentialBlockSize method consistent with hcc path and the CUDA API.
2019-07-04 00:43:22 +00:00
Yaxun (Sam) Liu 85c3b903df Add __hip_pinned_shadow__ for hip-clang and fix texture reference tests (#1200) 2019-07-03 08:51:59 +00:00
Alex Voicu 67abac1365 Put 3-wide vector types on a ketogenic diet. (#1180)
* Put 3-wide vector types on a ketogenic diet.

* Remove needless include.

* Do not be narrow-minded.

* Do not be narrow-minded.

* Put the C people on a diet too.
2019-06-25 06:32:09 +05:30
Aryan Salmanpour 96dc74897d [hip] implement the hipExtLaunchMultiKernelMultiDevice API (#1165)
* [hip] implement the hipExtLaunchMultiKernelMultiDevice API

* add a guard to check the HCC version for acquire_locked_hsa_queue() API which was introdued in HCC for ROCm 2.5

* modified code based on the requested changes

* changes to lock all streams before launching kernels for each device and unlock them after the dispatches

* check each stream to be valid before starting to lock all the streams
2019-06-20 05:59:05 +05:30
wkwchau d492f1fd6b Implement the hipOccupancyMaxPotentialBlockSize function (#1162)
* Implement the hipOccupancyMaxPotentialBlockSize function

* Replaced hipGetDeviceProperties() call by ihipGetDeviceProperties() in ihipOccupancyMaxPotentialBlockSize()

* Add test for hipOccupancyMaxPotentialBlockSize in Module API

* Added extern declaration for ihipGetDeviceProperties() to be accessed inside ihipOccupancyMaxPotentialBlockSize()

* fixed hipOccupancyMaxPotentialBlockSize test build issue

* Fix hipOccupancyMaxPotentialBlockSize dtest

* Add BUILD_CMD in hipOccupancyMaxPotentialBlockSize dtest

* Revert "Add BUILD_CMD in hipOccupancyMaxPotentialBlockSize dtest"

This reverts commit 0480ff56f1441fc515d2c26ce33783e303423938.

* Disable hipOccupancyMaxPotentialBlockSize dtest on NVCC

* move extern declaration of ihipGetDeviceProperties to hip_module.cpp

* Update the limiation of 32 wavefronts per CU and 800/512 SGPRs for VI/pre-VI chips to calculate the occupancy
2019-06-20 05:58:29 +05:30
Evgeny 19ca7a2a08 prof layer includes refactoring 2019-06-11 20:13:29 -05:00
Michael LIAO 52a8f2fed4 [hip] Replace implicit conversions with explicit ones. 2019-06-05 10:00:35 -04:00
Maneesh Gupta 7013f87885 Merge pull request #1140 from scchan/program_state_stage_2-rebase-20190524
migrate more program_state logic from header into shared library (phase II)
2019-06-05 16:09:01 +05:30
Maneesh Gupta 0982bb52f2 Merge pull request #1159 from mhbliao/hliao/master/explicit
[hip] Make vector type's scalar conversion explicit.
2019-06-05 12:13:56 +05:30
Michael LIAO 9bb4ecfcfe [hip] Make vector type's scalar conversion explicit. 2019-06-04 15:15:26 -04:00
Michael LIAO 498fe92734 [hip] Minor fix to silence compilation warnings.
- Add parenthese to silence repeative compilation warnings across
  projects built against hip.
2019-06-04 10:39:31 -04:00
Maneesh Gupta 0564d8af57 Merge pull request #1154 from yxsamliu/tex
Add device_builtin_texture_type attribute to texture type for hip-clang
2019-06-04 13:21:29 +05:30
Maneesh Gupta 49e254f5e8 Merge pull request #1152 from asalmanp/hip_as_b
Header change for new hip API hipExtLaunchMultiKernelMultiDevice
2019-06-04 13:21:13 +05:30
Maneesh Gupta b0d90ac2ed Merge pull request #1149 from zuhaib27/SWDEV-185448
Structured hipFloatComplex as typedef of float2, and hipDoubleComplex as typedef of double2.
2019-06-04 13:21:02 +05:30
Yaxun Sam Liu 80bcf0785a Add device_builtin_texture_type attribute to texture type for hip-clang
This is required to support texture type for hip-clang.
2019-05-31 12:07:58 -04:00
Aryan Salmanpour f34654f8aa Header change for new hip API hipExtLaunchMultiKernelMultiDevice 2019-05-30 18:04:05 -04:00
Siu Chi Chan 72e51f3ad0 fix compilation error when host compiler is clang (#1147)
* fix compilation error when host compiler is clang

* use a macro specifically for hcc && hip-clang
2019-05-29 12:34:48 +05:30
Zuhaib Khan cc0f3445bb Structured hipFloatComplex as typedef of float2, and hipDoubleComplex as typedef of double2. 2019-05-28 16:57:51 -04:00
Maneesh Gupta 21a9e05867 Header changes for cooperative groups
Change-Id: I5f3acca94275d74adc97adcb168aed9f74951189
2019-05-28 16:58:55 +05:30
Siu Chi Chan 54f94ed02f remove code_object_bundle.hpp, clean up the old Agent_global 2019-05-24 18:13:05 -04:00
Siu Chi Chan cec0926924 fix breakage due to compiling in C++17 2019-05-24 17:27:58 -04:00
Siu Chi Chan fc08f29735 replace std::vector for kernarg 2019-05-24 17:27:43 -04:00
Siu Chi Chan 00824be34c move executable_cache into program_state.cpp 2019-05-24 17:27:25 -04:00
Siu Chi Chan 80fec2b477 remove executables() from program_state 2019-05-24 17:27:01 -04:00
Siu Chi Chan 4239cfcf02 moving agent_globals_impl into hip_module 2019-05-24 16:43:38 -04:00
Maneesh Gupta 68ae4a0a3e Merge pull request #1128 from aaronenyeshi/fix-smid-func
Fix bug in __smid not setting correct size
2019-05-24 14:16:12 +05:30
Aaron Enye Shi ff74babedf Fix bug in __smid not setting correct size
The SZ field should minus by 1 since SIZE range is 1..32. Also add comments that results may vary.
2019-05-22 19:20:09 +00:00
Evgeny Mankov eede980a39 [HIP] fix typo in #1127 2019-05-22 20:48:18 +03:00
Evgeny Mankov 47bbb9c605 [HIP] fix nvcc path break in #1127 2019-05-22 20:04:45 +03:00
Evgeny Mankov 8f059b0ee9 [HIP][HIPIFY] Make hipMemcpyParam2D coherent with cuMemcpy2D
+ Makes hip_Memcpy2D struct compatible with CUDA_MEMCPY2D struct
+ Add hipMemcpyParam2D support in nvcc fallback path
+ Update hipify-clang, tests and docs accordingly
2019-05-22 18:31:39 +03:00
Alex Voicu ccfb764a59 Add HIPRTC, glorious ersatz for NVRTC (#1097)
* Add ersatz for NVRTC.

* Fix extraneous paren and use correct namespace.

* Use lowerCamelCase (yuck, yuck) consistently.

* Link against FS when building hiprtc lib.

* Correctly mark Manipulators. Fix dual compile.

* Add unit tests. Extend HIT to accept linker options.

* Make sure the HIPRTC library is installed.

* Better logging. Try to auto-detect the target.

* Stop specifying the target explicitly.

* Add missing flavour of `hipModuleLaunchKernel`.

* Program was already destroyed.

* Don't use `--genco`. Fix mangled name trimming.

* Fix HIPRTC breakage due to upstream noise.

* [dtests] Replace RUN -> TEST in hiprtc tests

Change-Id: Ie499e92dfe4e5c94634b1c2b76cf52d241bcfea3

* [hit] Set HIP_PATH to HIP_ROOT_DIR for all tests

Change-Id: Ib0ad1f99bc71c03e363e055dd508a7a4a210680a
2019-05-16 18:28:54 +05:30
Wen-Heng (Jack) Chung 9b9257f9b0 Revert "HACK for SWDEV-173477" (#1004)
* Revert "HACK for SWDEV-173477"

This reverts commit d941f19399.
2019-05-13 14:42:05 +05:30