Commit Graph

1023 Commits

Author SHA1 Message Date
Michael LIAO ef6ca030cb [hip] Replace implicit conversions with explicit ones.
[ROCm/clr commit: 4d5cf71c48]
2019-06-05 10:00:35 -04:00
Maneesh Gupta b4fb2b0ab4 Merge pull request #1140 from scchan/program_state_stage_2-rebase-20190524
migrate more program_state logic from header into shared library (phase II)

[ROCm/clr commit: 1d5d923d36]
2019-06-05 16:09:01 +05:30
Maneesh Gupta f616eed25e Merge pull request #1159 from mhbliao/hliao/master/explicit
[hip] Make vector type's scalar conversion explicit.

[ROCm/clr commit: b2c6817df0]
2019-06-05 12:13:56 +05:30
Michael LIAO 0067b36618 [hip] Make vector type's scalar conversion explicit.
[ROCm/clr commit: 62d70cb2c8]
2019-06-04 15:15:26 -04:00
Michael LIAO a56332edc7 [hip] Minor fix to silence compilation warnings.
- Add parenthese to silence repeative compilation warnings across
  projects built against hip.


[ROCm/clr commit: 2655003fde]
2019-06-04 10:39:31 -04:00
Maneesh Gupta 0853a2dcfb Merge pull request #1154 from yxsamliu/tex
Add device_builtin_texture_type attribute to texture type for hip-clang

[ROCm/clr commit: 8a859f8020]
2019-06-04 13:21:29 +05:30
Maneesh Gupta 601bd522af Merge pull request #1152 from asalmanp/hip_as_b
Header change for new hip API hipExtLaunchMultiKernelMultiDevice

[ROCm/clr commit: ef87f7eaef]
2019-06-04 13:21:13 +05:30
Maneesh Gupta 5ca1fc546e Merge pull request #1149 from zuhaib27/SWDEV-185448
Structured hipFloatComplex as typedef of float2, and hipDoubleComplex as typedef of double2.

[ROCm/clr commit: 98aa6cf895]
2019-06-04 13:21:02 +05:30
Yaxun Sam Liu 7a4dfd23d5 Add device_builtin_texture_type attribute to texture type for hip-clang
This is required to support texture type for hip-clang.


[ROCm/clr commit: 7c20081f8d]
2019-05-31 12:07:58 -04:00
Aryan Salmanpour aab9b5a13b Header change for new hip API hipExtLaunchMultiKernelMultiDevice
[ROCm/clr commit: d8e94fd5b5]
2019-05-30 18:04:05 -04:00
Siu Chi Chan 339a048377 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


[ROCm/clr commit: b2ffd6afc2]
2019-05-29 12:34:48 +05:30
Zuhaib Khan d030730c70 Structured hipFloatComplex as typedef of float2, and hipDoubleComplex as typedef of double2.
[ROCm/clr commit: 6aa704e7b9]
2019-05-28 16:57:51 -04:00
Maneesh Gupta b70b2c4e9d Header changes for cooperative groups
Change-Id: I5f3acca94275d74adc97adcb168aed9f74951189


[ROCm/clr commit: 4af81134ba]
2019-05-28 16:58:55 +05:30
Siu Chi Chan cf2e2c5077 remove code_object_bundle.hpp, clean up the old Agent_global
[ROCm/clr commit: d40a4627c3]
2019-05-24 18:13:05 -04:00
Siu Chi Chan f2cdf536f6 fix breakage due to compiling in C++17
[ROCm/clr commit: 2bc9455928]
2019-05-24 17:27:58 -04:00
Siu Chi Chan bafd29662c replace std::vector for kernarg
[ROCm/clr commit: 1fb9ab2d44]
2019-05-24 17:27:43 -04:00
Siu Chi Chan 304a1e2dbe move executable_cache into program_state.cpp
[ROCm/clr commit: 1a2d332e76]
2019-05-24 17:27:25 -04:00
Siu Chi Chan 305eb4239e remove executables() from program_state
[ROCm/clr commit: e2c0122892]
2019-05-24 17:27:01 -04:00
Siu Chi Chan 0cae3e06c1 moving agent_globals_impl into hip_module
[ROCm/clr commit: 6852be819f]
2019-05-24 16:43:38 -04:00
Maneesh Gupta d1bc228f25 Merge pull request #1128 from aaronenyeshi/fix-smid-func
Fix bug in __smid not setting correct size

[ROCm/clr commit: f03a8cc1b0]
2019-05-24 14:16:12 +05:30
Aaron Enye Shi 2fd8de1749 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.


