From 8a25b239bca21dbcd46fcc07974b505c8f5d972d Mon Sep 17 00:00:00 2001 From: Gopesh Bhardwaj Date: Tue, 6 Feb 2024 23:25:07 +0530 Subject: [PATCH] Fixing counter collection in tools and enabling tests (#436) * Fixing coutner colleciton in tools and enabling tests * fixing tests * improving coverage on test * Adding vector operations app * Fixing tools bug for counter collection * removing roctx linking --- samples/counter_collection/CMakeLists.txt | 2 +- source/lib/rocprofiler-sdk-tool/config.cpp | 4 - source/lib/rocprofiler-sdk-tool/tool.cpp | 5 +- tests/apps/CMakeLists.txt | 1 + tests/apps/vector-operations/CMakeLists.txt | 46 ++++ tests/apps/vector-operations/vector-ops.cpp | 258 ++++++++++++++++++ .../counter-collection/input1/CMakeLists.txt | 26 +- .../counter-collection/input1/validate.py | 6 +- .../counter-collection/input2/CMakeLists.txt | 25 +- .../counter-collection/input2/validate.py | 4 +- .../rocprofv3/tracing-plus-cc/CMakeLists.txt | 22 +- tests/rocprofv3/tracing-plus-cc/validate.py | 2 +- 12 files changed, 358 insertions(+), 43 deletions(-) create mode 100644 tests/apps/vector-operations/CMakeLists.txt create mode 100644 tests/apps/vector-operations/vector-ops.cpp diff --git a/samples/counter_collection/CMakeLists.txt b/samples/counter_collection/CMakeLists.txt index 59a8987d71..9ad7aa5c53 100644 --- a/samples/counter_collection/CMakeLists.txt +++ b/samples/counter_collection/CMakeLists.txt @@ -46,7 +46,7 @@ set_tests_properties( counter-collection-buffer PROPERTIES TIMEOUT - 300 + 600 LABELS "samples" ENVIRONMENT diff --git a/source/lib/rocprofiler-sdk-tool/config.cpp b/source/lib/rocprofiler-sdk-tool/config.cpp index 8463755cd0..1e1b7e67de 100644 --- a/source/lib/rocprofiler-sdk-tool/config.cpp +++ b/source/lib/rocprofiler-sdk-tool/config.cpp @@ -191,10 +191,6 @@ parse_counters(std::string line) { counters.emplace(counter); } - else - { - LOG(ERROR) << "invalid counter: " << counter; - } } } diff --git a/source/lib/rocprofiler-sdk-tool/tool.cpp b/source/lib/rocprofiler-sdk-tool/tool.cpp index f1a836fbb4..e679c93f46 100644 --- a/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -687,8 +687,11 @@ dispatch_callback(rocprofiler_queue_id_t queue_id, kernel_data.rlock([](const kernel_symbol_data_map_t& kdata, uint64_t kernel_id_v) { return kdata.at(kernel_id_v); }, kernel_id); - auto is_targeted_kernel = [&kernel_info]() { + // if kernel name is provided by user then by default all kernels in the application are + // targeted + if(tool::get_config().kernel_names.empty()) return true; + for(const auto& name : tool::get_config().kernel_names) { if(name == kernel_info.truncated_kernel_name) diff --git a/tests/apps/CMakeLists.txt b/tests/apps/CMakeLists.txt index 9b16f9d734..e478c4db94 100644 --- a/tests/apps/CMakeLists.txt +++ b/tests/apps/CMakeLists.txt @@ -10,6 +10,7 @@ set(CMAKE_BUILD_RPATH "\$ORIGIN:\$ORIGIN/../lib") # applications used by integration tests which DO NOT link to rocprofiler-sdk-roctx add_subdirectory(simple-transpose) add_subdirectory(multistream) +add_subdirectory(vector-operations) set(CMAKE_BUILD_RPATH "\$ORIGIN:\$ORIGIN/../lib:$" diff --git a/tests/apps/vector-operations/CMakeLists.txt b/tests/apps/vector-operations/CMakeLists.txt new file mode 100644 index 0000000000..5fb0a9c010 --- /dev/null +++ b/tests/apps/vector-operations/CMakeLists.txt @@ -0,0 +1,46 @@ +# +# +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +if(NOT CMAKE_HIP_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 NO_CACHE) + mark_as_advanced(amdclangpp_EXECUTABLE) + + if(amdclangpp_EXECUTABLE) + set(CMAKE_HIP_COMPILER "${amdclangpp_EXECUTABLE}") + endif() +endif() + +project(rocprofiler-tool-test-app-transpose LANGUAGES CXX HIP) + +foreach(_TYPE DEBUG MINSIZEREL RELEASE RELWITHDEBINFO) + if("${CMAKE_HIP_FLAGS_${_TYPE}}" STREQUAL "") + set(CMAKE_HIP_FLAGS_${_TYPE} "${CMAKE_CXX_FLAGS_${_TYPE}}") + endif() +endforeach() + +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_EXTENSIONS OFF) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_HIP_STANDARD 17) +set(CMAKE_HIP_EXTENSIONS OFF) +set(CMAKE_HIP_STANDARD_REQUIRED ON) + +set_source_files_properties(vector-ops.cpp PROPERTIES LANGUAGE HIP) +add_executable(vector-ops) +target_sources(vector-ops PRIVATE vector-ops.cpp) +target_compile_options(vector-ops PRIVATE -W -Wall -Wextra -Wpedantic -Wshadow -Werror) + +find_package(Threads REQUIRED) +target_link_libraries(vector-ops PRIVATE Threads::Threads) + +install( + TARGETS vector-ops + DESTINATION bin + COMPONENT tests) diff --git a/tests/apps/vector-operations/vector-ops.cpp b/tests/apps/vector-operations/vector-ops.cpp new file mode 100644 index 0000000000..71f7050df4 --- /dev/null +++ b/tests/apps/vector-operations/vector-ops.cpp @@ -0,0 +1,258 @@ +/* +Copyright (c) 2015-2016 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 +#include +#include +#include +#include +#include +#include +#include + +#define HIP_API_CALL(CALL) \ + { \ + hipError_t error_ = (CALL); \ + if(error_ != hipSuccess) \ + { \ + auto _hip_api_print_lk = auto_lock_t{print_lock}; \ + fprintf(stderr, \ + "%s:%d :: HIP error : %s\n", \ + __FILE__, \ + __LINE__, \ + hipGetErrorString(error_)); \ + throw std::runtime_error("hip_api_call"); \ + } \ + } + +namespace +{ +using auto_lock_t = std::unique_lock; +auto print_lock = std::mutex{}; +} // namespace + +#define WIDTH (1024) +#define HEIGHT (1024) + +#define NUM (WIDTH * HEIGHT) + +#define THREADS_PER_BLOCK_X 64 +#define THREADS_PER_BLOCK_Y 1 +#define THREADS_PER_BLOCK_Z 1 + +// Computes vectorAdd with matrix-multiply +template +__global__ void +addition_kernel(T* __restrict__ a, + const float* __restrict__ b, + const float* __restrict__ c, + int width, + [[maybe_unused]] int height) +{ + // printf("addition kernel\n"); + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; + + if(x >= WIDTH || y >= HEIGHT) return; + int index = y * width + x; + + a[index] = b[index] + c[index]; +} + +__global__ void +subtract_kernel(float* __restrict__ a, + const float* __restrict__ b, + const float* __restrict__ c, + int width, + [[maybe_unused]] int height) +{ + // printf("subtract kernel\n"); + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; + + if(x >= WIDTH || y >= HEIGHT) return; + int index = y * width + x; + + a[index] = abs(b[index] - c[index]); +} + +__global__ void +multiply_kernel(float* __restrict__ a, + const float* __restrict__ b, + const float* __restrict__ c, + int width, + [[maybe_unused]] int height) +{ + // printf("multiply kernel\n"); + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; + + if(x >= WIDTH || y >= HEIGHT) return; + int index = y * width + x; + + a[index] = (b[index] - 1) * (c[index] - 1) + 1; +} + +__global__ void +divide_kernel(float* __restrict__ a, + const float* __restrict__ b, + const float* __restrict__ c, + int width, + [[maybe_unused]] int height) +{ + // printf("divide kernel\n"); + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; + + if(x >= WIDTH || y >= HEIGHT) return; + int index = y * width + x; + + a[index] = (b[index] - c[index]) / abs(c[index] + b[index]) + 1; +} + +using namespace std; + +void +run(int NUM_QUEUE) +{ + std::vector hostA(NUM_QUEUE); + std::vector hostB(NUM_QUEUE); + std::vector hostC(NUM_QUEUE); + + std::vector deviceA(NUM_QUEUE); + std::vector deviceB(NUM_QUEUE); + std::vector deviceC(NUM_QUEUE); + + std::vector streams(NUM_QUEUE); + + hipDeviceProp_t devProp; + HIP_API_CALL(hipGetDeviceProperties(&devProp, 0)); + + int i; + + for(int q = 0; q < NUM_QUEUE; q++) + { + HIP_API_CALL(hipStreamCreateWithFlags(&streams[q], hipStreamNonBlocking)); + + HIP_API_CALL(hipHostMalloc(&hostA[q], NUM * sizeof(float), 0)); + HIP_API_CALL(hipHostMalloc(&hostB[q], NUM * sizeof(float), 0)); + HIP_API_CALL(hipHostMalloc(&hostC[q], NUM * sizeof(float), 0)); + + // initialize the input data + for(i = 0; i < NUM; i++) + { + hostB[q][i] = (float) i; + hostC[q][i] = (float) i * 100.0f; + } + + HIP_API_CALL(hipMalloc((void**) (&deviceA[q]), NUM * sizeof(float))); + HIP_API_CALL(hipMalloc((void**) (&deviceB[q]), NUM * sizeof(float))); + HIP_API_CALL(hipMalloc((void**) (&deviceC[q]), NUM * sizeof(float))); + + HIP_API_CALL(hipMemcpyAsync( + deviceB[q], hostB[q], NUM * sizeof(float), hipMemcpyHostToDevice, streams[q])); + HIP_API_CALL(hipMemcpyAsync( + deviceC[q], hostC[q], NUM * sizeof(float), hipMemcpyHostToDevice, streams[q])); + } + HIP_API_CALL(hipDeviceSynchronize()); + + for(int RUN_I = 0; RUN_I < 2; RUN_I++) + { + int q = (4 * RUN_I + 0) % NUM_QUEUE; + hipLaunchKernelGGL(addition_kernel, + dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), + 0, + streams[q], + deviceA[q], + deviceB[q], + deviceC[q], + WIDTH, + HEIGHT); + + HIP_API_CALL(hipDeviceSynchronize()); + q = (4 * RUN_I + 1) % NUM_QUEUE; + hipLaunchKernelGGL(subtract_kernel, + dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), + 0, + streams[q], + deviceA[q], + deviceB[q], + deviceC[q], + WIDTH, + HEIGHT); + + HIP_API_CALL(hipDeviceSynchronize()); + q = (4 * RUN_I + 2) % NUM_QUEUE; + hipLaunchKernelGGL(multiply_kernel, + dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), + 0, + streams[q], + deviceA[q], + deviceB[q], + deviceC[q], + WIDTH, + HEIGHT); + + HIP_API_CALL(hipDeviceSynchronize()); + q = (4 * RUN_I + 3) % NUM_QUEUE; + hipLaunchKernelGGL(divide_kernel, + dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), + 0, + streams[q], + deviceB[q], + deviceA[q], + deviceC[q], + WIDTH, + HEIGHT); + HIP_API_CALL(hipDeviceSynchronize()); + } + + for(int q = 0; q < NUM_QUEUE; q++) + HIP_API_CALL(hipMemcpyAsync( + hostA[q], deviceA[q], NUM * sizeof(float), hipMemcpyDeviceToHost, streams[q])); + + for(int q = 0; q < NUM_QUEUE; q++) + { + HIP_API_CALL(hipMemcpy(hostA[q], deviceA[q], NUM * sizeof(float), hipMemcpyDeviceToHost)); + HIP_API_CALL(hipDeviceSynchronize()); + + HIP_API_CALL(hipFree(deviceA[q])); + HIP_API_CALL(hipFree(deviceB[q])); + HIP_API_CALL(hipFree(deviceC[q])); + + HIP_API_CALL(hipHostFree(hostA[q])); + HIP_API_CALL(hipHostFree(hostB[q])); + HIP_API_CALL(hipHostFree(hostC[q])); + HIP_API_CALL(hipStreamDestroy(streams[q])); + } +} + +int +main() +{ + run(1); + return 0; +} diff --git a/tests/rocprofv3/counter-collection/input1/CMakeLists.txt b/tests/rocprofv3/counter-collection/input1/CMakeLists.txt index 5122c328b1..886860bf59 100644 --- a/tests/rocprofv3/counter-collection/input1/CMakeLists.txt +++ b/tests/rocprofv3/counter-collection/input1/CMakeLists.txt @@ -1,5 +1,5 @@ # -# +# rocprofv3 tool test # cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) @@ -23,11 +23,19 @@ add_test( ${CMAKE_CURRENT_BINARY_DIR}/input.txt -d ${CMAKE_CURRENT_BINARY_DIR}/out_cc_1 -o pmc1 $) +string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV + "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}") + +if(ROCPROFILER_MEMCHECK STREQUAL "LeakSanitizer") + set(LOG_LEVEL "warning") # info produces memory leak +else() + set(LOG_LEVEL "info") +endif() + set(cc-env-pmc1 - "${PRELOAD_ENV}" - "HSA_TOOLS_LIB=$" - "LD_LIBRARY_PATH=$:$ENV{LD_LIBRARY_PATH}" - ) + "${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}" + "ROCPROFILER_LOG_LEVEL=${LOG_LEVEL}" + "HSA_TOOLS_LIB=$") set_tests_properties( rocprofv3-test-counter-collection-pmc1-execute @@ -36,16 +44,10 @@ set_tests_properties( add_test(NAME rocprofv3-test-counter-collection-pmc1-validate COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --input - "${CMAKE_CURRENT_BINARY_DIR}/out_cc_1/pmc_1/pmc1_counter_collection.csv") + ${CMAKE_CURRENT_BINARY_DIR}/out_cc_1/pmc_1/pmc1_counter_collection.csv) set_tests_properties( rocprofv3-test-counter-collection-pmc1-validate PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS rocprofv3-test-counter-collection-pmc1-execute FAIL_REGULAR_EXPRESSION "threw an exception") - -# Needs to be enabled once counter colelction tool support is merged -set_tests_properties(rocprofv3-test-counter-collection-pmc1-execute PROPERTIES DISABLED - TRUE) -set_tests_properties(rocprofv3-test-counter-collection-pmc1-validate PROPERTIES DISABLED - TRUE) diff --git a/tests/rocprofv3/counter-collection/input1/validate.py b/tests/rocprofv3/counter-collection/input1/validate.py index ac4514d916..128644ddb2 100644 --- a/tests/rocprofv3/counter-collection/input1/validate.py +++ b/tests/rocprofv3/counter-collection/input1/validate.py @@ -7,8 +7,10 @@ def test_validate_counter_collection_pmc1(input_data: pd.DataFrame): df = input_data assert df.empty == False - assert df["agent-id"].map(type).eq(int).all() - assert len(df["kernel-name"]) > 0 + assert df["Agent_Id"].map(type).eq(int).all() + assert len(df["Kernel-Name"]) > 0 + assert df["Kernel-Name"].str.contains("matrixTranspose").all() + assert df["Counter_Name"].str.contains("SQ_WAVES").all() if __name__ == "__main__": diff --git a/tests/rocprofv3/counter-collection/input2/CMakeLists.txt b/tests/rocprofv3/counter-collection/input2/CMakeLists.txt index fa8a1c812d..5daf371a94 100644 --- a/tests/rocprofv3/counter-collection/input2/CMakeLists.txt +++ b/tests/rocprofv3/counter-collection/input2/CMakeLists.txt @@ -1,5 +1,5 @@ # -# +# rocprofv3 tool test # cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) @@ -22,11 +22,20 @@ add_test( $ -i ${CMAKE_CURRENT_BINARY_DIR}/input.txt -d ${CMAKE_CURRENT_BINARY_DIR}/out_cc_2 -o pmc2 $) + +string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV + "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}") + +if(ROCPROFILER_MEMCHECK STREQUAL "LeakSanitizer") + set(LOG_LEVEL "warning") # info produces memory leak +else() + set(LOG_LEVEL "info") +endif() + set(cc-env-pmc2 - "${PRELOAD_ENV}" - "HSA_TOOLS_LIB=$" - "LD_LIBRARY_PATH=$:$ENV{LD_LIBRARY_PATH}" - ) + "${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}" + "ROCPROFILER_LOG_LEVEL=${LOG_LEVEL}" + "HSA_TOOLS_LIB=$") set_tests_properties( rocprofv3-test-counter-collection-pmc2-execute @@ -42,9 +51,3 @@ set_tests_properties( PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS rocprofv3-test-counter-collection-pmc2-execute FAIL_REGULAR_EXPRESSION "threw an exception") - -# Needs to be enabled once counter colelction tool support is merged -set_tests_properties(rocprofv3-test-counter-collection-pmc2-execute PROPERTIES DISABLED - TRUE) -set_tests_properties(rocprofv3-test-counter-collection-pmc2-validate PROPERTIES DISABLED - TRUE) diff --git a/tests/rocprofv3/counter-collection/input2/validate.py b/tests/rocprofv3/counter-collection/input2/validate.py index 4a9793d198..33290f8d89 100644 --- a/tests/rocprofv3/counter-collection/input2/validate.py +++ b/tests/rocprofv3/counter-collection/input2/validate.py @@ -34,9 +34,9 @@ def test_validate_counter_collection_pmc2(input_dir: pd.DataFrame): with open(file_path, "r") as file: df = pd.read_csv(file) # check if kernel-name is present - assert len(df["kernel-name"]) > 0 + assert len(df["Kernel-Name"]) > 0 # check if counter value is positive - assert len(df["counter_value"]) > 0 + assert len(df["Counter_Value"]) > 0 if __name__ == "__main__": diff --git a/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt b/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt index fc6a6ec2eb..99c5298b6d 100644 --- a/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt +++ b/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt @@ -1,5 +1,5 @@ # -# +# rocprofv3 tool # cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) @@ -23,11 +23,19 @@ add_test( ${CMAKE_CURRENT_BINARY_DIR}/input.txt -d ${CMAKE_CURRENT_BINARY_DIR}/out_cc_trace -o pmc3 $) +string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV + "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}") + +if(ROCPROFILER_MEMCHECK STREQUAL "LeakSanitizer") + set(LOG_LEVEL "warning") # info produces memory leak +else() + set(LOG_LEVEL "info") +endif() + set(cc-tracing-env - "${PRELOAD_ENV}" - "HSA_TOOLS_LIB=$" - "LD_LIBRARY_PATH=$:$ENV{LD_LIBRARY_PATH}" - ) + "${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}" + "ROCPROFILER_LOG_LEVEL=${LOG_LEVEL}" + "HSA_TOOLS_LIB=$") set_tests_properties( rocprofv3-test-tracing-plus-cc-execute @@ -43,7 +51,3 @@ set_tests_properties( PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS rocprofv3-test-tracing-plus-cc-execute FAIL_REGULAR_EXPRESSION "threw an exception") - -# Needs to be enabled once counter colelction tool support is merged -set_tests_properties(rocprofv3-test-tracing-plus-cc-execute PROPERTIES DISABLED TRUE) -set_tests_properties(rocprofv3-test-tracing-plus-cc-validate PROPERTIES DISABLED TRUE) diff --git a/tests/rocprofv3/tracing-plus-cc/validate.py b/tests/rocprofv3/tracing-plus-cc/validate.py index ef54f9f781..43548661da 100644 --- a/tests/rocprofv3/tracing-plus-cc/validate.py +++ b/tests/rocprofv3/tracing-plus-cc/validate.py @@ -34,7 +34,7 @@ def test_validate_counter_collection_plus_tracing(input_dir: pd.DataFrame): with open(file_path, "r") as file: df = pd.read_csv(file) # check if either kernel-name/FUNCTION is present - assert "kernel-name" in df.columns or "FUNCTION" in df.columns + assert "Kernel-Name" in df.columns or "Function" in df.columns if __name__ == "__main__":