Merge 'master' into 'amd-master'
Change-Id: I4189feaa83580761c6fb2d0b636040a40964cbeb
Этот коммит содержится в:
Jenkins
@@ -253,7 +253,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
|
||||
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")
|
||||
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS}")
|
||||
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx803 --amdgpu-target=gfx900 --amdgpu-target=gfx906")
|
||||
if(COMPILE_HIP_ATP_MARKER)
|
||||
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L/opt/rocm/profiler/CXLActivityLogger/bin/x86_64 -lCXLActivityLogger")
|
||||
@@ -263,6 +263,12 @@ if(HIP_PLATFORM STREQUAL "hcc")
|
||||
if(HIP_COMPILER STREQUAL "hcc")
|
||||
target_link_libraries(hip_hcc PRIVATE hc_am)
|
||||
target_link_libraries(hip_hcc_static PRIVATE hc_am)
|
||||
|
||||
add_library(hiprtc SHARED src/hiprtc.cpp)
|
||||
target_include_directories(
|
||||
hiprtc SYSTEM
|
||||
PRIVATE ${PROJECT_SOURCE_DIR}/include ${HSA_PATH}/include)
|
||||
target_link_libraries(hiprtc PUBLIC stdc++fs)
|
||||
endif()
|
||||
|
||||
string(REPLACE " " ";" HCC_CXX_FLAGS_LIST ${HCC_CXX_FLAGS})
|
||||
@@ -297,7 +303,7 @@ endif()
|
||||
#############################
|
||||
# Install hip_hcc if platform is hcc
|
||||
if(HIP_PLATFORM STREQUAL "hcc")
|
||||
install(TARGETS hip_hcc_static hip_hcc DESTINATION lib)
|
||||
install(TARGETS hip_hcc_static hip_hcc hiprtc DESTINATION lib)
|
||||
|
||||
# Install .hipInfo
|
||||
install(FILES ${PROJECT_BINARY_DIR}/.hipInfo DESTINATION lib)
|
||||
|
||||
@@ -293,7 +293,7 @@
|
||||
|`cudnnScaleTensor` |`hipdnnScaleTensor` |
|
||||
|`cudnnCreateFilterDescriptor` |`hipdnnCreateFilterDescriptor` |
|
||||
|`cudnnSetFilter4dDescriptor` |`hipdnnSetFilter4dDescriptor` |
|
||||
|`cudnnGetFilter4dDescriptor` | |
|
||||
|`cudnnGetFilter4dDescriptor` |`hipdnnGetFilter4dDescriptor` |
|
||||
|`cudnnSetFilterNdDescriptor` |`hipdnnSetFilterNdDescriptor` |
|
||||
|`cudnnGetFilterNdDescriptor` |`hipdnnGetFilterNdDescriptor` |
|
||||
|`cudnnDestroyFilterDescriptor` |`hipdnnDestroyFilterDescriptor` |
|
||||
|
||||
@@ -103,8 +103,8 @@ To run it:
|
||||
mkdir build dist
|
||||
cd build
|
||||
```
|
||||
- **Linux**:
|
||||
|
||||
- **Linux**:
|
||||
```shell
|
||||
cmake \
|
||||
-DCMAKE_INSTALL_PREFIX=../dist \
|
||||
@@ -113,9 +113,9 @@ To run it:
|
||||
../llvm
|
||||
make -j install
|
||||
```
|
||||
- **Windows**:
|
||||
|
||||
```shell
|
||||
- **Windows**:
|
||||
```shell
|
||||
cmake \
|
||||
-G "Visual Studio 15 2017 Win64" \
|
||||
-DCMAKE_INSTALL_PREFIX=../dist \
|
||||
@@ -123,7 +123,7 @@ To run it:
|
||||
-DCMAKE_BUILD_TYPE=Release \
|
||||
-Thost=x64 \
|
||||
../llvm
|
||||
```
|
||||
```
|
||||
|
||||
Run `Visual Studio 15 2017`, open the generated `LLVM.sln`, build all, build project `INSTALL`.
|
||||
|
||||
@@ -242,7 +242,7 @@ make test-hipify
|
||||
```
|
||||
*A corresponding successful output:*
|
||||
```shell
|
||||
[100%] Running HIPify regression tests
|
||||
Running HIPify regression tests
|
||||
========================================
|
||||
CUDA 10.0 - will be used for testing
|
||||
LLVM 8.0.0 - will be used for testing
|
||||
@@ -251,56 +251,60 @@ Linux 4.13.0-32-generic - Platform OS
|
||||
64 - hipify-clang binary bitness
|
||||
64 - python 2.7.12 binary bitness
|
||||
========================================
|
||||
-- Testing: 44 tests, 12 threads --
|
||||
PASS: hipify :: unit_tests/headers/headers_test_03.cu (1 of 44)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_02.cu (2 of 44)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_05.cu (3 of 44)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_01.cu (4 of 44)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_06.cu (5 of 44)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_07.cu (6 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (7 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (8 of 44)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_04.cu (9 of 44)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_08.cu (10 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (11 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (12 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (13 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (14 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (15 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (16 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (17 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (18 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (19 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (20 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (21 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (22 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (23 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (24 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (25 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (26 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (27 of 44)
|
||||
PASS: hipify :: unit_tests/samples/allocators.cu (28 of 44)
|
||||
PASS: hipify :: unit_tests/samples/coalescing.cu (29 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (30 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (31 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (32 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (33 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (34 of 44)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_09.cu (35 of 44)
|
||||
PASS: hipify :: unit_tests/samples/axpy.cu (36 of 44)
|
||||
PASS: hipify :: unit_tests/samples/cudaRegister.cu (37 of 44)
|
||||
PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (38 of 44)
|
||||
PASS: hipify :: unit_tests/samples/vec_add.cu (39 of 44)
|
||||
PASS: hipify :: unit_tests/samples/intro.cu (40 of 44)
|
||||
PASS: hipify :: unit_tests/samples/square.cu (41 of 44)
|
||||
PASS: hipify :: unit_tests/samples/static_shared_memory.cu (42 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (43 of 44)
|
||||
PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (44 of 44)
|
||||
Testing Time: 2.51s
|
||||
Expected Passes : 44
|
||||
-- Testing: 49 tests, 12 threads --
|
||||
PASS: hipify :: unit_tests/headers/headers_test_01.cu (1 of 49)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_03.cu (2 of 49)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_02.cu (3 of 49)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_05.cu (4 of 49)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_06.cu (5 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (6 of 49)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_07.cu (7 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (8 of 49)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_04.cu (9 of 49)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_08.cu (10 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (11 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (12 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (13 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (14 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (15 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (16 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (17 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (18 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (19 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (20 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (21 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (22 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (23 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (24 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (25 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (26 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (27 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (28 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (29 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (30 of 49)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_09.cu (31 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (32 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (33 of 49)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (34 of 49)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (35 of 49)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (36 of 49)
|
||||
PASS: hipify :: unit_tests/samples/allocators.cu (37 of 49)
|
||||
PASS: hipify :: unit_tests/samples/coalescing.cu (38 of 49)
|
||||
PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (39 of 49)
|
||||
PASS: hipify :: unit_tests/samples/intro.cu (40 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (41 of 49)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (42 of 49)
|
||||
PASS: hipify :: unit_tests/samples/square.cu (43 of 49)
|
||||
PASS: hipify :: unit_tests/samples/static_shared_memory.cu (44 of 49)
|
||||
PASS: hipify :: unit_tests/samples/axpy.cu (45 of 49)
|
||||
PASS: hipify :: unit_tests/samples/vec_add.cu (46 of 49)
|
||||
PASS: hipify :: unit_tests/samples/cudaRegister.cu (47 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (48 of 49)
|
||||
PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (49 of 49)
|
||||
Testing Time: 2.82s
|
||||
Expected Passes : 49
|
||||
[100%] Built target test-hipify
|
||||
```
|
||||
|
||||
### <a name="windows"></a >Windows
|
||||
|
||||
On Windows 10 the following configurations are tested:
|
||||
|
||||
@@ -76,7 +76,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DNN_FUNCTION_MAP{
|
||||
// cuDNN Filter functions
|
||||
{"cudnnCreateFilterDescriptor", {"hipdnnCreateFilterDescriptor", "", CONV_LIB_FUNC, API_DNN}},
|
||||
{"cudnnSetFilter4dDescriptor", {"hipdnnSetFilter4dDescriptor", "", CONV_LIB_FUNC, API_DNN}},
|
||||
{"cudnnGetFilter4dDescriptor", {"hipdnnGetFilter4dDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
|
||||
{"cudnnGetFilter4dDescriptor", {"hipdnnGetFilter4dDescriptor", "", CONV_LIB_FUNC, API_DNN}},
|
||||
{"cudnnSetFilterNdDescriptor", {"hipdnnSetFilterNdDescriptor", "", CONV_LIB_FUNC, API_DNN}},
|
||||
{"cudnnGetFilterNdDescriptor", {"hipdnnGetFilterNdDescriptor", "", CONV_LIB_FUNC, API_DNN}},
|
||||
{"cudnnDestroyFilterDescriptor", {"hipdnnDestroyFilterDescriptor", "", CONV_LIB_FUNC, API_DNN}},
|
||||
|
||||
@@ -352,7 +352,7 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match
|
||||
}
|
||||
|
||||
if (!typeName.empty()) {
|
||||
clang::SourceLocation slStart = llcompat::getBeginLoc(sharedVar->getTypeSourceInfo()->getTypeLoc());
|
||||
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;
|
||||
|
||||
@@ -0,0 +1,111 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
#include <cstddef>
|
||||
#include <string>
|
||||
|
||||
enum hiprtcResult {
|
||||
HIPRTC_SUCCESS = 0,
|
||||
HIPRTC_ERROR_OUT_OF_MEMORY = 1,
|
||||
HIPRTC_ERROR_PROGRAM_CREATION_FAILURE = 2,
|
||||
HIPRTC_ERROR_INVALID_INPUT = 3,
|
||||
HIPRTC_ERROR_INVALID_PROGRAM = 4,
|
||||
HIPRTC_ERROR_INVALID_OPTION = 5,
|
||||
HIPRTC_ERROR_COMPILATION = 6,
|
||||
HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE = 7,
|
||||
HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 8,
|
||||
HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 9,
|
||||
HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 10,
|
||||
HIPRTC_ERROR_INTERNAL_ERROR = 11
|
||||
};
|
||||
|
||||
const char* hiprtcGetErrorString(hiprtcResult result);
|
||||
|
||||
inline
|
||||
hiprtcResult hiprtcVersion(int* major, int* minor) noexcept
|
||||
{ // TODO: NVRTC versioning is somewhat unclear.
|
||||
if (!major || !minor) return HIPRTC_ERROR_INVALID_INPUT;
|
||||
|
||||
// TODO: this should be generic / set by the build infrastructure.
|
||||
*major = 9;
|
||||
*minor = 0;
|
||||
|
||||
return HIPRTC_SUCCESS;
|
||||
}
|
||||
|
||||
struct _hiprtcProgram;
|
||||
using hiprtcProgram = _hiprtcProgram*;
|
||||
|
||||
hiprtcResult hiprtcAddNameExpression(hiprtcProgram prog,
|
||||
const char* name_expression);
|
||||
|
||||
hiprtcResult hiprtcCompileProgram(hiprtcProgram prog, int numOptions,
|
||||
const char** options);
|
||||
|
||||
hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog, const char* src,
|
||||
const char* name, int numHeaders,
|
||||
const char** headers,
|
||||
const char** includeNames);
|
||||
|
||||
hiprtcResult hiprtcDestroyProgram(hiprtcProgram* prog);
|
||||
|
||||
hiprtcResult hiprtcGetLoweredName(hiprtcProgram prog,
|
||||
const char* name_expression,
|
||||
const char** lowered_name);
|
||||
|
||||
hiprtcResult hiprtcGetProgramLog(hiprtcProgram prog, char* log);
|
||||
|
||||
hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram prog,
|
||||
std::size_t* logSizeRet);
|
||||
|
||||
hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* code);
|
||||
|
||||
hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, std::size_t* codeSizeRet);
|
||||
|
||||
namespace hip_impl
|
||||
{
|
||||
std::string demangle(const char* mangled_expression);
|
||||
}
|
||||
|
||||
#if defined(HIPRTC_GET_TYPE_NAME)
|
||||
#include <typeinfo>
|
||||
|
||||
#if defined(_WIN32)
|
||||
#include <dbghelp.h>
|
||||
|
||||
template<typename>
|
||||
hiprtcResult hiprtcGetTypeName(std::string*) = delete;
|
||||
#else
|
||||
template<typename T>
|
||||
inline
|
||||
hiprtcResult hiprtcGetTypeName(std::string* result)
|
||||
{
|
||||
if (!result) return HIPRTC_ERROR_INVALID_INPUT;
|
||||
|
||||
*result = hip_impl::demangle(typeid(T).name());
|
||||
|
||||
return (result->empty()) ? HIPRTC_ERROR_INTERNAL_ERROR :
|
||||
HIPRTC_SUCCESS;
|
||||
}
|
||||
#endif
|
||||
#endif
|
||||
@@ -75,7 +75,7 @@ public:
|
||||
|
||||
hipFunction_t kernel_descriptor(std::uintptr_t,
|
||||
hsa_agent_t);
|
||||
|
||||
|
||||
kernargs_size_align get_kernargs_size_align(std::uintptr_t);
|
||||
hsa_executable_t load_executable(const char*, const size_t,
|
||||
hsa_executable_t,
|
||||
|
||||
@@ -0,0 +1,32 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
#include <hip/hip_common.h>
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) && !defined(__HIP_PLATFORM_NVCC__)
|
||||
#include <hip/hcc_detail/hiprtc.h>
|
||||
#elif defined(__HIP_PLATFORM_NVCC__) && !defined(__HIP_PLATFORM_HCC__)
|
||||
#include <hip/nvcc_detail/nvrtc.h>
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
@@ -3,6 +3,7 @@ project(hip_hcc)
|
||||
|
||||
install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib)
|
||||
install(FILES @PROJECT_BINARY_DIR@/libhip_hcc_static.a DESTINATION lib)
|
||||
install(FILES @PROJECT_BINARY_DIR@/libhiprtc.so DESTINATION lib)
|
||||
install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib)
|
||||
install(FILES @PROJECT_BINARY_DIR@/hip-config.cmake @PROJECT_BINARY_DIR@/hip-config-version.cmake DESTINATION lib/cmake/hip)
|
||||
install(FILES @hip_SOURCE_DIR@/packaging/hip-targets.cmake @hip_SOURCE_DIR@/packaging/hip-targets-release.cmake DESTINATION lib/cmake/hip)
|
||||
|
||||
@@ -574,9 +574,10 @@ hipError_t hipExtGetLinkTypeAndHopCount(int device1, int device2, uint32_t* link
|
||||
HSA_ERROR_CHECK(err);
|
||||
*linktype = link_info.link_type;
|
||||
|
||||
err = hsa_amd_agent_memory_pool_get_info(device1Handle->_hsaAgent, pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_NUM_LINK_HOPS, hopcount);
|
||||
HSA_ERROR_CHECK(err);
|
||||
|
||||
if (link_info.numa_distance < 30)
|
||||
*hopcount = 1;
|
||||
else
|
||||
*hopcount = 2;
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -371,6 +371,8 @@ struct ihipModule_t {
|
||||
hsa_executable_t executable = {};
|
||||
hsa_code_object_reader_t coReader = {};
|
||||
std::string hash;
|
||||
std::unordered_map<
|
||||
std::string, std::vector<std::pair<std::size_t, std::size_t>>> kernargs;
|
||||
|
||||
~ihipModule_t() {
|
||||
if (executable.handle) hsa_executable_destroy(executable);
|
||||
|
||||
+28
-17
@@ -91,6 +91,7 @@ struct ihipModuleSymbol_t {
|
||||
uint64_t _object{}; // The kernel object.
|
||||
amd_kernel_code_t const* _header{};
|
||||
string _name; // TODO - review for performance cost. Name is just used for debug.
|
||||
vector<pair<size_t, size_t>> _kernarg_layout{};
|
||||
};
|
||||
|
||||
template <>
|
||||
@@ -132,6 +133,8 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
|
||||
uint32_t localWorkSizeZ, size_t sharedMemBytes,
|
||||
hipStream_t hStream, void** kernelParams, void** extra,
|
||||
hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags) {
|
||||
using namespace hip_impl;
|
||||
|
||||
auto ctx = ihipGetTlsDefaultCtx();
|
||||
hipError_t ret = hipSuccess;
|
||||
|
||||
@@ -146,19 +149,26 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
|
||||
void* config[5] = {0};
|
||||
size_t kernArgSize;
|
||||
|
||||
if (kernelParams != NULL) {
|
||||
std::string name = f->_name;
|
||||
struct ihipKernArgInfo pl = kernelArguments[name];
|
||||
char* argBuf = (char*)malloc(pl.totalSize);
|
||||
memset(argBuf, 0, pl.totalSize);
|
||||
int index = 0;
|
||||
for (int i = 0; i < pl.Size.size(); i++) {
|
||||
memcpy(argBuf + index, kernelParams[i], pl.Size[i]);
|
||||
index += pl.Align[i];
|
||||
std::vector<char> tmp{};
|
||||
if (kernelParams) {
|
||||
if (extra) return hipErrorInvalidValue;
|
||||
|
||||
for (auto&& x : f->_kernarg_layout) {
|
||||
const auto p{static_cast<const char*>(*kernelParams)};
|
||||
|
||||
tmp.insert(
|
||||
tmp.cend(),
|
||||
round_up_to_next_multiple_nonnegative(
|
||||
tmp.size(), x.second) - tmp.size(),
|
||||
'\0');
|
||||
tmp.insert(tmp.cend(), p, p + x.first);
|
||||
|
||||
++kernelParams;
|
||||
}
|
||||
config[1] = (void*)argBuf;
|
||||
kernArgSize = pl.totalSize;
|
||||
} else if (extra != NULL) {
|
||||
config[1] = static_cast<void*>(tmp.data());
|
||||
|
||||
kernArgSize = tmp.size();
|
||||
} else if (extra) {
|
||||
memcpy(config, extra, sizeof(size_t) * 5);
|
||||
if (config[0] == HIP_LAUNCH_PARAM_BUFFER_POINTER &&
|
||||
config[2] == HIP_LAUNCH_PARAM_BUFFER_SIZE && config[4] == HIP_LAUNCH_PARAM_END) {
|
||||
@@ -236,10 +246,6 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
|
||||
stopEvent->attachToCompletionFuture(&cf, hStream, hipEventTypeStopCommand);
|
||||
}
|
||||
|
||||
|
||||
if (kernelParams != NULL) {
|
||||
free(config[1]);
|
||||
}
|
||||
ihipPostLaunchKernel(f->_name.c_str(), hStream, lp);
|
||||
}
|
||||
|
||||
@@ -461,7 +467,7 @@ hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const ch
|
||||
// TODO: refactor the whole ihipThisThat, which is a mess and yields the
|
||||
// below, due to hipFunction_t being a pointer to ihipModuleSymbol_t.
|
||||
func[0][0] = *static_cast<hipFunction_t>(
|
||||
Kernel_descriptor{kernel_object(kernel), name});
|
||||
Kernel_descriptor{kernel_object(kernel), name, hmod->kernargs[name]});
|
||||
|
||||
return hipSuccess;
|
||||
}
|
||||
@@ -549,6 +555,11 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) {
|
||||
(*module)->executable = get_program_state().load_executable(
|
||||
content.data(), content.size(), (*module)->executable,
|
||||
this_agent());
|
||||
istringstream elf{content};
|
||||
ELFIO::elfio reader;
|
||||
if (reader.load(elf)) {
|
||||
program_state_impl::read_kernarg_metadata(reader, (*module)->kernargs);
|
||||
}
|
||||
|
||||
// compute the hash of the code object
|
||||
(*module)->hash = checksum(content.length(), content.data());
|
||||
|
||||
@@ -0,0 +1,601 @@
|
||||
/*
|
||||
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 "../include/hip/hiprtc.h"
|
||||
#include "../include/hip/hcc_detail/code_object_bundle.hpp"
|
||||
#include "../include/hip/hcc_detail/elfio/elfio.hpp"
|
||||
#include "../include/hip/hcc_detail/program_state.hpp"
|
||||
|
||||
#include "../lpl_ca/pstreams/pstream.h"
|
||||
|
||||
#include <hsa/hsa.h>
|
||||
|
||||
#include <cxxabi.h>
|
||||
|
||||
#include <algorithm>
|
||||
#include <cassert>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <experimental/filesystem>
|
||||
#include <fstream>
|
||||
#include <future>
|
||||
#include <iterator>
|
||||
#include <memory>
|
||||
#include <mutex>
|
||||
#include <stdexcept>
|
||||
#include <string>
|
||||
#include <unordered_map>
|
||||
#include <utility>
|
||||
#include <vector>
|
||||
|
||||
#include <iostream>
|
||||
|
||||
const char* hiprtcGetErrorString(hiprtcResult x)
|
||||
{
|
||||
switch (x) {
|
||||
case HIPRTC_SUCCESS:
|
||||
return "HIPRTC_SUCCESS";
|
||||
case HIPRTC_ERROR_OUT_OF_MEMORY:
|
||||
return "HIPRTC_ERROR_OUT_OF_MEMORY";
|
||||
case HIPRTC_ERROR_PROGRAM_CREATION_FAILURE:
|
||||
return "HIPRTC_ERROR_PROGRAM_CREATION_FAILURE";
|
||||
case HIPRTC_ERROR_INVALID_INPUT:
|
||||
return "HIPRTC_ERROR_INVALID_INPUT";
|
||||
case HIPRTC_ERROR_INVALID_PROGRAM:
|
||||
return "HIPRTC_ERROR_INVALID_PROGRAM";
|
||||
case HIPRTC_ERROR_INVALID_OPTION:
|
||||
return "HIPRTC_ERROR_INVALID_OPTION";
|
||||
case HIPRTC_ERROR_COMPILATION:
|
||||
return "HIPRTC_ERROR_COMPILATION";
|
||||
case HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE:
|
||||
return "HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE";
|
||||
case HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION:
|
||||
return "HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION";
|
||||
case HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION:
|
||||
return "HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION";
|
||||
case HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID:
|
||||
return "HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID";
|
||||
case HIPRTC_ERROR_INTERNAL_ERROR:
|
||||
return "HIPRTC_ERROR_INTERNAL_ERROR";
|
||||
default: throw std::logic_error{"Invalid HIPRTC result."};
|
||||
};
|
||||
}
|
||||
|
||||
namespace
|
||||
{
|
||||
struct Symbol {
|
||||
std::string name;
|
||||
ELFIO::Elf64_Addr value = 0;
|
||||
ELFIO::Elf_Xword size = 0;
|
||||
ELFIO::Elf_Half sect_idx = 0;
|
||||
std::uint8_t bind = 0;
|
||||
std::uint8_t type = 0;
|
||||
std::uint8_t other = 0;
|
||||
};
|
||||
|
||||
inline
|
||||
Symbol read_symbol(const ELFIO::symbol_section_accessor& section,
|
||||
unsigned int idx) {
|
||||
assert(idx < section.get_symbols_num());
|
||||
|
||||
Symbol r;
|
||||
section.get_symbol(
|
||||
idx, r.name, r.value, r.size, r.bind, r.type, r.sect_idx, r.other);
|
||||
|
||||
return r;
|
||||
}
|
||||
} // Unnamed namespace.
|
||||
|
||||
struct _hiprtcProgram {
|
||||
// DATA - STATICS
|
||||
static std::vector<std::unique_ptr<_hiprtcProgram>> programs;
|
||||
static std::mutex mtx;
|
||||
|
||||
// DATA
|
||||
std::vector<std::pair<std::string, std::string>> headers;
|
||||
std::vector<std::pair<std::string, std::string>> names;
|
||||
std::vector<std::string> loweredNames;
|
||||
std::vector<char> elf;
|
||||
std::string source;
|
||||
std::string name;
|
||||
std::string log;
|
||||
bool compiled;
|
||||
|
||||
// STATICS
|
||||
static
|
||||
hiprtcResult destroy(_hiprtcProgram* p)
|
||||
{
|
||||
using namespace std;
|
||||
|
||||
lock_guard<mutex> lck{mtx};
|
||||
|
||||
const auto it{find_if(programs.cbegin(), programs.cend(),
|
||||
[=](const unique_ptr<_hiprtcProgram>& x) {
|
||||
return x.get() == p;
|
||||
})};
|
||||
|
||||
if (it == programs.cend()) return HIPRTC_ERROR_INVALID_PROGRAM;
|
||||
|
||||
return HIPRTC_SUCCESS;
|
||||
}
|
||||
|
||||
static
|
||||
std::string handleMangledName(std::string name)
|
||||
{
|
||||
using namespace std;
|
||||
|
||||
name = hip_impl::demangle(name.c_str());
|
||||
|
||||
if (name.empty()) return name;
|
||||
|
||||
if (name.find("void ") == 0) name.erase(0, strlen("void "));
|
||||
|
||||
auto dx{name.find_first_of("(<")};
|
||||
|
||||
if (dx == string::npos) return name;
|
||||
|
||||
if (name[dx] == '<') {
|
||||
auto cnt{1u};
|
||||
do {
|
||||
++dx;
|
||||
cnt += (name[dx] == '<') ? 1 : ((name[dx] == '>') ? -1 : 0);
|
||||
} while (cnt);
|
||||
|
||||
name.erase(++dx);
|
||||
}
|
||||
else name.erase(dx);
|
||||
|
||||
return name;
|
||||
}
|
||||
|
||||
static
|
||||
_hiprtcProgram* make(std::string s, std::string n,
|
||||
std::vector<std::pair<std::string, std::string>> h)
|
||||
{
|
||||
using namespace std;
|
||||
|
||||
unique_ptr<_hiprtcProgram> tmp{new _hiprtcProgram{move(h), {}, {}, {},
|
||||
move(s), move(n), {},
|
||||
false}};
|
||||
|
||||
lock_guard<mutex> lck{mtx};
|
||||
|
||||
programs.push_back(move(tmp));
|
||||
|
||||
return programs.back().get();
|
||||
}
|
||||
|
||||
static
|
||||
bool isValid(_hiprtcProgram* p) noexcept
|
||||
{
|
||||
return std::find_if(programs.cbegin(), programs.cend(),
|
||||
[=](const std::unique_ptr<_hiprtcProgram>& x) {
|
||||
return x.get() == p;
|
||||
}) != programs.cend();
|
||||
}
|
||||
|
||||
// MANIPULATORS
|
||||
bool compile(const std::vector<std::string>& args,
|
||||
const std::experimental::filesystem::path& program_folder)
|
||||
{
|
||||
using namespace ELFIO;
|
||||
using namespace redi;
|
||||
using namespace std;
|
||||
|
||||
ipstream compile{args.front(), args, pstreambuf::pstderr};
|
||||
|
||||
constexpr const auto tmp_size{1024u};
|
||||
char tmp[tmp_size]{};
|
||||
while (!compile.eof()) {
|
||||
log.append(tmp, tmp + compile.readsome(tmp, tmp_size));
|
||||
}
|
||||
|
||||
compile.close();
|
||||
|
||||
if (compile.rdbuf()->exited() &&
|
||||
compile.rdbuf()->status() != EXIT_SUCCESS) return false;
|
||||
|
||||
elfio reader;
|
||||
if (!reader.load(args.back())) return false;
|
||||
|
||||
const auto it{find_if(reader.sections.begin(), reader.sections.end(),
|
||||
[](const section* x) {
|
||||
return x->get_name() == ".kernel";
|
||||
})};
|
||||
|
||||
if (it == reader.sections.end()) return false;
|
||||
|
||||
hip_impl::Bundled_code_header h{(*it)->get_data()};
|
||||
|
||||
if (bundles(h).empty()) return false;
|
||||
|
||||
elf.assign(bundles(h).back().blob.cbegin(),
|
||||
bundles(h).back().blob.cend());
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool readLoweredNames()
|
||||
{
|
||||
using namespace ELFIO;
|
||||
using namespace hip_impl;
|
||||
using namespace std;
|
||||
|
||||
if (names.empty()) return true;
|
||||
|
||||
istringstream blob{string{elf.cbegin(), elf.cend()}};
|
||||
|
||||
elfio reader;
|
||||
|
||||
if (!reader.load(blob)) return false;
|
||||
|
||||
const auto it{find_if(reader.sections.begin(), reader.sections.end(),
|
||||
[](const section* x) {
|
||||
return x->get_type() == SHT_SYMTAB;
|
||||
})};
|
||||
|
||||
ELFIO::symbol_section_accessor symbols{reader, *it};
|
||||
|
||||
auto n{symbols.get_symbols_num()};
|
||||
|
||||
if (n < loweredNames.size()) return false;
|
||||
|
||||
while (n--) {
|
||||
const auto tmp{read_symbol(symbols, n)};
|
||||
|
||||
auto it{find_if(names.cbegin(), names.cend(),
|
||||
[&](const pair<string, string>& x) {
|
||||
return x.second == tmp.name;
|
||||
})};
|
||||
|
||||
if (it == names.cend()) {
|
||||
const auto name{handleMangledName(tmp.name)};
|
||||
|
||||
if (name.empty()) continue;
|
||||
|
||||
it = find_if(names.cbegin(), names.cend(),
|
||||
[&](const pair<string, string>& x) {
|
||||
return x.second == name;
|
||||
});
|
||||
|
||||
if (it == names.cend()) continue;
|
||||
}
|
||||
|
||||
loweredNames[distance(names.cbegin(), it)] = tmp.name;
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
// ACCESSORS
|
||||
std::experimental::filesystem::path writeTemporaryFiles(
|
||||
const std::experimental::filesystem::path& programFolder) const
|
||||
{
|
||||
using namespace std;
|
||||
|
||||
vector<future<void>> fut{headers.size()};
|
||||
transform(headers.cbegin(), headers.cend(), begin(fut),
|
||||
[&](const pair<string, string>& x) {
|
||||
return async([&]() {
|
||||
ofstream h{programFolder / x.first};
|
||||
h.write(x.second.data(), x.second.size());
|
||||
});
|
||||
});
|
||||
|
||||
auto tmp{(programFolder / name).replace_extension(".cpp")};
|
||||
ofstream{tmp}.write(source.data(), source.size());
|
||||
|
||||
return tmp;
|
||||
}
|
||||
|
||||
|
||||
};
|
||||
std::vector<std::unique_ptr<_hiprtcProgram>> _hiprtcProgram::programs{};
|
||||
std::mutex _hiprtcProgram::mtx{};
|
||||
|
||||
namespace
|
||||
{
|
||||
inline
|
||||
bool isValidProgram(const hiprtcProgram p)
|
||||
{
|
||||
if (!p) return false;
|
||||
|
||||
std::lock_guard<std::mutex> lck{_hiprtcProgram::mtx};
|
||||
|
||||
return _hiprtcProgram::isValid(p);
|
||||
}
|
||||
} // Unnamed namespace.
|
||||
|
||||
hiprtcResult hiprtcAddNameExpression(hiprtcProgram p, const char* n)
|
||||
{
|
||||
if (!n) return HIPRTC_ERROR_INVALID_INPUT;
|
||||
if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
|
||||
if (p->compiled) return HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION;
|
||||
|
||||
const auto id{p->names.size()};
|
||||
|
||||
p->names.emplace_back(n, n);
|
||||
p->loweredNames.emplace_back();
|
||||
|
||||
if (p->names.back().second.back() == ')') {
|
||||
p->names.back().second.pop_back();
|
||||
p->names.back().second.erase(0, p->names.back().second.find('('));
|
||||
}
|
||||
if (p->names.back().second.front() == '&') {
|
||||
p->names.back().second.erase(0, 1);
|
||||
}
|
||||
|
||||
const auto var{"__hiprtc_" + std::to_string(id)};
|
||||
p->source.append("\nextern \"C\" constexpr auto " + var + " = " + n + ';');
|
||||
|
||||
return HIPRTC_SUCCESS;
|
||||
}
|
||||
|
||||
namespace
|
||||
{
|
||||
class Unique_temporary_path {
|
||||
// DATA
|
||||
std::experimental::filesystem::path path_{};
|
||||
public:
|
||||
// CREATORS
|
||||
Unique_temporary_path() : path_{std::tmpnam(nullptr)}
|
||||
{
|
||||
while (std::experimental::filesystem::exists(path_)) {
|
||||
path_ = std::tmpnam(nullptr);
|
||||
}
|
||||
}
|
||||
Unique_temporary_path(const std::string& extension)
|
||||
: Unique_temporary_path{}
|
||||
{
|
||||
path_.replace_extension(extension);
|
||||
}
|
||||
|
||||
Unique_temporary_path(const Unique_temporary_path&) = default;
|
||||
Unique_temporary_path(Unique_temporary_path&&) = default;
|
||||
|
||||
~Unique_temporary_path() noexcept
|
||||
{
|
||||
std::experimental::filesystem::remove_all(path_);
|
||||
}
|
||||
|
||||
// MANIPULATORS
|
||||
Unique_temporary_path& operator=(
|
||||
const Unique_temporary_path&) = default;
|
||||
Unique_temporary_path& operator=(Unique_temporary_path&&) = default;
|
||||
|
||||
// ACCESSORS
|
||||
const std::experimental::filesystem::path& path() const noexcept
|
||||
{
|
||||
return path_;
|
||||
}
|
||||
};
|
||||
} // Unnamed namespace.
|
||||
|
||||
namespace hip_impl
|
||||
{
|
||||
inline
|
||||
std::string demangle(const char* x)
|
||||
{
|
||||
if (!x) return {};
|
||||
|
||||
int s{};
|
||||
std::unique_ptr<char, decltype(std::free)*> tmp{
|
||||
abi::__cxa_demangle(x, nullptr, nullptr, &s), std::free};
|
||||
|
||||
if (s != 0) return {};
|
||||
|
||||
return tmp.get();
|
||||
}
|
||||
} // Namespace hip_impl.
|
||||
|
||||
namespace
|
||||
{
|
||||
const std::string& defaultTarget()
|
||||
{
|
||||
using namespace std;
|
||||
|
||||
static string r{"gfx900"};
|
||||
static once_flag f{};
|
||||
|
||||
call_once(f, []() {
|
||||
static hsa_agent_t a{};
|
||||
hsa_iterate_agents([](hsa_agent_t x, void*) {
|
||||
hsa_device_type_t t{};
|
||||
hsa_agent_get_info(x, HSA_AGENT_INFO_DEVICE, &t);
|
||||
|
||||
if (t != HSA_DEVICE_TYPE_GPU) return HSA_STATUS_SUCCESS;
|
||||
|
||||
a = x;
|
||||
|
||||
return HSA_STATUS_INFO_BREAK;
|
||||
}, nullptr);
|
||||
|
||||
if (!a.handle) return;
|
||||
|
||||
hsa_agent_iterate_isas(a, [](hsa_isa_t x, void*){
|
||||
uint32_t n{};
|
||||
hsa_isa_get_info_alt(x, HSA_ISA_INFO_NAME_LENGTH, &n);
|
||||
|
||||
if (n == 0) return HSA_STATUS_SUCCESS;
|
||||
|
||||
r.resize(n);
|
||||
hsa_isa_get_info_alt(x, HSA_ISA_INFO_NAME, &r[0]);
|
||||
|
||||
r.erase(0, r.find("gfx"));
|
||||
|
||||
return HSA_STATUS_INFO_BREAK;
|
||||
}, nullptr);
|
||||
});
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
inline
|
||||
void handleTarget(std::vector<std::string>& args)
|
||||
{
|
||||
using namespace std;
|
||||
|
||||
bool hasTarget{false};
|
||||
for (auto&& x : args) {
|
||||
const auto dx{x.find("--gpu-architecture")};
|
||||
const auto dy{(dx == string::npos) ? x.find("-arch")
|
||||
: string::npos};
|
||||
|
||||
if (dx == dy) continue;
|
||||
|
||||
x.replace(0, x.find('=', min(dx, dy)), "--amdgpu-target");
|
||||
hasTarget = true;
|
||||
|
||||
break;
|
||||
}
|
||||
if (!hasTarget) args.push_back("--amdgpu-target=" + defaultTarget());
|
||||
}
|
||||
} // Unnamed namespace.
|
||||
|
||||
hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o)
|
||||
{
|
||||
using namespace std;
|
||||
|
||||
if (n && !o) return HIPRTC_ERROR_INVALID_INPUT;
|
||||
if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
|
||||
if (p->compiled) return HIPRTC_ERROR_COMPILATION;
|
||||
|
||||
static const string hipcc{
|
||||
getenv("HIP_PATH") ? (getenv("HIP_PATH") + string{"/bin/hipcc"})
|
||||
: "/opt/rocm/bin/hipcc"};
|
||||
|
||||
if (!experimental::filesystem::exists(hipcc)) {
|
||||
return HIPRTC_ERROR_INTERNAL_ERROR;
|
||||
}
|
||||
|
||||
Unique_temporary_path tmp{};
|
||||
experimental::filesystem::create_directory(tmp.path());
|
||||
|
||||
const auto src{p->writeTemporaryFiles(tmp.path())};
|
||||
|
||||
vector<string> args{hipcc, "-shared"};
|
||||
if (n) args.insert(args.cend(), o, o + n);
|
||||
|
||||
handleTarget(args);
|
||||
|
||||
args.emplace_back(src);
|
||||
args.emplace_back("-o");
|
||||
args.emplace_back(tmp.path() / "hiprtc.out");
|
||||
|
||||
if (!p->compile(args, tmp.path())) return HIPRTC_ERROR_INTERNAL_ERROR;
|
||||
if (!p->readLoweredNames()) return HIPRTC_ERROR_INTERNAL_ERROR;
|
||||
|
||||
p->compiled = true;
|
||||
|
||||
return HIPRTC_SUCCESS;
|
||||
}
|
||||
|
||||
hiprtcResult hiprtcCreateProgram(hiprtcProgram* p, const char* src,
|
||||
const char* name, int n, const char** hdrs,
|
||||
const char** incs)
|
||||
{
|
||||
using namespace std;
|
||||
|
||||
if (!p) return HIPRTC_ERROR_INVALID_PROGRAM;
|
||||
if (n < 0) return HIPRTC_ERROR_INVALID_INPUT;
|
||||
if (n && (!hdrs || !incs)) return HIPRTC_ERROR_INVALID_INPUT;
|
||||
|
||||
vector<pair<string, string>> h;
|
||||
for (auto i = 0; i != n; ++i) h.emplace_back(incs[i], hdrs[i]);
|
||||
|
||||
*p = _hiprtcProgram::make(src, name ? name : "default_name", move(h));
|
||||
|
||||
return HIPRTC_SUCCESS;
|
||||
}
|
||||
|
||||
hiprtcResult hiprtcDestroyProgram(hiprtcProgram* p)
|
||||
{
|
||||
if (!p) return HIPRTC_SUCCESS;
|
||||
|
||||
return _hiprtcProgram::destroy(*p);
|
||||
}
|
||||
|
||||
hiprtcResult hiprtcGetLoweredName(hiprtcProgram p, const char* n,
|
||||
const char** ln)
|
||||
{
|
||||
using namespace std;
|
||||
|
||||
if (!n || !ln) return HIPRTC_ERROR_INVALID_INPUT;
|
||||
if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
|
||||
if (!p->compiled) return HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION;
|
||||
|
||||
const auto it{find_if(p->names.cbegin(), p->names.cend(),
|
||||
[=](const pair<string, string>& x) {
|
||||
return x.first == n;
|
||||
})};
|
||||
|
||||
if (it == p->names.cend()) return HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID;
|
||||
|
||||
*ln = p->loweredNames[distance(p->names.cbegin(), it)].c_str();
|
||||
|
||||
return HIPRTC_SUCCESS;
|
||||
}
|
||||
|
||||
hiprtcResult hiprtcGetProgramLog(hiprtcProgram p, char* l)
|
||||
{
|
||||
if (!l) return HIPRTC_ERROR_INVALID_INPUT;
|
||||
if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
|
||||
if (!p->compiled) return HIPRTC_ERROR_INVALID_PROGRAM;
|
||||
|
||||
l = std::copy_n(p->log.data(), p->log.size(), l);
|
||||
*l = '\0';
|
||||
|
||||
return HIPRTC_SUCCESS;
|
||||
}
|
||||
|
||||
hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram p, std::size_t* sz)
|
||||
{
|
||||
if (!sz) return HIPRTC_ERROR_INVALID_INPUT;
|
||||
if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
|
||||
if (!p->compiled) return HIPRTC_ERROR_INVALID_PROGRAM;
|
||||
|
||||
*sz = p->log.empty() ? 0 : p->log.size() + 1;
|
||||
|
||||
return HIPRTC_SUCCESS;
|
||||
}
|
||||
|
||||
hiprtcResult hiprtcGetCode(hiprtcProgram p, char* c)
|
||||
{
|
||||
if (!c) return HIPRTC_ERROR_INVALID_INPUT;
|
||||
if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
|
||||
if (!p->compiled) return HIPRTC_ERROR_INVALID_PROGRAM;
|
||||
|
||||
std::copy_n(p->elf.data(), p->elf.size(), c);
|
||||
|
||||
return HIPRTC_SUCCESS;
|
||||
}
|
||||
|
||||
hiprtcResult hiprtcGetCodeSize(hiprtcProgram p, std::size_t* sz)
|
||||
{
|
||||
if (!sz) return HIPRTC_ERROR_INVALID_INPUT;
|
||||
if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
|
||||
if (!p->compiled) return HIPRTC_ERROR_INVALID_PROGRAM;
|
||||
|
||||
*sz = p->elf.size();
|
||||
|
||||
return HIPRTC_SUCCESS;
|
||||
}
|
||||
@@ -66,10 +66,17 @@ class Kernel_descriptor {
|
||||
std::uint64_t kernel_object_{};
|
||||
amd_kernel_code_t const* kernel_header_{nullptr};
|
||||
std::string name_{};
|
||||
std::vector<std::pair<std::size_t, std::size_t>> kernarg_layout_{};
|
||||
public:
|
||||
Kernel_descriptor() = default;
|
||||
Kernel_descriptor(std::uint64_t kernel_object, const std::string& name)
|
||||
: kernel_object_{kernel_object}, name_{name}
|
||||
Kernel_descriptor(
|
||||
std::uint64_t kernel_object,
|
||||
const std::string& name,
|
||||
std::vector<std::pair<std::size_t, std::size_t>> kernarg_layout = {})
|
||||
:
|
||||
kernel_object_{kernel_object},
|
||||
name_{name},
|
||||
kernarg_layout_{std::move(kernarg_layout)}
|
||||
{
|
||||
bool supported{false};
|
||||
std::uint16_t min_v{UINT16_MAX};
|
||||
@@ -548,6 +555,7 @@ public:
|
||||
return functions[agent].second;
|
||||
}
|
||||
|
||||
static
|
||||
std::size_t parse_args(
|
||||
const std::string& metadata,
|
||||
std::size_t f,
|
||||
@@ -576,6 +584,7 @@ public:
|
||||
} while (true);
|
||||
}
|
||||
|
||||
static
|
||||
void read_kernarg_metadata(
|
||||
ELFIO::elfio& reader,
|
||||
std::unordered_map<
|
||||
|
||||
@@ -176,6 +176,19 @@ int main() {
|
||||
out_desc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT,
|
||||
out_n, out_c, out_h, out_w));
|
||||
|
||||
|
||||
cudnnDataType_t *dataType = nullptr;
|
||||
cudnnTensorFormat_t *tensorFormat = nullptr;
|
||||
int *p_filt_k = nullptr;
|
||||
int *p_filt_c = nullptr;
|
||||
int *p_filt_h = nullptr;
|
||||
int *p_filt_w = nullptr;
|
||||
|
||||
// CHECK: CUDNN_CALL(hipdnnGetFilter4dDescriptor(
|
||||
CUDNN_CALL(cudnnGetFilter4dDescriptor(
|
||||
filt_desc, dataType, tensorFormat,
|
||||
p_filt_k, p_filt_c, p_filt_h, p_filt_w));
|
||||
|
||||
float *out_data;
|
||||
// CHECK: CUDA_CALL(hipMalloc(
|
||||
CUDA_CALL(cudaMalloc(
|
||||
|
||||
+1
-1
@@ -1,6 +1,6 @@
|
||||
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
|
||||
/*
|
||||
Copyright (c) 2015-2019 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -0,0 +1,153 @@
|
||||
// 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.
|
||||
*/
|
||||
|
||||
#include <iostream>
|
||||
// CHECK: #include <hip/hip_runtime.h>
|
||||
#include <cuda.h>
|
||||
|
||||
#define WIDTH 32
|
||||
|
||||
#define NUM (WIDTH * WIDTH)
|
||||
|
||||
#define THREADS_PER_BLOCK_X 4
|
||||
#define THREADS_PER_BLOCK_Y 4
|
||||
#define THREADS_PER_BLOCK_Z 1
|
||||
|
||||
using namespace std;
|
||||
|
||||
__global__ void matrixTranspose_static_shared(float* out, float* in,
|
||||
const int width) {
|
||||
// CHECK-NOT: HIP_DYNAMIC_SHARED(int, sharedMem);
|
||||
// CHECK: __shared__ float sharedMem[WIDTH * WIDTH];
|
||||
__shared__ float sharedMem[WIDTH * WIDTH];
|
||||
|
||||
int x = blockDim.x * blockIdx.x + threadIdx.x;
|
||||
int y = blockDim.y * blockIdx.y + threadIdx.y;
|
||||
|
||||
sharedMem[y * width + x] = in[x * width + y];
|
||||
|
||||
__syncthreads();
|
||||
|
||||
out[y * width + x] = sharedMem[y * width + x];
|
||||
}
|
||||
|
||||
__global__ void matrixTranspose_dynamic_shared(float* out, float* in,
|
||||
const int width) {
|
||||
// declare dynamic shared memory
|
||||
// CHECK-NOT: extern __shared__
|
||||
// CHECK: HIP_DYNAMIC_SHARED(int, sharedMem)
|
||||
extern __shared__ int sharedMem[];
|
||||
|
||||
int x = blockDim.x * blockIdx.x + threadIdx.x;
|
||||
int y = blockDim.y * blockIdx.y + threadIdx.y;
|
||||
|
||||
sharedMem[y * width + x] = in[x * width + y];
|
||||
|
||||
__syncthreads();
|
||||
|
||||
out[y * width + x] = sharedMem[y * width + x];
|
||||
}
|
||||
|
||||
void MultipleStream(float** data, float* randArray, float** gpuTransposeMatrix,
|
||||
float** TransposeMatrix, int width) {
|
||||
const int num_streams = 2;
|
||||
// CHECK: hipStream_t streams[num_streams];
|
||||
cudaStream_t streams[num_streams];
|
||||
|
||||
// CHECK: for (int i = 0; i < num_streams; i++) hipStreamCreate(&streams[i]);
|
||||
for (int i = 0; i < num_streams; i++) cudaStreamCreate(&streams[i]);
|
||||
|
||||
for (int i = 0; i < num_streams; i++) {
|
||||
// CHECK: hipMalloc((void**)&data[i], NUM * sizeof(float));
|
||||
cudaMalloc((void**)&data[i], NUM * sizeof(float));
|
||||
// CHECK: hipMemcpyAsync(data[i], randArray, NUM * sizeof(float), hipMemcpyHostToDevice, streams[i]);
|
||||
cudaMemcpyAsync(data[i], randArray, NUM * sizeof(float), cudaMemcpyHostToDevice, streams[i]);
|
||||
}
|
||||
|
||||
dim3 dimGrid(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y);
|
||||
dim3 dimBlock(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y);
|
||||
// CHECK: hipLaunchKernelGGL(matrixTranspose_static_shared, dim3(dimGrid), dim3(dimBlock), 0, streams[0], gpuTransposeMatrix[0], data[0], width);
|
||||
matrixTranspose_static_shared <<<dimGrid, dimBlock, 0, streams[0]>>>(gpuTransposeMatrix[0], data[0], width);
|
||||
// CHECK: hipLaunchKernelGGL(matrixTranspose_dynamic_shared, dim3(dimGrid), dim3(dimBlock), 0, streams[1], gpuTransposeMatrix[1], data[1], width);
|
||||
matrixTranspose_dynamic_shared <<<dimGrid, dimBlock, 0, streams[1]>>>(gpuTransposeMatrix[1], data[1], width);
|
||||
|
||||
for (int i = 0; i < num_streams; i++)
|
||||
// CHECK: hipMemcpyAsync(TransposeMatrix[i], gpuTransposeMatrix[i], NUM * sizeof(float), hipMemcpyDeviceToHost, streams[i]);
|
||||
cudaMemcpyAsync(TransposeMatrix[i], gpuTransposeMatrix[i], NUM * sizeof(float), cudaMemcpyDeviceToHost, streams[i]);
|
||||
}
|
||||
|
||||
int main() {
|
||||
// CHECK: hipSetDevice(0);
|
||||
cudaSetDevice(0);
|
||||
|
||||
float *data[2], *TransposeMatrix[2], *gpuTransposeMatrix[2], *randArray;
|
||||
|
||||
int width = WIDTH;
|
||||
|
||||
randArray = (float*)malloc(NUM * sizeof(float));
|
||||
|
||||
TransposeMatrix[0] = (float*)malloc(NUM * sizeof(float));
|
||||
TransposeMatrix[1] = (float*)malloc(NUM * sizeof(float));
|
||||
|
||||
// CHECK: hipMalloc((void**)&gpuTransposeMatrix[0], NUM * sizeof(float));
|
||||
cudaMalloc((void**)&gpuTransposeMatrix[0], NUM * sizeof(float));
|
||||
// CHECK: hipMalloc((void**)&gpuTransposeMatrix[1], NUM * sizeof(float));
|
||||
cudaMalloc((void**)&gpuTransposeMatrix[1], NUM * sizeof(float));
|
||||
|
||||
for (int i = 0; i < NUM; i++) {
|
||||
randArray[i] = (float)i * 1.0f;
|
||||
}
|
||||
|
||||
MultipleStream(data, randArray, gpuTransposeMatrix, TransposeMatrix, width);
|
||||
|
||||
// CHECK: hipDeviceSynchronize();
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
// verify the results
|
||||
int errors = 0;
|
||||
double eps = 1.0E-6;
|
||||
for (int i = 0; i < NUM; i++) {
|
||||
if (std::abs(TransposeMatrix[0][i] - TransposeMatrix[1][i]) > eps) {
|
||||
printf("%d stream0: %f stream1 %f\n", i, TransposeMatrix[0][i], TransposeMatrix[1][i]);
|
||||
errors++;
|
||||
}
|
||||
}
|
||||
if (errors != 0) {
|
||||
printf("FAILED: %d errors\n", errors);
|
||||
} else {
|
||||
printf("stream PASSED!\n");
|
||||
}
|
||||
|
||||
free(randArray);
|
||||
for (int i = 0; i < 2; i++) {
|
||||
// CHECK: hipFree(data[i]);
|
||||
cudaFree(data[i]);
|
||||
// CHECK: hipFree(gpuTransposeMatrix[i]);
|
||||
cudaFree(gpuTransposeMatrix[i]);
|
||||
free(TransposeMatrix[i]);
|
||||
}
|
||||
|
||||
// CHECK: hipDeviceReset();
|
||||
cudaDeviceReset();
|
||||
return 0;
|
||||
}
|
||||
@@ -9,6 +9,7 @@
|
||||
__global__ void dynamicReverse(int *d, int n)
|
||||
{
|
||||
// Dynamic shared memory
|
||||
// CHECK-NOT: extern __shared__
|
||||
// CHECK: HIP_DYNAMIC_SHARED(int, s);
|
||||
extern __shared__ int s[];
|
||||
int t = threadIdx.x;
|
||||
|
||||
@@ -8,6 +8,8 @@
|
||||
|
||||
__global__ void staticReverse(int *d, int n)
|
||||
{
|
||||
// CHECK-NOT: HIP_DYNAMIC_SHARED(int, s);
|
||||
// CHECK: __shared__ int s[64];
|
||||
__shared__ int s[64];
|
||||
int t = threadIdx.x;
|
||||
int tr = n-t-1;
|
||||
|
||||
@@ -3,17 +3,19 @@ find_package(HIP REQUIRED)
|
||||
|
||||
#-------------------------------------------------------------------------------
|
||||
# Helper macro to parse BUILD instructions
|
||||
macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_options _exclude_platforms _dir)
|
||||
macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_options _link_options _exclude_platforms _dir)
|
||||
set(${_target})
|
||||
set(${_sources})
|
||||
set(${_hipcc_options})
|
||||
set(${_hcc_options})
|
||||
set(${_nvcc_options})
|
||||
set(${_link_options})
|
||||
set(${_exclude_platforms})
|
||||
set(_target_found FALSE)
|
||||
set(_hipcc_options_found FALSE)
|
||||
set(_hcc_options_found FALSE)
|
||||
set(_nvcc_options_found FALSE)
|
||||
set(_link_options_found FALSE)
|
||||
set(_exclude_platforms_found FALSE)
|
||||
foreach(arg ${ARGN})
|
||||
if(NOT _target_found)
|
||||
@@ -23,21 +25,31 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt
|
||||
set(_hipcc_options_found TRUE)
|
||||
set(_hcc_options_found FALSE)
|
||||
set(_nvcc_options_found FALSE)
|
||||
set(_link_options_found FALSE)
|
||||
set(_exclude_platforms_found FALSE)
|
||||
elseif("x${arg}" STREQUAL "xHCC_OPTIONS")
|
||||
set(_hipcc_options_found FALSE)
|
||||
set(_hcc_options_found TRUE)
|
||||
set(_nvcc_options_found FALSE)
|
||||
set(_link_options_found FALSE)
|
||||
set(_exclude_platforms_found FALSE)
|
||||
elseif("x${arg}" STREQUAL "xNVCC_OPTIONS")
|
||||
set(_hipcc_options_found FALSE)
|
||||
set(_hcc_options_found FALSE)
|
||||
set(_nvcc_options_found TRUE)
|
||||
set(_link_options_found FALSE)
|
||||
set(_exclude_platforms_found FALSE)
|
||||
elseif("x${arg}" STREQUAL "xLINK_OPTIONS")
|
||||
set(_hipcc_options_found FALSE)
|
||||
set(_hcc_options_found FALSE)
|
||||
set(_nvcc_options_found FALSE)
|
||||
set(_link_options_found TRUE)
|
||||
set(_exclude_platforms_found FALSE)
|
||||
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_PLATFORM")
|
||||
set(_hipcc_options_found FALSE)
|
||||
set(_hcc_options_found FALSE)
|
||||
set(_nvcc_options_found FALSE)
|
||||
set(_link_options_found FALSE)
|
||||
set(_exclude_platforms_found TRUE)
|
||||
else()
|
||||
if(_hipcc_options_found)
|
||||
@@ -46,6 +58,8 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt
|
||||
list(APPEND ${_hcc_options} ${arg})
|
||||
elseif(_nvcc_options_found)
|
||||
list(APPEND ${_nvcc_options} ${arg})
|
||||
elseif(_link_options_found)
|
||||
list(APPEND ${_link_options} ${arg})
|
||||
elseif(_exclude_platforms_found)
|
||||
set(${_exclude_platforms} ${arg})
|
||||
else()
|
||||
@@ -143,12 +157,12 @@ macro(MAKE_TEST exe)
|
||||
string(REPLACE " " "" smush_args ${ARGN})
|
||||
set(testname ${exe}${smush_args}.tst)
|
||||
add_test(NAME ${testname} COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN})
|
||||
set_tests_properties(${testname} PROPERTIES PASS_REGULAR_EXPRESSION "PASSED")
|
||||
set_tests_properties(${testname} PROPERTIES PASS_REGULAR_EXPRESSION "PASSED" ENVIRONMENT HIP_PATH=${HIP_ROOT_DIR})
|
||||
endmacro()
|
||||
|
||||
macro(MAKE_NAMED_TEST exe testname)
|
||||
add_test(NAME ${testname} COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN})
|
||||
set_tests_properties(${testname} PROPERTIES PASS_REGULAR_EXPRESSION "PASSED")
|
||||
set_tests_properties(${testname} PROPERTIES PASS_REGULAR_EXPRESSION "PASSED" ENVIRONMENT HIP_PATH=${HIP_ROOT_DIR})
|
||||
endmacro()
|
||||
#-------------------------------------------------------------------------------
|
||||
|
||||
@@ -165,7 +179,7 @@ macro(HIT_ADD_FILES _dir _label _parent)
|
||||
string(REGEX REPLACE "\n" ";" _contents "${_contents}")
|
||||
foreach(_cmd ${_contents})
|
||||
string(REGEX REPLACE " " ";" _cmd "${_cmd}")
|
||||
parse_build_command(_target _sources _hipcc_options _hcc_options _nvcc_options _exclude_platforms ${_dir} ${_cmd})
|
||||
parse_build_command(_target _sources _hipcc_options _hcc_options _nvcc_options _link_options _exclude_platforms ${_dir} ${_cmd})
|
||||
string(REGEX REPLACE "/" "." target ${_label}/${_target})
|
||||
insert_into_map("_exclude" "${target}" "${_exclude_platforms}")
|
||||
if(_exclude_platforms STREQUAL "all" OR _exclude_platforms STREQUAL ${HIP_PLATFORM})
|
||||
@@ -173,6 +187,7 @@ macro(HIT_ADD_FILES _dir _label _parent)
|
||||
set_source_files_properties(${_sources} PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1)
|
||||
hip_reset_flags()
|
||||
hip_add_executable(${target} ${_sources} HIPCC_OPTIONS ${_hipcc_options} HCC_OPTIONS ${_hcc_options} NVCC_OPTIONS ${_nvcc_options} EXCLUDE_FROM_ALL)
|
||||
target_link_libraries(${target} PRIVATE ${_link_options})
|
||||
set_target_properties(${target} PROPERTIES OUTPUT_NAME ${_target} RUNTIME_OUTPUT_DIRECTORY ${_label} LINK_DEPENDS "${HIP_LIB_FILES}")
|
||||
add_dependencies(${_parent} ${target})
|
||||
endif()
|
||||
|
||||
@@ -0,0 +1,152 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
#include <test_common.h>
|
||||
|
||||
#include <hip/hiprtc.h>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
#include <iostream>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
|
||||
static constexpr const char gpu_program[]{
|
||||
R"(
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
__device__ int V1; // set from host code
|
||||
static __global__ void f1(int *result) { *result = V1 + 10; }
|
||||
namespace N1 {
|
||||
namespace N2 {
|
||||
__constant__ int V2; // set from host code
|
||||
__global__ void f2(int *result) { *result = V2 + 20; }
|
||||
}
|
||||
}
|
||||
template<typename T>
|
||||
__global__ void f3(int *result) { *result = sizeof(T); }
|
||||
)"};
|
||||
|
||||
int main()
|
||||
{
|
||||
using namespace std;
|
||||
|
||||
hiprtcProgram prog;
|
||||
hiprtcCreateProgram(&prog, gpu_program, "prog.cu", 0, nullptr, nullptr);
|
||||
|
||||
vector<string> kernel_name_vec;
|
||||
vector<string> variable_name_vec;
|
||||
vector<int> variable_initial_value;
|
||||
vector<int> expected_result;
|
||||
|
||||
kernel_name_vec.push_back("&f1");
|
||||
expected_result.push_back(10 + 100);
|
||||
kernel_name_vec.push_back("N1::N2::f2");
|
||||
expected_result.push_back(20 + 200);
|
||||
kernel_name_vec.push_back("f3<int>");
|
||||
expected_result.push_back(sizeof(int));
|
||||
kernel_name_vec.push_back("f3<double>");
|
||||
expected_result.push_back(sizeof(double));
|
||||
|
||||
for (auto&& x : kernel_name_vec) hiprtcAddNameExpression(prog, x.c_str());
|
||||
|
||||
variable_name_vec.push_back("&V1");
|
||||
variable_initial_value.push_back(100);
|
||||
variable_name_vec.push_back("&N1::N2::V2");
|
||||
variable_initial_value.push_back(200);
|
||||
|
||||
for (auto&& x : variable_name_vec) hiprtcAddNameExpression(prog, x.c_str());
|
||||
|
||||
hiprtcResult compileResult = hiprtcCompileProgram(prog, 0, nullptr);
|
||||
|
||||
// Obtain compilation log from the program.
|
||||
size_t logSize;
|
||||
hiprtcGetProgramLogSize(prog, &logSize);
|
||||
|
||||
if (logSize) {
|
||||
string log(logSize, '\0');
|
||||
hiprtcGetProgramLog(prog, &log[0]);
|
||||
|
||||
cout << log << '\n';
|
||||
}
|
||||
|
||||
if (compileResult != HIPRTC_SUCCESS) { failed("Compilation failed."); }
|
||||
|
||||
size_t codeSize;
|
||||
hiprtcGetCodeSize(prog, &codeSize);
|
||||
|
||||
vector<char> code(codeSize);
|
||||
hiprtcGetCode(prog, code.data());
|
||||
|
||||
hipModule_t module;
|
||||
hipModuleLoadData(&module, code.data());
|
||||
|
||||
hipDeviceptr_t dResult;
|
||||
int hResult = 0;
|
||||
hipMalloc(&dResult, sizeof(hResult));
|
||||
hipMemcpyHtoD(dResult, &hResult, sizeof(hResult));
|
||||
|
||||
for (decltype(variable_name_vec.size()) i = 0; i != variable_name_vec.size(); ++i) {
|
||||
const char* name;
|
||||
hiprtcGetLoweredName(prog, variable_name_vec[i].c_str(), &name);
|
||||
|
||||
int initial_value = variable_initial_value[i];
|
||||
|
||||
hipDeviceptr_t variable_addr;
|
||||
size_t bytes{};
|
||||
hipModuleGetGlobal(&variable_addr, &bytes, module, name);
|
||||
hipMemcpyHtoD(variable_addr, &initial_value, sizeof(initial_value));
|
||||
}
|
||||
|
||||
for (decltype(kernel_name_vec.size()) i = 0; i != kernel_name_vec.size(); ++i) {
|
||||
const char* name;
|
||||
hiprtcGetLoweredName(prog, kernel_name_vec[i].c_str(), &name);
|
||||
|
||||
hipFunction_t kernel;
|
||||
hipModuleGetFunction(&kernel, module, name);
|
||||
|
||||
struct { hipDeviceptr_t a_; } args{dResult};
|
||||
|
||||
auto size = sizeof(args);
|
||||
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
|
||||
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
||||
HIP_LAUNCH_PARAM_END};
|
||||
|
||||
hipModuleLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr,
|
||||
config);
|
||||
|
||||
hipMemcpyDtoH(&hResult, dResult, sizeof(hResult));
|
||||
|
||||
if (expected_result[i] != hResult) { failed("Validation failed."); }
|
||||
}
|
||||
|
||||
hipFree(dResult);
|
||||
hipModuleUnload(module);
|
||||
|
||||
hiprtcDestroyProgram(&prog);
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,138 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <test_common.h>
|
||||
|
||||
#define HIPRTC_GET_TYPE_NAME
|
||||
#include <hip/hiprtc.h>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
#include <iostream>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
static constexpr auto gpu_program{
|
||||
R"(
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
namespace N1 { struct S1_t { int i; double d; }; }
|
||||
template<typename T>
|
||||
__global__ void f3(int *result) { *result = sizeof(T); }
|
||||
)"};
|
||||
|
||||
// note: this structure is also defined in GPU code string. Should ideally
|
||||
// be in a header file included by both GPU code string and by CPU code.
|
||||
namespace N1 { struct S1_t { int i; double d; }; };
|
||||
|
||||
template <typename T>
|
||||
std::string getKernelNameForType(void)
|
||||
{
|
||||
std::string type_name;
|
||||
hiprtcGetTypeName<T>(&type_name);
|
||||
return std::string{"f3<"} + type_name + '>';
|
||||
}
|
||||
|
||||
int main()
|
||||
{
|
||||
using namespace std;
|
||||
|
||||
hiprtcProgram prog;
|
||||
hiprtcCreateProgram(&prog, gpu_program, "gpu_program.cu", 0, nullptr,
|
||||
nullptr);
|
||||
|
||||
vector<string> name_vec;
|
||||
vector<int> expected_result;
|
||||
|
||||
name_vec.push_back(getKernelNameForType<int>());
|
||||
expected_result.push_back(sizeof(int));
|
||||
name_vec.push_back(getKernelNameForType<double>());
|
||||
expected_result.push_back(sizeof(double));
|
||||
name_vec.push_back(getKernelNameForType<N1::S1_t>());
|
||||
expected_result.push_back(sizeof(N1::S1_t));
|
||||
|
||||
for (auto&& x : name_vec) hiprtcAddNameExpression(prog, x.c_str());
|
||||
|
||||
hiprtcResult compileResult = hiprtcCompileProgram(prog, 0, nullptr);
|
||||
|
||||
size_t logSize;
|
||||
hiprtcGetProgramLogSize(prog, &logSize);
|
||||
|
||||
if (logSize) {
|
||||
string log(logSize, '\0');
|
||||
hiprtcGetProgramLog(prog, &log[0]);
|
||||
|
||||
cout << log << '\n';
|
||||
}
|
||||
|
||||
if (compileResult != HIPRTC_SUCCESS) { failed("Compilation failed."); }
|
||||
|
||||
size_t codeSize;
|
||||
hiprtcGetCodeSize(prog, &codeSize);
|
||||
|
||||
vector<char> code(codeSize);
|
||||
hiprtcGetCode(prog, code.data());
|
||||
|
||||
hipModule_t module;
|
||||
hipModuleLoadDataEx(&module, code.data(), 0, nullptr, nullptr);
|
||||
|
||||
hipDeviceptr_t dResult;
|
||||
int hResult = 0;
|
||||
hipMalloc(&dResult, sizeof(hResult));
|
||||
hipMemcpyHtoD(dResult, &hResult, sizeof(hResult));
|
||||
|
||||
for (size_t i = 0; i < name_vec.size(); ++i) {
|
||||
const char *name;
|
||||
hiprtcGetLoweredName(prog, name_vec[i].c_str(), &name);
|
||||
|
||||
hipFunction_t kernel;
|
||||
hipModuleGetFunction(&kernel, module, name);
|
||||
|
||||
struct { hipDeviceptr_t a_; } args{dResult};
|
||||
|
||||
auto size = sizeof(args);
|
||||
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
|
||||
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
||||
HIP_LAUNCH_PARAM_END};
|
||||
|
||||
hipModuleLaunchKernel(kernel,
|
||||
1, 1, 1,
|
||||
1, 1, 1,
|
||||
0, nullptr,
|
||||
nullptr, config);
|
||||
|
||||
hipMemcpyDtoH(&hResult, dResult, sizeof(hResult));
|
||||
|
||||
if (expected_result[i] != hResult) { failed("Validation failed."); }
|
||||
}
|
||||
|
||||
hipFree(dResult);
|
||||
hipModuleUnload(module);
|
||||
|
||||
hiprtcDestroyProgram(&prog);
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,146 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <test_common.h>
|
||||
|
||||
#include <hip/hiprtc.h>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
#include <cassert>
|
||||
#include <cstddef>
|
||||
#include <memory>
|
||||
#include <iostream>
|
||||
#include <iterator>
|
||||
|
||||
static constexpr auto NUM_THREADS{128};
|
||||
static constexpr auto NUM_BLOCKS{32};
|
||||
|
||||
static constexpr auto saxpy{
|
||||
R"(
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
extern "C"
|
||||
__global__
|
||||
void saxpy(float a, float* x, float* y, float* out, size_t n)
|
||||
{
|
||||
size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (tid < n) {
|
||||
out[tid] = a * x[tid] + y[tid];
|
||||
}
|
||||
}
|
||||
)"};
|
||||
|
||||
int main()
|
||||
{
|
||||
using namespace std;
|
||||
|
||||
hiprtcProgram prog;
|
||||
hiprtcCreateProgram(&prog, // prog
|
||||
saxpy, // buffer
|
||||
"saxpy.cu", // name
|
||||
0, // numHeaders
|
||||
nullptr, // headers
|
||||
nullptr); // includeNames
|
||||
|
||||
hiprtcResult compileResult{hiprtcCompileProgram(prog, 0, nullptr)};
|
||||
|
||||
size_t logSize;
|
||||
hiprtcGetProgramLogSize(prog, &logSize);
|
||||
|
||||
if (logSize) {
|
||||
string log(logSize, '\0');
|
||||
hiprtcGetProgramLog(prog, &log[0]);
|
||||
|
||||
cout << log << '\n';
|
||||
}
|
||||
|
||||
if (compileResult != HIPRTC_SUCCESS) { failed("Compilation failed."); }
|
||||
|
||||
size_t codeSize;
|
||||
hiprtcGetCodeSize(prog, &codeSize);
|
||||
|
||||
vector<char> code(codeSize);
|
||||
hiprtcGetCode(prog, code.data());
|
||||
|
||||
hiprtcDestroyProgram(&prog);
|
||||
|
||||
hipModule_t module;
|
||||
hipFunction_t kernel;
|
||||
|
||||
hipModuleLoadData(&module, code.data());
|
||||
hipModuleGetFunction(&kernel, module, "saxpy");
|
||||
|
||||
size_t n = NUM_THREADS * NUM_BLOCKS;
|
||||
size_t bufferSize = n * sizeof(float);
|
||||
|
||||
float a = 5.1f;
|
||||
unique_ptr<float[]> hX{new float[n]};
|
||||
unique_ptr<float[]> hY{new float[n]};
|
||||
unique_ptr<float[]> hOut{new float[n]};
|
||||
|
||||
for (size_t i = 0; i < n; ++i) {
|
||||
hX[i] = static_cast<float>(i);
|
||||
hY[i] = static_cast<float>(i * 2);
|
||||
}
|
||||
|
||||
hipDeviceptr_t dX, dY, dOut;
|
||||
hipMalloc(&dX, bufferSize);
|
||||
hipMalloc(&dY, bufferSize);
|
||||
hipMalloc(&dOut, bufferSize);
|
||||
hipMemcpyHtoD(dX, hX.get(), bufferSize);
|
||||
hipMemcpyHtoD(dY, hY.get(), bufferSize);
|
||||
|
||||
struct {
|
||||
float a_;
|
||||
hipDeviceptr_t b_;
|
||||
hipDeviceptr_t c_;
|
||||
hipDeviceptr_t d_;
|
||||
size_t e_;
|
||||
} args{a, dX, dY, dOut, n};
|
||||
|
||||
auto size = sizeof(args);
|
||||
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
|
||||
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
||||
HIP_LAUNCH_PARAM_END};
|
||||
|
||||
hipModuleLaunchKernel(kernel, NUM_BLOCKS, 1, 1, NUM_THREADS, 1, 1,
|
||||
0, nullptr, nullptr, config);
|
||||
|
||||
hipMemcpyDtoH(hOut.get(), dOut, bufferSize);
|
||||
|
||||
for (size_t i = 0; i < n; ++i) {
|
||||
if (a * hX[i] + hY[i] != hOut[i]) { failed("Validation failed."); }
|
||||
}
|
||||
|
||||
hipFree(dX);
|
||||
hipFree(dY);
|
||||
hipFree(dOut);
|
||||
|
||||
hipModuleUnload(module);
|
||||
|
||||
passed();
|
||||
}
|
||||
Исполняемый файл
Двоичные данные
Двоичный файл не отображается.
Ссылка в новой задаче
Block a user