forked from ROCm/rocprofiler-systems
-
Notifications
You must be signed in to change notification settings - Fork 0
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
* Porting from ROCm/omnitrace#411 * Improve OMPT support * Add OpenMP target example to testing * Update Timemory submodule to use ROCm/Timemory rather than NERSC/Timemory Signed-off-by: David Galiffi <[email protected]>
- Loading branch information
1 parent
21d9ab7
commit 99db079
Showing
13 changed files
with
757 additions
and
81 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,108 @@ | ||
| # | ||
| # | ||
| # | ||
| cmake_minimum_required(VERSION 3.18.4 FATAL_ERROR) | ||
|
|
||
| # try to find a compatible HIP version | ||
| if(ROCmVersion_MAJOR_VERSION AND ROCmVersion_MAJOR_VERSION GREATER_EQUAL 6) | ||
| find_package(hip ${ROCmVersion_MAJOR_VERSION}.0.0) | ||
| else() | ||
| find_package(hip) | ||
| endif() | ||
|
|
||
| if(NOT hip_FOUND) | ||
| message(WARNING "ROCm >= 5.6 not found. Skipping OpenMP target example.") | ||
| return() | ||
| elseif(hip_FOUND AND hip_VERSION VERSION_LESS 5.6.0) | ||
| message( | ||
| WARNING | ||
| "ROCm >= 5.6 not found (found ${hip_VERSION}). Skipping OpenMP target example." | ||
| ) | ||
| return() | ||
| endif() | ||
|
|
||
| if(NOT OMP_TARGET_COMPILER) | ||
| find_program( | ||
| amdclangpp_EXECUTABLE | ||
| NAMES amdclang++ | ||
| HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm | ||
| PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm | ||
| PATH_SUFFIXES bin llvm/bin) | ||
| mark_as_advanced(amdclangpp_EXECUTABLE) | ||
|
|
||
| if(amdclangpp_EXECUTABLE) | ||
| set(OMP_TARGET_COMPILER | ||
| "${amdclangpp_EXECUTABLE}" | ||
| CACHE FILEPATH "OpenMP target compiler") | ||
| else() | ||
| message(WARNING "OpenMP target compiler not found. Skipping this example.") | ||
| return() | ||
| endif() | ||
| endif() | ||
|
|
||
| project(rocprofiler-systems-example-openmp-target-lib LANGUAGES CXX) | ||
|
|
||
| set(CMAKE_BUILD_TYPE "RelWithDebInfo") | ||
|
|
||
| set(DEFAULT_GPU_TARGETS | ||
| "gfx900" | ||
| "gfx906" | ||
| "gfx908" | ||
| "gfx90a" | ||
| "gfx940" | ||
| "gfx941" | ||
| "gfx942" | ||
| "gfx1030" | ||
| "gfx1010" | ||
| "gfx1100" | ||
| "gfx1101" | ||
| "gfx1102") | ||
|
|
||
| set(GPU_TARGETS | ||
| "${DEFAULT_GPU_TARGETS}" | ||
| CACHE STRING "GPU targets to compile for") | ||
|
|
||
| find_package(Threads REQUIRED) | ||
|
|
||
| add_library(openmp-target-lib SHARED) | ||
| target_sources(openmp-target-lib PRIVATE library.cpp) | ||
| target_link_libraries(openmp-target-lib PUBLIC Threads::Threads) | ||
| target_compile_options(openmp-target-lib PRIVATE -fopenmp -ggdb) | ||
| target_link_options(openmp-target-lib PUBLIC -fopenmp) | ||
|
|
||
| foreach(_TARGET ${GPU_TARGETS}) | ||
| target_compile_options(openmp-target-lib PRIVATE --offload-arch=${_TARGET}) | ||
| target_link_options(openmp-target-lib PUBLIC --offload-arch=${_TARGET}) | ||
| endforeach() | ||
|
|
||
| message(STATUS "Using OpenMP target compiler: ${OMP_TARGET_COMPILER}") | ||
|
|
||
| get_filename_component(OMP_TARGET_COMPILER_DIR ${OMP_TARGET_COMPILER} PATH) | ||
| get_filename_component(OMP_TARGET_COMPILER_DIR ${OMP_TARGET_COMPILER_DIR} PATH) | ||
|
|
||
| message(STATUS "Using OpemMP target compiler directory: ${OMP_TARGET_COMPILER_DIR}") | ||
|
|
||
| if(NOT EXISTS ${OMP_TARGET_COMPILER_DIR}/llvm/lib) | ||
| message(FATAL_ERROR "${OMP_TARGET_COMPILER_DIR}/llvm/lib does not exist") | ||
| endif() | ||
| set_target_properties( | ||
| openmp-target-lib | ||
| PROPERTIES BUILD_RPATH | ||
| "${OMP_TARGET_COMPILER_DIR}/llvm/lib:${OMP_TARGET_COMPILER_DIR}/lib" | ||
| OUTPUT_NAME "openmp-target" | ||
| POSITION_INDEPENDENT_CODE ON) | ||
|
|
||
| rocprofiler_systems_custom_compilation(TARGET openmp-target-lib COMPILER ${OMP_TARGET_COMPILER}) | ||
|
|
||
| add_executable(openmp-target) | ||
| target_sources(openmp-target PRIVATE main.cpp) | ||
| target_link_libraries(openmp-target PRIVATE openmp-target-lib) | ||
| target_compile_options(openmp-target PRIVATE -ggdb) | ||
|
|
||
| set_target_properties( | ||
| openmp-target | ||
| PROPERTIES BUILD_RPATH | ||
| "${OMP_TARGET_COMPILER_DIR}/llvm/lib:${OMP_TARGET_COMPILER_DIR}/lib" | ||
| POSITION_INDEPENDENT_CODE ON) | ||
|
|
||
| rocprofiler_systems_custom_compilation(TARGET openmp-target COMPILER ${OMP_TARGET_COMPILER}) |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,149 @@ | ||
| // MIT License | ||
| // | ||
| // Copyright (c) 2023 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 <cstdlib> | ||
| #include <fstream> | ||
| #include <iostream> | ||
| #include <math.h> | ||
| #include <sstream> | ||
| #include <stdio.h> | ||
| #include <string> | ||
| #include <thread> | ||
| #include <unistd.h> | ||
|
|
||
| constexpr float EPS_FLOAT = 1.0e-7f; | ||
| constexpr double EPS_DOUBLE = 1.0e-15; | ||
|
|
||
| #pragma omp declare target | ||
| template <typename T> | ||
| T | ||
| mul(T a, T b) | ||
| { | ||
| T c; | ||
| c = a * b; | ||
| return c; | ||
| } | ||
| #pragma omp end declare target | ||
|
|
||
| template <typename T> | ||
| void | ||
| vmul(T* a, T* b, T* c, int N) | ||
| { | ||
| #pragma omp target map(to : a [0:N], b [0:N]) map(from : c [0:N]) | ||
| #pragma omp teams distribute parallel for | ||
| for(int i = 0; i < N; i++) | ||
| { | ||
| c[i] = mul(a[i], b[i]); | ||
| } | ||
| } | ||
|
|
||
| int | ||
| run_impl() | ||
| { | ||
| std::this_thread::sleep_for(std::chrono::milliseconds{ 50 }); | ||
|
|
||
| constexpr int N = 100000; | ||
| constexpr int Nc = N / 100; | ||
| int a_i[N], b_i[N], c_i[N], validate_i[N]; | ||
| float a_f[N], b_f[N], c_f[N], validate_f[N]; | ||
| double a_d[N], b_d[N], c_d[N], validate_d[N]; | ||
| int N_errors = 0; | ||
| bool flag = false; | ||
|
|
||
| #pragma omp parallel for schedule(dynamic, Nc) | ||
| for(int i = 0; i < N; ++i) | ||
| { | ||
| a_f[i] = a_i[i] = i + 1; | ||
| b_f[i] = b_i[i] = i + 2; | ||
| a_d[i] = a_i[i]; | ||
| b_d[i] = b_i[i]; | ||
| validate_i[i] = a_i[i] * b_i[i]; | ||
| validate_f[i] = a_f[i] * b_f[i]; | ||
| validate_d[i] = a_d[i] * b_d[i]; | ||
| } | ||
|
|
||
| for(int i = 0; i < 2; ++i) | ||
| { | ||
| vmul(a_i, b_i, c_i, N); | ||
| vmul(a_f, b_f, c_f, N); | ||
| vmul(a_d, b_d, c_d, N); | ||
| } | ||
|
|
||
| for(int i = 0; i < N; i++) | ||
| { | ||
| if(c_i[i] != validate_i[i]) | ||
| { | ||
| ++N_errors; | ||
| // print 1st bad index | ||
| if(!flag) | ||
| { | ||
| printf("First fail: c_i[%d](%d) != validate_i[%d](%d)\n", i, c_i[i], i, | ||
| validate_i[i]); | ||
| flag = true; | ||
| } | ||
| } | ||
| } | ||
| flag = false; | ||
| for(int i = 0; i < N; i++) | ||
| { | ||
| if(fabs(c_f[i] - validate_f[i]) > EPS_FLOAT) | ||
| { | ||
| ++N_errors; | ||
| // print 1st bad index | ||
| if(!flag) | ||
| { | ||
| printf("First fail: c_f[%d](%f) != validate_f[%d](%f)\n", i, | ||
| static_cast<double>(c_f[i]), i, | ||
| static_cast<double>(validate_f[i])); | ||
| flag = true; | ||
| } | ||
| } | ||
| } | ||
| flag = false; | ||
| for(int i = 0; i < N; i++) | ||
| { | ||
| if(fabs(c_d[i] - validate_d[i]) > EPS_DOUBLE) | ||
| { | ||
| ++N_errors; | ||
| // print 1st bad index | ||
| if(!flag) | ||
| { | ||
| printf("First fail: c_d[%d](%f) != validate_d[%d](%f)\n", i, c_d[i], i, | ||
| validate_d[i]); | ||
| flag = true; | ||
| } | ||
| } | ||
| } | ||
|
|
||
| return N_errors; | ||
| } | ||
|
|
||
| int | ||
| run() | ||
| { | ||
| #pragma omp parallel | ||
| { | ||
| run_impl(); | ||
| } | ||
|
|
||
| return 0; | ||
| } |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,52 @@ | ||
| // MIT License | ||
| // | ||
| // Copyright (c) 2023 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 <cstdlib> | ||
| #include <fstream> | ||
| #include <iostream> | ||
| #include <math.h> | ||
| #include <sstream> | ||
| #include <stdio.h> | ||
| #include <string> | ||
| #include <unistd.h> | ||
|
|
||
| extern int | ||
| run(); | ||
|
|
||
| int | ||
| main() | ||
| { | ||
| auto N_errors = run(); | ||
| auto _ec = EXIT_SUCCESS; | ||
| if(N_errors == 0) | ||
| { | ||
| printf("Success\n"); | ||
| } | ||
| else | ||
| { | ||
| printf("Total %d failures\n", N_errors); | ||
| printf("Fail\n"); | ||
| _ec = EXIT_FAILURE; | ||
| } | ||
|
|
||
| return _ec; | ||
| } |
Submodule timemory
updated
9 files
Oops, something went wrong.