diff --git a/CMakeLists.txt b/CMakeLists.txt index 18a57f4f..3adabde1 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -5,6 +5,8 @@ option(USE_CCACHE "Attempt using CCache to wrap the compilation" ON) option(USE_CXX11_ABI "Use the new C++-11 ABI, which is not backwards compatible." ON) option(USE_MANYLINUX "Build for manylinux" OFF) +option(BUILD_NVBENCH "Build the nvbench binary" OFF) + set(CMAKE_CXX_STANDARD 17) set(CMAKE_CXX_STANDARD_REQUIRED ON) set(CMAKE_CXX_EXTENSIONS OFF) diff --git a/cmake/nvbench_binary.cmake b/cmake/nvbench_binary.cmake new file mode 100644 index 00000000..a79a2091 --- /dev/null +++ b/cmake/nvbench_binary.cmake @@ -0,0 +1,40 @@ +include(CMakeParseArguments) + +function(nvbench_binary) + if(NOT BUILD_NVBENCH) + return() + endif() + + cmake_parse_arguments( + NV_BINARY # prefix + "" # options + "NAME" # one value args + "HDRS;SRCS;COPTS;DEFINES;LINKOPTS;DEPS" # multi value args + ${ARGN} + ) + + add_executable(${NV_BINARY_NAME} "") + target_sources(${NV_BINARY_NAME} + PRIVATE ${NV_BINARY_SRCS} ${NV_BINARY_HDRS} + ) + target_link_libraries(${NV_BINARY_NAME} + PUBLIC + ${NV_BINARY_DEPS} + nvbench::nvbench + nvbench::main + PRIVATE + ${NV_BINARY_LINKOPTS} + ) + target_include_directories(${NV_BINARY_NAME} + PUBLIC + "$" + ) + target_compile_options(${NV_BINARY_NAME} + PRIVATE + ${NV_BINARY_COPTS} + -lineinfo + ) + target_compile_definitions(${NV_BINARY_NAME} PUBLIC ${NV_BINARY_DEFINES}) + + add_executable(:${NV_BINARY_NAME} ALIAS ${NV_BINARY_NAME}) +endfunction() diff --git a/src/kernels/CMakeLists.txt b/src/kernels/CMakeLists.txt index 76c20ebe..406eb40a 100644 --- a/src/kernels/CMakeLists.txt +++ b/src/kernels/CMakeLists.txt @@ -28,7 +28,6 @@ cc_library( add_subdirectory(attention) add_subdirectory(quantization) -add_subdirectory(bench) add_subdirectory(playground) add_subdirectory(triton) diff --git a/src/kernels/attention/CMakeLists.txt b/src/kernels/attention/CMakeLists.txt index 934f051e..a72a05cb 100644 --- a/src/kernels/attention/CMakeLists.txt +++ b/src/kernels/attention/CMakeLists.txt @@ -1,4 +1,5 @@ include(cc_binary) +include(nvbench_binary) include(cc_library) include(cc_test) @@ -66,31 +67,23 @@ cc_test( torch ) -cc_binary( +nvbench_binary( NAME attention_sm80_bench SRCS attention_sm80_bench.cu DEPS - nvbench::nvbench - nvbench::main - :attention.template - COPTS - -lineinfo + :attention.template ) -cc_binary( +nvbench_binary( NAME attention_sm80_pagedkv_bench SRCS attention_sm80_pagedkv_bench.cu DEPS absl::random_random - nvbench::nvbench - nvbench::main :attention.template - COPTS - -lineinfo ) add_subdirectory(tools) \ No newline at end of file diff --git a/src/kernels/bench/CMakeLists.txt b/src/kernels/bench/CMakeLists.txt deleted file mode 100644 index ea62671e..00000000 --- a/src/kernels/bench/CMakeLists.txt +++ /dev/null @@ -1,11 +0,0 @@ -include(cc_binary) - -cc_binary( - NAME - kernel_bench - SRCS - bench_demo.cu - DEPS - nvbench::nvbench - nvbench::main -) diff --git a/src/kernels/bench/bench_demo.cu b/src/kernels/bench/bench_demo.cu deleted file mode 100644 index f50806b2..00000000 --- a/src/kernels/bench/bench_demo.cu +++ /dev/null @@ -1,25 +0,0 @@ -#include - -#include -#include - -__global__ void sleep_kernel(nvbench::int64_t microseconds) { - const auto start = cuda::std::chrono::high_resolution_clock::now(); - const auto target_duration = cuda::std::chrono::microseconds(microseconds); - const auto finish = start + target_duration; - - auto now = cuda::std::chrono::high_resolution_clock::now(); - while (now < finish) { - now = cuda::std::chrono::high_resolution_clock::now(); - } -} - -void sleep_benchmark(nvbench::state& state) { - const auto duration_us = state.get_int64("Duration (us)"); - state.exec([&duration_us](nvbench::launch& launch) { - sleep_kernel<<<1, 1, 0, launch.get_stream()>>>(duration_us); - }); -} -NVBENCH_BENCH(sleep_benchmark) - .add_int64_axis("Duration (us)", nvbench::range(0, 100, 5)) - .set_timeout(1); // Limit to one second per measurement. \ No newline at end of file diff --git a/third_party/CMakeLists.txt b/third_party/CMakeLists.txt index 7be466a1..066a46c1 100644 --- a/third_party/CMakeLists.txt +++ b/third_party/CMakeLists.txt @@ -17,6 +17,8 @@ cc_library( ) add_subdirectory(sentencepiece) -add_subdirectory(nvbench) +if (BUILD_NVBENCH) + add_subdirectory(nvbench) +endif() # add_subdirectory(jinja2cpp)