diff --git a/CMakeLists.txt b/CMakeLists.txt index a249f0efbc..10a5393755 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -240,7 +240,8 @@ if(HIP_PLATFORM STREQUAL "hcc") src/hip_surface.cpp src/hip_intercept.cpp src/env.cpp - src/program_state.cpp) + src/program_state.cpp + src/h2f.cpp) execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS) set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS} -Wl,-Bsymbolic") diff --git a/bin/hipcc b/bin/hipcc index b649bb6e7f..4c6a5ded07 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -424,8 +424,14 @@ foreach $arg (@ARGV) $swallowArg = 1; } + if (($trimarg eq '-mno-code-object-v3') and ($HIP_PLATFORM eq 'clang') ) { + $useCodeObjectV3 = 0; + $swallowArg = 1; + } + if (($trimarg eq '-mcode-object-v3') and ($HIP_PLATFORM eq 'clang') ) { $useCodeObjectV3 = 1; + $swallowArg = 1; } if (($arg =~ /--genco/) and $HIP_PLATFORM eq 'clang' ) { diff --git a/hipify-clang/README.md b/hipify-clang/README.md index deb951efbb..ee25b23f2d 100644 --- a/hipify-clang/README.md +++ b/hipify-clang/README.md @@ -1,6 +1,6 @@ # hipify-clang -`hipify-clang` is a clang-based tool to automatically translate CUDA source code into portable HIP C++. +`hipify-clang` is a clang-based tool to translate CUDA source code into portable HIP C++ automatically. ## Table of Contents @@ -32,25 +32,27 @@ ## Dependencies `hipify-clang` requires: -1. LLVM+CLANG of at least version 3.8.0, latest stable and recommended release: 6.0.1 (linux and windows). +1. LLVM+CLANG of at least version 3.8.0, the latest stable and recommended release: 6.0.1 (Linux and windows). -2. CUDA at least version 7.0, latest supported version is 9.0. +2. CUDA at least version 7.0, the latest supported version is 9.0. -| **LLVM release version** | **CUDA latest supported version** | **Comments** | -|:------------------------:|:---------------------------------:|:------------:| -| 3.8.0 | 7.5 | -| 3.8.1 | 7.5 | -| 3.9.0 | 7.5 | -| 3.9.1 | 7.5 | -| 4.0.0 | 8.0 | -| 4.0.1 | 8.0 | -| 5.0.0 | 8.0 | -| 5.0.1 | 8.0 | -| 5.0.2 | 8.0 | -| 6.0.0 | 9.0 | -| **6.0.1** | **9.0** | **LATEST STABLE RELEASE** | -| 7.0.0 | 9.2 | windows is not supported, on linux there is a clang bug: https://bugs.llvm.org/show_bug.cgi?id=36384 | -| | 10.0 | not yet supported | +| **LLVM release version** | **CUDA latest supported version** | **Windows** | **Linux** | +|:------------------------:|:---------------------------------:|:------------:|:---------:| +| 3.8.0 | 7.5 | + | + | +| 3.8.1 | 7.5 | + | + | +| 3.9.0 | 7.5 | + | + | +| 3.9.1 | 7.5 | + | + | +| 4.0.0 | 8.0 | + | + | +| 4.0.1 | 8.0 | + | + | +| 5.0.0 | 8.0 | + | + | +| 5.0.1 | 8.0 | + | + | +| 5.0.2 | 8.0 | + | + | +| 6.0.0 | 9.0 | + | + | +| **6.0.1** | **9.0** | +
**LATEST STABLE RELEASE** | +
**LATEST STABLE RELEASE** | +| 7.0.0 | 9.2 | -
not working due to
the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | -
not working due to
the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) | +| 7.0.1 | 9.2 | -
not working due to
the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | -
not working due to
the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) | +| 7.1.0 | 9.2 (?) | -
LLVM 7.1.0
is not yet released | -
LLVM 7.1.0
is not yet released | +| 8.0.0 | 10.0 (?) | -
LLVM 8.0.0
is not yet released | -
LLVM 8.0.0
is not yet released | In most cases, you can get a suitable version of LLVM+CLANG with your package manager. @@ -78,7 +80,7 @@ make -j install On Windows, the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"`; the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make.` Debug build type `-DCMAKE_BUILD_TYPE=Debug` is also supported and tested; `LLVM+CLANG` should be built in `Debug` mode as well. -64 bit build mode `-Thost=x64` is supported as well; `LLVM+CLANG` should be built in 64bit mode as well. +64-bit build mode `-Thost=x64` is supported as well; `LLVM+CLANG` should be built in 64-bit mode as well. The binary can then be found at `./dist/bin/hipify-clang`. @@ -131,7 +133,7 @@ To run it: `-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0"` -4. Ensure [`cuDNN`](https://developer.nvidia.com/rdp/cudnn-archive) of version corresponding to CUDA's version is installed. +4. Ensure [`cuDNN`](https://developer.nvidia.com/rdp/cudnn-archive) of the version corresponding to CUDA's version is installed. * Path to cuDNN should be specified by the `CUDA_DNN_ROOT_DIR` option: @@ -234,37 +236,45 @@ make test-hipify *A corresponding successful output:* ```shell [100%] Running HIPify regression tests --- Testing: 28 tests, 12 threads -- -PASS: hipify :: allocators.cu (1 of 28) -PASS: hipify :: coalescing.cu (2 of 28) -PASS: hipify :: cuDNN/cudnn_softmax.cu (3 of 28) -PASS: hipify :: cuFFT/simple_cufft.cu (4 of 28) -PASS: hipify :: cuComplex/cuComplex_Julia.cu (5 of 28) -PASS: hipify :: cuBLAS/cublas_sgemm_matrix_multiplication.cu (6 of 28) -PASS: hipify :: cuBLAS/cublas_1_based_indexing.cu (7 of 28) -PASS: hipify :: cuBLAS/cublas_0_based_indexing.cu (8 of 28) -PASS: hipify :: axpy.cu (9 of 28) -PASS: hipify :: dynamic_shared_memory.cu (10 of 28) -PASS: hipify :: headers_test_01.cu (11 of 28) -PASS: hipify :: headers_test_02.cu (12 of 28) -PASS: hipify :: headers_test_03.cu (13 of 28) -PASS: hipify :: headers_test_05.cu (14 of 28) -PASS: hipify :: cuDNN/cudnn_convolution_forward.cu (15 of 28) -PASS: hipify :: cuRAND/poisson_api_example.cu (16 of 28) -PASS: hipify :: cudaRegister.cu (17 of 28) -PASS: hipify :: headers_test_06.cu (18 of 28) -PASS: hipify :: headers_test_04.cu (19 of 28) -PASS: hipify :: intro.cu (20 of 28) -PASS: hipify :: headers_test_07.cu (21 of 28) -PASS: hipify :: square.cu (22 of 28) -PASS: hipify :: static_shared_memory.cu (23 of 28) -PASS: hipify :: vec_add.cu (24 of 28) -PASS: hipify :: headers_test_08.cu (25 of 28) -PASS: hipify :: cuRAND/benchmark_curand_generate.cpp (26 of 28) -PASS: hipify :: cuRAND/benchmark_curand_kernel.cpp (27 of 28) -PASS: hipify :: headers_test_09.cu (28 of 28) -Testing Time: 1.71s - Expected Passes : 28 +CUDA 8.0 will be used for testing. +-- Testing: 35 tests, 12 threads -- +PASS: hipify :: unit_tests/headers/headers_test_05.cu (1 of 35) +PASS: hipify :: unit_tests/headers/headers_test_01.cu (2 of 35) +PASS: hipify :: unit_tests/headers/headers_test_02.cu (3 of 35) +PASS: hipify :: unit_tests/headers/headers_test_03.cu (4 of 35) +PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (5 of 35) +PASS: hipify :: unit_tests/headers/headers_test_06.cu (6 of 35) +PASS: hipify :: unit_tests/headers/headers_test_07.cu (7 of 35) +PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (8 of 35) +PASS: hipify :: unit_tests/headers/headers_test_04.cu (9 of 35) +PASS: hipify :: unit_tests/headers/headers_test_08.cu (10 of 35) +PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (11 of 35) +PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (12 of 35) +PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (13 of 35) +PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (14 of 35) +PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (15 of 35) +PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (16 of 35) +PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (17 of 35) +PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (18 of 35) +PASS: hipify :: unit_tests/samples/allocators.cu (19 of 35) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (20 of 35) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (21 of 35) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (22 of 35) +PASS: hipify :: unit_tests/samples/coalescing.cu (23 of 35) +PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (24 of 35) +PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (25 of 35) +PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (26 of 35) +PASS: hipify :: unit_tests/samples/intro.cu (27 of 35) +PASS: hipify :: unit_tests/samples/static_shared_memory.cu (28 of 35) +PASS: hipify :: unit_tests/samples/axpy.cu (29 of 35) +PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (30 of 35) +PASS: hipify :: unit_tests/samples/square.cu (31 of 35) +PASS: hipify :: unit_tests/samples/cudaRegister.cu (32 of 35) +PASS: hipify :: unit_tests/samples/vec_add.cu (33 of 35) +PASS: hipify :: unit_tests/headers/headers_test_09.cu (34 of 35) +PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (35 of 35) +Testing Time: 1.81s + Expected Passes : 35 [100%] Built target test-hipify ``` @@ -318,10 +328,7 @@ To process a file, `hipify-clang` needs access to the same headers that would be For example: ```shell -./hipify-clang \ - square.cu \ - --cuda-path=/usr/local/cuda-8.0 \ - -I /usr/local/cuda-8.0/samples/common/inc +./hipify-clang square.cu --cuda-path=/usr/local/cuda-8.0 -I /usr/local/cuda-8.0/samples/common/inc ``` `hipify-clang` arguments are given first, followed by a separator, and then the arguments you'd pass to `clang` if you diff --git a/hipify-clang/src/CUDA2HIP.cpp b/hipify-clang/src/CUDA2HIP.cpp index 7536c5c0f7..b9e97e874d 100644 --- a/hipify-clang/src/CUDA2HIP.cpp +++ b/hipify-clang/src/CUDA2HIP.cpp @@ -64,6 +64,17 @@ const std::map CUDA_INCLUDE_MAP{ // cuBLAS includes {"cusparse.h", {"hipsparse.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_SPARSE}}, {"cusparse_v2.h", {"hipsparse.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_SPARSE}}, + // CAFFE2 includes + {"caffe2/core/common_gpu.h", {"caffe2/core/hip/common_gpu.h", "", CONV_INCLUDE, API_CAFFE2, UNSUPPORTED}}, + {"caffe2/core/context_gpu.h", {"caffe2/core/hip/context_gpu.h", "", CONV_INCLUDE, API_CAFFE2, UNSUPPORTED}}, + {"caffe2/operators/operator_fallback_gpu.h", {"", "", CONV_INCLUDE, API_CAFFE2, UNSUPPORTED}}, + {"caffe2/operators/spatial_batch_norm_op.h", {"caffe2/operators/hip/spatial_batch_norm_op_miopen.hip", "", CONV_INCLUDE, API_CAFFE2}}, + {"caffe2/operators/generate_proposals_op_util_nms_gpu.h", {"", "", CONV_INCLUDE, API_CAFFE2, UNSUPPORTED}}, + {"caffe2/operators/max_pool_with_index_gpu.h", {"", "", CONV_INCLUDE, API_CAFFE2, UNSUPPORTED}}, + {"caffe2/operators/rnn/recurrent_network_executor_gpu.h", {"", "", CONV_INCLUDE, API_CAFFE2, UNSUPPORTED}}, + {"caffe2/utils/math/reduce.cuh", {"caffe2/utils/math/hip/reduce.cuh", "", CONV_INCLUDE, API_CAFFE2, UNSUPPORTED}}, + {"caffe2/operators/gather_op.cuh", {"caffe2/operators/math/gather_op.cuh", "", CONV_INCLUDE, API_CAFFE2, UNSUPPORTED}}, + {"caffe2/core/common_cudnn.h", {"caffe2/core/hip/common_miopen.h", "", CONV_INCLUDE, API_CAFFE2}}, }; const std::map& CUDA_RENAMES_MAP() { @@ -88,5 +99,7 @@ const std::map& CUDA_RENAMES_MAP() { ret.insert(CUDA_FFT_FUNCTION_MAP.begin(), CUDA_FFT_FUNCTION_MAP.end()); ret.insert(CUDA_SPARSE_TYPE_NAME_MAP.begin(), CUDA_SPARSE_TYPE_NAME_MAP.end()); ret.insert(CUDA_SPARSE_FUNCTION_MAP.begin(), CUDA_SPARSE_FUNCTION_MAP.end()); + ret.insert(CUDA_CAFFE2_TYPE_NAME_MAP.begin(), CUDA_CAFFE2_TYPE_NAME_MAP.end()); + ret.insert(CUDA_CAFFE2_FUNCTION_MAP.begin(), CUDA_CAFFE2_FUNCTION_MAP.end()); return ret; }; diff --git a/hipify-clang/src/CUDA2HIP.h b/hipify-clang/src/CUDA2HIP.h index d95d4fea30..b8961097b3 100644 --- a/hipify-clang/src/CUDA2HIP.h +++ b/hipify-clang/src/CUDA2HIP.h @@ -61,6 +61,10 @@ extern const std::map CUDA_FFT_FUNCTION_MAP; extern const std::map CUDA_SPARSE_TYPE_NAME_MAP; // Maps the names of CUDA SPARSE API functions to the corresponding HIP functions extern const std::map CUDA_SPARSE_FUNCTION_MAP; +// Maps the names of CUDA CAFFE2 API types to the corresponding HIP types +extern const std::map CUDA_CAFFE2_TYPE_NAME_MAP; +// Maps the names of CUDA CAFFE2 API functions to the corresponding HIP functions +extern const std::map CUDA_CAFFE2_FUNCTION_MAP; /** * The union of all the above maps, except includes. diff --git a/hipify-clang/src/CUDA2HIP_CAFFE2_API_functions.cpp b/hipify-clang/src/CUDA2HIP_CAFFE2_API_functions.cpp new file mode 100644 index 0000000000..63860de262 --- /dev/null +++ b/hipify-clang/src/CUDA2HIP_CAFFE2_API_functions.cpp @@ -0,0 +1,28 @@ +/* +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. +*/ + +#include "CUDA2HIP.h" + +// Maps the names of CUDA SPARSE API functions to the corresponding HIP functions +const std::map CUDA_CAFFE2_FUNCTION_MAP{ + {"cuda_stream", {"hip_stream", "", CONV_LIB_FUNC, API_CAFFE2}}, +}; \ No newline at end of file diff --git a/hipify-clang/src/CUDA2HIP_CAFFE2_API_types.cpp b/hipify-clang/src/CUDA2HIP_CAFFE2_API_types.cpp new file mode 100644 index 0000000000..4791cffeee --- /dev/null +++ b/hipify-clang/src/CUDA2HIP_CAFFE2_API_types.cpp @@ -0,0 +1,34 @@ +/* +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. +*/ + +#include "CUDA2HIP.h" + +// Map of all types +const std::map CUDA_CAFFE2_TYPE_NAME_MAP{ + + // 5. Defines + {"REGISTER_CUDA_OPERATOR", {"REGISTER_HIP_OPERATOR", "", CONV_DEFINE, API_CAFFE2}}, + {"REGISTER_CUDA_OPERATOR_CREATOR", {"REGISTER_HIP_OPERATOR_CREATOR", "", CONV_DEFINE, API_CAFFE2}}, + + // 6. Classes + {"CUDAContext", {"HIPContext", "", CONV_TYPE, API_CAFFE2}}, +}; diff --git a/hipify-clang/src/CUDA2HIP_Complex_API_functions.cpp b/hipify-clang/src/CUDA2HIP_Complex_API_functions.cpp index af012e27cb..6e0c1a54e7 100644 --- a/hipify-clang/src/CUDA2HIP_Complex_API_functions.cpp +++ b/hipify-clang/src/CUDA2HIP_Complex_API_functions.cpp @@ -22,7 +22,7 @@ THE SOFTWARE. #include "CUDA2HIP.h" -// Maps the names of CUDA Complex API types to the corresponding HIP types +// Maps the names of CUDA Complex API functions to the corresponding HIP functions const std::map CUDA_COMPLEX_FUNCTION_MAP{ {"cuCrealf", {"hipCrealf", "", CONV_COMPLEX, API_COMPLEX}}, {"cuCimagf", {"hipCimagf", "", CONV_COMPLEX, API_COMPLEX}}, diff --git a/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp b/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp index d89e292852..210ab1ee06 100644 --- a/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp +++ b/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp @@ -22,7 +22,7 @@ THE SOFTWARE. #include "CUDA2HIP.h" -// Maps the names of CUDA SPARSE API types to the corresponding HIP types +// Maps the names of CUDA SPARSE API functions to the corresponding HIP functions const std::map CUDA_SPARSE_FUNCTION_MAP{ // 5. cuSPARSE Helper Function Reference {"cusparseCreate", {"hipsparseCreate", "", CONV_LIB_FUNC, API_SPARSE}}, diff --git a/hipify-clang/src/Statistics.cpp b/hipify-clang/src/Statistics.cpp index 39f70e9d8a..e673957458 100644 --- a/hipify-clang/src/Statistics.cpp +++ b/hipify-clang/src/Statistics.cpp @@ -73,7 +73,8 @@ const char *apiNames[NUM_API_TYPES] = { "cuRAND API", "cuDNN API", "cuFFT API", - "cuSPARSE API" + "cuSPARSE API", + "CAFFE2 API" }; namespace { diff --git a/hipify-clang/src/Statistics.h b/hipify-clang/src/Statistics.h index 6d8986bc52..8305f0de93 100644 --- a/hipify-clang/src/Statistics.h +++ b/hipify-clang/src/Statistics.h @@ -131,6 +131,7 @@ enum ApiTypes { API_DNN, API_FFT, API_SPARSE, + API_CAFFE2, API_LAST }; constexpr int NUM_API_TYPES = (int) ApiTypes::API_LAST; diff --git a/include/hip/hcc_detail/host_defines.h b/include/hip/hcc_detail/host_defines.h index 5d1c3d8f68..689a4530da 100644 --- a/include/hip/hcc_detail/host_defines.h +++ b/include/hip/hcc_detail/host_defines.h @@ -85,7 +85,7 @@ THE SOFTWARE. #define __global__ #define __noinline__ -#define __forceinline__ +#define __forceinline__ inline #define __shared__ #define __constant__ diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index 8ac87425b9..c384c5573d 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/include/hip/hcc_detail/math_functions.h @@ -918,9 +918,10 @@ double nan(const char* tagp) return tmp.val; #else + static_assert(sizeof(uint64_t)==sizeof(double)); uint64_t val = __make_mantissa(tagp); val |= 0xFFF << 51; - return reinterpret_cast(val); + return *reinterpret_cast(&val); #endif } __DEVICE__ diff --git a/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index bbeddf2f48..6181c49afe 100644 --- a/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -205,7 +205,7 @@ void RunBenchmark_H2D(ResultDatabase& resultDB) { } double speed = - (double(sizeToBytes(thisSize) * p_beatsperiteration) / (1000 * 1000)) / t; + (double(double(sizeToBytes(thisSize)/1000) * p_beatsperiteration) / 1000) / t; char sizeStr[256]; if (p_beatsperiteration > 1) { sprintf(sizeStr, "%9sx%d", sizeToString(thisSize).c_str(), p_beatsperiteration); @@ -377,7 +377,8 @@ void RunBenchmark_D2H(ResultDatabase& resultDB) { std::cerr << "size " << sizeToString(thisSize) << " took " << t << " ms\n"; } - double speed = (double(sizeToBytes(thisSize)) / (1000 * 1000)) / t; + double speed = + (double(double(sizeToBytes(thisSize)/1000) * p_beatsperiteration) / 1000) / t; char sizeStr[256]; sprintf(sizeStr, "%9s", sizeToString(thisSize).c_str()); if (p_beatsperiteration > 1) { @@ -744,7 +745,7 @@ void RunBenchmark_P2P_Unidir(ResultDatabase& resultDB) { } double speed = - (double(sizeToBytes(thisSize) * p_beatsperiteration) / (1000 * 1000)) / t; + (double(double(sizeToBytes(thisSize)/1000) * p_beatsperiteration) / 1000) / t; char sizeStr[256]; if (p_beatsperiteration > 1) { sprintf(sizeStr, "%9sx%d", sizeToString(thisSize).c_str(), @@ -812,6 +813,7 @@ void RunBenchmark_P2P_Bidir(ResultDatabase& resultDB) { hipSetDevice(currentGpu); hipMalloc((void**)¤tGpuMem[0], sizeof(float) * numMaxFloats); hipMalloc((void**)¤tGpuMem[1], sizeof(float) * numMaxFloats); + enablePeer2Peer(peerGpu,currentGpu); hipSetDevice(peerGpu); hipMalloc((void**)&peerGpuMem[0], sizeof(float) * numMaxFloats); @@ -867,7 +869,7 @@ void RunBenchmark_P2P_Bidir(ResultDatabase& resultDB) { } double speed = - (double(sizeToBytes(2 * thisSize) * p_beatsperiteration) / (1000 * 1000)) / + (double(double(sizeToBytes(2 * thisSize)/1000) * p_beatsperiteration) / 1000) / t; char sizeStr[256]; if (p_beatsperiteration > 1) { @@ -899,6 +901,7 @@ void RunBenchmark_P2P_Bidir(ResultDatabase& resultDB) { } disablePeer2Peer(currentGpu, peerGpu); + disablePeer2Peer(peerGpu, currentGpu); hipEventDestroy(start); hipEventDestroy(stop); diff --git a/src/h2f.cpp b/src/h2f.cpp new file mode 100644 index 0000000000..a8c60e7c48 --- /dev/null +++ b/src/h2f.cpp @@ -0,0 +1,68 @@ +/* +Copyright (c) 2018 - 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. +*/ + +#include +#include + +// conversion routines between float and half precision +static inline std::uint32_t f32_as_u32(float f) { union { float f; std::uint32_t u; } v; v.f = f; return v.u; } +static inline float u32_as_f32(std::uint32_t u) { union { float f; std::uint32_t u; } v; v.u = u; return v.f; } +static inline int clamp_int(int i, int l, int h) { return std::min(std::max(i, l), h); } + +// half to float, the f16 is in the low 16 bits of the input argument a +static inline float __convert_half_to_float(std::uint32_t a) noexcept { + std::uint32_t u = ((a << 13) + 0x70000000U) & 0x8fffe000U; + std::uint32_t v = f32_as_u32(u32_as_f32(u) * 0x1.0p+112f) + 0x38000000U; + u = (a & 0x7fff) != 0 ? v : u; + return u32_as_f32(u) * 0x1.0p-112f; +} + +// float to half with nearest even rounding +// The lower 16 bits of the result is the bit pattern for the f16 +static inline std::uint32_t __convert_float_to_half(float a) noexcept { + std::uint32_t u = f32_as_u32(a); + int e = static_cast((u >> 23) & 0xff) - 127 + 15; + std::uint32_t m = ((u >> 11) & 0xffe) | ((u & 0xfff) != 0); + std::uint32_t i = 0x7c00 | (m != 0 ? 0x0200 : 0); + std::uint32_t n = ((std::uint32_t)e << 12) | m; + std::uint32_t s = (u >> 16) & 0x8000; + int b = clamp_int(1-e, 0, 13); + std::uint32_t d = (0x1000 | m) >> b; + d |= (d << b) != (0x1000 | m); + std::uint32_t v = e < 1 ? d : n; + v = (v >> 2) + (((v & 0x7) == 3) | ((v & 0x7) > 5)); + v = e > 30 ? 0x7c00 : v; + v = e == 143 ? i : v; + return s | v; +} + +// On machines without fp16 instructions, clang lowers llvm.convert.from.fp16 +// to call of this function. +extern "C" float __gnu_h2f_ieee(unsigned short h){ + return __convert_half_to_float((std::uint32_t) h); +} + +// On machines without fp16 instructions, clang lowers llvm.convert.to.fp16 +// to call of this function. +extern "C" unsigned short __gnu_f2h_ieee(float f){ + return (unsigned short)__convert_float_to_half(f); +} diff --git a/tests/hipify-clang/lit.cfg b/tests/hipify-clang/lit.cfg index 0fd0aef275..00a3eeb873 100644 --- a/tests/hipify-clang/lit.cfg +++ b/tests/hipify-clang/lit.cfg @@ -14,6 +14,8 @@ lit_config.load_config(config, site_cfg) print("CUDA " + config.cuda_version + " will be used for testing.") config.excludes = ['cmdparser.hpp'] +config.excludes.append('spatial_batch_norm_op.h') +config.excludes.append('common_cudnn.h') if config.cuda_version_major == 7 and config.cuda_version_minor == 0: config.excludes.append('headers_test_09.cu') diff --git a/tests/hipify-clang/unit_tests/libraries/CAFFE2/caffe2/core/common_cudnn.h b/tests/hipify-clang/unit_tests/libraries/CAFFE2/caffe2/core/common_cudnn.h new file mode 100644 index 0000000000..e9437c11f5 --- /dev/null +++ b/tests/hipify-clang/unit_tests/libraries/CAFFE2/caffe2/core/common_cudnn.h @@ -0,0 +1,7 @@ +#ifndef CAFFE2_CORE_COMMON_CUDNN_H_ +#define CAFFE2_CORE_COMMON_CUDNN_H_ + +#include +#include + +#endif // CAFFE2_CORE_COMMON_CUDNN_H_ diff --git a/tests/hipify-clang/unit_tests/libraries/CAFFE2/caffe2/operators/spatial_batch_norm_op.h b/tests/hipify-clang/unit_tests/libraries/CAFFE2/caffe2/operators/spatial_batch_norm_op.h new file mode 100644 index 0000000000..7b8a13788a --- /dev/null +++ b/tests/hipify-clang/unit_tests/libraries/CAFFE2/caffe2/operators/spatial_batch_norm_op.h @@ -0,0 +1,14 @@ +#ifndef CAFFE2_OPERATORS_SPATIAL_BATCH_NORM_OP_H_ +#define CAFFE2_OPERATORS_SPATIAL_BATCH_NORM_OP_H_ + +#include +#include +#include +#include +#include + +namespace caffe2 { + +} // namespace caffe2 + +#endif // CAFFE2_OPERATORS_SPATIAL_BATCH_NORM_OP_H_ diff --git a/tests/hipify-clang/unit_tests/libraries/CAFFE2/caffe2_01.cu b/tests/hipify-clang/unit_tests/libraries/CAFFE2/caffe2_01.cu new file mode 100644 index 0000000000..3c82045d30 --- /dev/null +++ b/tests/hipify-clang/unit_tests/libraries/CAFFE2/caffe2_01.cu @@ -0,0 +1,12 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args "-roc" %clang_args + +// NOTE: Nonworking code just for conversion testing + +// CHECK: #include +#include +#include +#include +// CHECK: #include "caffe2/operators/hip/spatial_batch_norm_op_miopen.hip" +#include "caffe2/operators/spatial_batch_norm_op.h" +// CHECK: #include "caffe2/core/hip/common_miopen.h" +#include "caffe2/core/common_cudnn.h" diff --git a/tests/hipify-clang/unit_tests/libraries/CAFFE2/caffe2_02.cu b/tests/hipify-clang/unit_tests/libraries/CAFFE2/caffe2_02.cu new file mode 100644 index 0000000000..7f29cfe25c --- /dev/null +++ b/tests/hipify-clang/unit_tests/libraries/CAFFE2/caffe2_02.cu @@ -0,0 +1,102 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args "-roc" %clang_args + +// NOTE: Nonworking code just for conversion testing + +// CHECK: #include +#include +#include +#include +#include + +namespace caffe2 { + +// Operator Definition. +struct OperatorDef { + int input = 1; + int output = 2; + int name = 3; +}; + +class OperatorBase; +class Workspace; + +template +class Observable { + public: + Observable() = default; + + Observable(Observable&&) = default; + Observable& operator =(Observable&&) = default; + + virtual ~Observable() = default; +}; + +template +class ObserverBase { + public: + explicit ObserverBase(T* subject) : subject_(subject) {} + + virtual void Start() {} + virtual void Stop() {} + + virtual std::string debugInfo() { + return "Not implemented."; + } + + virtual ~ObserverBase() noexcept {}; + + T* subject() const { + return subject_; + } + + protected: + T* subject_; +}; + +typedef ObserverBase OperatorObserver; + +class OperatorBase : public Observable { + public: + explicit OperatorBase(const OperatorDef& operator_def, Workspace* ws); + virtual ~OperatorBase() noexcept {} +}; + +template +class Operator : public OperatorBase { + public: + explicit Operator(const OperatorDef& operator_def, Workspace* ws) + : OperatorBase(operator_def, ws) { + } + ~Operator() noexcept override {} +}; + +template +class DummyEmptyOp : public Operator { + public: + DummyEmptyOp(const OperatorDef& def, Workspace* ws) + : Operator(def, ws) {} + + bool RunOnDevice() final { return true; } +}; + + +class CUDAContext { +public: + CUDAContext(); + virtual ~CUDAContext() noexcept {} +}; + +#define REGISTER_CUDA_OPERATOR(name, ...) \ + void CAFFE2_PLEASE_ADD_OPERATOR_SCHEMA_FOR_##name(); \ + static void CAFFE_ANONYMOUS_VARIABLE_CUDA##name() { \ + CAFFE2_PLEASE_ADD_OPERATOR_SCHEMA_FOR_##name(); \ + } + +#define REGISTER_CUDA_OPERATOR_CREATOR(key, ...) + +// CHECK: REGISTER_HIP_OPERATOR(Operator, DummyEmptyOp); +REGISTER_CUDA_OPERATOR(Operator, DummyEmptyOp); +// CHECK: REGISTER_HIP_OPERATOR_CREATOR(Operator, DummyEmptyOp); +REGISTER_CUDA_OPERATOR_CREATOR(Operator, DummyEmptyOp); + +}