Code Coverage Support (#46)

* Code-coverage support

* Examples update

- code-coverage example
- tweak transpose and parallel-overhead

* Coverage output + testing

- config::get_setting value(...)
- REGULAR_EXPRESSION -> REGEX in cmake func args
- coverage.hpp header
- coverage JSON
- coverage tests

* cmake formatting

* Library instrumentation w/o main + more

- fixed library instrumentation w/o main
- use TIMEMORY_PROJECT_NAME in output messages
- removed '--driver' option from omnitrace exe
- support coverage in trace mode
- OMNITRACE_KOKKOS_KERNEL_LOGGER
- support multiple calls to omnitrace_set_env after init if already called
- support multiple calls to omnitrace_set_mpi after init if same args
- support multiple calls to omnitrace_init if same mode
- unique_ptr_t for thread_data which calls finalize when thread_data is destroyed
- tweaked openmp tests
- improved finalization

* Replace CI --output-on-failure with -V

* Fix to OMNITRACE_DL_INVOKE

* omnitrace-exe and testing updates

- omnitrace::omnitrace-timemory interface library
- support for configs in omnitrace exe
- print-{available,instrumented,...} opts no longer exit w/o --simulate
- all tests apply --print-instrumented functions
- tweaked coverage tests
- print-* options print instructions not address range

* Remove OMNITRACE_DEBUG_FINALIZE=ON from CI

* Python cmake tweaks

* Tweak test ordering

* Upload CI artifacts if fail or success

* CI Python tweaks

- Use OMNITRACE_PYTHON_PREFIX and OMNITRACE_PYTHON_ENVS

* CI ELFULTILS_DOWNLOAD_VERSION

* test tweaks

- labels and more coverage tests

* tweak to omnitrace --config handling

* Update module/function constraint handling + PP

- tweak pre-processor definition handling
- removed free-standing module_constraint
- remove free-standing routine_constraint
- remove module_name.find("omnitrace") module constraint
- fully handle the output path of omnitrace *-instr files
- get_use_code_coverage config option
- print-coverage option
- coverage_module_functions

* use github.job not github.name

* Re-enable HSA_ENABLE_INTERRUPT

- remove coverage address report

examples/CMakeLists.txt

@@ -7,6 +7,7 @@ set(CMAKE_CXX_VISIBILITY_PRESET "default")
 
 add_subdirectory(transpose)
 add_subdirectory(parallel-overhead)
+add_subdirectory(code-coverage)
 add_subdirectory(user-api)
 add_subdirectory(openmp)
 

examples/code-coverage/CMakeLists.txt

@@ -0,0 +1,16 @@
+cmake_minimum_required(VERSION 3.15 FATAL_ERROR)
+
+project(omnitrace-code-coverage LANGUAGES CXX)
+
+set(CMAKE_BUILD_TYPE "RelWithDebInfo")
+string(REPLACE " " ";" _FLAGS "${CMAKE_CXX_FLAGS_DEBUG}")
+
+find_package(Threads REQUIRED)
+add_executable(code-coverage code-coverage.cpp)
+target_link_libraries(code-coverage Threads::Threads)
+target_compile_options(code-coverage PRIVATE ${_FLAGS})
+
+if(NOT CMAKE_PROJECT_NAME STREQUAL PROJECT_NAME)
+    set_target_properties(code-coverage PROPERTIES RUNTIME_OUTPUT_DIRECTORY
+                                                   ${CMAKE_BINARY_DIR})
+endif()

examples/code-coverage/code-coverage.cpp

@@ -0,0 +1,88 @@
+
+#include <atomic>
+#include <cstdio>
+#include <cstdlib>
+#include <string>
+#include <thread>
+#include <vector>
+
+#define NOINLINE __attribute__((noinline))
+
+std::atomic<long> total{ 0 };
+
+long
+fib(long n) NOINLINE;
+
+void
+run_real(size_t nitr, long) NOINLINE;
+
+void
+run_fake(size_t nitr, long) NOINLINE;
+
+int
+main(int argc, char** argv)
+{
+    using exec_t = void (*)(size_t, long);
+
+    std::string _name = argv[0];
+    auto        _pos  = _name.find_last_of('/');
+    if(_pos != std::string::npos) _name = _name.substr(_pos + 1);
+
+    size_t nthread = std::min<size_t>(16, std::thread::hardware_concurrency());
+    size_t nitr    = 5000;
+    long   nfib    = 10;
+
+    if(argc > 1) nfib = atol(argv[1]);
+    if(argc > 2) nthread = atol(argv[2]);
+    if(argc > 3) nitr = atol(argv[3]);
+
+    exec_t _exec = &run_real;
+
+    // ensure that compiler cannot optimize run_fake away
+    if(std::getenv("CODE_COVERAGE_USE_FAKE") != nullptr) _exec = &run_fake;
+
+    printf("[%s] Threads: %zu\n[%s] Iterations: %zu\n[%s] fibonacci(%li)...\n",
+           _name.c_str(), nthread, _name.c_str(), nitr, _name.c_str(), nfib);
+
+    std::vector<std::thread> threads{};
+    for(size_t i = 0; i < nthread; ++i)
+    {
+        size_t _nitr = ((i % 2) == 1) ? (nitr - (0.1 * nitr)) : (nitr + (0.1 * nitr));
+        _nitr        = std::max<size_t>(_nitr, 1);
+        threads.emplace_back(_exec, _nitr, nfib);
+    }
+
+    auto _nitr = std::max<size_t>(nitr - 0.25 * nitr, 1);
+    (*_exec)(_nitr, nfib - 0.1 * nfib);
+    for(auto& itr : threads)
+        itr.join();
+
+    printf("[%s] fibonacci(%li) x %lu = %li\n", _name.c_str(), nfib, nthread,
+           total.load());
+
+    return 0;
+}
+
+long
+fib(long n)
+{
+    return (n < 2) ? n : fib(n - 1) + fib(n - 2);
+}
+
+void
+run_real(size_t nitr, long n)
+{
+    long local = 0;
+    for(size_t i = 0; i < nitr; ++i)
+        local += fib(n);
+    total += local;
+}
+
+void
+run_fake(size_t nitr, long n)
+{
+    long local = 0;
+    for(size_t i = 0; i < nitr; ++i)
+        local += fib(n);
+    total += local;
+}

examples/parallel-overhead/parallel-overhead.cpp

@@ -2,6 +2,7 @@
 #include <atomic>
 #include <cstdio>
 #include <cstdlib>
+#include <string>
 #include <thread>
 #include <vector>
 
@@ -29,28 +30,36 @@ run(size_t nitr, long n)
 int
 main(int argc, char** argv)
 {
+    std::string _name = argv[0];
+    auto        _pos  = _name.find_last_of('/');
+    if(_pos != std::string::npos) _name = _name.substr(_pos + 1);
+
     size_t nthread = std::min<size_t>(16, std::thread::hardware_concurrency());
     size_t nitr    = 50000;
     long   nfib    = 10;
+
     if(argc > 1) nfib = atol(argv[1]);
     if(argc > 2) nthread = atol(argv[2]);
     if(argc > 3) nitr = atol(argv[3]);
 
-    printf("[%s] Threads: %zu\n[%s] Iterations: %zu\n[%s] fibonacci(%li)...\n", argv[0],
-           nthread, argv[0], nitr, argv[0], nfib);
+    printf("[%s] Threads: %zu\n[%s] Iterations: %zu\n[%s] fibonacci(%li)...\n",
+           _name.c_str(), nthread, _name.c_str(), nitr, _name.c_str(), nfib);
 
     std::vector<std::thread> threads{};
     for(size_t i = 0; i < nthread; ++i)
     {
         size_t _nitr = ((i % 2) == 1) ? (nitr - (0.1 * nitr)) : (nitr + (0.1 * nitr));
+        _nitr        = std::max<size_t>(_nitr, 1);
         threads.emplace_back(&run, _nitr, nfib);
     }
 
-    run(nitr - 0.25 * nitr, nfib - 0.1 * nfib);
+    auto _nitr = std::max<size_t>(nitr - 0.25 * nitr, 1);
+    run(_nitr, nfib - 0.1 * nfib);
     for(auto& itr : threads)
         itr.join();
 
-    printf("[%s] fibonacci(%li) x %lu = %li\n", argv[0], nfib, nthread, total.load());
+    printf("[%s] fibonacci(%li) x %lu = %li\n", _name.c_str(), nfib, nthread,
+           total.load());
 
     return 0;
 }

examples/transpose/transpose.cpp

@@ -104,33 +104,36 @@ run(int rank, int tid, hipStream_t stream, int argc, char** argv)
     std::cout << "[" << rank << "][" << tid << "] M: " << M << " N: " << N << std::endl;
     _lk.unlock();
 
-    size_t size   = sizeof(int) * M * N;
-    int*   matrix = new int[size];
+    size_t size       = sizeof(int) * M * N;
+    int*   inp_matrix = new int[size];
+    int*   out_matrix = new int[size];
     for(size_t i = 0; i < M * N; i++)
-        matrix[i] = rand() % 1002;
+    {
+        inp_matrix[i] = rand() % 1002;
+        out_matrix[i] = 0;
+    }
     int* in  = nullptr;
     int* out = nullptr;
 
-    std::chrono::high_resolution_clock::time_point t1, t2;
-
     HIP_API_CALL(hipMalloc(&in, size));
     HIP_API_CALL(hipMalloc(&out, size));
-    HIP_API_CALL(hipMemset(in, 0, size));
-    HIP_API_CALL(hipMemset(out, 0, size));
-    HIP_API_CALL(hipMemcpy(in, matrix, size, hipMemcpyHostToDevice));
-    HIP_API_CALL(hipDeviceSynchronize());
+    HIP_API_CALL(hipMemsetAsync(in, 0, size, stream));
+    HIP_API_CALL(hipMemsetAsync(out, 0, size, stream));
+    HIP_API_CALL(hipMemcpyAsync(in, inp_matrix, size, hipMemcpyHostToDevice, stream));
+    HIP_API_CALL(hipStreamSynchronize(stream));
 
     dim3 grid(M / 32, N / 32, 1);
     dim3 block(32, 32, 1);  // transpose_a
 
-    t1 = std::chrono::high_resolution_clock::now();
+    auto t1 = std::chrono::high_resolution_clock::now();
     for(size_t i = 0; i < nitr; i++)
     {
         transpose_a<<<grid, block, 0, stream>>>(in, out, M, N);
         check_hip_error();
     }
+    auto t2 = std::chrono::high_resolution_clock::now();
     HIP_API_CALL(hipStreamSynchronize(stream));
-    t2 = std::chrono::high_resolution_clock::now();
+    HIP_API_CALL(hipMemcpyAsync(out_matrix, out, size, hipMemcpyDeviceToHost, stream));
     double time =
         std::chrono::duration_cast<std::chrono::duration<double>>(t2 - t1).count();
     float GB = (float) size * nitr * 2 / (1 << 30);
@@ -142,18 +145,15 @@ run(int rank, int tid, hipStream_t stream, int argc, char** argv)
               << std::endl;
     print_lock.unlock();
 
-    HIP_API_CALL(hipDeviceSynchronize());
-
-    int* out_matrix = new int[size];
-    HIP_API_CALL(hipMemcpy(out_matrix, out, size, hipMemcpyDeviceToHost));
+    HIP_API_CALL(hipStreamSynchronize(stream));
 
     // cpu_transpose(matrix, out_matrix, M, N);
-    verify(matrix, out_matrix, M, N);
+    verify(inp_matrix, out_matrix, M, N);
 
     HIP_API_CALL(hipFree(in));
     HIP_API_CALL(hipFree(out));
 
-    delete[] matrix;
+    delete[] inp_matrix;
     delete[] out_matrix;
 }