Commit graph

21 Commits

Autor SHA1 Nachricht Datum
Evgeny Mankov 9ae76609a1 [HIP] Unify hipError_t (Step 2)
Step 2. Make a few hipError codes deprecated
Update hipify-clang, hipify-perl, docs and samples accordingly
2019-12-22 02:05:31 +03:00
Evgeny Mankov e2191e23e6 [HIPIFY][#1569] Fix 2019-10-22 11:08:37 +03:00
Evgeny Mankov e3319acbf8 [HIPIFY][test] Update allocators.cu test 2019-10-10 17:20:41 +03:00
Evgeny Mankov fee2f0fa78 [HIPIFY] Return to wrapping with HIP_KERNEL_NAME(...) macro of a template instantiation kernel launch
[REASON]

1. hip-clang is fine with the templated kernel launch, brackets are unneeded: HIP_KERNEL_NAME(...) __VA_ARGS__
2. HCC is not, thus: HIP_KERNEL_NAME(...) (__VA_ARGS__)

[TODO] Clean-up entirely kernel name wrapping when HCC is finally obsolete.

+ Update perl generation, hipify-perl, and affected tests accordingly.
2019-10-02 16:01:07 +03:00
Evgeny Mankov 1db7c1d694 [HIPIFY][#1439] Add reinterpret_cast to args of some functions
+ Perl part of [#1458]
+ Affected functions: hipFuncSetCacheConfig, hipFuncGetAttributes
+ Implement function generateHostFunctions() in hipify-clang for that purposes
+ Update hipify-perl accordingly
2019-09-25 18:53:17 +03:00
Evgeny Mankov cf467b4550 [HIPIFY][#1439] Add reinterpret_cast to args of some functions
+ Affected functions: hipFuncSetCacheConfig, hipFuncGetAttributes
+ Add a corresponding Matcher cudaReinterpretCastArgFuncCall
+ Add reinterpret_cast.cu test

TODO: Do the same for hipify-perl
2019-09-24 09:44:17 +03:00
Evgeny Mankov 4e3845be83 [HIPIFY][#1400] Fix Template Instantiation kernel launch (clang & perl)
+ Enclose template instantiation kernel calls into round brackets, leave regular kernel names unchanged (hipify-perl doesn't handle cases with macros).
+ Fix corresponding tests.

PS. hipify-perl couldn't handle correctly the following cases due to macros expansion disability, thus hipify-clang should be used instead:

#define KERNEL_NAME_MACRO axpy<float>
#define KERNEL_CALL_MACRO axpy<float><<<1, 2>>>
#define KERNEL_ARG_LIST_MACRO a, x, y

// CUDA:
KERNEL_NAME_MACRO<<<1, 2>>>(KERNEL_ARG_LIST_MACRO);
KERNEL_CALL_MACRO(KERNEL_ARG_LIST_MACRO);

// hipify-perl:
hipLaunchKernelGGL(KERNEL_NAME_MACRO, dim3(1), dim3(2), 0, 0, KERNEL_ARG_LIST_MACRO);
KERNEL_CALL_MACRO(KERNEL_ARG_LIST_MACRO);

// hipify-clang:
hipLaunchKernelGGL((KERNEL_NAME_MACRO), dim3(1), dim3(2), 0, 0, KERNEL_ARG_LIST_MACRO);
hipLaunchKernelGGL((axpy<float>), dim3(1), dim3(2), 0, 0, KERNEL_ARG_LIST_MACRO);
2019-09-10 15:59:06 +03:00
Evgeny Mankov e26ec02a99 [HIPIFY][perl][#259] Fix
empty<<<1, 2>>> ( );     >>  hipLaunchKernelGGL(empty, dim3(1), dim3(2), 0, 0);
empty<<<1, 2, 0>>>();    >>
empty<<<1, 2, 0, 0>>>(); >>

instead of erroneous:    >> hipLaunchKernelGGL((empty), dim3(1), dim3(2), 0, 0, );
2019-09-03 16:44:20 +03:00
Evgeny Mankov 608331250e [HIPIFY][tests] Add occupancy test 2019-09-02 17:54:06 +03:00
Evgeny Mankov 9166b1554e [HIPIFY] Add cudaMallocManaged -> hipMallocManaged
+ Add mapping for corresponding data types
+ Add a test
+ Update docs
2019-08-13 17:56:06 +03:00
Evgeny Mankov f0832fd968 [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
Evgeny Mankov 7db992cf44 [HIPIFY] Add device texture fetch functions support
+ Add a corresponding reverse engineered sample tex2dKernel with texture template
2019-05-23 12:47:08 +03:00
Evgeny Mankov 5c205b75f5 [HIPIFY][tests] Add reverse engineered HIP sample "peer2peer"
+ Fix typo in sample "stream"
2019-05-20 17:16:11 +03:00
Evgeny Mankov 7cc12df514 [HIPIFY][tests] Add reverse engineered HIP sample "stream"
+ Add additional checks for extern __shared__ due to [#1109]
2019-05-15 20:17:03 +03:00
Evgeny Mankov bf65120156 [HIPIFY][fix][#1109] Do not preserve extern __shared__ for IncompleteArrayType
+ Update tests accordingly
2019-05-15 20:05:56 +03:00
Evgeny Mankov 5e49c25faa [HIPIFY][tests] Add reverse engineered HIP sample Profiler
+ Add missing cuda_profiler_api.h to hip/hip_profile.h transformation.
NOTE: HIP Profiler API is under development. This is NOT WORKING example.
TODO: Find out a way to generate HIP_SCOPED_MARKER, HIP_BEGIN_MARKER, HIP_END_MARKER, declared in hip/hip_profile.h in particular place (signatures are to obtain).
2019-05-14 16:43:44 +03:00
Evgeny Mankov 9860dac7fa [HIPIFY][tests] Add reverse engineered HIP sample hipEvent 2019-05-13 22:12:43 +03:00
emankov cdc76af186 [HIPIFY][tests] Add reverse engineered HIP sample MatrixTranspose 2019-05-13 19:37:18 +03:00
Evgeny Mankov c4268855e3 [HIPIFY] Introduce CUDA installation path option '-cuda-path'
Repeats clang's '--cuda-path' option.

[Reason]
In case of absence of any other clang's options setting '-cuda-path' allows not to specify separator '--' before clang's '--cuda-path'.

+ Tests and scripts are updated accordingly.
2019-01-09 20:18:36 +03:00
emankov 410f8f0b92 [HIPIFY][tests] Reduce cudaRegister.cu test 2019-01-03 17:17:21 +03:00
Evgeny Mankov b76cf94562 [HIPIFY][tests] Reorganization 2018-12-04 20:47:34 +03:00