AI/rocm-systems
2
0
Rozštěpit 0
Zdrojový kód Úkoly Pull requesty Akce Balíčky Projekty Vydání Wiki Aktivita

Merge 'master' into 'amd-master'

Procházet zdrojové kódy 
Change-Id: I27cef7b297b1a06b20be5e4d7f76ec4ecd958eb5
Tento commit je obsažen v:
Jenkins
2019-09-27 05:58:37 -04:00
rodič 82f3b6f818 85d702ad2a
revize 2bd24bb9c1
11 změnil soubory, kde provedl 544 přidání a 118 odebrání
Stáhněte soubor záplaty Stáhněte rozdílový soubor Rozbalit všechny soubory Sbalit všechny soubory
bin/hipify-perl
+179 -8
Zobrazit soubor
@@ -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/(?<!\/\/ CHECK: )($func)\s*\(\s*([^,]+)\s*,\s*([^,\)]+)\s*(,\s*|\))\s*/$func\($2, HIP_SYMBOL\($3\)$4/g;
}
foreach $func (
"hipFuncSetCacheConfig"
)
{
$m += s/(?<!\/\/ CHECK: )($func)\s*\(\s*([^,]+)\s*,/$func\(reinterpret_cast<const void*>\($2\),/g
}
foreach $func (
"hipFuncGetAttributes"
)
{
$m += s/(?<!\/\/ CHECK: )($func)\s*\(\s*([^,]+)\s*,\s*([^,\)]+)\s*(,\s*|\))\s*/$func\($2, reinterpret_cast<const void*>\($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",
hipify-clang/README.md
+67 -65
Zobrazit soubor
@@ -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
```
### <a name="windows"></a >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`:
hipify-clang/src/CUDA2HIP_Device_functions.cpp
+168 -5
Zobrazit soubor
@@ -432,13 +432,176 @@ const std::map<llvm::StringRef, hipCounter> 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}},
};
hipify-clang/src/CUDA2HIP_Perl.cpp
+23 -7
Zobrazit soubor
@@ -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<std::ostream>& perlStreamPtr) {
*perlStreamPtr.get() << "\n" << sSub << " transformSymbolFunctions\n" << "{\n" << tab << sMy;
void generateHostFunctions(std::unique_ptr<std::ostream>& perlStreamPtr) {
*perlStreamPtr.get() << "\n" << sSub << " transformHostFunctions\n" << "{\n" << tab << sMy;
std::set<std::string> &funcSet = DeviceSymbolFunctions0;
for (int i = 0; i < 2; ++i) {
const std::string s0 = "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,/$func\\(";
const std::string s1 = "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,\\s*([^,\\)]+)\\s*(,\\s*|\\))\\s*/$func\\($2, ";
for (int i = 0; i < 4; ++i) {
*perlStreamPtr.get() << tab + sForeach;
if (i == 1) funcSet = DeviceSymbolFunctions1;
switch (i) {
case 1: funcSet = DeviceSymbolFunctions1; break;
case 2: funcSet = ReinterpretFunctions0; break;
case 3: funcSet = ReinterpretFunctions1; break;
default: funcSet = DeviceSymbolFunctions0;
}
unsigned int count = 0;
for (auto& f : funcSet) {
const auto found = CUDA_RUNTIME_FUNCTION_MAP.find(f);
@@ -60,8 +67,17 @@ namespace perl {
}
}
*perlStreamPtr.get() << "\n" << tab << ")\n" << tab << "{\n" << double_tab;
if (i ==0) *perlStreamPtr.get() << "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,/$func\\(HIP_SYMBOL\\($2\\),/g\n";
else *perlStreamPtr.get() << "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,\\s*([^,\\)]+)\\s*(,\\s*|\\))\\s*/$func\\($2, HIP_SYMBOL\\($3\\)$4/g;\n";
switch (i) {
case 0:
default:
*perlStreamPtr.get() << s0 << sHIP_SYMBOL << "\\($2\\),/g\n"; break;
case 1:
*perlStreamPtr.get() << s1 << sHIP_SYMBOL << "\\($3\\)$4/g;\n"; break;
case 2:
*perlStreamPtr.get() << s0 << s_reinterpret_cast << "\\($2\\),/g\n"; break;
case 3:
*perlStreamPtr.get() << s1 << s_reinterpret_cast << "\\($3\\)$4/g;\n"; break;
}
*perlStreamPtr.get() << tab << "}\n";
}
*perlStreamPtr.get() << tab << sReturn_m << "}\n";
@@ -164,7 +180,7 @@ namespace perl {
}
}
}
generateSymbolFunctions(perlStreamPtr);
generateHostFunctions(perlStreamPtr);
generateDeviceFunctions(perlStreamPtr);
perlStreamPtr.get()->flush();
bool ret = true;
hipify-clang/src/CUDA2HIP_Scripting.h
+5
Zobrazit soubor
@@ -24,6 +24,11 @@ THE SOFTWARE.
extern std::set<std::string> DeviceSymbolFunctions0;
extern std::set<std::string> DeviceSymbolFunctions1;
extern std::set<std::string> ReinterpretFunctions0;
extern std::set<std::string> ReinterpretFunctions1;
extern std::string sHIP_SYMBOL;
extern std::string s_reinterpret_cast;
namespace perl {
hipify-clang/src/HipifyAction.cpp
+45 -32
Zobrazit soubor
@@ -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 void*>";
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<std::string> DeviceSymbolFunctions0 {
{sCudaMemcpyToSymbol},
@@ -55,6 +63,19 @@ std::set<std::string> DeviceSymbolFunctions1 {
{sCudaMemcpyFromSymbolAsync}
};
std::set<std::string> ReinterpretFunctions{
{sCudaFuncSetCacheConfig},
{sCudaFuncGetAttributes}
};
std::set<std::string> ReinterpretFunctions0{
{sCudaFuncSetCacheConfig}
};
std::set<std::string> 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<clang::CallExpr>("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<clang::CallExpr>("cudaSymbolFuncCall")) {
bool HipifyAction::cudaHostFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
if (const clang::CallExpr * call = Result.Nodes.getNodeAs<clang::CallExpr>("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<clang::ASTConsumer> 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<PPCallbackProxy>(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;
}
hipify-clang/src/HipifyAction.h
+1 -1
Zobrazit soubor
@@ -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,
include/hip/hip_runtime_api.h
+2
Zobrazit soubor
@@ -126,6 +126,8 @@ typedef struct hipDeviceProp_t {
size_t textureAlignment; ///<Alignment requirement for textures
int kernelExecTimeoutEnabled; ///<Run time limit for kernels executed on the device
int ECCEnabled; ///<Device has ECC support enabled
int tccDriver; ///< 1:If device is Tesla device using TCC driver, else 0
} hipDeviceProp_t;
include/hip/nvcc_detail/hip_runtime_api.h
+1
Zobrazit soubor
@@ -846,6 +846,7 @@ inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t* p_prop, int dev
p_prop->textureAlignment = cdprop.textureAlignment;
p_prop->kernelExecTimeoutEnabled = cdprop.kernelExecTimeoutEnabled;
p_prop->ECCEnabled = cdprop.ECCEnabled;
p_prop->tccDriver = cdprop.tccDriver;
return hipCUDAErrorTohipError(cerror);
}
src/hip_hcc.cpp
+1
Zobrazit soubor
@@ -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;
}
tests/hipify-clang/unit_tests/casts/reinterpret_cast.cu
+52
Zobrazit soubor
@@ -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 <hip/hip_runtime.h>
#include <cuda_runtime.h>
__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<const void*>(func), cacheConfig);
cudaFuncSetCacheConfig(func, cacheConfig);
// CHECK: hipFuncAttributes attr{};
cudaFuncAttributes attr{};
// CHECK: auto r = hipFuncGetAttributes(&attr, reinterpret_cast<const void*>(&fn));
auto r = cudaFuncGetAttributes(&attr, &fn);
// CHECK: if (r != hipSuccess || attr.maxThreadsPerBlock == 0) {
if (r != cudaSuccess || attr.maxThreadsPerBlock == 0) {
return 1;
}
return 0;
}