diff --git a/bin/hipify-perl b/bin/hipify-perl index 02a540b133..2887402b4e 100755 --- a/bin/hipify-perl +++ b/bin/hipify-perl @@ -1738,12 +1738,6 @@ while (@ARGV) { if ((/(\bcuda[A-Z]\w+)/) or (/<<<.*>>>/)) { # flag any remaining code that look like cuda API calls, may want to add these to hipify $tag = (defined $1) ? $1 : "Launch"; - } elsif (/(\bcublas[A-Z]\w+)/) { - $warningsCublas++; - $tag = $1; - } elsif (/(\bcurand[A-Z]\w+)/) { - $warningsCurand++; - $tag = $1; } if (defined $tag) { $warnings++; @@ -1763,7 +1757,7 @@ while (@ARGV) { $ft{'kernel_func'} += countSupportedDeviceFunctions(); } - $ft{'memory'} += transformSymbolFunctions(); + transformHostFunctions(); # Print it! # TODO - would like to move this code outside loop but it uses $_ which contains the whole file. @@ -1819,7 +1813,7 @@ if ($count_conversions) { } } -sub transformSymbolFunctions +sub transformHostFunctions { my $m = 0; foreach $func ( @@ -1838,6 +1832,18 @@ sub transformSymbolFunctions { $m += s/(?\($2\),/g + } + foreach $func ( + "hipFuncGetAttributes" + ) + { + $m += s/(?\($3\)$4/g; + } return $m; } @@ -1868,6 +1874,13 @@ sub countSupportedDeviceFunctions "__fdividef", "__ffs", "__ffsll", + "__float22half2_rn", + "__float2half", + "__float2half2_rn", + "__float2half_rd", + "__float2half_rn", + "__float2half_ru", + "__float2half_rz", "__float2int_rd", "__float2int_rn", "__float2int_ru", @@ -1886,6 +1899,7 @@ sub countSupportedDeviceFunctions "__float2ull_rz", "__float_as_int", "__float_as_uint", + "__floats2half2_rn", "__fmaf_rd", "__fmaf_rn", "__fmaf_ru", @@ -1907,19 +1921,127 @@ sub countSupportedDeviceFunctions "__fsub_rn", "__fsub_ru", "__fsub_rz", + "__h2div", "__hadd", + "__hadd2", + "__hadd2_sat", + "__hadd_sat", + "__half22float2", + "__half2float", + "__half2half2", + "__half2int_rd", + "__half2int_rn", + "__half2int_ru", + "__half2int_rz", + "__half2ll_rd", + "__half2ll_rn", + "__half2ll_ru", + "__half2ll_rz", + "__half2short_rd", + "__half2short_rn", + "__half2short_ru", + "__half2short_rz", + "__half2uint_rd", + "__half2uint_rn", + "__half2uint_ru", + "__half2uint_rz", + "__half2ull_rd", + "__half2ull_rn", + "__half2ull_ru", + "__half2ull_rz", + "__half2ushort_rd", + "__half2ushort_rn", + "__half2ushort_ru", + "__half2ushort_rz", + "__half_as_short", + "__half_as_ushort", + "__halves2half2", + "__hbeq2", + "__hbequ2", + "__hbge2", + "__hbgeu2", + "__hbgt2", + "__hbgtu2", + "__hble2", + "__hbleu2", + "__hblt2", + "__hbltu2", + "__hbne2", + "__hbneu2", + "__hdiv", + "__heq", + "__heq2", + "__hequ", + "__hequ2", + "__hfma", + "__hfma2", + "__hfma2_sat", + "__hfma_sat", + "__hge", + "__hge2", + "__hgeu", + "__hgeu2", + "__hgt", + "__hgt2", + "__hgtu", + "__hgtu2", + "__high2float", + "__high2half", + "__high2half2", + "__highs2half2", + "__hisinf", + "__hisnan", + "__hisnan2", + "__hle", + "__hle2", + "__hleu", + "__hleu2", + "__hlt", + "__hlt2", + "__hltu", + "__hltu2", + "__hmul", + "__hmul2", + "__hmul2_sat", + "__hmul_sat", + "__hne", + "__hne2", + "__hneg", + "__hneg2", + "__hneu", + "__hneu2", + "__hsub", + "__hsub2", + "__hsub2_sat", + "__hsub_sat", "__int2float_rd", "__int2float_rn", "__int2float_ru", "__int2float_rz", + "__int2half_rn", + "__int2half_ru", + "__int2half_rz", "__int_as_float", + "__ldca", + "__ldcg", + "__ldcs", + "__ldg", "__ll2float_rd", "__ll2float_rn", "__ll2float_ru", "__ll2float_rz", + "__ll2half_rd", + "__ll2half_rn", + "__ll2half_ru", + "__ll2half_rz", "__log10f", "__log2f", "__logf", + "__low2float", + "__low2half", + "__low2half2", + "__lowhigh2highlow", + "__lows2half2", "__mul24", "__mul64hi", "__mulhi", @@ -1929,6 +2051,11 @@ sub countSupportedDeviceFunctions "__rhadd", "__sad", "__saturatef", + "__short2half_rd", + "__short2half_rn", + "__short2half_ru", + "__short2half_rz", + "__short_as_half", "__sincosf", "__sinf", "__syncthreads", @@ -1940,16 +2067,29 @@ sub countSupportedDeviceFunctions "__uint2float_rn", "__uint2float_ru", "__uint2float_rz", + "__uint2half_rd", + "__uint2half_rn", + "__uint2half_ru", + "__uint2half_rz", "__uint_as_float", "__ull2float_rd", "__ull2float_rn", "__ull2float_ru", "__ull2float_rz", + "__ull2half_rd", + "__ull2half_rn", + "__ull2half_ru", + "__ull2half_rz", "__umul24", "__umul64hi", "__umulhi", "__urhadd", "__usad", + "__ushort2half_rd", + "__ushort2half_rn", + "__ushort2half_ru", + "__ushort2half_rz", + "__ushort_as_half", "abs", "acos", "acosf", @@ -2017,6 +2157,36 @@ sub countSupportedDeviceFunctions "fmodf", "frexp", "frexpf", + "h2ceil", + "h2cos", + "h2exp", + "h2exp10", + "h2exp2", + "h2floor", + "h2log", + "h2log10", + "h2log2", + "h2rcp", + "h2rint", + "h2rsqrt", + "h2sin", + "h2sqrt", + "h2trunc", + "hceil", + "hcos", + "hexp", + "hexp10", + "hexp2", + "hfloor", + "hlog", + "hlog10", + "hlog2", + "hrcp", + "hrint", + "hrsqrt", + "hsin", + "hsqrt", + "htrunc", "hypot", "hypotf", "ilogb", @@ -2251,6 +2421,7 @@ sub warnUnsupportedDeviceFunctions "__vsubus4", "_fdsign", "_ldsign", + "atomicAdd", "float2int", "float_as_int", "float_as_uint", diff --git a/hipify-clang/README.md b/hipify-clang/README.md index eb64d94367..4631b0de40 100644 --- a/hipify-clang/README.md +++ b/hipify-clang/README.md @@ -109,7 +109,7 @@ To run it: cmake \ -DCMAKE_INSTALL_PREFIX=../dist \ -DLLVM_SOURCE_DIR=../llvm \ - -DLLVM_TARGETS_TO_BUILD="NVPTX" \ + -DLLVM_TARGETS_TO_BUILD="X86;NVPTX" \ -DCMAKE_BUILD_TYPE=Release \ ../llvm make -j install @@ -121,7 +121,7 @@ To run it: -G "Visual Studio 16 2019 Win64" \ -DCMAKE_INSTALL_PREFIX=../dist \ -DLLVM_SOURCE_DIR=../llvm \ - -DLLVM_TARGETS_TO_BUILD="NVPTX" \ + -DLLVM_TARGETS_TO_BUILD="X86;NVPTX" \ -DCMAKE_BUILD_TYPE=Release \ -Thost=x64 \ ../llvm @@ -253,68 +253,70 @@ Linux 5.2.0 - Platform OS 64 - hipify-clang binary bitness 64 - python 2.7.12 binary bitness ======================================== --- Testing: 59 tests, 12 threads -- -PASS: hipify :: unit_tests/headers/headers_test_01.cu (1 of 59) -PASS: hipify :: unit_tests/headers/headers_test_05.cu (2 of 59) -PASS: hipify :: unit_tests/headers/headers_test_03.cu (3 of 59) -PASS: hipify :: unit_tests/headers/headers_test_10.cu (4 of 59) -PASS: hipify :: unit_tests/headers/headers_test_11.cu (5 of 59) -PASS: hipify :: unit_tests/headers/headers_test_02.cu (6 of 59) -PASS: hipify :: unit_tests/headers/headers_test_06.cu (7 of 59) -PASS: hipify :: unit_tests/headers/headers_test_07.cu (8 of 59) -PASS: hipify :: unit_tests/headers/headers_test_04.cu (9 of 59) -PASS: hipify :: unit_tests/headers/headers_test_08.cu (10 of 59) -PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (11 of 59) -PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (12 of 59) -PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (13 of 59) -PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (14 of 59) -PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (15 of 59) -PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (16 of 59) -PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (17 of 59) -PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (18 of 59) -PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (19 of 59) -PASS: hipify :: unit_tests/device/math_functions.cu (20 of 59) -PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (21 of 59) -PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (22 of 59) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (23 of 59) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (24 of 59) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (25 of 59) -PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (26 of 59) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (27 of 59) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (28 of 59) -PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (29 of 59) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (30 of 59) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (31 of 59) -PASS: hipify :: unit_tests/headers/headers_test_09.cu (32 of 59) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (33 of 59) -PASS: hipify :: unit_tests/pp/pp_if_else_conditionals_01.cu (34 of 59) -PASS: hipify :: unit_tests/pp/pp_if_else_conditionals.cu (35 of 59) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (36 of 59) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (37 of 59) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (38 of 59) -PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp (39 of 59) -PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (40 of 59) -PASS: hipify :: unit_tests/samples/MallocManaged.cpp (41 of 59) -PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (42 of 59) -PASS: hipify :: unit_tests/samples/allocators.cu (43 of 59) -PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp (44 of 59) -PASS: hipify :: unit_tests/samples/2_Cookbook/13_occupancy/occupancy.cpp (45 of 59) -PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (46 of 59) -PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (47 of 59) -PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (48 of 59) -PASS: hipify :: unit_tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp (49 of 59) -PASS: hipify :: unit_tests/samples/coalescing.cu (50 of 59) -PASS: hipify :: unit_tests/samples/intro.cu (51 of 59) -PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (52 of 59) -PASS: hipify :: unit_tests/samples/axpy.cu (53 of 59) -PASS: hipify :: unit_tests/samples/static_shared_memory.cu (54 of 59) -PASS: hipify :: unit_tests/samples/vec_add.cu (55 of 59) -PASS: hipify :: unit_tests/samples/square.cu (56 of 59) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (57 of 59) -PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (58 of 59) -PASS: hipify :: unit_tests/samples/cudaRegister.cu (59 of 59) -Testing Time: 3.23s - Expected Passes : 59 +-- Testing: 61 tests, 12 threads -- +PASS: hipify :: unit_tests/casts/reinterpret_cast.cu (1 of 61) +PASS: hipify :: unit_tests/headers/headers_test_03.cu (2 of 61) +PASS: hipify :: unit_tests/headers/headers_test_05.cu (3 of 61) +PASS: hipify :: unit_tests/headers/headers_test_02.cu (4 of 61) +PASS: hipify :: unit_tests/headers/headers_test_01.cu (5 of 61) +PASS: hipify :: unit_tests/device/math_functions.cu (6 of 61) +PASS: hipify :: unit_tests/headers/headers_test_06.cu (7 of 61) +PASS: hipify :: unit_tests/headers/headers_test_07.cu (8 of 61) +PASS: hipify :: unit_tests/device/device_symbols.cu (9 of 61) +PASS: hipify :: unit_tests/headers/headers_test_04.cu (10 of 61) +PASS: hipify :: unit_tests/headers/headers_test_10.cu (11 of 61) +PASS: hipify :: unit_tests/headers/headers_test_11.cu (12 of 61) +PASS: hipify :: unit_tests/headers/headers_test_08.cu (13 of 61) +PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (14 of 61) +PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (15 of 61) +PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (16 of 61) +PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (17 of 61) +PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (18 of 61) +PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (19 of 61) +PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (20 of 61) +PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (21 of 61) +PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (22 of 61) +PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (23 of 61) +PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (24 of 61) +PASS: hipify :: unit_tests/headers/headers_test_09.cu (25 of 61) +PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (26 of 61) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (27 of 61) +PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (28 of 61) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (29 of 61) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (30 of 61) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (31 of 61) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (32 of 61) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (33 of 61) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (34 of 61) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (35 of 61) +PASS: hipify :: unit_tests/pp/pp_if_else_conditionals.cu (36 of 61) +PASS: hipify :: unit_tests/pp/pp_if_else_conditionals_01.cu (37 of 61) +PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (38 of 61) +PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp (39 of 61) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (40 of 61) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (41 of 61) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (42 of 61) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (43 of 61) +PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (44 of 61) +PASS: hipify :: unit_tests/samples/MallocManaged.cpp (45 of 61) +PASS: hipify :: unit_tests/samples/allocators.cu (46 of 61) +PASS: hipify :: unit_tests/samples/2_Cookbook/13_occupancy/occupancy.cpp (47 of 61) +PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp (48 of 61) +PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (49 of 61) +PASS: hipify :: unit_tests/samples/coalescing.cu (50 of 61) +PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (51 of 61) +PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (52 of 61) +PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (53 of 61) +PASS: hipify :: unit_tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp (54 of 61) +PASS: hipify :: unit_tests/samples/intro.cu (55 of 61) +PASS: hipify :: unit_tests/samples/axpy.cu (56 of 61) +PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (57 of 61) +PASS: hipify :: unit_tests/samples/square.cu (58 of 61) +PASS: hipify :: unit_tests/samples/static_shared_memory.cu (59 of 61) +PASS: hipify :: unit_tests/samples/vec_add.cu (60 of 61) +PASS: hipify :: unit_tests/samples/cudaRegister.cu (61 of 61) +Testing Time: 2.81s + Expected Passes : 61 [100%] Built target test-hipify ``` ### Windows @@ -329,7 +331,7 @@ LLVM 7.0.0 - 9.0.0, CUDA 7.5 - 10.1, cudnn-7.0.5.15 - cudnn-7.6.3.30 Build system for the above configurations: -Python 3.6 (min), cmake 3.12.3 (min), Visual Studio 2017 (15.5.2) - 2019 (16.2.5). +Python 3.6 - 3.7.4, cmake 3.12.3 - 3.15.3, Visual Studio 2017 (15.5.2) - 2019 (16.3.1). Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 16 2019`: diff --git a/hipify-clang/src/CUDA2HIP_Device_functions.cpp b/hipify-clang/src/CUDA2HIP_Device_functions.cpp index 791e477ecb..b61e4bd73c 100644 --- a/hipify-clang/src/CUDA2HIP_Device_functions.cpp +++ b/hipify-clang/src/CUDA2HIP_Device_functions.cpp @@ -432,13 +432,176 @@ const std::map CUDA_DEVICE_FUNC_MAP{ {"__vabsdiffs4", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, {"__vsads4", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, // fp16 functions + {"__float2half", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__float2half_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__float2half_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__float2half_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__float2half_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2float", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__float2half2_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__floats2half2_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__low2float", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__high2float", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__float22half2_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half22float2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2int_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2int_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2int_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2int_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__int2half_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__int2half_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__int2half_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2short_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2short_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2short_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2short_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__short2half_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__short2half_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__short2half_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__short2half_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2uint_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2uint_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2uint_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2uint_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__uint2half_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__uint2half_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__uint2half_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__uint2half_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2ushort_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2ushort_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2ushort_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2ushort_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ushort2half_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ushort2half_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ushort2half_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ushort2half_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2ull_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2ull_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2ull_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2ull_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ull2half_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ull2half_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ull2half_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ull2half_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2ll_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2ll_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2ll_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2ll_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ll2half_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ll2half_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ll2half_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ll2half_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"htrunc", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"hceil", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"hfloor", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"hrint", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"h2trunc", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"h2ceil", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"h2floor", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"h2rint", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half2half2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__lowhigh2highlow", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__lows2half2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__highs2half2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__high2half", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__low2half", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hisinf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__halves2half2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__low2half2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__high2half2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half_as_short", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__half_as_ushort", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__short_as_half", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ushort_as_half", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ldg", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ldg", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ldcg", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ldca", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__ldcs", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__heq2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hne2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hle2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hge2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hge2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hlt2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hgt2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hequ2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hneu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hleu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hleu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hgeu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hltu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hgtu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hisnan2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hadd2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hsub2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hmul2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__h2div", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hadd2_sat", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hsub2_sat", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hmul2_sat", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hfma2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hfma2_sat", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hneg2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hsub", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hmul", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hdiv", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hadd_sat", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hsub_sat", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hmul_sat", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hfma", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hfma_sat", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hneg", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hbeq2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hbne2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hble2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hbge2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hblt2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hbgt2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hbequ2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hbneu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hbleu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hbgeu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hbltu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hbgtu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__heq", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hne", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hle", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hge", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hlt", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hgt", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hequ", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hneu", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hleu", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hgeu", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hltu", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hgtu", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__hisnan", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"hsqrt", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"hrsqrt", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"hrcp", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"hlog", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"hlog2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"hlog10", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"hexp", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"hexp2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"hexp10", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"hcos", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"hsin", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"h2sqrt", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"h2rsqrt", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"h2rcp", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"h2log", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"h2log2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"h2log10", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"h2exp", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"h2exp2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"h2exp10", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"h2cos", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"h2sin", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__shfl_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, {"__shfl_up_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, {"__shfl_down_sync",{"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, {"__shfl_xor_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, - {"__shfl_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, - {"__shfl_up_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, - {"__shfl_down_sync",{"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, - {"__shfl_xor_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, - {"__shfl_xor_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, + {"atomicAdd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, }; diff --git a/hipify-clang/src/CUDA2HIP_Perl.cpp b/hipify-clang/src/CUDA2HIP_Perl.cpp index 26a72886c7..5407cae35c 100644 --- a/hipify-clang/src/CUDA2HIP_Perl.cpp +++ b/hipify-clang/src/CUDA2HIP_Perl.cpp @@ -45,12 +45,19 @@ namespace perl { const std::string sForeach = "foreach $func (\n"; const std::string sMy = "my $m = 0;\n"; - void generateSymbolFunctions(std::unique_ptr& perlStreamPtr) { - *perlStreamPtr.get() << "\n" << sSub << " transformSymbolFunctions\n" << "{\n" << tab << sMy; + void generateHostFunctions(std::unique_ptr& perlStreamPtr) { + *perlStreamPtr.get() << "\n" << sSub << " transformHostFunctions\n" << "{\n" << tab << sMy; std::set &funcSet = DeviceSymbolFunctions0; - for (int i = 0; i < 2; ++i) { + const std::string s0 = "$m += s/(?flush(); bool ret = true; diff --git a/hipify-clang/src/CUDA2HIP_Scripting.h b/hipify-clang/src/CUDA2HIP_Scripting.h index 9469f25845..76b103735b 100644 --- a/hipify-clang/src/CUDA2HIP_Scripting.h +++ b/hipify-clang/src/CUDA2HIP_Scripting.h @@ -24,6 +24,11 @@ THE SOFTWARE. extern std::set DeviceSymbolFunctions0; extern std::set DeviceSymbolFunctions1; +extern std::set ReinterpretFunctions0; +extern std::set ReinterpretFunctions1; + +extern std::string sHIP_SYMBOL; +extern std::string s_reinterpret_cast; namespace perl { diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index 4acbb7e78e..e331cfbec5 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/hipify-clang/src/HipifyAction.cpp @@ -36,12 +36,20 @@ THE SOFTWARE. namespace ct = clang::tooling; namespace mat = clang::ast_matchers; +const std::string sHIP_DYNAMIC_SHARED = "HIP_DYNAMIC_SHARED"; +std::string sHIP_SYMBOL = "HIP_SYMBOL"; +std::string s_reinterpret_cast = "reinterpret_cast"; +const std::string sHipLaunchKernelGGL = "hipLaunchKernelGGL("; +const std::string sDim3 = "dim3("; + const std::string sCudaMemcpyToSymbol = "cudaMemcpyToSymbol"; const std::string sCudaMemcpyToSymbolAsync = "cudaMemcpyToSymbolAsync"; const std::string sCudaGetSymbolSize = "cudaGetSymbolSize"; const std::string sCudaGetSymbolAddress = "cudaGetSymbolAddress"; const std::string sCudaMemcpyFromSymbol = "cudaMemcpyFromSymbol"; const std::string sCudaMemcpyFromSymbolAsync = "cudaMemcpyFromSymbolAsync"; +const std::string sCudaFuncSetCacheConfig = "cudaFuncSetCacheConfig"; +const std::string sCudaFuncGetAttributes = "cudaFuncGetAttributes"; std::set DeviceSymbolFunctions0 { {sCudaMemcpyToSymbol}, @@ -55,6 +63,19 @@ std::set DeviceSymbolFunctions1 { {sCudaMemcpyFromSymbolAsync} }; +std::set ReinterpretFunctions{ + {sCudaFuncSetCacheConfig}, + {sCudaFuncGetAttributes} +}; + +std::set ReinterpretFunctions0{ + {sCudaFuncSetCacheConfig} +}; + +std::set ReinterpretFunctions1{ + {sCudaFuncGetAttributes} +}; + void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) { clang::SourceManager& SM = getCompilerInstance().getSourceManager(); size_t begin = 0; @@ -136,10 +157,8 @@ namespace { clang::SourceRange getReadRange(clang::SourceManager& SM, const clang::SourceRange& exprRange) { clang::SourceLocation begin = exprRange.getBegin(); clang::SourceLocation end = exprRange.getEnd(); - bool beginSafe = !SM.isMacroBodyExpansion(begin) || clang::Lexer::isAtStartOfMacroExpansion(begin, SM, clang::LangOptions{}); bool endSafe = !SM.isMacroBodyExpansion(end) || clang::Lexer::isAtEndOfMacroExpansion(end, SM, clang::LangOptions{}); - if (beginSafe && endSafe) { return {SM.getFileLoc(begin), SM.getFileLoc(end)}; } else { @@ -268,7 +287,6 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, DE.Report(sl, DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "Unsupported CUDA header.")); return; } - clang::StringRef newInclude; // Keep the same include type that the user gave. if (!exclude) { @@ -329,15 +347,13 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc llvm::raw_svector_ostream OS(XStr); clang::LangOptions DefaultLangOptions; clang::SourceManager* SM = Result.SourceManager; - OS << "hipLaunchKernelGGL("; + OS << sHipLaunchKernelGGL; if (caleeDecl->isTemplateInstantiation()) OS << "("; OS << readSourceText(*SM, calleeExpr->getSourceRange()); if (caleeDecl->isTemplateInstantiation()) OS << ")"; OS << ", "; - // Next up are the four kernel configuration parameters, the last two of which are optional and default to zero. // Copy the two dimensional arguments verbatim. - std::string sDim3 = "dim3("; for (unsigned int i = 0; i < 2; ++i) { const std::string sArg = readSourceText(*SM, config->getArg(i)->getSourceRange()).str(); bool bDim3 = std::equal(sDim3.begin(), sDim3.end(), sArg.c_str()); @@ -346,7 +362,6 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc // The stream/memory arguments default to zero if omitted. OS << stringifyZeroDefaultedArg(*SM, config->getArg(2)) << ", "; OS << stringifyZeroDefaultedArg(*SM, config->getArg(3)); - // If there are ordinary arguments to the kernel, just copy them verbatim into our new call. int numArgs = launchKernel->getNumArgs(); if (numArgs > 0) { @@ -358,7 +373,6 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc OS << readSourceText(*SM, {argStart, argEnd}); } OS << ")"; - clang::SourceRange replacementRange = getWriteRange(*SM, {llcompat::getBeginLoc(launchKernel), llcompat::getEndLoc(launchKernel)}); clang::SourceLocation launchStart = replacementRange.getBegin(); clang::SourceLocation launchEnd = replacementRange.getEnd(); @@ -381,7 +395,6 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match if (!sharedVar->hasExternalFormalLinkage()) { return false; } - clang::QualType QT = sharedVar->getType(); std::string typeName; if (QT->isIncompleteArrayType()) { @@ -400,58 +413,59 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match typeName = QT.getAsString(); } } - if (!typeName.empty()) { clang::SourceLocation slStart = sharedVar->getOuterLocStart(); clang::SourceLocation slEnd = llcompat::getEndLoc(sharedVar->getTypeSourceInfo()->getTypeLoc()); clang::SourceManager* SM = Result.SourceManager; size_t repLength = SM->getCharacterData(slEnd) - SM->getCharacterData(slStart) + 1; std::string varName = sharedVar->getNameAsString(); - std::string repName = "HIP_DYNAMIC_SHARED(" + typeName + ", " + varName + ")"; + std::string repName = sHIP_DYNAMIC_SHARED + "(" + typeName + ", " + varName + ")"; ct::Replacement Rep(*SM, slStart, repLength, repName); clang::FullSourceLoc fullSL(slStart, *SM); insertReplacement(Rep, fullSL); - hipCounter counter = {"HIP_DYNAMIC_SHARED", "", ConvTypes::CONV_MEMORY, ApiTypes::API_RUNTIME}; + hipCounter counter = {sHIP_DYNAMIC_SHARED, "", ConvTypes::CONV_MEMORY, ApiTypes::API_RUNTIME}; Statistics::current().incrementCounter(counter, refName.str()); + return true; } - return true; + return false; } bool HipifyAction::cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) { if (const clang::CallExpr *call = Result.Nodes.getNodeAs("cudaDeviceFuncCall")) { const clang::FunctionDecl *funcDcl = call->getDirectCallee(); if (!funcDcl) { - return true; + return false; } FindAndReplace(funcDcl->getDeclName().getAsString(), llcompat::getBeginLoc(call), CUDA_DEVICE_FUNC_MAP, false); + return true; } - return true; + return false; } -bool HipifyAction::cudaSymbolFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) { - if (const clang::CallExpr * call = Result.Nodes.getNodeAs("cudaSymbolFuncCall")) { +bool HipifyAction::cudaHostFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) { + if (const clang::CallExpr * call = Result.Nodes.getNodeAs("cudaHostFuncCall")) { if (!call->getNumArgs()) { - return true; + return false; } const clang::FunctionDecl* funcDcl = call->getDirectCallee(); if (!funcDcl) { - return true; + return false; } std::string sName = funcDcl->getDeclName().getAsString(); unsigned int argNum = 0; - if (DeviceSymbolFunctions0.find(sName) != DeviceSymbolFunctions0.end()) { + bool b_reinterpret = (ReinterpretFunctions.find(sName) != ReinterpretFunctions.end()) ? true : false; + if (DeviceSymbolFunctions0.find(sName) != DeviceSymbolFunctions0.end() || sCudaFuncSetCacheConfig == sName) { argNum = 0; - } else if (call->getNumArgs() > 1 && DeviceSymbolFunctions1.find(sName) != DeviceSymbolFunctions1.end()) { + } else if (call->getNumArgs() > 1 && (DeviceSymbolFunctions1.find(sName) != DeviceSymbolFunctions1.end() || sCudaFuncGetAttributes == sName)) { argNum = 1; } else { - return true; + return false; } clang::SmallString<40> XStr; llvm::raw_svector_ostream OS(XStr); clang::SourceRange sr = call->getArg(argNum)->getSourceRange(); clang::SourceManager* SM = Result.SourceManager; - const std::string sSymbol = "HIP_SYMBOL"; - OS << sSymbol << "(" << readSourceText(*SM, sr) << ")"; + OS << (b_reinterpret ? s_reinterpret_cast : sHIP_SYMBOL) << "(" << readSourceText(*SM, sr) << ")"; clang::SourceRange replacementRange = getWriteRange(*SM, { sr.getBegin(), sr.getEnd() }); clang::SourceLocation s = replacementRange.getBegin(); clang::SourceLocation e = replacementRange.getEnd(); @@ -460,8 +474,9 @@ bool HipifyAction::cudaSymbolFuncCall(const clang::ast_matchers::MatchFinder::Ma ct::Replacement Rep(*SM, s, length, OS.str()); clang::FullSourceLoc fullSL(s, *SM); insertReplacement(Rep, fullSL); + return true; } - return true; + return false; } void HipifyAction::insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL) { @@ -498,11 +513,13 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi sCudaMemcpyFromSymbol, sCudaMemcpyFromSymbolAsync, sCudaMemcpyToSymbol, - sCudaMemcpyToSymbolAsync + sCudaMemcpyToSymbolAsync, + sCudaFuncSetCacheConfig, + sCudaFuncGetAttributes ) ) ) - ).bind("cudaSymbolFuncCall"), + ).bind("cudaHostFuncCall"), this ); Finder->addMatcher( @@ -606,7 +623,6 @@ public: hipifyAction.Ifndef(Loc, MacroNameTok, MD); } }; - } bool HipifyAction::BeginInvocation(clang::CompilerInstance &CI) { @@ -617,12 +633,10 @@ bool HipifyAction::BeginInvocation(clang::CompilerInstance &CI) { void HipifyAction::ExecuteAction() { clang::Preprocessor& PP = getCompilerInstance().getPreprocessor(); clang::SourceManager& SM = getCompilerInstance().getSourceManager(); - // Start lexing the specified input file. const llvm::MemoryBuffer* FromFile = SM.getBuffer(SM.getMainFileID()); clang::Lexer RawLex(SM.getMainFileID(), FromFile, SM, PP.getLangOpts()); RawLex.SetKeepWhitespaceMode(true); - // Perform a token-level rewrite of CUDA identifiers to hip ones. The raw-mode lexer gives us enough // information to tell the difference between identifiers, string literals, and "other stuff". It also // ignores preprocessor directives, so this transformation will operate inside preprocessor-deleted code. @@ -632,7 +646,6 @@ void HipifyAction::ExecuteAction() { RewriteToken(RawTok); RawLex.LexFromRawLexer(RawTok); } - // Register yourself as the preprocessor callback, by proxy. PP.addPPCallbacks(std::unique_ptr(new PPCallbackProxy(*this))); // Now we're done futzing with the lexer, have the subclass proceeed with Sema and AST matching. @@ -642,6 +655,6 @@ void HipifyAction::ExecuteAction() { void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Result) { if (cudaLaunchKernel(Result)) return; if (cudaSharedIncompleteArrayVar(Result)) return; - if (cudaSymbolFuncCall(Result)) return; + if (cudaHostFuncCall(Result)) return; if (cudaDeviceFuncCall(Result)) return; } diff --git a/hipify-clang/src/HipifyAction.h b/hipify-clang/src/HipifyAction.h index 6c04e2c0cc..3c85604ced 100644 --- a/hipify-clang/src/HipifyAction.h +++ b/hipify-clang/src/HipifyAction.h @@ -71,7 +71,7 @@ public: bool cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result); bool cudaSharedIncompleteArrayVar(const clang::ast_matchers::MatchFinder::MatchResult& Result); bool cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result); - bool cudaSymbolFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result); + bool cudaHostFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result); // Called by the preprocessor for each include directive during the non-raw lexing pass. void InclusionDirective(clang::SourceLocation hash_loc, const clang::Token &include_token, diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index e3e5a14aad..f25ab72255 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -126,6 +126,8 @@ typedef struct hipDeviceProp_t { size_t textureAlignment; ///textureAlignment = cdprop.textureAlignment; p_prop->kernelExecTimeoutEnabled = cdprop.kernelExecTimeoutEnabled; p_prop->ECCEnabled = cdprop.ECCEnabled; + p_prop->tccDriver = cdprop.tccDriver; return hipCUDAErrorTohipError(cerror); } diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index 2a8610db5e..093e7ee292 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -927,6 +927,7 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) { else prop->ECCEnabled = 0; //Device has ECC support disabled + prop->tccDriver = 0; // valid only for nvcc platform return e; } diff --git a/tests/hipify-clang/unit_tests/casts/reinterpret_cast.cu b/tests/hipify-clang/unit_tests/casts/reinterpret_cast.cu new file mode 100644 index 0000000000..6b6f4f5dde --- /dev/null +++ b/tests/hipify-clang/unit_tests/casts/reinterpret_cast.cu @@ -0,0 +1,52 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args + +/* +Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +// CHECK: #include +#include + +__global__ +void fn(float* px, float* py) { + bool a[42]; + __shared__ double b[69]; + for (auto&& x : b) x = *py++; + for (auto&& x : a) x = *px++ > 0.0; + for (auto&& x : a) if (x) *--py = *--px; +} + +int main() { + // CHECK: hipFuncCache_t cacheConfig; + cudaFuncCache cacheConfig; + void* func; + // CHECK: hipFuncSetCacheConfig(reinterpret_cast(func), cacheConfig); + cudaFuncSetCacheConfig(func, cacheConfig); + // CHECK: hipFuncAttributes attr{}; + cudaFuncAttributes attr{}; + // CHECK: auto r = hipFuncGetAttributes(&attr, reinterpret_cast(&fn)); + auto r = cudaFuncGetAttributes(&attr, &fn); + // CHECK: if (r != hipSuccess || attr.maxThreadsPerBlock == 0) { + if (r != cudaSuccess || attr.maxThreadsPerBlock == 0) { + return 1; + } + return 0; +}