From c123102aa571ae91695db6ca7a3f9d51b7d98d7c Mon Sep 17 00:00:00 2001 From: pciolkosz Date: Wed, 24 Jul 2024 12:53:11 -0700 Subject: [PATCH 1/4] Add a header to interact with driver APIs --- .../experimental/__utility/driver_api.cuh | 74 +++++++++++++++++++ 1 file changed, 74 insertions(+) create mode 100644 cudax/include/cuda/experimental/__utility/driver_api.cuh diff --git a/cudax/include/cuda/experimental/__utility/driver_api.cuh b/cudax/include/cuda/experimental/__utility/driver_api.cuh new file mode 100644 index 00000000000..2262a41dc70 --- /dev/null +++ b/cudax/include/cuda/experimental/__utility/driver_api.cuh @@ -0,0 +1,74 @@ +//===----------------------------------------------------------------------===// +// +// Part of CUDA Experimental in CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDAX__UTILITY_DRIVER_API +#define _CUDAX__UTILITY_DRIVER_API + +#include + +#include + +#define CUDAX_GET_DRIVER_FUNCTION(function_name) \ + reinterpret_cast(getDriverEntryPoint(#function_name)) + +namespace cuda::experimental::detail::driver +{ +inline void* getDriverEntryPoint(const char* name) +{ + void* fn; + cudaDriverEntryPointQueryResult result; + cudaGetDriverEntryPoint(name, &fn, cudaEnableDefault, &result); + if (result != cudaDriverEntryPointSuccess) { + if (result == cudaDriverEntryPointVersionNotSufficent) { + ::cuda::__throw_cuda_error(cudaErrorNotSupported, "Driver does not support this API"); + } + else { + ::cuda::__throw_cuda_error(cudaErrorUnknown, "Failed to access driver API"); + } + } + return fn; +} + +inline void ctxPush(CUcontext ctx) +{ + static auto driver_fn = CUDAX_GET_DRIVER_FUNCTION(cuCtxPushCurrent); + CUresult status = driver_fn(ctx); + if (status != CUDA_SUCCESS) + { + ::cuda::__throw_cuda_error(static_cast(status), "Failed to push context"); + } +} + +inline void ctxPop() +{ + CUcontext dummy; + static auto driver_fn = CUDAX_GET_DRIVER_FUNCTION(cuCtxPopCurrent); + CUresult status = driver_fn(&dummy); + if (status != CUDA_SUCCESS) + { + ::cuda::__throw_cuda_error(static_cast(status), "Failed to pop context"); + } +} + +inline CUcontext streamGetCtx(CUstream stream) +{ + CUcontext result; + static auto driver_fn = CUDAX_GET_DRIVER_FUNCTION(cuStreamGetCtx); + CUresult status = driver_fn(stream, &result); + if (status != CUDA_SUCCESS) + { + ::cuda::__throw_cuda_error(static_cast(status), "Failed to get context from a stream"); + } + return result; +} +} // namespace cuda::experimental::detail::driver + +#undef CUDAX_GET_DRIVER_FUNCTION +#endif \ No newline at end of file From 8357b5aa21ef030e22dd730852d1038ec924b96c Mon Sep 17 00:00:00 2001 From: pciolkosz Date: Wed, 24 Jul 2024 15:15:41 -0700 Subject: [PATCH 2/4] Add a test for the driver API interaction --- .../experimental/__utility/driver_api.cuh | 42 ++++++++++-------- cudax/test/CMakeLists.txt | 4 ++ cudax/test/utility/driver_api.cu | 44 +++++++++++++++++++ 3 files changed, 71 insertions(+), 19 deletions(-) create mode 100644 cudax/test/utility/driver_api.cu diff --git a/cudax/include/cuda/experimental/__utility/driver_api.cuh b/cudax/include/cuda/experimental/__utility/driver_api.cuh index 2262a41dc70..5b32be62fe7 100644 --- a/cudax/include/cuda/experimental/__utility/driver_api.cuh +++ b/cudax/include/cuda/experimental/__utility/driver_api.cuh @@ -15,12 +15,13 @@ #include +// Get the driver function by name using this macro #define CUDAX_GET_DRIVER_FUNCTION(function_name) \ - reinterpret_cast(getDriverEntryPoint(#function_name)) + reinterpret_cast(get_driver_entry_point(#function_name)) namespace cuda::experimental::detail::driver { -inline void* getDriverEntryPoint(const char* name) +inline void* get_driver_entry_point(const char* name) { void* fn; cudaDriverEntryPointQueryResult result; @@ -36,36 +37,39 @@ inline void* getDriverEntryPoint(const char* name) return fn; } +template +inline void call_driver_fn(Fn fn, const char* err_msg, Args... args) { + CUresult status = fn(args...); + if (status != CUDA_SUCCESS) { + ::cuda::__throw_cuda_error(static_cast(status), err_msg); + } +} + inline void ctxPush(CUcontext ctx) { static auto driver_fn = CUDAX_GET_DRIVER_FUNCTION(cuCtxPushCurrent); - CUresult status = driver_fn(ctx); - if (status != CUDA_SUCCESS) - { - ::cuda::__throw_cuda_error(static_cast(status), "Failed to push context"); - } + call_driver_fn(driver_fn, "Failed to push context", ctx); } inline void ctxPop() { - CUcontext dummy; static auto driver_fn = CUDAX_GET_DRIVER_FUNCTION(cuCtxPopCurrent); - CUresult status = driver_fn(&dummy); - if (status != CUDA_SUCCESS) - { - ::cuda::__throw_cuda_error(static_cast(status), "Failed to pop context"); - } + CUcontext dummy; + call_driver_fn(driver_fn, "Failed to pop context", &dummy); +} + +inline CUcontext ctxGetCurrent() { + static auto driver_fn = CUDAX_GET_DRIVER_FUNCTION(cuCtxGetCurrent); + CUcontext result; + call_driver_fn(driver_fn, "Failed to get current context", &result); + return result; } inline CUcontext streamGetCtx(CUstream stream) { - CUcontext result; static auto driver_fn = CUDAX_GET_DRIVER_FUNCTION(cuStreamGetCtx); - CUresult status = driver_fn(stream, &result); - if (status != CUDA_SUCCESS) - { - ::cuda::__throw_cuda_error(static_cast(status), "Failed to get context from a stream"); - } + CUcontext result; + call_driver_fn(driver_fn, "Failed to get context from a stream", stream, &result); return result; } } // namespace cuda::experimental::detail::driver diff --git a/cudax/test/CMakeLists.txt b/cudax/test/CMakeLists.txt index 5cf3b67c843..0f8464f2bbc 100644 --- a/cudax/test/CMakeLists.txt +++ b/cudax/test/CMakeLists.txt @@ -57,4 +57,8 @@ foreach(cn_target IN LISTS cudax_TARGETS) launch/configuration.cu ) target_compile_options(${test_target} PRIVATE $<$:--extended-lambda>) + + Cudax_add_catch2_test(test_target misc_tests ${cn_target} + utility/driver_api.cu + ) endforeach() diff --git a/cudax/test/utility/driver_api.cu b/cudax/test/utility/driver_api.cu new file mode 100644 index 00000000000..ada726bf60b --- /dev/null +++ b/cudax/test/utility/driver_api.cu @@ -0,0 +1,44 @@ +//===----------------------------------------------------------------------===// +// +// Part of CUDA Experimental in CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// +#define LIBCUDACXX_ENABLE_EXCEPTIONS + +#include + +#include "../hierarchy/testing_common.cuh" + +TEST_CASE("Call each one", "[driver api]") { + cudaStream_t stream; + // Assumes the ctx stack was empty or had one ctx, should be the case unless some other + // test leaves 2+ ctxs on the stack + + // Pushes the primary context if the stack is empty + CUDART(cudaStreamCreate(&stream)); + + auto ctx = cuda::experimental::detail::driver::ctxGetCurrent(); + CUDAX_REQUIRE(ctx != nullptr); + + cuda::experimental::detail::driver::ctxPop(); + CUDAX_REQUIRE(cuda::experimental::detail::driver::ctxGetCurrent() == nullptr); + + cuda::experimental::detail::driver::ctxPush(ctx); + CUDAX_REQUIRE(cuda::experimental::detail::driver::ctxGetCurrent() == ctx); + + cuda::experimental::detail::driver::ctxPush(ctx); + CUDAX_REQUIRE(cuda::experimental::detail::driver::ctxGetCurrent() == ctx); + + cuda::experimental::detail::driver::ctxPop(); + CUDAX_REQUIRE(cuda::experimental::detail::driver::ctxGetCurrent() == ctx); + + + auto stream_ctx = cuda::experimental::detail::driver::streamGetCtx(stream); + CUDAX_REQUIRE(ctx == stream_ctx); + + CUDART(cudaStreamDestroy(stream)); +} \ No newline at end of file From 242e135c050002265b0c27b6f261512f474cb800 Mon Sep 17 00:00:00 2001 From: pciolkosz Date: Wed, 24 Jul 2024 16:44:20 -0700 Subject: [PATCH 3/4] Format --- .../experimental/__utility/driver_api.cuh | 20 +++++---- cudax/test/utility/driver_api.cu | 42 +++++++++---------- 2 files changed, 34 insertions(+), 28 deletions(-) diff --git a/cudax/include/cuda/experimental/__utility/driver_api.cuh b/cudax/include/cuda/experimental/__utility/driver_api.cuh index 5b32be62fe7..21b8c4d7425 100644 --- a/cudax/include/cuda/experimental/__utility/driver_api.cuh +++ b/cudax/include/cuda/experimental/__utility/driver_api.cuh @@ -26,11 +26,14 @@ inline void* get_driver_entry_point(const char* name) void* fn; cudaDriverEntryPointQueryResult result; cudaGetDriverEntryPoint(name, &fn, cudaEnableDefault, &result); - if (result != cudaDriverEntryPointSuccess) { - if (result == cudaDriverEntryPointVersionNotSufficent) { + if (result != cudaDriverEntryPointSuccess) + { + if (result == cudaDriverEntryPointVersionNotSufficent) + { ::cuda::__throw_cuda_error(cudaErrorNotSupported, "Driver does not support this API"); } - else { + else + { ::cuda::__throw_cuda_error(cudaErrorUnknown, "Failed to access driver API"); } } @@ -38,9 +41,11 @@ inline void* get_driver_entry_point(const char* name) } template -inline void call_driver_fn(Fn fn, const char* err_msg, Args... args) { +inline void call_driver_fn(Fn fn, const char* err_msg, Args... args) +{ CUresult status = fn(args...); - if (status != CUDA_SUCCESS) { + if (status != CUDA_SUCCESS) + { ::cuda::__throw_cuda_error(static_cast(status), err_msg); } } @@ -58,7 +63,8 @@ inline void ctxPop() call_driver_fn(driver_fn, "Failed to pop context", &dummy); } -inline CUcontext ctxGetCurrent() { +inline CUcontext ctxGetCurrent() +{ static auto driver_fn = CUDAX_GET_DRIVER_FUNCTION(cuCtxGetCurrent); CUcontext result; call_driver_fn(driver_fn, "Failed to get current context", &result); @@ -75,4 +81,4 @@ inline CUcontext streamGetCtx(CUstream stream) } // namespace cuda::experimental::detail::driver #undef CUDAX_GET_DRIVER_FUNCTION -#endif \ No newline at end of file +#endif diff --git a/cudax/test/utility/driver_api.cu b/cudax/test/utility/driver_api.cu index ada726bf60b..513d6476eb5 100644 --- a/cudax/test/utility/driver_api.cu +++ b/cudax/test/utility/driver_api.cu @@ -13,32 +13,32 @@ #include "../hierarchy/testing_common.cuh" -TEST_CASE("Call each one", "[driver api]") { - cudaStream_t stream; - // Assumes the ctx stack was empty or had one ctx, should be the case unless some other - // test leaves 2+ ctxs on the stack +TEST_CASE("Call each one", "[driver api]") +{ + cudaStream_t stream; + // Assumes the ctx stack was empty or had one ctx, should be the case unless some other + // test leaves 2+ ctxs on the stack - // Pushes the primary context if the stack is empty - CUDART(cudaStreamCreate(&stream)); + // Pushes the primary context if the stack is empty + CUDART(cudaStreamCreate(&stream)); - auto ctx = cuda::experimental::detail::driver::ctxGetCurrent(); - CUDAX_REQUIRE(ctx != nullptr); + auto ctx = cuda::experimental::detail::driver::ctxGetCurrent(); + CUDAX_REQUIRE(ctx != nullptr); - cuda::experimental::detail::driver::ctxPop(); - CUDAX_REQUIRE(cuda::experimental::detail::driver::ctxGetCurrent() == nullptr); + cuda::experimental::detail::driver::ctxPop(); + CUDAX_REQUIRE(cuda::experimental::detail::driver::ctxGetCurrent() == nullptr); - cuda::experimental::detail::driver::ctxPush(ctx); - CUDAX_REQUIRE(cuda::experimental::detail::driver::ctxGetCurrent() == ctx); + cuda::experimental::detail::driver::ctxPush(ctx); + CUDAX_REQUIRE(cuda::experimental::detail::driver::ctxGetCurrent() == ctx); - cuda::experimental::detail::driver::ctxPush(ctx); - CUDAX_REQUIRE(cuda::experimental::detail::driver::ctxGetCurrent() == ctx); + cuda::experimental::detail::driver::ctxPush(ctx); + CUDAX_REQUIRE(cuda::experimental::detail::driver::ctxGetCurrent() == ctx); - cuda::experimental::detail::driver::ctxPop(); - CUDAX_REQUIRE(cuda::experimental::detail::driver::ctxGetCurrent() == ctx); + cuda::experimental::detail::driver::ctxPop(); + CUDAX_REQUIRE(cuda::experimental::detail::driver::ctxGetCurrent() == ctx); + auto stream_ctx = cuda::experimental::detail::driver::streamGetCtx(stream); + CUDAX_REQUIRE(ctx == stream_ctx); - auto stream_ctx = cuda::experimental::detail::driver::streamGetCtx(stream); - CUDAX_REQUIRE(ctx == stream_ctx); - - CUDART(cudaStreamDestroy(stream)); -} \ No newline at end of file + CUDART(cudaStreamDestroy(stream)); +} From 4b7bb4c081cfb1533236f720e6e2fd1cf9eba005 Mon Sep 17 00:00:00 2001 From: pciolkosz Date: Tue, 30 Jul 2024 15:59:51 -0700 Subject: [PATCH 4/4] Fix formatting --- cudax/test/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cudax/test/CMakeLists.txt b/cudax/test/CMakeLists.txt index da42e678dcf..e0c73a7b76e 100644 --- a/cudax/test/CMakeLists.txt +++ b/cudax/test/CMakeLists.txt @@ -71,7 +71,7 @@ foreach(cn_target IN LISTS cudax_TARGETS) cudax_add_catch2_test(test_target stream_tests ${cn_target} stream/stream_smoke.cu ) - + cudax_add_catch2_test(test_target misc_tests ${cn_target} utility/driver_api.cu )