[ROCm/clr commit: 2b11a8bf0c]
2019-05-22 19:20:09 +00:00
Evgeny Mankov 3afaf0d2de [HIP] fix typo in #1127
[ROCm/clr commit: 49b9df7a9e]
2019-05-22 20:48:18 +03:00
Evgeny Mankov a0e1887ff3 [HIP] fix nvcc path break in #1127
[ROCm/clr commit: 6806ab6745]
2019-05-22 20:04:45 +03:00
Evgeny Mankov 204043c6e0 [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


[ROCm/clr commit: 9cb3e9aa5e]
2019-05-22 18:31:39 +03:00
Alex Voicu a4a3132c64 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


[ROCm/clr commit: a538eb705a]
2019-05-16 18:28:54 +05:30
Wen-Heng (Jack) Chung e92ffd2261 Revert "HACK for SWDEV-173477" (#1004)
* Revert "HACK for SWDEV-173477"

This reverts commit 86379d694f.

[ROCm/clr commit: a4db991cbf]
2019-05-13 14:42:05 +05:30
Rahul Garg d44e800a17 Add fine grained host memory lock support (#1095)
* Add fine grained host memory lock support

* Fix default flag check


[ROCm/clr commit: e1f3dc0c80]
2019-05-13 11:48:26 +05:30
Siu Chi Chan 76f535b4ce migrate program_state logic from header into shared library (phase I) (#1077)
* Revert "Revert "Use COMgr to read Kernel Args Metadata (#1006)""

This reverts commit f8d108a815.

* Revert "Use COMgr to read Kernel Args Metadata (#1006)"

This reverts commit 10048a5631.

* Revert "improve program state commentary"

This reverts commit 5233d41c6c.

* Revert "load program state once per agent"

This reverts commit 9cee2c5311.

* start moving function_names() into the hip shared lib

* start moving code_object_blobs to a new "state" object

* Consolidate various program state related static objects into a
single program_state object

* minor clean up

* move more stuffs from functional_grid_launch into program_state

* debug make_kernarg

* moving lookup for kernargs size_align into program_state

* clean up old code for kernarg size and alignment

* update hip_module to use newer api in program_state

* Create public member functions for program_state

* move most program state functions into shared library

* Pass the data buffer size to load_executable
Otherwise, it can't figure what the data size is
just from the char* (since the data is not really a string)

* turning free functions in program state into members of program_state_impl

* change the free function globals() into a member of program_state_impl

* replace the static mutex used for populating globals

* moving associate_code_object_symbols_with_host_allocation into
program_state_impl

* move load_code_object_and_freeze_executable into program_state_impl

* moving executables and functions_names into program_state_impl

* moving kernels() into program_state_impl

* moving functions() into program_state_impl

* move get_kernargs into program_state_impl

* moving kernel_descriptor into program_state_impl

* moving kernargs_size_align calculation into program_state_impl

* Changing the handle to program_state_impl to a pointer

* moving program_state_impl into a separate inline source file

* fixing/cleaning up some header file includes

* moving member function for kernargs_size_align into program_state.cpp

* moving Kernel_descriptor into program_state.inl

* add a new class to manage agent globals

* moving all agent globals processing functions into agent_globals_impl

* load program state once per agent

re-merging PR991 against other program state changes

* fix per-agent program state member initialization

* cache executables based on elf name, isa, and agent.

This avoids program state reloading executables after a shared library is dlopened.

re-merging PR1057 against other program state changes

* protect executables cache by a global mutex

* return ref to executables cache

* adapt PR#981 Make hipModuleGetGlobal be in HIP runtime


[ROCm/clr commit: 05a1b696da]
2019-05-12 19:24:03 +05:30
Maneesh Gupta 54f932e569 Merge pull request #1084 from mhbliao/hliao/master/api_ext
[hip] Add API `hipExtModuleLaunchKernel` in HIP runtime

[ROCm/clr commit: e78a09c041]
2019-05-09 18:26:31 +05:30
Maneesh Gupta 301a9292ff Merge pull request #1082 from gargrahul/fix_hipmemcpy_symbol_nvcc
Fix symbol address issue on NVCC path

[ROCm/clr commit: c6cf2a9e26]
2019-05-07 16:17:01 +05:30
Maneesh Gupta 30c7ed3e28 Merge pull request #1081 from mangupta/swdev-181624
Implement hipExtGetLinkTypeAndHopCount for ROCm devices

[ROCm/clr commit: c6c5e4cee8]
2019-05-07 16:15:41 +05:30
Maneesh Gupta 46d0385435 Merge pull request #1068 from mhbliao/hliao/master/dev_vec_func
[devfunc] Add necessary `__device__` and `__host__` attributes.

[ROCm/clr commit: 11972049c6]
2019-05-07 16:01:48 +05:30
Michael LIAO d94d566410 [hip] Add API hipExtModuleLaunchKernel in HIP runtime
[ROCm/clr commit: de768c22ae]
2019-05-06 21:20:28 -04:00
Rahul Garg bf3bafb9f5 Fix symbol address issue on NVCC path
[ROCm/clr commit: 6cbc70d238]
2019-05-07 03:59:43 +05:30
Maneesh Gupta f657eba4a5 Implement hipExtGetLinkTypeAndHopCount for ROCm devices
Change-Id: Ie5bb4f640ac6d189c7fceeab22627a7494fd10bd


[ROCm/clr commit: 2f43f110d9]
2019-05-06 15:54:31 +05:30
Maneesh Gupta 13b13c3493 Merge pull request #1062 from mhbliao/hliao/master/icmp
[hip] Re-implement ballot using AMDGCN builtins

[ROCm/clr commit: 2eafa5dcf9]
2019-05-03 17:48:19 +05:30
Michael LIAO 08fa23f774 [devfunc] Add necessary __device__ and __host__ attributes.
- Minor clean up to keep consistent function declaration.


[ROCm/clr commit: a9f90713f3]
2019-05-01 22:26:35 -04:00
Michael LIAO eb43303d0b [Device Function] Fix implementation of __bitinsert_u64
- It's a common mistake by assuming 1 << shamt would be promoted to
  64-bit, if shamt is a 64-bit integer. That's not the case. Replace
  that left shift to a 64-bit one to ensure it won't fall into undefined
  behavior.
- Fix the host-side implementation as well for device function testing.


[ROCm/clr commit: 2380eb8ecc]
2019-04-30 08:59:13 -04:00
Michael LIAO e8de293fc5 [devfunc] Re-implement ballot using AMDGCN builtins
- As the signature of `amdgcn.icmp` is changed for next-gen chip, using
  clang builtins is portable way to hide that details.


[ROCm/clr commit: a7a4d80f54]
2019-04-29 17:21:25 -04:00
Aaron Enye Shi f8d108a815 Revert "Use COMgr to read Kernel Args Metadata (#1006)"
This reverts commit 10048a5631.


[ROCm/clr commit: 235c6877c8]
2019-04-26 16:04:56 -04:00
Maneesh Gupta dad6abcd7a Merge pull request #1043 from mhbliao/hliao/master/fp16
[hip] Fix including of hip_fp16.h

[ROCm/clr commit: 7f81c72f1c]
2019-04-24 16:50:46 +05:30
Maneesh Gupta f8f49d57dc Merge pull request #1042 from mhbliao/hliao/master/ldg
[hip] Fix use of `__HIP_CLANG_ONLY__` in `hip_ldg.h`.

[ROCm/clr commit: 63ab2ea945]
2019-04-24 16:50:37 +05:30
Maneesh Gupta b41d81d74e Merge pull request #1040 from eshcherb/roctracer-hip-frontend-190422
hip_prof_api.h include under __cplusplus

[ROCm/clr commit: 54cdeabe6e]
2019-04-24 16:50:27 +05:30
Maneesh Gupta a709855e9d Merge pull request #1039 from gargrahul/fix_ptrgetattr_nvcc
Fix hipPointerGetAttributes for NVCC

[ROCm/clr commit: 7edb43bc83]
2019-04-24 16:50:18 +05:30
Rahul Garg d69edbbb7f Add hipMallocManaged default functional support (#1036)
* Add hipMallocManaged default functional support

* Fix build error

* Add dtest


[ROCm/clr commit: 94769fc8dd]
2019-04-24 16:50:03 +05:30
Michael LIAO 59e6127969 [hip] Fix including of hip_fp16.h
- Separate the definition of `__HCC_OR_HIP_CLANG__`, `__HCC_ONLY__`, and
  `__HIP_CLANG_ONLY__` into hip_common.h so that it could be included in
  hip_fp16.h, which may be included separately in app.


[ROCm/clr commit: d086dbd0e5]
2019-04-23 09:16:00 -04:00
Michael LIAO 619050ae96 [hip] Fix use of __HIP_CLANG_ONLY__ in hip_ldg.h.
- Check its value instead of whether it's defined or not.


[ROCm/clr commit: ca6a5c07eb]
2019-04-22 23:22:32 -04:00
Evgeny 79df39e3c3 hip_prof_api.h include under __cplusplus
[ROCm/clr commit: 165c42483b]
2019-04-22 21:14:18 -05:00
Rahul Garg 0198199780 Fix hipPointerGetAttributes for NVCC
[ROCm/clr commit: c0e0f0b7fd]
2019-04-23 03:22:25 +05:30
Konstantin Pyzhov a525cc8f47 Fix for __popcll() device function implementation.
[ROCm/clr commit: f6fbf8751d]
2019-04-19 08:53:22 -04:00