diff --git a/CMakeLists.txt b/CMakeLists.txt index 30aa071515..30a699c7ce 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -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) diff --git a/docs/markdown/CUDNN_API_supported_by_HIP.md b/docs/markdown/CUDNN_API_supported_by_HIP.md index 43e7ca419e..2a3fa1a67a 100644 --- a/docs/markdown/CUDNN_API_supported_by_HIP.md +++ b/docs/markdown/CUDNN_API_supported_by_HIP.md @@ -293,7 +293,7 @@ |`cudnnScaleTensor` |`hipdnnScaleTensor` | |`cudnnCreateFilterDescriptor` |`hipdnnCreateFilterDescriptor` | |`cudnnSetFilter4dDescriptor` |`hipdnnSetFilter4dDescriptor` | -|`cudnnGetFilter4dDescriptor` | | +|`cudnnGetFilter4dDescriptor` |`hipdnnGetFilter4dDescriptor` | |`cudnnSetFilterNdDescriptor` |`hipdnnSetFilterNdDescriptor` | |`cudnnGetFilterNdDescriptor` |`hipdnnGetFilterNdDescriptor` | |`cudnnDestroyFilterDescriptor` |`hipdnnDestroyFilterDescriptor` | diff --git a/hipify-clang/README.md b/hipify-clang/README.md index 86e05a9567..393fd4fc3d 100644 --- a/hipify-clang/README.md +++ b/hipify-clang/README.md @@ -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 ``` - ### Windows On Windows 10 the following configurations are tested: diff --git a/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp b/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp index 85e138a8fd..a52c392b72 100644 --- a/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp +++ b/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp @@ -76,7 +76,7 @@ const std::map 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}}, diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index baa66e064c..5ea83c8376 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/hipify-clang/src/HipifyAction.cpp @@ -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; diff --git a/include/hip/hcc_detail/hiprtc.h b/include/hip/hcc_detail/hiprtc.h new file mode 100644 index 0000000000..d4174e4ce1 --- /dev/null +++ b/include/hip/hcc_detail/hiprtc.h @@ -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 +#include + +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 + + #if defined(_WIN32) + #include + + template + hiprtcResult hiprtcGetTypeName(std::string*) = delete; + #else + template + 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 \ No newline at end of file diff --git a/include/hip/hcc_detail/program_state.hpp b/include/hip/hcc_detail/program_state.hpp index ca8b6dcd88..19db8f9e0d 100644 --- a/include/hip/hcc_detail/program_state.hpp +++ b/include/hip/hcc_detail/program_state.hpp @@ -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, diff --git a/include/hip/hiprtc.h b/include/hip/hiprtc.h new file mode 100644 index 0000000000..22d78d2656 --- /dev/null +++ b/include/hip/hiprtc.h @@ -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 + +#if defined(__HIP_PLATFORM_HCC__) && !defined(__HIP_PLATFORM_NVCC__) + #include +#elif defined(__HIP_PLATFORM_NVCC__) && !defined(__HIP_PLATFORM_HCC__) + #include +#else + #error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); +#endif \ No newline at end of file diff --git a/packaging/hip_hcc.txt b/packaging/hip_hcc.txt index fe866e47f9..5aebc6c36d 100644 --- a/packaging/hip_hcc.txt +++ b/packaging/hip_hcc.txt @@ -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) diff --git a/src/hip_device.cpp b/src/hip_device.cpp index 6570b73316..746eaa2ddf 100644 --- a/src/hip_device.cpp +++ b/src/hip_device.cpp @@ -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); } } diff --git a/src/hip_hcc_internal.h b/src/hip_hcc_internal.h index b40fac93a5..769d6b7914 100644 --- a/src/hip_hcc_internal.h +++ b/src/hip_hcc_internal.h @@ -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>> kernargs; ~ihipModule_t() { if (executable.handle) hsa_executable_destroy(executable); diff --git a/src/hip_module.cpp b/src/hip_module.cpp index e029ec6e78..e8a8801e98 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -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> _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 tmp{}; + if (kernelParams) { + if (extra) return hipErrorInvalidValue; + + for (auto&& x : f->_kernarg_layout) { + const auto p{static_cast(*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(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( - 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()); diff --git a/src/hiprtc.cpp b/src/hiprtc.cpp new file mode 100644 index 0000000000..aadf48c3ed --- /dev/null +++ b/src/hiprtc.cpp @@ -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 + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +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> programs; + static std::mutex mtx; + + // DATA + std::vector> headers; + std::vector> names; + std::vector loweredNames; + std::vector elf; + std::string source; + std::string name; + std::string log; + bool compiled; + + // STATICS + static + hiprtcResult destroy(_hiprtcProgram* p) + { + using namespace std; + + lock_guard 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> h) + { + using namespace std; + + unique_ptr<_hiprtcProgram> tmp{new _hiprtcProgram{move(h), {}, {}, {}, + move(s), move(n), {}, + false}}; + + lock_guard 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& 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& 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& 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> fut{headers.size()}; + transform(headers.cbegin(), headers.cend(), begin(fut), + [&](const pair& 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> _hiprtcProgram::programs{}; +std::mutex _hiprtcProgram::mtx{}; + +namespace +{ + inline + bool isValidProgram(const hiprtcProgram p) + { + if (!p) return false; + + std::lock_guard 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 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& 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 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> 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& 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; +} \ No newline at end of file diff --git a/src/program_state.inl b/src/program_state.inl index 9729da8115..f1397b3fe9 100644 --- a/src/program_state.inl +++ b/src/program_state.inl @@ -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> 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> 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< diff --git a/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu b/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu index 6aae90e403..5a2a993d54 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu @@ -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( diff --git a/tests/hipify-clang/unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp b/tests/hipify-clang/unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp index 2898587143..2bf196c0cd 100644 --- a/tests/hipify-clang/unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp +++ b/tests/hipify-clang/unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp @@ -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 diff --git a/tests/hipify-clang/unit_tests/samples/2_Cookbook/7_streams/stream.cpp b/tests/hipify-clang/unit_tests/samples/2_Cookbook/7_streams/stream.cpp new file mode 100644 index 0000000000..9da1dca9af --- /dev/null +++ b/tests/hipify-clang/unit_tests/samples/2_Cookbook/7_streams/stream.cpp @@ -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 +// CHECK: #include +#include + +#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 <<>>(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 <<>>(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; +} diff --git a/tests/hipify-clang/unit_tests/samples/dynamic_shared_memory.cu b/tests/hipify-clang/unit_tests/samples/dynamic_shared_memory.cu index 9e00a9b8e3..7bdd6cafbe 100644 --- a/tests/hipify-clang/unit_tests/samples/dynamic_shared_memory.cu +++ b/tests/hipify-clang/unit_tests/samples/dynamic_shared_memory.cu @@ -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; diff --git a/tests/hipify-clang/unit_tests/samples/static_shared_memory.cu b/tests/hipify-clang/unit_tests/samples/static_shared_memory.cu index f53932bf74..e8108ead3b 100644 --- a/tests/hipify-clang/unit_tests/samples/static_shared_memory.cu +++ b/tests/hipify-clang/unit_tests/samples/static_shared_memory.cu @@ -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; diff --git a/tests/hit/HIT.cmake b/tests/hit/HIT.cmake index f62ba65657..d4c162f4b9 100644 --- a/tests/hit/HIT.cmake +++ b/tests/hit/HIT.cmake @@ -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() diff --git a/tests/src/hiprtc/hiprtcGetLoweredName.cpp b/tests/src/hiprtc/hiprtcGetLoweredName.cpp new file mode 100644 index 0000000000..e3fa057a81 --- /dev/null +++ b/tests/src/hiprtc/hiprtcGetLoweredName.cpp @@ -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 + +#include +#include + +#include +#include +#include + + +static constexpr const char gpu_program[]{ +R"( +#include + +__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 +__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 kernel_name_vec; + vector variable_name_vec; + vector variable_initial_value; + vector 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"); + expected_result.push_back(sizeof(int)); + kernel_name_vec.push_back("f3"); + 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 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(); +} diff --git a/tests/src/hiprtc/hiprtcGetTypeName.cpp b/tests/src/hiprtc/hiprtcGetTypeName.cpp new file mode 100644 index 0000000000..b0348408f3 --- /dev/null +++ b/tests/src/hiprtc/hiprtcGetTypeName.cpp @@ -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 + +#define HIPRTC_GET_TYPE_NAME +#include +#include + +#include +#include +#include + +static constexpr auto gpu_program{ +R"( +#include + +namespace N1 { struct S1_t { int i; double d; }; } +template +__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 +std::string getKernelNameForType(void) +{ + std::string type_name; + hiprtcGetTypeName(&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 name_vec; + vector expected_result; + + name_vec.push_back(getKernelNameForType()); + expected_result.push_back(sizeof(int)); + name_vec.push_back(getKernelNameForType()); + expected_result.push_back(sizeof(double)); + name_vec.push_back(getKernelNameForType()); + 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 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(); +} diff --git a/tests/src/hiprtc/saxpy.cpp b/tests/src/hiprtc/saxpy.cpp new file mode 100644 index 0000000000..5f9dc7a125 --- /dev/null +++ b/tests/src/hiprtc/saxpy.cpp @@ -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 + +#include +#include + +#include +#include +#include +#include +#include + +static constexpr auto NUM_THREADS{128}; +static constexpr auto NUM_BLOCKS{32}; + +static constexpr auto saxpy{ +R"( +#include + +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 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 hX{new float[n]}; + unique_ptr hY{new float[n]}; + unique_ptr hOut{new float[n]}; + + for (size_t i = 0; i < n; ++i) { + hX[i] = static_cast(i); + hY[i] = static_cast(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(); +} diff --git a/tests/src/runtimeApi/module/vcpy_kernel.code b/tests/src/runtimeApi/module/vcpy_kernel.code new file mode 100755 index 0000000000..737ee1ce9b Binary files /dev/null and b/tests/src/runtimeApi/module/vcpy_kernel.code differ