Merge 'master' into 'amd-master'
Change-Id: Iba89017b25273c611bee5a15fa619c7ddf5ae0aa
This commit is contained in:
Jenkins
+2
-1
@@ -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")
|
||||
|
||||
@@ -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' ) {
|
||||
|
||||
+62
-55
@@ -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 @@
|
||||
## <a name="dependencies"></a> 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** | + <br/> **LATEST STABLE RELEASE** | + <br/> **LATEST STABLE RELEASE** |
|
||||
| 7.0.0 | 9.2 | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
|
||||
| 7.0.1 | 9.2 | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
|
||||
| 7.1.0 | 9.2 (?) | - <br/> LLVM 7.1.0 <br/> is not yet released | - <br/> LLVM 7.1.0 <br/> is not yet released |
|
||||
| 8.0.0 | 10.0 (?) | - <br/> LLVM 8.0.0 <br/> is not yet released | - <br/> LLVM 8.0.0 <br/> 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
|
||||
|
||||
@@ -64,6 +64,17 @@ const std::map <llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter>& CUDA_RENAMES_MAP() {
|
||||
@@ -88,5 +99,7 @@ const std::map<llvm::StringRef, hipCounter>& 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;
|
||||
};
|
||||
|
||||
@@ -61,6 +61,10 @@ extern const std::map<llvm::StringRef, hipCounter> CUDA_FFT_FUNCTION_MAP;
|
||||
extern const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_TYPE_NAME_MAP;
|
||||
// Maps the names of CUDA SPARSE API functions to the corresponding HIP functions
|
||||
extern const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP;
|
||||
// Maps the names of CUDA CAFFE2 API types to the corresponding HIP types
|
||||
extern const std::map<llvm::StringRef, hipCounter> CUDA_CAFFE2_TYPE_NAME_MAP;
|
||||
// Maps the names of CUDA CAFFE2 API functions to the corresponding HIP functions
|
||||
extern const std::map<llvm::StringRef, hipCounter> CUDA_CAFFE2_FUNCTION_MAP;
|
||||
|
||||
/**
|
||||
* The union of all the above maps, except includes.
|
||||
|
||||
@@ -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<llvm::StringRef, hipCounter> CUDA_CAFFE2_FUNCTION_MAP{
|
||||
{"cuda_stream", {"hip_stream", "", CONV_LIB_FUNC, API_CAFFE2}},
|
||||
};
|
||||
@@ -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<llvm::StringRef, hipCounter> 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}},
|
||||
};
|
||||
@@ -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<llvm::StringRef, hipCounter> CUDA_COMPLEX_FUNCTION_MAP{
|
||||
{"cuCrealf", {"hipCrealf", "", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCimagf", {"hipCimagf", "", CONV_COMPLEX, API_COMPLEX}},
|
||||
|
||||
@@ -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<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP{
|
||||
// 5. cuSPARSE Helper Function Reference
|
||||
{"cusparseCreate", {"hipsparseCreate", "", CONV_LIB_FUNC, API_SPARSE}},
|
||||
|
||||
@@ -73,7 +73,8 @@ const char *apiNames[NUM_API_TYPES] = {
|
||||
"cuRAND API",
|
||||
"cuDNN API",
|
||||
"cuFFT API",
|
||||
"cuSPARSE API"
|
||||
"cuSPARSE API",
|
||||
"CAFFE2 API"
|
||||
};
|
||||
|
||||
namespace {
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -85,7 +85,7 @@ THE SOFTWARE.
|
||||
#define __global__
|
||||
|
||||
#define __noinline__
|
||||
#define __forceinline__
|
||||
#define __forceinline__ inline
|
||||
|
||||
#define __shared__
|
||||
#define __constant__
|
||||
|
||||
@@ -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<double>(val);
|
||||
return *reinterpret_cast<double*>(&val);
|
||||
#endif
|
||||
}
|
||||
__DEVICE__
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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 <cstdint>
|
||||
#include <algorithm>
|
||||
|
||||
// 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<int>((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);
|
||||
}
|
||||
@@ -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')
|
||||
|
||||
@@ -0,0 +1,7 @@
|
||||
#ifndef CAFFE2_CORE_COMMON_CUDNN_H_
|
||||
#define CAFFE2_CORE_COMMON_CUDNN_H_
|
||||
|
||||
#include <array>
|
||||
#include <mutex>
|
||||
|
||||
#endif // CAFFE2_CORE_COMMON_CUDNN_H_
|
||||
+14
@@ -0,0 +1,14 @@
|
||||
#ifndef CAFFE2_OPERATORS_SPATIAL_BATCH_NORM_OP_H_
|
||||
#define CAFFE2_OPERATORS_SPATIAL_BATCH_NORM_OP_H_
|
||||
|
||||
#include <algorithm>
|
||||
#include <array>
|
||||
#include <functional>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
namespace caffe2 {
|
||||
|
||||
} // namespace caffe2
|
||||
|
||||
#endif // CAFFE2_OPERATORS_SPATIAL_BATCH_NORM_OP_H_
|
||||
@@ -0,0 +1,12 @@
|
||||
// RUN: %run_test hipify "%s" "%t" %hipify_args "-roc" %clang_args
|
||||
|
||||
// NOTE: Nonworking code just for conversion testing
|
||||
|
||||
// CHECK: #include <hip/hip_runtime.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <math.h>
|
||||
// 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"
|
||||
@@ -0,0 +1,102 @@
|
||||
// RUN: %run_test hipify "%s" "%t" %hipify_args "-roc" %clang_args
|
||||
|
||||
// NOTE: Nonworking code just for conversion testing
|
||||
|
||||
// CHECK: #include <hip/hip_runtime.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <math.h>
|
||||
#include <string>
|
||||
|
||||
namespace caffe2 {
|
||||
|
||||
// Operator Definition.
|
||||
struct OperatorDef {
|
||||
int input = 1;
|
||||
int output = 2;
|
||||
int name = 3;
|
||||
};
|
||||
|
||||
class OperatorBase;
|
||||
class Workspace;
|
||||
|
||||
template <class T>
|
||||
class Observable {
|
||||
public:
|
||||
Observable() = default;
|
||||
|
||||
Observable(Observable&&) = default;
|
||||
Observable& operator =(Observable&&) = default;
|
||||
|
||||
virtual ~Observable() = default;
|
||||
};
|
||||
|
||||
template <class T>
|
||||
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<OperatorBase> OperatorObserver;
|
||||
|
||||
class OperatorBase : public Observable<OperatorBase> {
|
||||
public:
|
||||
explicit OperatorBase(const OperatorDef& operator_def, Workspace* ws);
|
||||
virtual ~OperatorBase() noexcept {}
|
||||
};
|
||||
|
||||
template <class Context>
|
||||
class Operator : public OperatorBase {
|
||||
public:
|
||||
explicit Operator(const OperatorDef& operator_def, Workspace* ws)
|
||||
: OperatorBase(operator_def, ws) {
|
||||
}
|
||||
~Operator() noexcept override {}
|
||||
};
|
||||
|
||||
template <class Context>
|
||||
class DummyEmptyOp : public Operator<Context> {
|
||||
public:
|
||||
DummyEmptyOp(const OperatorDef& def, Workspace* ws)
|
||||
: Operator<Context>(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<HIPContext>);
|
||||
REGISTER_CUDA_OPERATOR(Operator, DummyEmptyOp<CUDAContext>);
|
||||
// CHECK: REGISTER_HIP_OPERATOR_CREATOR(Operator, DummyEmptyOp<HIPContext>);
|
||||
REGISTER_CUDA_OPERATOR_CREATOR(Operator, DummyEmptyOp<CUDAContext>);
|
||||
|
||||
}
|
||||
مرجع در شماره جدید
Block a user