Wykres commitów

147 Commity

Autor SHA1 Wiadomość Data
Maneesh Gupta 0ea6697192 Merge branch 'amd-master' into amd-master-next
Change-Id: I3094c15008093f2072bcd38aca4ea90aeae2d97b
2020-04-09 06:31:00 -04:00
Michael LIAO a14695d4eb [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
2020-03-31 00:26:53 -04:00
Sameer Sahasrabuddhe c014c733c6 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
2020-03-28 00:18:21 -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
Sameer Sahasrabuddhe d48738856c SWDEV-204784: separate printf declaration for vdi/clang
There are now two implementations of printf in HIP:

1. The implemenation for HCC is controlled by the HC_FEATURE_PRINTF
   macro, and it works only with the HCC compiler used in combination
   with the HCC runtime.

2. The implementation for hip-clang requires the VDI runtime, and is
   always enabled with that combination.

Change-Id: Ibaeda7900ffe2ce602ca0094aafed0f1147ac2b6
2020-03-16 04:00:39 -04:00
Maneesh Gupta c091635f5a Revert PRs that break ROCm builds (#1781)
Fixes SWDEV-218626 and SWDEV-218629

Changes:
- Revert "`static inline` in a header, just like excess sugar in a diet, causes bloat (#1692)"
   This reverts commit cf526f5484.
- Revert "Fix rocFFT build failure (#1777)"
   This reverts commit cbd8f604ec.
2020-01-08 15:11:58 +05:30
Yaxun (Sam) Liu 35988729f3 Add macro __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__ (#1761)
This is to allow force enable cuda wrapper for OpenMP for flexibility
2019-12-23 19:24:54 +05:30
Alex Voicu cf526f5484 static inline in a header, just like excess sugar in a diet, causes bloat (#1692) 2019-12-23 19:09:38 +05:30
Michael LIAO d28ad401c9 Remove redundant declarations.
- The revised `hip/hip_ext.h` have that declarations.
2019-11-07 10:11:22 -05:00
Alex Voicu 33c40ee510 Optimise the gridDim.n * blockDim.m idiom (#1468) 2019-09-30 10:39:23 +05:30
Yaxun (Sam) Liu e72346dc24 Do not include cuda wappers for OMP for hip-clang (#1382) 2019-09-03 05:13:59 +00:00
Sarbojit2019 4d270775c8 Removed hipLaunchKernel macro got missed in Merge (#1374) 2019-09-03 05:13:07 +00:00
Yaxun (Sam) Liu 68a59ddbe8 Fix missing decl for hip-clang
Add back decl for hipHccModuleLaunchKernel and hipExtModuleLaunchKernel for HIP/VDI only
2019-08-19 18:27:13 -04:00
Rahul Garg 7f9de881cb Fix undefined identifier issue for hipExtModuleLaunchKernel 2019-08-14 16:46:32 -04:00
Jatin Chaudhary 46553d719a Adding fix to initialize a dim3 variable from blockIdx. (#1212)
Adding gitignore for ctags
2019-07-11 09:16:03 +00:00
Aryan Salmanpour d6ad690cb6 [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
Michael LIAO de768c22ae [hip] Add API hipExtModuleLaunchKernel in HIP runtime 2019-05-06 21:20:28 -04:00
Michael LIAO d086dbd0e5 [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.
2019-04-23 09:16:00 -04:00
Yaxun (Sam) Liu e200ece4da hip-clang: Add __align__
CUDA has __align__. Define eqivalent for hip-clang.
2019-04-10 14:17:18 -04:00
Michael LIAO 13655df76e [hip] Fix typo in macro hipLaunchKernel 2019-03-25 12:06:46 -04:00
Maneesh Gupta e7453483e2 Merge pull request #965 from mhbliao/hliao/master/immarg
[Device Function] Support immediate argument.
2019-03-19 18:41:31 +05:30
Michael LIAO b74b4500c4 [Device Function] Support immediate argument.
- `immarg`, immediate argument, is enabled on all AMDGPU intrinsics.
  Revise device functions using these intrinsics with immediate
  arguments.
2019-03-15 12:38:04 -04:00
Michael Liao f46d6cf249 [hip] Re-implement hipLaunchKernelGGL as macros. 2019-03-08 09:10:51 -05:00
Aaron Enye Shi ba43d1ce1b Remove libhip_device.a static library
Move remaining function definitions from device_util.cpp to hip_runtime.h header. We can now remove the static library completely as device_util.cpp was the last file part of target hip::hip_device .
2018-10-25 19:52:35 +00:00
Alex Voicu 217fe345ca Minimal should mean minimal. 2018-10-11 00:21:41 +01:00
Alex Voicu 9ec697c620 Re-sync with upstream. 2018-10-10 11:43:49 +01:00
Yaxun Sam Liu 781f0d97b5 Fix build failure of hipTestHalf and hipTestIncludeMath for hip-clang 2018-09-18 21:00:15 -04:00
Yaxun Sam Liu 0471ce2da6 Fix hipLaunchKernelGGL for hip-clang
Do not decay function pointer type of the kernel argument passed to hipLaunchKernelGGL
and hipLaunchKernel, otherwise some type information is lost which may cause
type inference failure for the template.

This issue caused compilation error of FeatureLPPooling in Caffe2/PyTorch and this patch
fixes that.
2018-09-17 11:20:41 -04:00
Alex Voicu 26c26f63cb Align with HC Next. 2018-09-17 11:50:29 +03:00
Maneesh Gupta b23609ad42 Merge branch 'master' into support-malloc 2018-09-17 10:17:25 +05:30
Maneesh Gupta 9f553c0b40 Merge pull request #650 from ROCm-Developer-Tools/hip-clang-new
Support placement new in hip-clang
2018-09-15 11:21:01 +05:30
Maneesh Gupta 17d12bbfef Merge pull request #665 from aaronenyeshi/fix-min-funcs
Use templates for min to prevent ambiguity
2018-09-14 13:21:38 +05:30
Maneesh Gupta 2081cc9bd4 Merge pull request #661 from yxsamliu/add-empty-printf
Add empty printf for hip-clang
2018-09-13 07:54:03 +05:30
Aaron Enye Shi 5f0838300e Use templates for min to prevent ambiguity 2018-09-11 18:21:54 +00:00
Yaxun Sam Liu 6f2201475d Use template for hipLaunchKernelGGL for hip-clang 2018-09-07 16:20:00 -04:00
Yaxun Sam Liu be4a8f6859 Add empty printf for hip-clang 2018-09-07 16:15:08 -04:00
Yaxun (Sam) Liu cc0fbc0e2e Support placement new in hip-clang 2018-09-04 16:53:46 -04:00
Maneesh Gupta a5f9ae90d2 Merge pull request #613 from ROCm-Developer-Tools/hcc-compat
Add HCC compatibility mode for hip-clang
2018-08-10 09:32:31 +05:30
Maneesh Gupta b485957858 Merge pull request #603 from b-sumner/master
Drop use of NVCC_COMPAT since it is always set
2018-08-03 08:45:18 +05:30
Maneesh Gupta 0354367a07 Merge pull request #610 from aaronenyeshi/enable-hcc-module-launch
Enable hipHccModuleLaunchKernel on HIP-Clang
2018-08-02 08:08:46 +05:30
Yaxun Sam Liu 9c28758fa0 Add hc_get_workitem_absolute_id for hip-clang HCC compatibility mode 2018-08-01 15:14:38 -04:00
Yaxun Sam Liu b91691966f Add HCC compatibility mode 2018-08-01 15:14:38 -04:00
Aaron Enye Shi 0ad36dfa58 Enable hipHccModuleLaunchKernel on HIP-Clang 2018-07-31 20:38:20 +00:00
Yaxun Sam Liu d17dc21c9c Support malloc/free for hip-clang 2018-07-27 16:24:51 -04:00
Brian Sumner 8d80abb30f Drop use of NVCC_COMPAT since it is always set 2018-07-27 09:29:29 -07:00
Aaron Enye Shi bf6fdca1e8 Enable HIP texture functions with HIP vector types
Include texture_functions.h on HIP-Clang path. Use HIP vector types instead of hc::short_vectors for texture_functions, and remove HCC headers.
2018-07-25 19:52:19 +00:00
Aaron Enye Shi 98fca64220 [HIPClang] Switch ordering on math functions
Header math_functions.h should be included after including __clang_cuda_math_forward_declares.h to avoid warning: attribute declaration must precede definition.
2018-07-20 13:25:04 -04:00
Yaxun (Sam) Liu a8dc1257df Support std::complex for hip-clang 2018-07-18 00:08:04 -04:00
Maneesh Gupta 1fc3ef3cd4 Merge branch 'master' into move-memcpy 2018-07-17 10:51:42 +05:30
Yaxun (Sam) Liu 903dfa94eb Move __hip_hc_memcpy and __hip_hc_memset from device_utils.cpp to device_functions.h as inline functions 2018-07-10 18:12:41 -04:00