diff options
Diffstat (limited to 'Tests/HIP')
31 files changed, 597 insertions, 0 deletions
diff --git a/Tests/HIP/ArchitectureOff/CMakeLists.txt b/Tests/HIP/ArchitectureOff/CMakeLists.txt new file mode 100644 index 000000000..bccb3b4f8 --- /dev/null +++ b/Tests/HIP/ArchitectureOff/CMakeLists.txt @@ -0,0 +1,8 @@ +cmake_minimum_required(VERSION 3.18) +project(HIPArchitecture HIP) + +# Make sure CMake doesn't pass architectures if HIP_ARCHITECTURES is OFF. +string(APPEND CMAKE_HIP_FLAGS " --offload-arch=gfx908") + +add_executable(HIPOnlyArchitectureOff main.hip) +set_property(TARGET HIPOnlyArchitectureOff PROPERTY HIP_ARCHITECTURES OFF) diff --git a/Tests/HIP/ArchitectureOff/main.hip b/Tests/HIP/ArchitectureOff/main.hip new file mode 100644 index 000000000..9256318cc --- /dev/null +++ b/Tests/HIP/ArchitectureOff/main.hip @@ -0,0 +1,9 @@ +#ifdef __HIP_DEVICE_COMPILE__ +# ifndef __gfx908__ +# error "Passed architecture gfx908, but got something else." +# endif +#endif + +int main() +{ +} diff --git a/Tests/HIP/CMakeLists.txt b/Tests/HIP/CMakeLists.txt new file mode 100644 index 000000000..9499be8f6 --- /dev/null +++ b/Tests/HIP/CMakeLists.txt @@ -0,0 +1,15 @@ +macro (add_hip_test_macro name) + add_test_macro("${name}" ${ARGN}) + set_property(TEST "${name}" APPEND + PROPERTY LABELS "HIP") +endmacro () + +add_hip_test_macro(HIP.ArchitectureOff HIPOnlyArchitectureOff) +add_hip_test_macro(HIP.CompileFlags HIPOnlyCompileFlags) +add_hip_test_macro(HIP.EnableStandard HIPEnableStandard) +add_hip_test_macro(HIP.InferHipLang1 HIPInferHipLang1) +add_hip_test_macro(HIP.InferHipLang2 HIPInferHipLang2) +add_hip_test_macro(HIP.MathFunctions HIPOnlyMathFunctions) +add_hip_test_macro(HIP.MixedLanguage HIPMixedLanguage) +add_hip_test_macro(HIP.TryCompile HIPOnlyTryCompile) +add_hip_test_macro(HIP.WithDefs HIPOnlyWithDefs) diff --git a/Tests/HIP/CompileFlags/CMakeLists.txt b/Tests/HIP/CompileFlags/CMakeLists.txt new file mode 100644 index 000000000..c80831387 --- /dev/null +++ b/Tests/HIP/CompileFlags/CMakeLists.txt @@ -0,0 +1,8 @@ +cmake_minimum_required(VERSION 3.18) +project(CompileFlags HIP) + +add_executable(HIPOnlyCompileFlags main.hip) + +set_property(TARGET HIPOnlyCompileFlags PROPERTY HIP_ARCHITECTURES gfx803) + +target_compile_options(HIPOnlyCompileFlags PRIVATE -DALWAYS_DEFINE) diff --git a/Tests/HIP/CompileFlags/main.hip b/Tests/HIP/CompileFlags/main.hip new file mode 100644 index 000000000..3259f1d31 --- /dev/null +++ b/Tests/HIP/CompileFlags/main.hip @@ -0,0 +1,13 @@ +#ifdef __HIP_DEVICE_COMPILE__ +# ifndef __gfx803__ +# error "Passed architecture gfx803, but got something else." +# endif +#endif + +#ifndef ALWAYS_DEFINE +# error "ALWAYS_DEFINE not defined!" +#endif + +int main() +{ +} diff --git a/Tests/HIP/EnableStandard/CMakeLists.txt b/Tests/HIP/EnableStandard/CMakeLists.txt new file mode 100644 index 000000000..670172467 --- /dev/null +++ b/Tests/HIP/EnableStandard/CMakeLists.txt @@ -0,0 +1,20 @@ +cmake_minimum_required(VERSION 3.18) +project (EnableStandard HIP) + +set(CMAKE_CXX_COMPILER ${CMAKE_HIP_COMPILER}) +enable_language(CXX) + +#Goal for this example: +#build hip sources that require C++11 to be enabled. + +add_library(HIPStatic11 STATIC static.cxx) +set_source_files_properties(static.cxx PROPERTIES LANGUAGE HIP) + +add_library(HIPDynamic11 SHARED shared.hip) + +add_executable(HIPEnableStandard main.hip) +target_link_libraries(HIPEnableStandard PRIVATE HIPStatic11 HIPDynamic11) + +target_compile_features(HIPDynamic11 PRIVATE cxx_std_11) +set_target_properties(HIPStatic11 PROPERTIES HIP_STANDARD 11) +set_target_properties(HIPStatic11 PROPERTIES HIP_STANDARD_REQUIRED TRUE) diff --git a/Tests/HIP/EnableStandard/main.hip b/Tests/HIP/EnableStandard/main.hip new file mode 100644 index 000000000..7dc32e806 --- /dev/null +++ b/Tests/HIP/EnableStandard/main.hip @@ -0,0 +1,23 @@ + +#include <iostream> + +#ifdef _WIN32 +# define IMPORT __declspec(dllimport) +#else +# define IMPORT +#endif + +int static_hip11_func(int); +IMPORT int shared_hip11_func(int); + +void test_functions() +{ + static_hip11_func(int(42)); + shared_hip11_func(int(42)); +} + +int main(int argc, char** argv) +{ + test_functions(); + return 0; +} diff --git a/Tests/HIP/EnableStandard/shared.hip b/Tests/HIP/EnableStandard/shared.hip new file mode 100644 index 000000000..204225863 --- /dev/null +++ b/Tests/HIP/EnableStandard/shared.hip @@ -0,0 +1,15 @@ + +#include <type_traits> + +#ifdef _WIN32 +# define EXPORT __declspec(dllexport) +#else +# define EXPORT +#endif + +using tt = std::true_type; +using ft = std::false_type; +EXPORT int __host__ shared_hip11_func(int x) +{ + return x * x + std::integral_constant<int, 17>::value; +} diff --git a/Tests/HIP/EnableStandard/static.cxx b/Tests/HIP/EnableStandard/static.cxx new file mode 100644 index 000000000..29558945e --- /dev/null +++ b/Tests/HIP/EnableStandard/static.cxx @@ -0,0 +1,9 @@ + +#include <type_traits> + +using tt = std::true_type; +using ft = std::false_type; +int __host__ static_hip11_func(int x) +{ + return x * x + std::integral_constant<int, 17>::value; +} diff --git a/Tests/HIP/InferHipLang1/CMakeLists.txt b/Tests/HIP/InferHipLang1/CMakeLists.txt new file mode 100644 index 000000000..63d77fdaa --- /dev/null +++ b/Tests/HIP/InferHipLang1/CMakeLists.txt @@ -0,0 +1,12 @@ +cmake_minimum_required(VERSION 3.18) +project(InferHipLang C CXX HIP) + +#Goal for this example: +#make sure that we understand that HIP is the correct link language +add_library(InterfaceWithHIP INTERFACE) +target_sources(InterfaceWithHIP INTERFACE interface.hip main.cxx) +target_compile_features(InterfaceWithHIP INTERFACE hip_std_14) +target_compile_features(InterfaceWithHIP INTERFACE cxx_std_11) + +add_executable(HIPInferHipLang1 ) +target_link_libraries(HIPInferHipLang1 PRIVATE InterfaceWithHIP) diff --git a/Tests/HIP/InferHipLang1/interface.hip b/Tests/HIP/InferHipLang1/interface.hip new file mode 100644 index 000000000..6ac8641cc --- /dev/null +++ b/Tests/HIP/InferHipLang1/interface.hip @@ -0,0 +1,19 @@ +#include <system_error> +#include <type_traits> +#include <hip/hip_runtime_api.h> + +static __global__ void fake_hip_kernel() +{ +} + + +int __host__ interface_hip_func(int x) +{ + + fake_hip_kernel<<<1, 1>>>(); + bool valid = (hipSuccess == hipGetLastError()); + if (!valid) { + throw std::system_error(ENODEV, std::generic_category(), "no hip device"); + } + return x * x + std::integral_constant<int, 17>::value; +} diff --git a/Tests/HIP/InferHipLang1/main.cxx b/Tests/HIP/InferHipLang1/main.cxx new file mode 100644 index 000000000..987b6c651 --- /dev/null +++ b/Tests/HIP/InferHipLang1/main.cxx @@ -0,0 +1,19 @@ + +#include <iostream> + +#ifdef __HIP_PLATFORM_HCC__ +# error "__HIP_PLATFORM_HCC__ propagated to C++ compilation!" +#endif + +#ifdef __HIP_ROCclr__ +# error "__HIP_ROCclr__ propagated to C++ compilation!" +#endif + +int interface_hip_func(int); + +int main(int argc, char** argv) +{ + interface_hip_func(int(42)); + + return 0; +} diff --git a/Tests/HIP/InferHipLang2/CMakeLists.txt b/Tests/HIP/InferHipLang2/CMakeLists.txt new file mode 100644 index 000000000..0e69de38f --- /dev/null +++ b/Tests/HIP/InferHipLang2/CMakeLists.txt @@ -0,0 +1,12 @@ +cmake_minimum_required(VERSION 3.18) +project(InferHipLang C CXX HIP) + +#Goal for this example: +#make sure that we understand that HIP is the correct link language +add_library(InterfaceWithHIP OBJECT) +target_sources(InterfaceWithHIP PRIVATE interface.hip main.cxx) +target_compile_features(InterfaceWithHIP INTERFACE hip_std_14) +target_compile_features(InterfaceWithHIP INTERFACE cxx_std_11) + +add_executable(HIPInferHipLang2 ) +target_link_libraries(HIPInferHipLang2 PRIVATE InterfaceWithHIP) diff --git a/Tests/HIP/InferHipLang2/interface.hip b/Tests/HIP/InferHipLang2/interface.hip new file mode 100644 index 000000000..6ac8641cc --- /dev/null +++ b/Tests/HIP/InferHipLang2/interface.hip @@ -0,0 +1,19 @@ +#include <system_error> +#include <type_traits> +#include <hip/hip_runtime_api.h> + +static __global__ void fake_hip_kernel() +{ +} + + +int __host__ interface_hip_func(int x) +{ + + fake_hip_kernel<<<1, 1>>>(); + bool valid = (hipSuccess == hipGetLastError()); + if (!valid) { + throw std::system_error(ENODEV, std::generic_category(), "no hip device"); + } + return x * x + std::integral_constant<int, 17>::value; +} diff --git a/Tests/HIP/InferHipLang2/main.cxx b/Tests/HIP/InferHipLang2/main.cxx new file mode 100644 index 000000000..987b6c651 --- /dev/null +++ b/Tests/HIP/InferHipLang2/main.cxx @@ -0,0 +1,19 @@ + +#include <iostream> + +#ifdef __HIP_PLATFORM_HCC__ +# error "__HIP_PLATFORM_HCC__ propagated to C++ compilation!" +#endif + +#ifdef __HIP_ROCclr__ +# error "__HIP_ROCclr__ propagated to C++ compilation!" +#endif + +int interface_hip_func(int); + +int main(int argc, char** argv) +{ + interface_hip_func(int(42)); + + return 0; +} diff --git a/Tests/HIP/MathFunctions/CMakeLists.txt b/Tests/HIP/MathFunctions/CMakeLists.txt new file mode 100644 index 000000000..81e3ddb26 --- /dev/null +++ b/Tests/HIP/MathFunctions/CMakeLists.txt @@ -0,0 +1,18 @@ +cmake_minimum_required(VERSION 3.18) +project(MathFunctions HIP) + +# This test covers these major HIP language/runtime requirements: +# +# 1. This makes sure CMake properly specifies the internal clang header dirs +# that hold headers needed for overloads of device side functions +# +# 2. This makes sure that all HIP include directories are properly marked as +# system includes so we don't get the following warnings: +# replacement function 'operator delete' cannot be declared 'inline'# +# +# 3. This makes sure CMake properly links to all the built-in libraries +# that hip needs that inject support for __half support +# +add_executable(HIPOnlyMathFunctions main.hip) +target_compile_options(HIPOnlyMathFunctions PRIVATE -Werror) +target_compile_features(HIPOnlyMathFunctions PRIVATE hip_std_14) diff --git a/Tests/HIP/MathFunctions/main.hip b/Tests/HIP/MathFunctions/main.hip new file mode 100644 index 000000000..8a6e77f5d --- /dev/null +++ b/Tests/HIP/MathFunctions/main.hip @@ -0,0 +1,40 @@ + +#include <stdexcept> +#include <cmath> +#include <math.h> +#include <memory> + +#include <hip/hip_runtime.h> +#include <hip/hip_fp16.h> + +namespace { +template<class T, class F> +__global__ void global_entry_point(F f, T *out) { + *out = f(); +} + +template <class T, class F> +bool verify(F f, T expected) +{ + std::unique_ptr<T> cpu_T(new T); + T* gpu_T = nullptr; + hipMalloc((void**)&gpu_T, sizeof(T)); + hipLaunchKernelGGL(global_entry_point, 1, 1, 0, 0, f, gpu_T); + hipMemcpy(cpu_T.get(), gpu_T, sizeof(T), hipMemcpyDeviceToHost); + hipFree(gpu_T); + return (*cpu_T == expected); +} +} + +int main(int argc, char** argv) +{ + bool valid = verify([]__device__(){ return std::round(1.4f); }, 1.0f); + valid &= verify([]__device__(){ return max<_Float16>(1.0f, 2.0f); }, 2.0f); + valid &= verify([]__device__(){ return min<_Float16>(1.0f, 2.0f); }, 1.0f); + + if (valid) { + return 0; + } else { + return 1; + } +} diff --git a/Tests/HIP/MixedLanguage/CMakeLists.txt b/Tests/HIP/MixedLanguage/CMakeLists.txt new file mode 100644 index 000000000..4f6dd3b0f --- /dev/null +++ b/Tests/HIP/MixedLanguage/CMakeLists.txt @@ -0,0 +1,19 @@ +cmake_minimum_required(VERSION 3.18) +project (MixedLanguage C CXX HIP) + +set(CMAKE_HIP_STANDARD 14) +set(CMAKE_CXX_STANDARD 14) + +#Goal for this example: +#make sure that we can build multiple languages into targets +#and have the link language always be HIP +add_library(MixedSharedLib SHARED shared.c) +add_library(MixedObjectLib OBJECT shared.cxx shared.hip) +set_target_properties(MixedObjectLib PROPERTIES POSITION_INDEPENDENT_CODE ON) +target_link_libraries(MixedSharedLib PRIVATE MixedObjectLib) + +add_library(MixedStaticLib STATIC static.c static.cxx static.hip) +set_target_properties(MixedStaticLib PROPERTIES POSITION_INDEPENDENT_CODE ON) + +add_executable(HIPMixedLanguage main.cxx) +target_link_libraries(HIPMixedLanguage PRIVATE MixedStaticLib MixedSharedLib) diff --git a/Tests/HIP/MixedLanguage/main.cxx b/Tests/HIP/MixedLanguage/main.cxx new file mode 100644 index 000000000..003d18ac4 --- /dev/null +++ b/Tests/HIP/MixedLanguage/main.cxx @@ -0,0 +1,40 @@ + +#include <iostream> + +#ifdef __HIP_PLATFORM_HCC__ +# error "__HIP_PLATFORM_HCC__ propagated to C++ compilation!" +#endif + +#ifdef __HIP_ROCclr__ +# error "__HIP_ROCclr__ propagated to C++ compilation!" +#endif + +#ifdef _WIN32 +# define IMPORT __declspec(dllimport) +#else +# define IMPORT +#endif + +extern "C" { +IMPORT int shared_c_func(int); +int static_c_func(int); +} + +IMPORT int shared_cxx_func(int); +IMPORT int shared_hip_func(int); + +int static_cxx_func(int); +int static_hip_func(int); + +int main(int argc, char** argv) +{ + static_c_func(int(42)); + static_cxx_func(int(42)); + static_hip_func(int(42)); + + shared_c_func(int(42)); + shared_cxx_func(int(42)); + shared_hip_func(int(42)); + + return 0; +} diff --git a/Tests/HIP/MixedLanguage/shared.c b/Tests/HIP/MixedLanguage/shared.c new file mode 100644 index 000000000..347fba925 --- /dev/null +++ b/Tests/HIP/MixedLanguage/shared.c @@ -0,0 +1,12 @@ + + +#ifdef _WIN32 +# define EXPORT __declspec(dllexport) +#else +# define EXPORT +#endif + +EXPORT int shared_c_func(int x) +{ + return -x; +} diff --git a/Tests/HIP/MixedLanguage/shared.cxx b/Tests/HIP/MixedLanguage/shared.cxx new file mode 100644 index 000000000..8e6c1d35f --- /dev/null +++ b/Tests/HIP/MixedLanguage/shared.cxx @@ -0,0 +1,21 @@ + +#include <type_traits> + +#ifdef __HIP_PLATFORM_HCC__ +# error "__HIP_PLATFORM_HCC__ propagated to C++ compilation!" +#endif + +#ifdef __HIP_ROCclr__ +# error "__HIP_ROCclr__ propagated to C++ compilation!" +#endif + +#ifdef _WIN32 +# define EXPORT __declspec(dllexport) +#else +# define EXPORT +#endif + +EXPORT int shared_cxx_func(int x) +{ + return x * x + std::integral_constant<int, 14>::value; +} diff --git a/Tests/HIP/MixedLanguage/shared.hip b/Tests/HIP/MixedLanguage/shared.hip new file mode 100644 index 000000000..e6fea9fe2 --- /dev/null +++ b/Tests/HIP/MixedLanguage/shared.hip @@ -0,0 +1,26 @@ +#include <system_error> +#include <type_traits> +#include <hip/hip_runtime_api.h> + +#ifdef _WIN32 +# define EXPORT __declspec(dllexport) +#else +# define EXPORT +#endif + + +static __global__ void fake_hip_kernel() +{ +} + + +int __host__ shared_hip_func(int x) +{ + + fake_hip_kernel<<<1, 1>>>(); + bool valid = (hipSuccess == hipGetLastError()); + if (!valid) { + throw std::system_error(ENODEV, std::generic_category(), "no hip device"); + } + return x * x + std::integral_constant<int, 17>::value; +} diff --git a/Tests/HIP/MixedLanguage/static.c b/Tests/HIP/MixedLanguage/static.c new file mode 100644 index 000000000..06c33b4cc --- /dev/null +++ b/Tests/HIP/MixedLanguage/static.c @@ -0,0 +1,6 @@ + + +int static_c_func(int x) +{ + return -x; +} diff --git a/Tests/HIP/MixedLanguage/static.cxx b/Tests/HIP/MixedLanguage/static.cxx new file mode 100644 index 000000000..2c14fb101 --- /dev/null +++ b/Tests/HIP/MixedLanguage/static.cxx @@ -0,0 +1,7 @@ + +#include <type_traits> + +int static_cxx_func(int x) +{ + return x * x + std::integral_constant<int, 14>::value; +} diff --git a/Tests/HIP/MixedLanguage/static.hip b/Tests/HIP/MixedLanguage/static.hip new file mode 100644 index 000000000..359b9facd --- /dev/null +++ b/Tests/HIP/MixedLanguage/static.hip @@ -0,0 +1,21 @@ + +#include <type_traits> +#include <system_error> +#include <hip/hip_runtime_api.h> + + +static __global__ void fake_hip_kernel() +{ +} + + +int __host__ static_hip_func(int x) +{ + + fake_hip_kernel<<<1, 1>>>(); + bool valid = (hipSuccess == hipGetLastError()); + if (!valid) { + throw std::system_error(ENODEV, std::generic_category(), "no hip device"); + } + return x * x + std::integral_constant<int, 17>::value; +} diff --git a/Tests/HIP/TryCompile/CMakeLists.txt b/Tests/HIP/TryCompile/CMakeLists.txt new file mode 100644 index 000000000..92a834c61 --- /dev/null +++ b/Tests/HIP/TryCompile/CMakeLists.txt @@ -0,0 +1,15 @@ +cmake_minimum_required(VERSION 3.18) +project (TryCompile HIP) + +#Goal for this example: +# Verify try_compile with HIP language works +set(CMAKE_HIP_STANDARD 14) +set(CMAKE_HIP_ARCHITECTURES gfx803 gfx900) + +set(CMAKE_TRY_COMPILE_TARGET_TYPE STATIC_LIBRARY) +try_compile(result "${CMAKE_CURRENT_BINARY_DIR}" + SOURCES "${CMAKE_CURRENT_SOURCE_DIR}/device_function.hip" + COPY_FILE "${CMAKE_CURRENT_BINARY_DIR}/device_function.o") + +add_executable(HIPOnlyTryCompile main.hip) +target_link_libraries(HIPOnlyTryCompile "${CMAKE_CURRENT_BINARY_DIR}/device_function.o") diff --git a/Tests/HIP/TryCompile/device_function.hip b/Tests/HIP/TryCompile/device_function.hip new file mode 100644 index 000000000..7c1205e6d --- /dev/null +++ b/Tests/HIP/TryCompile/device_function.hip @@ -0,0 +1,17 @@ +#include <system_error> +#include <hip/hip_runtime_api.h> + +static __global__ void fake_hip_kernel() +{ +} + +int __host__ try_compile_hip_func(int x) +{ + + fake_hip_kernel<<<1, 1>>>(); + bool valid = (hipSuccess == hipGetLastError()); + if (!valid) { + throw std::system_error(ENODEV, std::generic_category(), "no hip device"); + } + return x * x; +} diff --git a/Tests/HIP/TryCompile/main.hip b/Tests/HIP/TryCompile/main.hip new file mode 100644 index 000000000..091dca3bc --- /dev/null +++ b/Tests/HIP/TryCompile/main.hip @@ -0,0 +1,8 @@ +int __host__ try_compile_hip_func(int x); + +int main(int argc, char** argv) +{ + try_compile_hip_func(int(42)); + + return 0; +} diff --git a/Tests/HIP/WithDefs/CMakeLists.txt b/Tests/HIP/WithDefs/CMakeLists.txt new file mode 100644 index 000000000..e2db1820a --- /dev/null +++ b/Tests/HIP/WithDefs/CMakeLists.txt @@ -0,0 +1,37 @@ + +cmake_minimum_required(VERSION 3.18) +project (WithDefs HIP) + +set(CMAKE_HIP_ARCHITECTURES OFF) +set(release_compile_defs DEFREL) + +#Goal for this example: +#build a executable that needs to be passed a complex define through add_definitions +#this verifies we can pass C++ style attributes to hipcc +add_definitions("-DPACKED_DEFINE=[[gnu::packed]]") + +add_executable(HIPOnlyWithDefs main.hip.cpp) +set_source_files_properties(main.hip.cpp PROPERTIES LANGUAGE HIP) + +target_compile_features(HIPOnlyWithDefs PRIVATE hip_std_17) + +target_compile_options(HIPOnlyWithDefs + PRIVATE + --offload-arch=gfx900 + -DFLAG_COMPILE_LANG_$<COMPILE_LANGUAGE> + $<$<HIP_COMPILER_ID:ROCMClang>:-DFLAG_LANG_IS_HIP=$<COMPILE_LANGUAGE:HIP>> # Host-only defines are possible only on NVCC. + ) + +target_compile_definitions(HIPOnlyWithDefs + PRIVATE + $<$<CONFIG:RELEASE>:$<BUILD_INTERFACE:${release_compile_defs}>> + -DDEF_COMPILE_LANG_$<COMPILE_LANGUAGE> + -DDEF_LANG_IS_HIP=$<COMPILE_LANGUAGE:HIP> + -DDEF_HIP_COMPILER=$<HIP_COMPILER_ID> + -DDEF_HIP_COMPILER_VERSION=$<HIP_COMPILER_VERSION> + ) + +target_include_directories(HIPOnlyWithDefs + PRIVATE + $<$<COMPILE_LANGUAGE:HIP>:${CMAKE_CURRENT_SOURCE_DIR}/inc_hip> +) diff --git a/Tests/HIP/WithDefs/inc_hip/inc_hip.h b/Tests/HIP/WithDefs/inc_hip/inc_hip.h new file mode 100644 index 000000000..fe2698a02 --- /dev/null +++ b/Tests/HIP/WithDefs/inc_hip/inc_hip.h @@ -0,0 +1 @@ +#define INC_HIP diff --git a/Tests/HIP/WithDefs/main.hip.cpp b/Tests/HIP/WithDefs/main.hip.cpp new file mode 100644 index 000000000..a8f2d185b --- /dev/null +++ b/Tests/HIP/WithDefs/main.hip.cpp @@ -0,0 +1,89 @@ +#include <iostream> + +#include <hip/hip_runtime_api.h> +#include <inc_hip.h> +#ifndef INC_HIP +# error "INC_HIP not defined!" +#endif + +#ifndef PACKED_DEFINE +# error "PACKED_DEFINE not defined!" +#endif + +#ifndef FLAG_COMPILE_LANG_HIP +# error "FLAG_COMPILE_LANG_HIP not defined!" +#endif + +#ifndef FLAG_LANG_IS_HIP +# error "FLAG_LANG_IS_HIP not defined!" +#endif + +#if !FLAG_LANG_IS_HIP +# error "Expected FLAG_LANG_IS_HIP" +#endif + +#ifndef DEF_COMPILE_LANG_HIP +# error "DEF_COMPILE_LANG_HIP not defined!" +#endif + +#ifndef DEF_LANG_IS_HIP +# error "DEF_LANG_IS_HIP not defined!" +#endif + +#if !DEF_LANG_IS_HIP +# error "Expected DEF_LANG_IS_HIP" +#endif + +#ifndef DEF_HIP_COMPILER +# error "DEF_HIP_COMPILER not defined!" +#endif + +#ifndef DEF_HIP_COMPILER_VERSION +# error "DEF_HIP_COMPILER_VERSION not defined!" +#endif + +static __global__ void DetermineIfValidHIPDevice() +{ +} + +#ifdef _MSC_VER +# pragma pack(push, 1) +# undef PACKED_DEFINE +# define PACKED_DEFINE +#endif +struct PACKED_DEFINE result_type +{ + bool valid; + int value; +#if defined(NDEBUG) && !defined(DEFREL) +# error missing DEFREL flag +#endif +}; +#ifdef _MSC_VER +# pragma pack(pop) +#endif + +result_type can_launch_kernel() +{ + result_type r; + DetermineIfValidHIPDevice<<<1, 1>>>(); + r.valid = (hipSuccess == hipGetLastError()); + if (r.valid) { + r.value = 1; + } else { + r.value = -1; + } + return r; +} + +int main(int argc, char** argv) +{ + hipError_t err; + int nDevices = 0; + err = hipGetDeviceCount(&nDevices); + if (err != hipSuccess) { + std::cerr << hipGetErrorString(err) << std::endl; + return 1; + } + return 0; +} |