diff --git a/.clang-format b/.clang-format new file mode 100644 index 00000000000..9093cdc5ca0 --- /dev/null +++ b/.clang-format @@ -0,0 +1,37 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: Apache-2.0 + +# A clang-format style derived from Python's PEP 7 (4-space indent, Stroustrup braces, +# return type on its own line), except that the column limit is 120 rather than 79. +# Based on https://gist.github.com/pganssle/0e3a5f828b4d07d79447f6ced8e7e4db +BasedOnStyle: Google +AlwaysBreakAfterReturnType: All +AllowShortIfStatementsOnASingleLine: false +BreakBeforeBraces: Stroustrup +ColumnLimit: 120 +DerivePointerAlignment: false +IndentWidth: 4 +Language: Cpp +PointerAlignment: Right +ReflowComments: true +SpaceBeforeParens: ControlStatements +SpacesInParentheses: false +TabWidth: 4 +UseTab: Never + +# Keep include order as written; reordering can break order-sensitive builds. +SortIncludes: Never + +# Ruff-like wrapping: break after "(", indent arguments, put ")" on its own line. +# This does not contradict PEP7, but extends it to the consistent style where +# calls either fit on a single line or they are broken across multiple lines with +# one argument per line (indented). +AlignAfterOpenBracket: BlockIndent +BinPackArguments: false +BinPackParameters: false +AllowAllArgumentsOnNextLine: false +AllowAllParametersOfDeclarationOnNextLine: false + +# Don't put short lambdas on a single line. +AllowShortLambdasOnASingleLine: None diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index 2cc4bb91d26..f875f324a35 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -32,6 +32,13 @@ repos: - id: ruff-format exclude: .*\.pyi$ + - repo: https://github.com/pre-commit/mirrors-clang-format + rev: 98ad4ddb9c057de2f5b9e93e44e6db959807d68c # frozen: v23.1.3 + hooks: + - id: clang-format + types_or: [c, c++, cuda] + exclude: ^cuda_core/(cuda/core/_include/(dlpack|aoti_shim)\.h|third_party/) + - repo: local hooks: - id: check-precommit-installed diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index 68e25ff0951..cebdf4bdd05 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -247,6 +247,15 @@ between commits, leaving stale headers or out-of-date stubs in the history. If the hook isn't installed, `pre-commit run` will print a visible warning reminding you to run `pre-commit install`. +C, C++ and CUDA sources are formatted with `clang-format` using the +[`.clang-format`](.clang-format) at the repository root, which follows PEP 7 +except for a 120-column limit. The +pre-commit hook installs its own pinned `clang-format` (from PyPI) into the +hook environment, so no system-wide installation is required. Vendored code +(for example `dlpack.h` and `cuda_core/third_party/`) is excluded in +`.pre-commit-config.yaml`. To format in your editor, install the same version +with `pip install clang-format==23.1.3`. + Windows contributors: see [Pre-commit lychee workaround](#pre-commit-lychee-workaround) under Development on Windows. ## Pixi lockfiles diff --git a/benchmarks/cuda_bindings/benchmarks/cpp/bench_ctx_device.cpp b/benchmarks/cuda_bindings/benchmarks/cpp/bench_ctx_device.cpp index 29358867186..3f96f0ea41b 100644 --- a/benchmarks/cuda_bindings/benchmarks/cpp/bench_ctx_device.cpp +++ b/benchmarks/cuda_bindings/benchmarks/cpp/bench_ctx_device.cpp @@ -9,18 +9,20 @@ #include #include - -static void check_cu(CUresult status, const char* message) { +static void +check_cu(CUresult status, const char *message) +{ if (status != CUDA_SUCCESS) { - const char* error_name = nullptr; + const char *error_name = nullptr; cuGetErrorName(status, &error_name); std::cerr << message << ": " << (error_name ? error_name : "unknown") << '\n'; std::exit(1); } } - -int main(int argc, char** argv) { +int +main(int argc, char **argv) +{ bench::Options options = bench::parse_args(argc, argv); // Setup: init CUDA and create a context @@ -80,10 +82,7 @@ int main(int argc, char** argv) { // --- device_primary_ctx_retain --- // Outer retain so the benchmarked retain/release pair just bumps the refcount. CUcontext primary_outer = nullptr; - check_cu( - cuDevicePrimaryCtxRetain(&primary_outer, device), - "cuDevicePrimaryCtxRetain (setup) failed" - ); + check_cu(cuDevicePrimaryCtxRetain(&primary_outer, device), "cuDevicePrimaryCtxRetain (setup) failed"); { CUcontext primary = nullptr; suite.run("ctx_device.device_primary_ctx_retain", [&]() { @@ -91,10 +90,7 @@ int main(int argc, char** argv) { check_cu(cuDevicePrimaryCtxRelease(device), "cuDevicePrimaryCtxRelease failed"); }); } - check_cu( - cuDevicePrimaryCtxRelease(device), - "cuDevicePrimaryCtxRelease (teardown) failed" - ); + check_cu(cuDevicePrimaryCtxRelease(device), "cuDevicePrimaryCtxRelease (teardown) failed"); // Cleanup check_cu(cuCtxDestroy(ctx), "cuCtxDestroy failed"); diff --git a/benchmarks/cuda_bindings/benchmarks/cpp/bench_event.cpp b/benchmarks/cuda_bindings/benchmarks/cpp/bench_event.cpp index c24aa983199..e2612e5d1b0 100644 --- a/benchmarks/cuda_bindings/benchmarks/cpp/bench_event.cpp +++ b/benchmarks/cuda_bindings/benchmarks/cpp/bench_event.cpp @@ -9,18 +9,20 @@ #include #include - -static void check_cu(CUresult status, const char* message) { +static void +check_cu(CUresult status, const char *message) +{ if (status != CUDA_SUCCESS) { - const char* error_name = nullptr; + const char *error_name = nullptr; cuGetErrorName(status, &error_name); std::cerr << message << ": " << (error_name ? error_name : "unknown") << '\n'; std::exit(1); } } - -int main(int argc, char** argv) { +int +main(int argc, char **argv) +{ bench::Options options = bench::parse_args(argc, argv); // Setup diff --git a/benchmarks/cuda_bindings/benchmarks/cpp/bench_launch.cpp b/benchmarks/cuda_bindings/benchmarks/cpp/bench_launch.cpp index 4897859a61a..89c2e6966d3 100644 --- a/benchmarks/cuda_bindings/benchmarks/cpp/bench_launch.cpp +++ b/benchmarks/cuda_bindings/benchmarks/cpp/bench_launch.cpp @@ -14,36 +14,47 @@ #include #include - -static void check_cu(CUresult status, const char* message) { +static void +check_cu(CUresult status, const char *message) +{ if (status != CUDA_SUCCESS) { - const char* error_name = nullptr; + const char *error_name = nullptr; cuGetErrorName(status, &error_name); std::cerr << message << ": " << (error_name ? error_name : "unknown") << '\n'; std::exit(1); } } -static void check_nvrtc(nvrtcResult status, const char* message) { +static void +check_nvrtc(nvrtcResult status, const char *message) +{ if (status != NVRTC_SUCCESS) { std::cerr << message << ": " << nvrtcGetErrorString(status) << '\n'; std::exit(1); } } -static CUmodule compile_and_load(const char* source, CUdevice device) { +static CUmodule +compile_and_load(const char *source, CUdevice device) +{ int major = 0, minor = 0; - check_cu(cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, device), - "cuDeviceGetAttribute failed"); - check_cu(cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, device), - "cuDeviceGetAttribute failed"); + check_cu( + cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, device), + "cuDeviceGetAttribute failed" + ); + check_cu( + cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, device), + "cuDeviceGetAttribute failed" + ); nvrtcProgram prog; - check_nvrtc(nvrtcCreateProgram(&prog, source, "benchmark_kernel.cu", 0, nullptr, nullptr), - "nvrtcCreateProgram failed"); + check_nvrtc( + nvrtcCreateProgram(&prog, source, "benchmark_kernel.cu", 0, nullptr, nullptr), + "nvrtcCreateProgram failed" + ); std::string arch = "--gpu-architecture=sm_" + std::to_string(major) + std::to_string(minor); - const char* opts[] = {"--fmad=false", arch.c_str()}; + const char *opts[] = {"--fmad=false", arch.c_str()}; nvrtcResult compile_result = nvrtcCompileProgram(prog, 2, opts); if (compile_result != NVRTC_SUCCESS) { @@ -66,8 +77,7 @@ static CUmodule compile_and_load(const char* source, CUdevice device) { return module; } - -static const char* KERNEL_SOURCE = R"( +static const char *KERNEL_SOURCE = R"( extern "C" __global__ void empty_kernel() { return; } extern "C" __global__ void small_kernel(float *f) { *f = 0.0f; } @@ -152,8 +162,9 @@ extern "C" __global__ void small_kernel_2048B(KernelFunctionParam<2048> param) {} )"; - -int main(int argc, char** argv) { +int +main(int argc, char **argv) +{ bench::Options options = bench::parse_args(argc, argv); // Setup @@ -169,23 +180,23 @@ int main(int argc, char** argv) { CUmodule module = compile_and_load(KERNEL_SOURCE, device); // Get all kernel handles - auto get_func = [&](const char* name) { + auto get_func = [&](const char *name) { CUfunction f; check_cu(cuModuleGetFunction(&f, module, name), "GetFunction failed"); return f; }; - CUfunction empty_kernel = get_func("empty_kernel"); - CUfunction small_kernel = get_func("small_kernel"); - CUfunction kernel_16_args = get_func("small_kernel_16_args"); - CUfunction kernel_256_args = get_func("small_kernel_256_args"); - CUfunction kernel_512_args = get_func("small_kernel_512_args"); - CUfunction kernel_512_bools = get_func("small_kernel_512_bools"); - CUfunction kernel_512_ints = get_func("small_kernel_512_ints"); + CUfunction empty_kernel = get_func("empty_kernel"); + CUfunction small_kernel = get_func("small_kernel"); + CUfunction kernel_16_args = get_func("small_kernel_16_args"); + CUfunction kernel_256_args = get_func("small_kernel_256_args"); + CUfunction kernel_512_args = get_func("small_kernel_512_args"); + CUfunction kernel_512_bools = get_func("small_kernel_512_bools"); + CUfunction kernel_512_ints = get_func("small_kernel_512_ints"); CUfunction kernel_512_doubles = get_func("small_kernel_512_doubles"); - CUfunction kernel_512_chars = get_func("small_kernel_512_chars"); + CUfunction kernel_512_chars = get_func("small_kernel_512_chars"); CUfunction kernel_512_longlongs = get_func("small_kernel_512_longlongs"); - CUfunction kernel_2048B = get_func("small_kernel_2048B"); + CUfunction kernel_2048B = get_func("small_kernel_2048B"); CUstream stream; check_cu(cuStreamCreate(&stream, CU_STREAM_NON_BLOCKING), "cuStreamCreate failed"); @@ -200,42 +211,62 @@ int main(int argc, char** argv) { } // Pre-pack pointer params - void* packed_16[16]; - for (int i = 0; i < 16; ++i) + void *packed_16[16]; + for (int i = 0; i < 16; ++i) { packed_16[i] = &int_ptrs[i]; + } - void* packed_256[256]; - for (int i = 0; i < 256; ++i) + void *packed_256[256]; + for (int i = 0; i < 256; ++i) { packed_256[i] = &int_ptrs[i]; + } - void* packed_512[512]; - for (int i = 0; i < 512; ++i) + void *packed_512[512]; + for (int i = 0; i < 512; ++i) { packed_512[i] = &int_ptrs[i]; + } // Typed args for 512-arg benchmarks bool bool_args[512]; - void* bool_params[512]; - for (int i = 0; i < 512; ++i) { bool_args[i] = true; bool_params[i] = &bool_args[i]; } + void *bool_params[512]; + for (int i = 0; i < 512; ++i) { + bool_args[i] = true; + bool_params[i] = &bool_args[i]; + } int int_args[512]; - void* int_params[512]; - for (int i = 0; i < 512; ++i) { int_args[i] = 123; int_params[i] = &int_args[i]; } + void *int_params[512]; + for (int i = 0; i < 512; ++i) { + int_args[i] = 123; + int_params[i] = &int_args[i]; + } double double_args[512]; - void* double_params[512]; - for (int i = 0; i < 512; ++i) { double_args[i] = 1.2345; double_params[i] = &double_args[i]; } + void *double_params[512]; + for (int i = 0; i < 512; ++i) { + double_args[i] = 1.2345; + double_params[i] = &double_args[i]; + } char char_args[512]; - void* char_params[512]; - for (int i = 0; i < 512; ++i) { char_args[i] = 127; char_params[i] = &char_args[i]; } + void *char_params[512]; + for (int i = 0; i < 512; ++i) { + char_args[i] = 127; + char_params[i] = &char_args[i]; + } long long ll_args[512]; - void* ll_params[512]; - for (int i = 0; i < 512; ++i) { ll_args[i] = 9223372036854775806LL; ll_params[i] = &ll_args[i]; } + void *ll_params[512]; + for (int i = 0; i < 512; ++i) { + ll_args[i] = 9223372036854775806LL; + ll_params[i] = &ll_args[i]; + } // 2048-byte struct - struct alignas(8) Struct2048B { unsigned char p[2048]; } struct_2048B = {}; - void* struct_params[] = {&struct_2048B}; + struct alignas(8) Struct2048B { + unsigned char p[2048]; + } struct_2048B = {}; + void *struct_params[] = {&struct_2048B}; bench::BenchmarkSuite suite(options); // After calibration, drain the persistent stream so the first measured @@ -246,83 +277,106 @@ int main(int argc, char** argv) { }); suite.run("launch.launch_empty_kernel", [&]() { - check_cu(cuLaunchKernel(empty_kernel, 1, 1, 1, 1, 1, 1, 0, stream, nullptr, nullptr), - "cuLaunchKernel failed"); + check_cu(cuLaunchKernel(empty_kernel, 1, 1, 1, 1, 1, 1, 0, stream, nullptr, nullptr), "cuLaunchKernel failed"); }); check_cu(cuStreamSynchronize(stream), "sync"); { - void* params[] = {&float_ptr}; + void *params[] = {&float_ptr}; suite.run("launch.launch_small_kernel", [&]() { - check_cu(cuLaunchKernel(small_kernel, 1, 1, 1, 1, 1, 1, 0, stream, params, nullptr), - "cuLaunchKernel failed"); + check_cu( + cuLaunchKernel(small_kernel, 1, 1, 1, 1, 1, 1, 0, stream, params, nullptr), + "cuLaunchKernel failed" + ); }); } check_cu(cuStreamSynchronize(stream), "sync"); suite.run("launch.launch_16_args", [&]() { - check_cu(cuLaunchKernel(kernel_16_args, 1, 1, 1, 1, 1, 1, 0, stream, packed_16, nullptr), - "cuLaunchKernel failed"); + check_cu( + cuLaunchKernel(kernel_16_args, 1, 1, 1, 1, 1, 1, 0, stream, packed_16, nullptr), + "cuLaunchKernel failed" + ); }); check_cu(cuStreamSynchronize(stream), "sync"); suite.run("launch.launch_16_args_pre_packed", [&]() { - check_cu(cuLaunchKernel(kernel_16_args, 1, 1, 1, 1, 1, 1, 0, stream, packed_16, nullptr), - "cuLaunchKernel failed"); + check_cu( + cuLaunchKernel(kernel_16_args, 1, 1, 1, 1, 1, 1, 0, stream, packed_16, nullptr), + "cuLaunchKernel failed" + ); }); check_cu(cuStreamSynchronize(stream), "sync"); suite.run("launch.launch_256_args", [&]() { - check_cu(cuLaunchKernel(kernel_256_args, 1, 1, 1, 1, 1, 1, 0, stream, packed_256, nullptr), - "cuLaunchKernel failed"); + check_cu( + cuLaunchKernel(kernel_256_args, 1, 1, 1, 1, 1, 1, 0, stream, packed_256, nullptr), + "cuLaunchKernel failed" + ); }); check_cu(cuStreamSynchronize(stream), "sync"); suite.run("launch.launch_512_args", [&]() { - check_cu(cuLaunchKernel(kernel_512_args, 1, 1, 1, 1, 1, 1, 0, stream, packed_512, nullptr), - "cuLaunchKernel failed"); + check_cu( + cuLaunchKernel(kernel_512_args, 1, 1, 1, 1, 1, 1, 0, stream, packed_512, nullptr), + "cuLaunchKernel failed" + ); }); check_cu(cuStreamSynchronize(stream), "sync"); suite.run("launch.launch_512_args_pre_packed", [&]() { - check_cu(cuLaunchKernel(kernel_512_args, 1, 1, 1, 1, 1, 1, 0, stream, packed_512, nullptr), - "cuLaunchKernel failed"); + check_cu( + cuLaunchKernel(kernel_512_args, 1, 1, 1, 1, 1, 1, 0, stream, packed_512, nullptr), + "cuLaunchKernel failed" + ); }); check_cu(cuStreamSynchronize(stream), "sync"); suite.run("launch.launch_512_bools", [&]() { - check_cu(cuLaunchKernel(kernel_512_bools, 1, 1, 1, 1, 1, 1, 0, stream, bool_params, nullptr), - "cuLaunchKernel failed"); + check_cu( + cuLaunchKernel(kernel_512_bools, 1, 1, 1, 1, 1, 1, 0, stream, bool_params, nullptr), + "cuLaunchKernel failed" + ); }); check_cu(cuStreamSynchronize(stream), "sync"); suite.run("launch.launch_512_ints", [&]() { - check_cu(cuLaunchKernel(kernel_512_ints, 1, 1, 1, 1, 1, 1, 0, stream, int_params, nullptr), - "cuLaunchKernel failed"); + check_cu( + cuLaunchKernel(kernel_512_ints, 1, 1, 1, 1, 1, 1, 0, stream, int_params, nullptr), + "cuLaunchKernel failed" + ); }); check_cu(cuStreamSynchronize(stream), "sync"); suite.run("launch.launch_512_doubles", [&]() { - check_cu(cuLaunchKernel(kernel_512_doubles, 1, 1, 1, 1, 1, 1, 0, stream, double_params, nullptr), - "cuLaunchKernel failed"); + check_cu( + cuLaunchKernel(kernel_512_doubles, 1, 1, 1, 1, 1, 1, 0, stream, double_params, nullptr), + "cuLaunchKernel failed" + ); }); check_cu(cuStreamSynchronize(stream), "sync"); suite.run("launch.launch_512_bytes", [&]() { - check_cu(cuLaunchKernel(kernel_512_chars, 1, 1, 1, 1, 1, 1, 0, stream, char_params, nullptr), - "cuLaunchKernel failed"); + check_cu( + cuLaunchKernel(kernel_512_chars, 1, 1, 1, 1, 1, 1, 0, stream, char_params, nullptr), + "cuLaunchKernel failed" + ); }); check_cu(cuStreamSynchronize(stream), "sync"); suite.run("launch.launch_512_longlongs", [&]() { - check_cu(cuLaunchKernel(kernel_512_longlongs, 1, 1, 1, 1, 1, 1, 0, stream, ll_params, nullptr), - "cuLaunchKernel failed"); + check_cu( + cuLaunchKernel(kernel_512_longlongs, 1, 1, 1, 1, 1, 1, 0, stream, ll_params, nullptr), + "cuLaunchKernel failed" + ); }); check_cu(cuStreamSynchronize(stream), "sync"); suite.run("launch.launch_2048b", [&]() { - check_cu(cuLaunchKernel(kernel_2048B, 1, 1, 1, 1, 1, 1, 0, stream, struct_params, nullptr), - "cuLaunchKernel failed"); + check_cu( + cuLaunchKernel(kernel_2048B, 1, 1, 1, 1, 1, 1, 0, stream, struct_params, nullptr), + "cuLaunchKernel failed" + ); }); check_cu(cuStreamSynchronize(stream), "sync"); diff --git a/benchmarks/cuda_bindings/benchmarks/cpp/bench_memory.cpp b/benchmarks/cuda_bindings/benchmarks/cpp/bench_memory.cpp index 803363be480..b85b762f20e 100644 --- a/benchmarks/cuda_bindings/benchmarks/cpp/bench_memory.cpp +++ b/benchmarks/cuda_bindings/benchmarks/cpp/bench_memory.cpp @@ -11,22 +11,23 @@ #include #include - -static void check_cu(CUresult status, const char* message) { +static void +check_cu(CUresult status, const char *message) +{ if (status != CUDA_SUCCESS) { - const char* error_name = nullptr; + const char *error_name = nullptr; cuGetErrorName(status, &error_name); std::cerr << message << ": " << (error_name ? error_name : "unknown") << '\n'; std::exit(1); } } - static constexpr size_t ALLOC_SIZE = 1024; static constexpr size_t COPY_SIZE = 8; - -int main(int argc, char** argv) { +int +main(int argc, char **argv) +{ bench::Options options = bench::parse_args(argc, argv); // Setup diff --git a/benchmarks/cuda_bindings/benchmarks/cpp/bench_module.cpp b/benchmarks/cuda_bindings/benchmarks/cpp/bench_module.cpp index 802025a6506..9dd84d46edd 100644 --- a/benchmarks/cuda_bindings/benchmarks/cpp/bench_module.cpp +++ b/benchmarks/cuda_bindings/benchmarks/cpp/bench_module.cpp @@ -12,36 +12,47 @@ #include #include - -static void check_cu(CUresult status, const char* message) { +static void +check_cu(CUresult status, const char *message) +{ if (status != CUDA_SUCCESS) { - const char* error_name = nullptr; + const char *error_name = nullptr; cuGetErrorName(status, &error_name); std::cerr << message << ": " << (error_name ? error_name : "unknown") << '\n'; std::exit(1); } } -static void check_nvrtc(nvrtcResult status, const char* message) { +static void +check_nvrtc(nvrtcResult status, const char *message) +{ if (status != NVRTC_SUCCESS) { std::cerr << message << ": " << nvrtcGetErrorString(status) << '\n'; std::exit(1); } } -static std::vector compile_cubin(const char* source, CUdevice device) { +static std::vector +compile_cubin(const char *source, CUdevice device) +{ int major = 0, minor = 0; - check_cu(cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, device), - "cuDeviceGetAttribute failed"); - check_cu(cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, device), - "cuDeviceGetAttribute failed"); + check_cu( + cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, device), + "cuDeviceGetAttribute failed" + ); + check_cu( + cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, device), + "cuDeviceGetAttribute failed" + ); nvrtcProgram prog; - check_nvrtc(nvrtcCreateProgram(&prog, source, "benchmark_kernel.cu", 0, nullptr, nullptr), - "nvrtcCreateProgram failed"); + check_nvrtc( + nvrtcCreateProgram(&prog, source, "benchmark_kernel.cu", 0, nullptr, nullptr), + "nvrtcCreateProgram failed" + ); std::string arch = "--gpu-architecture=sm_" + std::to_string(major) + std::to_string(minor); - const char* opts[] = {"--fmad=false", arch.c_str()}; + const char *opts[] = {"--fmad=false", arch.c_str()}; nvrtcResult compile_result = nvrtcCompileProgram(prog, 2, opts); if (compile_result != NVRTC_SUCCESS) { @@ -62,13 +73,13 @@ static std::vector compile_cubin(const char* source, CUdevice device) { return cubin; } - -static const char* KERNEL_SOURCE = R"( +static const char *KERNEL_SOURCE = R"( extern "C" __global__ void empty_kernel() { return; } )"; - -int main(int argc, char** argv) { +int +main(int argc, char **argv) +{ bench::Options options = bench::parse_args(argc, argv); // Setup @@ -88,8 +99,7 @@ int main(int argc, char** argv) { check_cu(cuModuleLoadData(&persistent_module, cubin.data()), "cuModuleLoadData failed"); CUfunction function; - check_cu(cuModuleGetFunction(&function, persistent_module, "empty_kernel"), - "cuModuleGetFunction failed"); + check_cu(cuModuleGetFunction(&function, persistent_module, "empty_kernel"), "cuModuleGetFunction failed"); bench::BenchmarkSuite suite(options); @@ -106,8 +116,7 @@ int main(int argc, char** argv) { { CUfunction f; suite.run("module.module_get_function", [&]() { - check_cu(cuModuleGetFunction(&f, persistent_module, "empty_kernel"), - "cuModuleGetFunction failed"); + check_cu(cuModuleGetFunction(&f, persistent_module, "empty_kernel"), "cuModuleGetFunction failed"); }); } @@ -115,8 +124,10 @@ int main(int argc, char** argv) { { int value = 0; suite.run("module.func_get_attribute", [&]() { - check_cu(cuFuncGetAttribute(&value, CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, function), - "cuFuncGetAttribute failed"); + check_cu( + cuFuncGetAttribute(&value, CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, function), + "cuFuncGetAttribute failed" + ); }); } diff --git a/benchmarks/cuda_bindings/benchmarks/cpp/bench_nvrtc.cpp b/benchmarks/cuda_bindings/benchmarks/cpp/bench_nvrtc.cpp index 842b7353cf4..80f174fd1fe 100644 --- a/benchmarks/cuda_bindings/benchmarks/cpp/bench_nvrtc.cpp +++ b/benchmarks/cuda_bindings/benchmarks/cpp/bench_nvrtc.cpp @@ -12,31 +12,34 @@ #include #include - -static void check_cu(CUresult status, const char* message) { +static void +check_cu(CUresult status, const char *message) +{ if (status != CUDA_SUCCESS) { - const char* error_name = nullptr; + const char *error_name = nullptr; cuGetErrorName(status, &error_name); std::cerr << message << ": " << (error_name ? error_name : "unknown") << '\n'; std::exit(1); } } -static void check_nvrtc(nvrtcResult status, const char* message) { +static void +check_nvrtc(nvrtcResult status, const char *message) +{ if (status != NVRTC_SUCCESS) { std::cerr << message << ": " << nvrtcGetErrorString(status) << '\n'; std::exit(1); } } - -static const char* KERNEL_SOURCE = R"(extern "C" __global__ void empty_kernel() { return; })"; -static const char* PROGRAM_NAME = "benchmark_kernel.cu"; +static const char *KERNEL_SOURCE = R"(extern "C" __global__ void empty_kernel() { return; })"; +static const char *PROGRAM_NAME = "benchmark_kernel.cu"; static constexpr int NUM_HEADERS = 100; - -int main(int argc, char** argv) { +int +main(int argc, char **argv) +{ bench::Options options = bench::parse_args(argc, argv); // Setup: need CUDA init to query compute capability for compile options @@ -46,19 +49,23 @@ int main(int argc, char** argv) { check_cu(cuDeviceGet(&device, 0), "cuDeviceGet failed"); int major = 0, minor = 0; - check_cu(cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, device), - "cuDeviceGetAttribute failed"); - check_cu(cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, device), - "cuDeviceGetAttribute failed"); + check_cu( + cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, device), + "cuDeviceGetAttribute failed" + ); + check_cu( + cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, device), + "cuDeviceGetAttribute failed" + ); std::string arch = "--gpu-architecture=sm_" + std::to_string(major) + std::to_string(minor); - const char* compile_opts[] = {"--fmad=false", arch.c_str()}; + const char *compile_opts[] = {"--fmad=false", arch.c_str()}; // Pre-build 100 empty headers std::vector header_name_strs(NUM_HEADERS); - std::vector header_names(NUM_HEADERS); - std::vector header_sources(NUM_HEADERS); - static const char* empty_header = "// empty"; + std::vector header_names(NUM_HEADERS); + std::vector header_sources(NUM_HEADERS); + static const char *empty_header = "// empty"; for (int i = 0; i < NUM_HEADERS; ++i) { header_name_strs[i] = "header_" + std::to_string(i) + ".cuh"; header_names[i] = header_name_strs[i].c_str(); @@ -84,8 +91,14 @@ int main(int argc, char** argv) { nvrtcProgram prog; suite.run("nvrtc.nvrtc_create_program_100_headers", [&]() { check_nvrtc( - nvrtcCreateProgram(&prog, KERNEL_SOURCE, PROGRAM_NAME, - NUM_HEADERS, header_sources.data(), header_names.data()), + nvrtcCreateProgram( + &prog, + KERNEL_SOURCE, + PROGRAM_NAME, + NUM_HEADERS, + header_sources.data(), + header_names.data() + ), "nvrtcCreateProgram failed" ); check_nvrtc(nvrtcDestroyProgram(&prog), "nvrtcDestroyProgram failed"); @@ -101,10 +114,7 @@ int main(int argc, char** argv) { nvrtcCreateProgram(&prog, KERNEL_SOURCE, PROGRAM_NAME, 0, nullptr, nullptr), "nvrtcCreateProgram failed" ); - check_nvrtc( - nvrtcCompileProgram(prog, 2, compile_opts), - "nvrtcCompileProgram failed" - ); + check_nvrtc(nvrtcCompileProgram(prog, 2, compile_opts), "nvrtcCompileProgram failed"); check_nvrtc(nvrtcDestroyProgram(&prog), "nvrtcDestroyProgram failed"); }); } diff --git a/benchmarks/cuda_bindings/benchmarks/cpp/bench_pointer_attributes.cpp b/benchmarks/cuda_bindings/benchmarks/cpp/bench_pointer_attributes.cpp index 6c6f42a58e1..80221fe8ddd 100644 --- a/benchmarks/cuda_bindings/benchmarks/cpp/bench_pointer_attributes.cpp +++ b/benchmarks/cuda_bindings/benchmarks/cpp/bench_pointer_attributes.cpp @@ -9,18 +9,20 @@ #include #include - -static void check_cu(CUresult status, const char* message) { +static void +check_cu(CUresult status, const char *message) +{ if (status != CUDA_SUCCESS) { - const char* error_name = nullptr; + const char *error_name = nullptr; cuGetErrorName(status, &error_name); std::cerr << message << ": " << (error_name ? error_name : "unknown") << '\n'; std::exit(1); } } - -int main(int argc, char** argv) { +int +main(int argc, char **argv) +{ bench::Options options = bench::parse_args(argc, argv); // Setup: init CUDA, allocate memory @@ -53,7 +55,7 @@ int main(int argc, char** argv) { { unsigned int memory_type = 0; CUdeviceptr dev_ptr_out = 0; - void* host_ptr_out = nullptr; + void *host_ptr_out = nullptr; unsigned long long buffer_id = 0; CUpointer_attribute attrs[4] = { @@ -62,13 +64,10 @@ int main(int argc, char** argv) { CU_POINTER_ATTRIBUTE_HOST_POINTER, CU_POINTER_ATTRIBUTE_BUFFER_ID, }; - void* data[4] = {&memory_type, &dev_ptr_out, &host_ptr_out, &buffer_id}; + void *data[4] = {&memory_type, &dev_ptr_out, &host_ptr_out, &buffer_id}; suite.run("pointer_attributes.pointer_get_attributes", [&]() { - check_cu( - cuPointerGetAttributes(4, attrs, data, ptr), - "cuPointerGetAttributes failed" - ); + check_cu(cuPointerGetAttributes(4, attrs, data, ptr), "cuPointerGetAttributes failed"); }); } diff --git a/benchmarks/cuda_bindings/benchmarks/cpp/bench_stream.cpp b/benchmarks/cuda_bindings/benchmarks/cpp/bench_stream.cpp index 95ad0790f9f..cc7e52c3f4d 100644 --- a/benchmarks/cuda_bindings/benchmarks/cpp/bench_stream.cpp +++ b/benchmarks/cuda_bindings/benchmarks/cpp/bench_stream.cpp @@ -9,18 +9,20 @@ #include #include - -static void check_cu(CUresult status, const char* message) { +static void +check_cu(CUresult status, const char *message) +{ if (status != CUDA_SUCCESS) { - const char* error_name = nullptr; + const char *error_name = nullptr; cuGetErrorName(status, &error_name); std::cerr << message << ": " << (error_name ? error_name : "unknown") << '\n'; std::exit(1); } } - -int main(int argc, char** argv) { +int +main(int argc, char **argv) +{ bench::Options options = bench::parse_args(argc, argv); // Setup diff --git a/benchmarks/cuda_bindings/benchmarks/cpp/bench_support.hpp b/benchmarks/cuda_bindings/benchmarks/cpp/bench_support.hpp index 1207ee7c4dc..ce638aa977d 100644 --- a/benchmarks/cuda_bindings/benchmarks/cpp/bench_support.hpp +++ b/benchmarks/cuda_bindings/benchmarks/cpp/bench_support.hpp @@ -42,7 +42,9 @@ struct RunResult { std::vector values; // seconds per loop }; -inline Options parse_args(int argc, char** argv) { +inline Options +parse_args(int argc, char **argv) +{ Options options; for (int i = 1; i < argc; ++i) { const std::string arg(argv[i]); @@ -102,7 +104,9 @@ inline Options parse_args(int argc, char** argv) { return options; } -inline std::string iso_now() { +inline std::string +iso_now() +{ const auto now = std::chrono::system_clock::now(); const std::time_t t = std::chrono::system_clock::to_time_t(now); std::tm tm{}; @@ -121,16 +125,20 @@ inline std::string iso_now() { // true when calibration reached `max_loops` before hitting `min_time_sec` // (meaning --min-time was NOT actually satisfied by the calibration). template -std::uint64_t calibrate_loops( - const Options& options, - Fn&& fn, - const std::function& post_calibrate = {}, - bool* capped_out = nullptr, - double* last_elapsed_out = nullptr -) { +std::uint64_t +calibrate_loops( + const Options &options, + Fn &&fn, + const std::function &post_calibrate = {}, + bool *capped_out = nullptr, + double *last_elapsed_out = nullptr +) +{ if (options.min_time_sec <= 0.0) { - if (capped_out) *capped_out = false; - if (last_elapsed_out) *last_elapsed_out = 0.0; + if (capped_out) + *capped_out = false; + if (last_elapsed_out) + *last_elapsed_out = 0.0; return options.loops; } @@ -175,7 +183,8 @@ std::uint64_t calibrate_loops( } if (loops > max_loops / 2) { loops = max_loops; - } else { + } + else { loops *= 2; } } @@ -187,15 +196,19 @@ std::uint64_t calibrate_loops( } } - if (capped_out) *capped_out = best_capped; - if (last_elapsed_out) *last_elapsed_out = best_elapsed; + if (capped_out) + *capped_out = best_capped; + if (last_elapsed_out) + *last_elapsed_out = best_elapsed; return best_loops; } // Run a benchmark function. The function signature is: void fn() — one call = one operation. // The harness calls fn() in a tight loop `loops` times per value. template -std::vector run_benchmark(const Options& options, Fn&& fn) { +std::vector +run_benchmark(const Options &options, Fn &&fn) +{ std::vector results; results.reserve(options.runs); @@ -234,10 +247,12 @@ std::vector run_benchmark(const Options& options, Fn&& fn) { return results; } -inline void print_summary(const std::string& name, const std::vector& results) { +inline void +print_summary(const std::string &name, const std::vector &results) +{ // Collect all timed values std::vector all_values; - for (const auto& run : results) { + for (const auto &run : results) { for (double v : run.values) { all_values.push_back(v); } @@ -246,8 +261,9 @@ inline void print_summary(const std::string& name, const std::vector& return; double sum = 0; - for (double v : all_values) + for (double v : all_values) { sum += v; + } double mean = sum / static_cast(all_values.size()); @@ -258,23 +274,25 @@ inline void print_summary(const std::string& name, const std::vector& } double stdev = std::sqrt(sq_sum / static_cast(all_values.size())); - std::cout << name << ": Mean +- std dev: " - << std::fixed << std::setprecision(0) - << (mean * 1e9) << " ns +- " + std::cout << name << ": Mean +- std dev: " << std::fixed << std::setprecision(0) << (mean * 1e9) << " ns +- " << (stdev * 1e9) << " ns\n"; } // Escape a JSON string (minimal — no control chars expected) -inline std::string json_str(const std::string& s) { +inline std::string +json_str(const std::string &s) +{ return "\"" + s + "\""; } -inline void write_pyperf_json( - const std::string& output_path, - const std::string& name, +inline void +write_pyperf_json( + const std::string &output_path, + const std::string &name, std::uint64_t loops, - const std::vector& results -) { + const std::vector &results +) +{ std::ofstream out(output_path); if (!out) { std::cerr << "Failed to open output file: " << output_path << '\n'; @@ -293,8 +311,9 @@ inline void write_pyperf_json( out << "\"benchmarks\": [{\"runs\": ["; for (std::size_t r = 0; r < results.size(); ++r) { - const auto& run = results[r]; - if (r > 0) out << ", "; + const auto &run = results[r]; + if (r > 0) + out << ", "; out << "{\"metadata\": {"; out << "\"date\": " << json_str(run.date) << ", "; @@ -304,7 +323,8 @@ inline void write_pyperf_json( // Warmups: array of [loops, value] pairs out << "\"warmups\": ["; for (std::size_t w = 0; w < run.warmup_values.size(); ++w) { - if (w > 0) out << ", "; + if (w > 0) + out << ", "; out << "[" << loops << ", " << run.warmup_values[w] << "]"; } out << "], "; @@ -312,7 +332,8 @@ inline void write_pyperf_json( // Values out << "\"values\": ["; for (std::size_t v = 0; v < run.values.size(); ++v) { - if (v > 0) out << ", "; + if (v > 0) + out << ", "; out << run.values[v]; } out << "]}"; @@ -331,7 +352,7 @@ struct BenchmarkEntry { // Collect multiple benchmarks from a single binary and write them all // to one pyperf-compatible JSON file. class BenchmarkSuite { -public: + public: explicit BenchmarkSuite(Options options) : options_(std::move(options)) {} // Post-calibration hook. If set, invoked between calibration probes so @@ -339,7 +360,9 @@ class BenchmarkSuite { // next one runs. The final probe leaves the benchmark in a drained state // before the first measured warmup/value. Can be overridden per-call via // the `post_calibrate` parameter on `run()`. - void set_post_calibrate(std::function hook) { + void + set_post_calibrate(std::function hook) + { post_calibrate_ = std::move(hook); } @@ -347,11 +370,9 @@ class BenchmarkSuite { // If --min-time is set, loop count is auto-calibrated. `post_calibrate`, // if provided, runs between calibration probes to reset async state. template - void run( - const std::string& name, - Fn&& fn, - std::function post_calibrate = {} - ) { + void + run(const std::string &name, Fn &&fn, std::function post_calibrate = {}) + { std::uint64_t loops = options_.loops; Options custom = options_; if (options_.min_time_sec > 0.0) { @@ -367,12 +388,9 @@ class BenchmarkSuite { // or a fixed count for slow ops like compilation). When --min-time is set, // calibration still runs but starts from `loops_override` as the minimum. template - void run( - const std::string& name, - std::uint64_t loops_override, - Fn&& fn, - std::function post_calibrate = {} - ) { + void + run(const std::string &name, std::uint64_t loops_override, Fn &&fn, std::function post_calibrate = {}) + { std::uint64_t loops = loops_override; Options custom = options_; custom.loops = loops_override; @@ -388,18 +406,22 @@ class BenchmarkSuite { } // Write all collected benchmarks to the output file (if -o was given). - void write() const { + void + write() const + { if (options_.output_path.empty() || entries_.empty()) return; write_multi_pyperf_json(options_.output_path, entries_); } -private: + private: Options options_; std::vector entries_; std::function post_calibrate_; - std::function select_post_calibrate(const std::function& per_call) const { + std::function + select_post_calibrate(const std::function &per_call) const + { if (per_call) { return per_call; } @@ -407,20 +429,19 @@ class BenchmarkSuite { } template - std::uint64_t calibrate_and_warn( - const std::string& name, - const Options& calib_opts, - Fn&& fn, - const std::function& post_calibrate - ) const { + std::uint64_t + calibrate_and_warn( + const std::string &name, + const Options &calib_opts, + Fn &&fn, + const std::function &post_calibrate + ) const + { bool capped = false; double last_elapsed = 0.0; - std::uint64_t loops = calibrate_loops( - calib_opts, std::forward(fn), post_calibrate, &capped, &last_elapsed - ); + std::uint64_t loops = calibrate_loops(calib_opts, std::forward(fn), post_calibrate, &capped, &last_elapsed); if (capped) { - std::cerr << "WARNING: " << name - << ": calibration hit --max-loops (" << calib_opts.max_loops + std::cerr << "WARNING: " << name << ": calibration hit --max-loops (" << calib_opts.max_loops << ") before reaching --min-time (" << calib_opts.min_time_sec << "s). Last sample: " << last_elapsed << "s. Raise --max-loops to satisfy --min-time for this benchmark.\n"; @@ -428,10 +449,9 @@ class BenchmarkSuite { return loops; } - static void write_multi_pyperf_json( - const std::string& output_path, - const std::vector& entries - ) { + static void + write_multi_pyperf_json(const std::string &output_path, const std::vector &entries) + { std::ofstream out(output_path); if (!out) { std::cerr << "Failed to open output file: " << output_path << '\n'; @@ -442,8 +462,9 @@ class BenchmarkSuite { out << "{\"version\": \"1.0\", \"benchmarks\": ["; for (std::size_t e = 0; e < entries.size(); ++e) { - const auto& entry = entries[e]; - if (e > 0) out << ", "; + const auto &entry = entries[e]; + if (e > 0) + out << ", "; out << "{\"metadata\": {"; out << "\"name\": " << json_str(entry.name) << ", "; @@ -452,8 +473,9 @@ class BenchmarkSuite { out << "}, \"runs\": ["; for (std::size_t r = 0; r < entry.results.size(); ++r) { - const auto& run = entry.results[r]; - if (r > 0) out << ", "; + const auto &run = entry.results[r]; + if (r > 0) + out << ", "; out << "{\"metadata\": {"; out << "\"date\": " << json_str(run.date) << ", "; @@ -462,14 +484,16 @@ class BenchmarkSuite { out << "\"warmups\": ["; for (std::size_t w = 0; w < run.warmup_values.size(); ++w) { - if (w > 0) out << ", "; + if (w > 0) + out << ", "; out << "[" << entry.loops << ", " << run.warmup_values[w] << "]"; } out << "], "; out << "\"values\": ["; for (std::size_t v = 0; v < run.values.size(); ++v) { - if (v > 0) out << ", "; + if (v > 0) + out << ", "; out << run.values[v]; } out << "]}"; diff --git a/benchmarks/cuda_bindings/benchmarks/cpp/bench_tensormap.cpp b/benchmarks/cuda_bindings/benchmarks/cpp/bench_tensormap.cpp index d2ea5d6b20c..5ef7b219bb9 100644 --- a/benchmarks/cuda_bindings/benchmarks/cpp/bench_tensormap.cpp +++ b/benchmarks/cuda_bindings/benchmarks/cpp/bench_tensormap.cpp @@ -9,18 +9,20 @@ #include #include - -static void check_cu(CUresult status, const char* message) { +static void +check_cu(CUresult status, const char *message) +{ if (status != CUDA_SUCCESS) { - const char* error_name = nullptr; + const char *error_name = nullptr; cuGetErrorName(status, &error_name); std::cerr << message << ": " << (error_name ? error_name : "unknown") << '\n'; std::exit(1); } } - -int main(int argc, char** argv) { +int +main(int argc, char **argv) +{ bench::Options options = bench::parse_args(argc, argv); // Setup: init CUDA, allocate device memory @@ -52,7 +54,7 @@ int main(int argc, char** argv) { &tensor_map, CU_TENSOR_MAP_DATA_TYPE_FLOAT32, 2u, - reinterpret_cast(ptr), + reinterpret_cast(ptr), global_dim, global_strides, box_dim, @@ -69,7 +71,7 @@ int main(int argc, char** argv) { &tensor_map, CU_TENSOR_MAP_DATA_TYPE_FLOAT32, 2u, - reinterpret_cast(ptr), + reinterpret_cast(ptr), global_dim, global_strides, box_dim, @@ -82,11 +84,11 @@ int main(int argc, char** argv) { "cuTensorMapEncodeTiled failed" ); }); - } else { - const char* err_name = nullptr; + } + else { + const char *err_name = nullptr; cuGetErrorName(probe, &err_name); - std::cerr << "Skipping tensormap.tensor_map_encode_tiled: " - << (err_name ? err_name : "unknown") << '\n'; + std::cerr << "Skipping tensormap.tensor_map_encode_tiled: " << (err_name ? err_name : "unknown") << '\n'; } } @@ -102,7 +104,7 @@ int main(int argc, char** argv) { &tensor_map, CU_TENSOR_MAP_DATA_TYPE_FLOAT16, 3u, - reinterpret_cast(ptr), + reinterpret_cast(ptr), global_dim, global_strides, pixel_box_lower, @@ -122,7 +124,7 @@ int main(int argc, char** argv) { &tensor_map, CU_TENSOR_MAP_DATA_TYPE_FLOAT16, 3u, - reinterpret_cast(ptr), + reinterpret_cast(ptr), global_dim, global_strides, pixel_box_lower, @@ -138,11 +140,11 @@ int main(int argc, char** argv) { "cuTensorMapEncodeIm2col failed" ); }); - } else { - const char* err_name = nullptr; + } + else { + const char *err_name = nullptr; cuGetErrorName(probe, &err_name); - std::cerr << "Skipping tensormap.tensor_map_encode_im2col: " - << (err_name ? err_name : "unknown") << '\n'; + std::cerr << "Skipping tensormap.tensor_map_encode_im2col: " << (err_name ? err_name : "unknown") << '\n'; } } @@ -157,7 +159,7 @@ int main(int argc, char** argv) { &tensor_map, CU_TENSOR_MAP_DATA_TYPE_FLOAT16, 3u, - reinterpret_cast(ptr), + reinterpret_cast(ptr), global_dim, global_strides, 0, @@ -179,7 +181,7 @@ int main(int argc, char** argv) { &tensor_map, CU_TENSOR_MAP_DATA_TYPE_FLOAT16, 3u, - reinterpret_cast(ptr), + reinterpret_cast(ptr), global_dim, global_strides, 0, @@ -196,11 +198,12 @@ int main(int argc, char** argv) { "cuTensorMapEncodeIm2colWide failed" ); }); - } else { - const char* err_name = nullptr; + } + else { + const char *err_name = nullptr; cuGetErrorName(probe, &err_name); - std::cerr << "Skipping tensormap.tensor_map_encode_im2col_wide: " - << (err_name ? err_name : "unknown") << '\n'; + std::cerr << "Skipping tensormap.tensor_map_encode_im2col_wide: " << (err_name ? err_name : "unknown") + << '\n'; } } #else diff --git a/cuda_core/cuda/core/_cpp/rt/api.hpp b/cuda_core/cuda/core/_cpp/rt/api.hpp index cd39a27c91e..399343c2707 100644 --- a/cuda_core/cuda/core/_cpp/rt/api.hpp +++ b/cuda_core/cuda/core/_cpp/rt/api.hpp @@ -17,45 +17,51 @@ namespace cuda_core::rt { // ============================================================================ // Function to create a non-owning context handle (references existing context). -ContextHandle create_context_handle_ref(CUcontext ctx); +ContextHandle +create_context_handle_ref(CUcontext ctx); // Create a context handle for the CUcontext view of the provided green context. // The returned ContextHandle keeps the green context alive, but the CUcontext // view is non-owning and is not destroyed independently. -ContextHandle create_context_handle_from_green_ctx(const GreenCtxHandle& h_green_ctx); +ContextHandle +create_context_handle_from_green_ctx(const GreenCtxHandle &h_green_ctx); // Return the green context dependency associated with a ContextHandle, if any. -GreenCtxHandle get_context_green_ctx(const ContextHandle& h) noexcept; +GreenCtxHandle +get_context_green_ctx(const ContextHandle &h) noexcept; // Create an owning green context handle from a list of device resources. -GreenCtxHandle create_green_ctx_handle(CUdevResource* resources, unsigned int nbResources, - CUdevice dev, unsigned int flags); +GreenCtxHandle +create_green_ctx_handle(CUdevResource *resources, unsigned int nbResources, CUdevice dev, unsigned int flags); // Create a non-owning green context handle. -GreenCtxHandle create_green_ctx_handle_ref(CUgreenCtx ctx); +GreenCtxHandle +create_green_ctx_handle_ref(CUgreenCtx ctx); // Get handle to the primary context for a device (with thread-local caching) // Returns empty handle on error (caller must check) -ContextHandle get_primary_context(int device_id); +ContextHandle +get_primary_context(int device_id); // Get handle to the current CUDA context // Returns empty handle if no context is current (caller must check) -ContextHandle get_current_context(); +ContextHandle +get_current_context(); // Synchronize the provided context. Releases the GIL around the driver call. // Returns CUDA_ERROR_INVALID_CONTEXT for an empty handle. -CUresult context_synchronize(const ContextHandle& h_context) noexcept; +CUresult +context_synchronize(const ContextHandle &h_context) noexcept; // Query the stream priority range for the provided context. // Returns CUDA_ERROR_INVALID_CONTEXT for an empty handle. -CUresult context_get_stream_priority_range( - const ContextHandle& h_context, - int* least_priority, - int* greatest_priority) noexcept; +CUresult +context_get_stream_priority_range(const ContextHandle &h_context, int *least_priority, int *greatest_priority) noexcept; // Query the device of the provided context. // Returns CUDA_ERROR_INVALID_CONTEXT for an empty handle. -CUresult context_get_device(const ContextHandle& h_context, CUdevice* device) noexcept; +CUresult +context_get_device(const ContextHandle &h_context, CUdevice *device) noexcept; // Call cuGraphNodeSetParams with h_context current (empty handle: the caller's // context). Returns the update status; *restore_status receives a failure to @@ -63,11 +69,13 @@ CUresult context_get_device(const ContextHandle& h_context, CUdevice* device) no // raises only after publishing the metadata that depends on the update. // Returns CUDA_ERROR_NOT_SUPPORTED when the driver lacks cuGraphNodeSetParams. // Implemented in graph.cpp -CUresult graph_node_set_params( +CUresult +graph_node_set_params( CUgraphNode node, - CUgraphNodeParams* params, - const ContextHandle& h_context, - CUresult* restore_status) noexcept; + CUgraphNodeParams *params, + const ContextHandle &h_context, + CUresult *restore_status +) noexcept; // ============================================================================ // Stream handle functions @@ -77,31 +85,38 @@ CUresult graph_node_set_params( // The stream structurally depends on the provided context handle. // When the last reference is released, cuStreamDestroy is called automatically. // Returns empty handle on error (caller must check). -StreamHandle create_stream_handle(const ContextHandle& h_ctx, unsigned int flags, int priority); +StreamHandle +create_stream_handle(const ContextHandle &h_ctx, unsigned int flags, int priority); // Create a non-owning stream handle (references existing stream). // Use for borrowed streams (from foreign code) or built-in streams. // The stream will NOT be destroyed when the handle is released. // Caller is responsible for keeping the stream's context alive. -StreamHandle create_stream_handle_ref(CUstream stream); +StreamHandle +create_stream_handle_ref(CUstream stream); // Initialize the process-lifetime CUDA user-object cleanup queue. Called once // from module initialization while Python is fully initialized. // Implemented in py_deferred_cleanup.cpp -void initialize_deferred_cleanup(); +void +initialize_deferred_cleanup(); // Implemented in py_deferred_cleanup.cpp -void retry_deferred_cleanup() noexcept; +void +retry_deferred_cleanup() noexcept; // Return the context dependency associated with a stream handle, if any. -ContextHandle get_stream_context(const StreamHandle& h) noexcept; +ContextHandle +get_stream_context(const StreamHandle &h) noexcept; // Get non-owning handle to the legacy default stream (CU_STREAM_LEGACY) // Note: Legacy stream has no specific context dependency. -StreamHandle get_legacy_stream(); +StreamHandle +get_legacy_stream(); // Get non-owning handle to the per-thread default stream (CU_STREAM_PER_THREAD) // Note: Per-thread stream has no specific context dependency. -StreamHandle get_per_thread_stream(); +StreamHandle +get_per_thread_stream(); // Wrap CU_STREAM_LEGACY with an explicit context, bypassing the "bind to // whatever is current" resolution that a bare default-stream token uses (see @@ -109,7 +124,8 @@ StreamHandle get_per_thread_stream(); // known context (e.g. a synchronous, non-pooled allocator) record a correct // deallocation context without requiring that context to be current when the // token is created. Returns an empty handle for an empty h_context. -StreamHandle create_context_bound_legacy_stream(const ContextHandle& h_context); +StreamHandle +create_context_bound_legacy_stream(const ContextHandle &h_context); // ============================================================================ // Event handle functions @@ -120,9 +136,15 @@ StreamHandle create_context_bound_legacy_stream(const ContextHandle& h_context); // Metadata fields are stored in the EventBox for later retrieval. // When the last reference is released, cuEventDestroy is called automatically. // Returns empty handle on error (caller must check). -EventHandle create_event_handle(const ContextHandle& h_ctx, unsigned int flags, - bool timing_enabled, bool is_blocking_sync, - bool ipc_enabled, int device_id); +EventHandle +create_event_handle( + const ContextHandle &h_ctx, + unsigned int flags, + bool timing_enabled, + bool is_blocking_sync, + bool ipc_enabled, + int device_id +); // Create an owning event in the context that owns `stream`, so it can be // recorded on that stream regardless of which context is current. Default- @@ -131,27 +153,34 @@ EventHandle create_event_handle(const ContextHandle& h_ctx, unsigned int flags, // same scope; the handle carries no device id. // When the last reference is released, cuEventDestroy is called automatically. // Returns empty handle on error (caller must check). -EventHandle create_event_handle_for_stream(CUstream stream, unsigned int flags); +EventHandle +create_event_handle_for_stream(CUstream stream, unsigned int flags); // Create an owning event handle from an IPC handle. // The originating process owns the event and its context. // When the last reference is released, cuEventDestroy is called automatically. // Returns empty handle on error (caller must check). -EventHandle create_event_handle_ipc(const CUipcEventHandle& ipc_handle, - bool is_blocking_sync); +EventHandle +create_event_handle_ipc(const CUipcEventHandle &ipc_handle, bool is_blocking_sync); // Create a non-owning event handle (references existing event). // Use for events that are managed by the CUDA graph or another owner. // The event will NOT be destroyed when the handle is released. // Metadata defaults to unknown (timing_enabled=false, device_id=-1). -EventHandle create_event_handle_ref(CUevent event); +EventHandle +create_event_handle_ref(CUevent event); // Event metadata accessors (read from EventBox via pointer arithmetic) -bool get_event_timing_enabled(const EventHandle& h) noexcept; -bool get_event_is_blocking_sync(const EventHandle& h) noexcept; -bool get_event_ipc_enabled(const EventHandle& h) noexcept; -int get_event_device_id(const EventHandle& h) noexcept; -ContextHandle get_event_context(const EventHandle& h) noexcept; +bool +get_event_timing_enabled(const EventHandle &h) noexcept; +bool +get_event_is_blocking_sync(const EventHandle &h) noexcept; +bool +get_event_ipc_enabled(const EventHandle &h) noexcept; +int +get_event_device_id(const EventHandle &h) noexcept; +ContextHandle +get_event_context(const EventHandle &h) noexcept; // ============================================================================ // Memory pool handle functions @@ -161,22 +190,26 @@ ContextHandle get_event_context(const EventHandle& h) noexcept; // Memory pools are device-scoped (not context-scoped). // When the last reference is released, cuMemPoolDestroy is called automatically. // Returns empty handle on error (caller must check). -MemoryPoolHandle create_mempool_handle(const CUmemPoolProps& props); +MemoryPoolHandle +create_mempool_handle(const CUmemPoolProps &props); // Create a non-owning memory pool handle (references existing pool). // Use for device default/current pools that are managed by the driver. // The pool will NOT be destroyed when the handle is released. -MemoryPoolHandle create_mempool_handle_ref(CUmemoryPool pool); +MemoryPoolHandle +create_mempool_handle_ref(CUmemoryPool pool); // Get non-owning handle to the current memory pool for a device. // Returns empty handle on error (caller must check). -MemoryPoolHandle get_device_mempool(int device_id); +MemoryPoolHandle +get_device_mempool(int device_id); // Create an owning memory pool handle from an IPC import. // The file descriptor is NOT owned by this handle (caller manages FD separately). // When the last reference is released, cuMemPoolDestroy is called automatically. // Returns empty handle on error (caller must check). -MemoryPoolHandle create_mempool_handle_ipc(int fd, CUmemAllocationHandleType handle_type); +MemoryPoolHandle +create_mempool_handle_ipc(int fd, CUmemAllocationHandleType handle_type); // ============================================================================ // Device pointer handle functions @@ -186,61 +219,63 @@ MemoryPoolHandle create_mempool_handle_ipc(int fd, CUmemAllocationHandleType han // The pointer structurally depends on the provided pool handle (captured in deleter). // When the last reference is released, cuMemFreeAsync is called on the stored stream. // Returns empty handle on error (caller must check). -DevicePtrHandle deviceptr_alloc_from_pool( - size_t size, - const MemoryPoolHandle& h_pool, - const StreamHandle& h_stream); +DevicePtrHandle +deviceptr_alloc_from_pool(size_t size, const MemoryPoolHandle &h_pool, const StreamHandle &h_stream); // Allocate device memory asynchronously via cuMemAllocAsync. // When the last reference is released, cuMemFreeAsync is called on the stored stream. // Returns empty handle on error (caller must check). -DevicePtrHandle deviceptr_alloc_async(size_t size, const StreamHandle& h_stream); +DevicePtrHandle +deviceptr_alloc_async(size_t size, const StreamHandle &h_stream); // Allocate device memory synchronously via cuMemAlloc with the provided // context current. The caller owns the pointer and releases it with cuMemFree. // Returns CUDA_ERROR_INVALID_CONTEXT for an empty handle. -CUresult deviceptr_alloc_raw(CUdeviceptr* ptr, size_t size, - const ContextHandle& h_context) noexcept; +CUresult +deviceptr_alloc_raw(CUdeviceptr *ptr, size_t size, const ContextHandle &h_context) noexcept; // Allocate pinned host memory via cuMemAllocHost. // When the last reference is released, cuMemFreeHost is called. // Returns empty handle on error (caller must check). -DevicePtrHandle deviceptr_alloc_host(size_t size); +DevicePtrHandle +deviceptr_alloc_host(size_t size); // Create a non-owning device pointer handle (references existing pointer). // Use for foreign pointers (e.g., from external libraries). // The pointer will NOT be freed when the handle is released. -DevicePtrHandle deviceptr_create_ref(CUdeviceptr ptr); +DevicePtrHandle +deviceptr_create_ref(CUdeviceptr ptr); // Create a device pointer handle for a mapped graphics resource. // The pointer structurally depends on the provided graphics resource handle. // When the last reference is released, cuGraphicsUnmapResources is called on // the stored stream, then the graphics resource may be unregistered when its // own handle is released. -DevicePtrHandle deviceptr_create_mapped_graphics( +DevicePtrHandle +deviceptr_create_mapped_graphics( CUdeviceptr ptr, - const GraphicsResourceHandle& h_resource, - const StreamHandle& h_stream); + const GraphicsResourceHandle &h_resource, + const StreamHandle &h_stream +); // Import a device pointer from IPC via cuMemPoolImportPointer. // When the last reference is released, cuMemFreeAsync is called on the stored stream. // Note: Does not yet implement reference counting for nvbug 5570902. // On error, returns empty handle and sets thread-local error (use get_last_error()). -DevicePtrHandle deviceptr_import_ipc( - const MemoryPoolHandle& h_pool, - const void* export_data, - const StreamHandle& h_stream); +DevicePtrHandle +deviceptr_import_ipc(const MemoryPoolHandle &h_pool, const void *export_data, const StreamHandle &h_stream); // Access the deallocation stream for a device pointer handle (read-only). // For non-owning handles, the stream is not used but can still be accessed. // Returns an empty handle for an empty device pointer handle. -StreamHandle deallocation_stream(const DevicePtrHandle& h) noexcept; +StreamHandle +deallocation_stream(const DevicePtrHandle &h) noexcept; // Set the deallocation stream for a device pointer handle. // Returns CUDA_ERROR_INVALID_CONTEXT when a default-stream token cannot be // bound because no CUDA context is current. -CUresult set_deallocation_stream( - const DevicePtrHandle& h, const StreamHandle& h_stream) noexcept; +CUresult +set_deallocation_stream(const DevicePtrHandle &h, const StreamHandle &h_stream) noexcept; // ============================================================================ // Virtual memory management (VMM_DESIGN.md) @@ -256,18 +291,20 @@ CUresult set_deallocation_stream( // applied to every mapping of this allocation. When the last reference is // released, cuMemRelease is called; the memory is freed once no mapping // remains. Returns empty handle on error (caller must check). -MemAllocationHandle create_mem_allocation_handle(size_t size, const CUmemAllocationProp& prop, - const CUmemAccessDesc* descs, size_t count); +MemAllocationHandle +create_mem_allocation_handle(size_t size, const CUmemAllocationProp &prop, const CUmemAccessDesc *descs, size_t count); // Size of the allocation; the only size cuMemMap accepts for it. -size_t mem_allocation_size(const MemAllocationHandle& h) noexcept; +size_t +mem_allocation_size(const MemAllocationHandle &h) noexcept; // Whether cuMemImportFromShareableHandle produced the allocation. The driver // exports only allocations this process created with a handle type, so an // imported allocation cannot be exported again; callers check this before // cuMemExportToShareableHandle instead of relying on the driver's error. The // mark travels with the allocation into every range that maps it. -bool mem_allocation_is_imported(const MemAllocationHandle& h) noexcept; +bool +mem_allocation_is_imported(const MemAllocationHandle &h) noexcept; // Import a physical allocation that another process exported with // cuMemExportToShareableHandle. `os_handle` is the shareable handle in the @@ -278,16 +315,24 @@ bool mem_allocation_is_imported(const MemAllocationHandle& h) noexcept; // a MemAllocationHandle like one from create_mem_allocation_handle: the // access descriptors are applied to every mapping, and the last reference // calls cuMemRelease. Returns empty handle on error (caller must check). -MemAllocationHandle import_mem_allocation_handle(void* os_handle, CUmemAllocationHandleType handle_type, - size_t size, const CUmemAccessDesc* descs, size_t count); +MemAllocationHandle +import_mem_allocation_handle( + void *os_handle, + CUmemAllocationHandleType handle_type, + size_t size, + const CUmemAccessDesc *descs, + size_t count +); // Reserve an address range via cuMemAddressReserve. Pass alignment 0 for the // driver default. When the last reference is released, cuMemAddressFree is // called with the exact reserved pair. Returns empty handle on error. -VaReservationHandle create_va_reservation_handle(size_t size, size_t alignment, CUdeviceptr hint); +VaReservationHandle +create_va_reservation_handle(size_t size, size_t alignment, CUdeviceptr hint); // Size of the reservation. -size_t va_reservation_size(const VaReservationHandle& h) noexcept; +size_t +va_reservation_size(const VaReservationHandle &h) noexcept; // Test-suite instrumentation: process-wide counts of live *owned* driver // resources. Owned memory pools (cuMemPoolCreate or imported) and VA @@ -295,36 +340,44 @@ size_t va_reservation_size(const VaReservationHandle& h) noexcept; // site and uncounted in their deleter. Borrowed pool references and // MemoryResource subclasses that hold no pool or reservation are not counted. // conftest.py reads these to gate the per-test gc.collect(). -long long live_owned_mempool_count() noexcept; -long long live_va_reservation_count() noexcept; +long long +live_owned_mempool_count() noexcept; +long long +live_va_reservation_count() noexcept; // Map the whole allocation at ptr inside the reservation via cuMemMap and // apply the allocation's access descriptors. The mapping structurally depends // on both handles. When the last reference is released, cuMemUnmap is called // first. Returns empty handle on error, including a range outside the // reservation; a failed cuMemSetAccess unmaps before returning. -VaMappingHandle create_va_mapping_handle(CUdeviceptr ptr, const MemAllocationHandle& h_alloc, - const VaReservationHandle& h_res); +VaMappingHandle +create_va_mapping_handle(CUdeviceptr ptr, const MemAllocationHandle &h_alloc, const VaReservationHandle &h_res); // Mapping accessors. -size_t va_mapping_size(const VaMappingHandle& h) noexcept; -MemAllocationHandle va_mapping_allocation(const VaMappingHandle& h) noexcept; +size_t +va_mapping_size(const VaMappingHandle &h) noexcept; +MemAllocationHandle +va_mapping_allocation(const VaMappingHandle &h) noexcept; // Build an immutable range from mappings in ascending, contiguous order. A // grow builds a new range for its result and never changes the input's. // May throw std::bad_alloc. -VmmRangeHandle create_vmm_range(const std::vector& mappings); +VmmRangeHandle +create_vmm_range(const std::vector &mappings); // The range of a device pointer handle created by deviceptr_create_vmm. Only // for such handles: the caller (VirtualMemoryBuffer) guarantees the origin. // Empty for an empty handle. -VmmRangeHandle vmm_range(const DevicePtrHandle& h) noexcept; +VmmRangeHandle +vmm_range(const DevicePtrHandle &h) noexcept; // Range accessors: a copy of the mapping list, which a grow extends and turns // into a new range, and the range total. Reads of an immutable range need no // synchronization. -std::vector vmm_range_mappings(const VmmRangeHandle& range); // may throw -size_t vmm_range_total(const VmmRangeHandle& range) noexcept; +std::vector +vmm_range_mappings(const VmmRangeHandle &range); // may throw +size_t +vmm_range_total(const VmmRangeHandle &range) noexcept; // Create a device pointer handle whose box owns a range (an empty range for // a size-zero buffer). The box records no deallocation stream; set one with @@ -332,7 +385,8 @@ size_t vmm_range_total(const VmmRangeHandle& range) noexcept; // stream is synchronized (skipped with a report when that would disturb a // capture) and the box freed, which unmaps every mapping this buffer was the // last to hold. Returns empty handle for a null range handle. -DevicePtrHandle deviceptr_create_vmm(CUdeviceptr base, const VmmRangeHandle& range); +DevicePtrHandle +deviceptr_create_vmm(CUdeviceptr base, const VmmRangeHandle &range); // ============================================================================ // Library handle functions @@ -341,18 +395,21 @@ DevicePtrHandle deviceptr_create_vmm(CUdeviceptr base, const VmmRangeHandle& ran // Create an owning library handle by loading from a file path. // When the last reference is released, cuLibraryUnload is called automatically. // Returns empty handle on error (caller must check). -LibraryHandle create_library_handle_from_file(const char* path); +LibraryHandle +create_library_handle_from_file(const char *path); // Create an owning library handle by loading from memory data. // The driver makes an internal copy of the data; caller can free it after return. // When the last reference is released, cuLibraryUnload is called automatically. // Returns empty handle on error (caller must check). -LibraryHandle create_library_handle_from_data(const void* data); +LibraryHandle +create_library_handle_from_data(const void *data); // Create a non-owning library handle (references existing library). // Use for borrowed libraries (e.g., from foreign code). // The library will NOT be unloaded when the handle is released. -LibraryHandle create_library_handle_ref(CUlibrary library); +LibraryHandle +create_library_handle_ref(CUlibrary library); // ============================================================================ // Kernel handle functions @@ -362,28 +419,32 @@ LibraryHandle create_library_handle_ref(CUlibrary library); // The kernel structurally depends on the provided library handle. // Kernels have no explicit destroy - their lifetime is tied to the library. // Returns empty handle on error (caller must check). -KernelHandle create_kernel_handle(const LibraryHandle& h_library, const char* name); +KernelHandle +create_kernel_handle(const LibraryHandle &h_library, const char *name); // Create a kernel handle from a raw CUkernel. // If the kernel is already managed (in the registry), returns the owning // handle with library dependency. Otherwise returns a non-owning ref. -KernelHandle create_kernel_handle_ref(CUkernel kernel); +KernelHandle +create_kernel_handle_ref(CUkernel kernel); // Get the library handle associated with a kernel (from KernelBox). // Returns empty handle if the kernel has no library dependency. -LibraryHandle get_kernel_library(const KernelHandle& h) noexcept; +LibraryHandle +get_kernel_library(const KernelHandle &h) noexcept; // ============================================================================ // Graph handle functions // ============================================================================ // Create the owning handle for a root graph and its hierarchy. -GraphHandle create_graph_handle(CUgraph graph); +GraphHandle +create_graph_handle(CUgraph graph); // Create the canonical handle for a graph whose CUDA lifetime is owned by a // node in h_parent. -GraphHandle create_child_graph_handle( - CUgraph child_graph, const GraphHandle& h_parent, CUgraphNode owner_node); +GraphHandle +create_child_graph_handle(CUgraph child_graph, const GraphHandle &h_parent, CUgraphNode owner_node); // ============================================================================ // Graph node attachments @@ -402,60 +463,59 @@ GraphHandle create_child_graph_handle( // ============================================================================ // Build an OpaqueHandle from a malloc'd buffer: std::free on release. -OpaqueHandle make_opaque_malloc(void* buf); +OpaqueHandle +make_opaque_malloc(void *buf); // Copy requested owners from node's current attachment. Pass nullptr to ignore // either owner; a missing attachment produces empty handles. -CUresult graph_get_attachment( - const GraphHandle& h_graph, - CUgraphNode node, - OpaqueHandle* owner0, - OpaqueHandle* owner1); +CUresult +graph_get_attachment(const GraphHandle &h_graph, CUgraphNode node, OpaqueHandle *owner0, OpaqueHandle *owner1); // Create and graph-retain a replacement attachment before a CUDA mutation. // Destruction rolls the prepared attachment back unless it is committed. -CUresult graph_prepare_attachment( - const GraphHandle& h_graph, +CUresult +graph_prepare_attachment( + const GraphHandle &h_graph, OpaqueHandle owner0, OpaqueHandle owner1, - PreparedAttachment* out_prepared); + PreparedAttachment *out_prepared +); // Publish a prepared attachment after the CUDA mutation succeeds. A null node // retains the attachment anonymously without publishing node metadata. -CUresult graph_commit_attachment( - PreparedAttachment& prepared, - CUgraphNode node); +CUresult +graph_commit_attachment(PreparedAttachment &prepared, CUgraphNode node); // Copy attachment metadata from a source graph hierarchy into its CUDA clone. -CUresult graph_clone_attachments( - const GraphHandle& h_clone, - const GraphHandle& h_source); +CUresult +graph_clone_attachments(const GraphHandle &h_clone, const GraphHandle &h_source); // Stage a complete metadata replacement before CUDA replaces an embedded // graph. Dropping the prepared state leaves the current hierarchy unchanged. -CUresult graph_prepare_child_graph_update( - const GraphHandle& h_parent, - const GraphHandle& h_old_child, +CUresult +graph_prepare_child_graph_update( + const GraphHandle &h_parent, + const GraphHandle &h_old_child, CUgraphNode owner_node, - const GraphHandle& h_source, - PreparedChildGraphUpdate* out_prepared); + const GraphHandle &h_source, + PreparedChildGraphUpdate *out_prepared +); // Rekey staged metadata to CUDA's replacement clone, retire the old embedded // hierarchy, and publish the replacement handle. -CUresult graph_commit_child_graph_update( - PreparedChildGraphUpdate& prepared, - GraphHandle* out_child); +CUresult +graph_commit_child_graph_update(PreparedChildGraphUpdate &prepared, GraphHandle *out_child); // Invalidate cuda.core state for child graphs CUDA destroyed with owner_node. -void invalidate_child_graph_state( - const GraphHandle& h_parent, - CUgraphNode owner_node) noexcept; +void +invalidate_child_graph_state(const GraphHandle &h_parent, CUgraphNode owner_node) noexcept; // Invalidate cuda.core state for a root graph that CUDA destroyed itself, such // as the graph of an invalidated capture ended by cuStreamEndCapture. The // owning handle then no longer calls cuGraphDestroy. No-op unless h_root is // the live root of its hierarchy. -void invalidate_root_graph_state(const GraphHandle& h_root) noexcept; +void +invalidate_root_graph_state(const GraphHandle &h_root) noexcept; // ============================================================================ // Graph exec handle functions @@ -470,30 +530,29 @@ void invalidate_root_graph_state(const GraphHandle& h_root) noexcept; // Returns empty handle on error (caller must check). The caller reads // params->result_out for the specific instantiation failure and // get_last_error() for a driver status. -GraphExecHandle create_graph_exec_handle( - const GraphHandle& h_source, - CUDA_GRAPH_INSTANTIATE_PARAMS* params); +GraphExecHandle +create_graph_exec_handle(const GraphHandle &h_source, CUDA_GRAPH_INSTANTIATE_PARAMS *params); // Update h_exec in place by calling cuGraphExecUpdate, and publish a fresh // accumulator when CUDA accepts the update. Writes result_info for the caller. -CUresult graph_exec_update( - const GraphExecHandle& h_exec, - const GraphHandle& h_source, - CUgraphExecUpdateResultInfo* result_info); +CUresult +graph_exec_update(const GraphExecHandle &h_exec, const GraphHandle &h_source, CUgraphExecUpdateResultInfo *result_info); // Append owners before an executable-node mutation. The accumulator grows // because CUDA cannot attach user objects to an exec after instantiation, so // old owners stay reachable. Dropping the transaction restores the accumulator // to its original size. -CUresult graph_prepare_exec_attachment( - const GraphExecHandle& h_exec, +CUresult +graph_prepare_exec_attachment( + const GraphExecHandle &h_exec, OpaqueHandle owner0, OpaqueHandle owner1, - PreparedExecAttachment* out_prepared); + PreparedExecAttachment *out_prepared +); // Keep the owners added by graph_prepare_exec_attachment. -void graph_commit_exec_attachment( - PreparedExecAttachment& prepared) noexcept; +void +graph_commit_exec_attachment(PreparedExecAttachment &prepared) noexcept; // ============================================================================ // Graph node handle functions @@ -502,13 +561,16 @@ void graph_commit_exec_attachment( // Create a node handle. Nodes are owned by their parent graph (not // independently destroyable). The GraphHandle dependency ensures the // graph outlives any node reference. -GraphNodeHandle create_graph_node_handle(CUgraphNode node, const GraphHandle& h_graph); +GraphNodeHandle +create_graph_node_handle(CUgraphNode node, const GraphHandle &h_graph); // Extract the owning graph handle from a node handle. -GraphHandle graph_node_get_graph(const GraphNodeHandle& h) noexcept; +GraphHandle +graph_node_get_graph(const GraphNodeHandle &h) noexcept; // Zero the CUgraphNode resource inside the handle, marking it invalid. -void invalidate_graph_node(const GraphNodeHandle& h) noexcept; +void +invalidate_graph_node(const GraphNodeHandle &h) noexcept; // ============================================================================ // Graphics resource handle functions @@ -517,7 +579,8 @@ void invalidate_graph_node(const GraphNodeHandle& h) noexcept; // Create an owning graphics resource handle. // When the last reference is released, cuGraphicsUnregisterResource is called automatically. // Use for CUgraphicsResource handles obtained from cuGraphicsGLRegisterBuffer etc. -GraphicsResourceHandle create_graphics_resource_handle(CUgraphicsResource resource); +GraphicsResourceHandle +create_graphics_resource_handle(CUgraphicsResource resource); // ============================================================================ // NVRTC Program handle functions @@ -526,11 +589,13 @@ GraphicsResourceHandle create_graphics_resource_handle(CUgraphicsResource resour // Create an owning NVRTC program handle. // When the last reference is released, nvrtcDestroyProgram is called. // Use this to wrap a program created via nvrtcCreateProgram. -NvrtcProgramHandle create_nvrtc_program_handle(nvrtcProgram prog); +NvrtcProgramHandle +create_nvrtc_program_handle(nvrtcProgram prog); // Create a non-owning NVRTC program handle (references existing program). // The program will NOT be destroyed when the handle is released. -NvrtcProgramHandle create_nvrtc_program_handle_ref(nvrtcProgram prog); +NvrtcProgramHandle +create_nvrtc_program_handle_ref(nvrtcProgram prog); // ============================================================================ // NVVM Program handle functions @@ -540,11 +605,13 @@ NvrtcProgramHandle create_nvrtc_program_handle_ref(nvrtcProgram prog); // When the last reference is released, nvvmDestroyProgram is called. // Use this to wrap a program created via nvvmCreateProgram. // Note: If NVVM is not available (p_nvvmDestroyProgram is null), the deleter is a no-op. -NvvmProgramHandle create_nvvm_program_handle(nvvmProgram prog); +NvvmProgramHandle +create_nvvm_program_handle(nvvmProgram prog); // Create a non-owning NVVM program handle (references existing program). // The program will NOT be destroyed when the handle is released. -NvvmProgramHandle create_nvvm_program_handle_ref(nvvmProgram prog); +NvvmProgramHandle +create_nvvm_program_handle_ref(nvvmProgram prog); // ============================================================================ // nvJitLink handle functions @@ -554,11 +621,13 @@ NvvmProgramHandle create_nvvm_program_handle_ref(nvvmProgram prog); // When the last reference is released, nvJitLinkDestroy is called. // Use this to wrap a handle created via nvJitLinkCreate. // Note: If nvJitLink is not available (p_nvJitLinkDestroy is null), the deleter is a no-op. -NvJitLinkHandle create_nvjitlink_handle(nvJitLink_t handle); +NvJitLinkHandle +create_nvjitlink_handle(nvJitLink_t handle); // Create a non-owning nvJitLink handle (references existing handle). // The handle will NOT be destroyed when the last reference is released. -NvJitLinkHandle create_nvjitlink_handle_ref(nvJitLink_t handle); +NvJitLinkHandle +create_nvjitlink_handle_ref(nvJitLink_t handle); // ============================================================================ // cuLink handle functions @@ -567,11 +636,13 @@ NvJitLinkHandle create_nvjitlink_handle_ref(nvJitLink_t handle); // Create an owning cuLink handle. // When the last reference is released, cuLinkDestroy is called. // Use this to wrap a CUlinkState created via cuLinkCreate. -CuLinkHandle create_culink_handle(CUlinkState state); +CuLinkHandle +create_culink_handle(CUlinkState state); // Create a non-owning cuLink handle (references existing CUlinkState). // The handle will NOT be destroyed when the last reference is released. -CuLinkHandle create_culink_handle_ref(CUlinkState state); +CuLinkHandle +create_culink_handle_ref(CUlinkState state); // ============================================================================ // File descriptor handle functions @@ -579,10 +650,12 @@ CuLinkHandle create_culink_handle_ref(CUlinkState state); // Create an owning file descriptor handle. // When the last reference is released, POSIX close() is called. -FileDescriptorHandle create_fd_handle(int fd); +FileDescriptorHandle +create_fd_handle(int fd); // Create a non-owning file descriptor handle (caller manages the fd). -FileDescriptorHandle create_fd_handle_ref(int fd); +FileDescriptorHandle +create_fd_handle_ref(int fd); // ============================================================================ // Array / mipmapped-array / texture / surface handle functions (PR #467) @@ -597,57 +670,78 @@ FileDescriptorHandle create_fd_handle_ref(int fd); // Create an owning CUDA array via cuArray3DCreate. // When the last reference is released, cuArrayDestroy is called automatically. // Returns empty handle on error (caller must check). -OpaqueArrayHandle create_array_handle(const ContextHandle& h_context, const CUDA_ARRAY3D_DESCRIPTOR& desc); +OpaqueArrayHandle +create_array_handle(const ContextHandle &h_context, const CUDA_ARRAY3D_DESCRIPTOR &desc); // Create a non-owning array handle (references an existing CUarray). // Use for arrays owned elsewhere (e.g. graphics interop). Never destroyed here. -OpaqueArrayHandle create_array_handle_ref(CUarray arr); +OpaqueArrayHandle +create_array_handle_ref(CUarray arr); // Create an owning array handle adopting an existing CUarray. // When the last reference is released, cuArrayDestroy is called automatically. -OpaqueArrayHandle create_array_handle_owning(CUarray arr); +OpaqueArrayHandle +create_array_handle_owning(CUarray arr); // Return the context dependency associated with an array, if known. -ContextHandle get_array_context(const OpaqueArrayHandle& h) noexcept; +ContextHandle +get_array_context(const OpaqueArrayHandle &h) noexcept; // Create a non-owning handle to a mipmap level via cuMipmappedArrayGetLevel. // The level CUarray is owned by the mipmap; the parent MipmappedArrayHandle is // embedded in the box so it outlives the level view. No destroy in the deleter. // Returns empty handle on error (caller must check). -OpaqueArrayHandle create_array_level_handle(const MipmappedArrayHandle& h_mip, unsigned int level); +OpaqueArrayHandle +create_array_level_handle(const MipmappedArrayHandle &h_mip, unsigned int level); // Create an owning mipmapped array via cuMipmappedArrayCreate. // When the last reference is released, cuMipmappedArrayDestroy is called. // Returns empty handle on error (caller must check). -MipmappedArrayHandle create_mipmapped_array_handle(const ContextHandle& h_context, - const CUDA_ARRAY3D_DESCRIPTOR& desc, - unsigned int num_levels); +MipmappedArrayHandle +create_mipmapped_array_handle( + const ContextHandle &h_context, + const CUDA_ARRAY3D_DESCRIPTOR &desc, + unsigned int num_levels +); // Return the context dependency associated with a mipmapped array, if known. -ContextHandle get_mipmapped_array_context(const MipmappedArrayHandle& h) noexcept; +ContextHandle +get_mipmapped_array_context(const MipmappedArrayHandle &h) noexcept; // Create an owning texture object via cuTexObjectCreate, embedding the backing // resource handle (array / mipmapped array / linear-or-pitch2d device pointer) // so the backing always outlives the texture. cuTexObjectDestroy runs in the // deleter. Returns empty handle on error (caller must check). -TexObjectHandle create_tex_object_handle_array(const ContextHandle& h_context, - const CUDA_RESOURCE_DESC& res, - const CUDA_TEXTURE_DESC& tex, - const OpaqueArrayHandle& h_backing); -TexObjectHandle create_tex_object_handle_mipmap(const ContextHandle& h_context, - const CUDA_RESOURCE_DESC& res, - const CUDA_TEXTURE_DESC& tex, - const MipmappedArrayHandle& h_backing); -TexObjectHandle create_tex_object_handle_linear(const ContextHandle& h_context, - const CUDA_RESOURCE_DESC& res, - const CUDA_TEXTURE_DESC& tex, - const DevicePtrHandle& h_backing); +TexObjectHandle +create_tex_object_handle_array( + const ContextHandle &h_context, + const CUDA_RESOURCE_DESC &res, + const CUDA_TEXTURE_DESC &tex, + const OpaqueArrayHandle &h_backing +); +TexObjectHandle +create_tex_object_handle_mipmap( + const ContextHandle &h_context, + const CUDA_RESOURCE_DESC &res, + const CUDA_TEXTURE_DESC &tex, + const MipmappedArrayHandle &h_backing +); +TexObjectHandle +create_tex_object_handle_linear( + const ContextHandle &h_context, + const CUDA_RESOURCE_DESC &res, + const CUDA_TEXTURE_DESC &tex, + const DevicePtrHandle &h_backing +); // Create an owning surface object via cuSurfObjectCreate, embedding the backing // array handle so it outlives the surface. cuSurfObjectDestroy runs in the // deleter. Returns empty handle on error (caller must check). -SurfObjectHandle create_surf_object_handle(const ContextHandle& h_context, - const CUDA_RESOURCE_DESC& res, - const OpaqueArrayHandle& h_backing); +SurfObjectHandle +create_surf_object_handle( + const ContextHandle &h_context, + const CUDA_RESOURCE_DESC &res, + const OpaqueArrayHandle &h_backing +); } // namespace cuda_core::rt diff --git a/cuda_core/cuda/core/_cpp/rt/context.cpp b/cuda_core/cuda/core/_cpp/rt/context.cpp index 47678502989..21a94dd517c 100644 --- a/cuda_core/cuda/core/_cpp/rt/context.cpp +++ b/cuda_core/cuda/core/_cpp/rt/context.cpp @@ -22,7 +22,9 @@ namespace { thread_local CUresult context_restore_fault = CUDA_SUCCESS; } // namespace -void set_context_restore_fault_for_testing(CUresult status) noexcept { +void +set_context_restore_fault_for_testing(CUresult status) noexcept +{ context_restore_fault = status; } @@ -33,7 +35,9 @@ namespace detail { // invoke_in_context_or_undo reject empty handles before getting here; only // graph_node_set_params relies on the no-op (pre-13.2 node updates run in the // caller's context). -CUresult enter_context(const ContextHandle& h_context, CUcontext* previous, int* changed) noexcept { +CUresult +enter_context(const ContextHandle &h_context, CUcontext *previous, int *changed) noexcept +{ *previous = nullptr; *changed = 0; clear_last_error_detail(); @@ -53,7 +57,9 @@ CUresult enter_context(const ContextHandle& h_context, CUcontext* previous, int* } // Restore the caller's context. Returns the restoration status. -CUresult restore_context(CUcontext previous) noexcept { +CUresult +restore_context(CUcontext previous) noexcept +{ if (context_restore_fault != CUDA_SUCCESS) { // Test hook: behave as if cuCtxSetCurrent(previous) failed, leaving the // target context current exactly as a real failure would. @@ -68,7 +74,9 @@ CUresult restore_context(CUcontext previous) noexcept { // operation error, if any, is returned; otherwise the restoration status is. // Either way a restoration failure is recorded as the detail of the returned // status, so the eventual CUDAError explains it (see take_last_error_detail()). -CUresult exit_context(CUcontext previous, int changed, CUresult operation_status) noexcept { +CUresult +exit_context(CUcontext previous, int changed, CUresult operation_status) noexcept +{ CUresult restore_status = changed ? restore_context(previous) : CUDA_SUCCESS; if (restore_status == CUDA_SUCCESS) { return operation_status; @@ -79,7 +87,9 @@ CUresult exit_context(CUcontext previous, int changed, CUresult operation_status } // namespace detail // Synchronize the provided context. -CUresult context_synchronize(const ContextHandle& h_context) noexcept { +CUresult +context_synchronize(const ContextHandle &h_context) noexcept +{ GILReleaseGuard gil; return invoke_in_context(h_context, []() noexcept { return DRIVER_FN(cuCtxSynchronize)(); @@ -87,9 +97,9 @@ CUresult context_synchronize(const ContextHandle& h_context) noexcept { } // Query the stream priority range for the provided context. -CUresult context_get_stream_priority_range(const ContextHandle& h_context, - int* least_priority, - int* greatest_priority) noexcept { +CUresult +context_get_stream_priority_range(const ContextHandle &h_context, int *least_priority, int *greatest_priority) noexcept +{ GILReleaseGuard gil; return invoke_in_context(h_context, [&]() noexcept { return DRIVER_CALL(cuCtxGetStreamPriorityRange, least_priority, greatest_priority); @@ -97,7 +107,9 @@ CUresult context_get_stream_priority_range(const ContextHandle& h_context, } // Query the device of the provided context. -CUresult context_get_device(const ContextHandle& h_context, CUdevice* device) noexcept { +CUresult +context_get_device(const ContextHandle &h_context, CUdevice *device) noexcept +{ return invoke_in_context(h_context, [&]() noexcept { return DRIVER_CALL(cuCtxGetDevice, device); }); @@ -117,42 +129,45 @@ struct GreenCtxBox { CUgreenCtx resource; }; -static const ContextBox* get_box(const ContextHandle& h) noexcept { - const CUcontext* p = h.get(); - return reinterpret_cast( - reinterpret_cast(p) - offsetof(ContextBox, resource) - ); +static const ContextBox * +get_box(const ContextHandle &h) noexcept +{ + const CUcontext *p = h.get(); + return reinterpret_cast(reinterpret_cast(p) - offsetof(ContextBox, resource)); } // See REGISTRY_DESIGN.md (Level 1: Driver Handle -> Resource Handle) static HandleRegistry context_registry; // Create a context handle reference, with optional green context as source. -ContextHandle create_context_handle_ref(CUcontext ctx, GreenCtxHandle h_green_ctx) { +ContextHandle +create_context_handle_ref(CUcontext ctx, GreenCtxHandle h_green_ctx) +{ if (!ctx) { return {}; } if (auto h = context_registry.lookup(ctx)) { return h; } - auto box = std::shared_ptr( - new ContextBox{ctx, std::move(h_green_ctx)}, - [](const ContextBox* b) { - context_registry.unregister_handle(b->resource); - delete b; - } - ); + auto box = std::shared_ptr(new ContextBox{ctx, std::move(h_green_ctx)}, [](const ContextBox *b) { + context_registry.unregister_handle(b->resource); + delete b; + }); ContextHandle h(box, &box->resource); context_registry.register_handle(ctx, h); return h; } } // namespace -ContextHandle create_context_handle_ref(CUcontext ctx) { +ContextHandle +create_context_handle_ref(CUcontext ctx) +{ return create_context_handle_ref(ctx, {}); } -ContextHandle create_context_handle_from_green_ctx(const GreenCtxHandle& h_green_ctx) { +ContextHandle +create_context_handle_from_green_ctx(const GreenCtxHandle &h_green_ctx) +{ GILReleaseGuard gil; if (!h_green_ctx) { return {}; @@ -165,15 +180,18 @@ ContextHandle create_context_handle_from_green_ctx(const GreenCtxHandle& h_green return create_context_handle_ref(ctx, h_green_ctx); } -GreenCtxHandle get_context_green_ctx(const ContextHandle& h) noexcept { +GreenCtxHandle +get_context_green_ctx(const ContextHandle &h) noexcept +{ if (!h) { return {}; } return get_box(h)->h_green_ctx; } -GreenCtxHandle create_green_ctx_handle(CUdevResource* resources, unsigned int nbResources, - CUdevice dev, unsigned int flags) { +GreenCtxHandle +create_green_ctx_handle(CUdevResource *resources, unsigned int nbResources, CUdevice dev, unsigned int flags) +{ GILReleaseGuard gil; CUdevResourceDesc desc = nullptr; if (CUDA_SUCCESS != (err = DRIVER_CALL(cuDevResourceGenerateDesc, &desc, resources, nbResources))) { @@ -185,18 +203,17 @@ GreenCtxHandle create_green_ctx_handle(CUdevResource* resources, unsigned int nb return {}; } - auto box = std::shared_ptr( - new GreenCtxBox{green_ctx}, - [](const GreenCtxBox* b) { - GILReleaseGuard gil; - pw_cuGreenCtxDestroy(b->resource); - delete b; - } - ); + auto box = std::shared_ptr(new GreenCtxBox{green_ctx}, [](const GreenCtxBox *b) { + GILReleaseGuard gil; + pw_cuGreenCtxDestroy(b->resource); + delete b; + }); return GreenCtxHandle(box, &box->resource); } -GreenCtxHandle create_green_ctx_handle_ref(CUgreenCtx green_ctx) { +GreenCtxHandle +create_green_ctx_handle_ref(CUgreenCtx green_ctx) +{ if (!green_ctx) { return {}; } @@ -207,7 +224,9 @@ GreenCtxHandle create_green_ctx_handle_ref(CUgreenCtx green_ctx) { // Thread-local cache of primary contexts indexed by device ID static thread_local std::vector primary_context_cache; -ContextHandle get_primary_context(int device_id) { +ContextHandle +get_primary_context(int device_id) +{ // Check thread-local cache if (static_cast(device_id) < primary_context_cache.size()) { if (auto cached = primary_context_cache[device_id]) { @@ -222,20 +241,17 @@ ContextHandle get_primary_context(int device_id) { return {}; } - auto box = std::shared_ptr( - new ContextBox{ctx, {}}, - [device_id](const ContextBox* b) { - context_registry.unregister_handle(b->resource); - // During interpreter shutdown, leave primary-context cleanup to - // process teardown. An unavailable table entry would need Python - // to report itself. - if (Py_IsInitialized() && !py_is_finalizing()) { - GILReleaseGuard gil; - DRIVER_CALL(cuDevicePrimaryCtxRelease, device_id); - } - delete b; + auto box = std::shared_ptr(new ContextBox{ctx, {}}, [device_id](const ContextBox *b) { + context_registry.unregister_handle(b->resource); + // During interpreter shutdown, leave primary-context cleanup to + // process teardown. An unavailable table entry would need Python + // to report itself. + if (Py_IsInitialized() && !py_is_finalizing()) { + GILReleaseGuard gil; + DRIVER_CALL(cuDevicePrimaryCtxRelease, device_id); } - ); + delete b; + }); auto h = ContextHandle(box, &box->resource); context_registry.register_handle(ctx, h); @@ -247,7 +263,9 @@ ContextHandle get_primary_context(int device_id) { return h; } -ContextHandle get_current_context() { +ContextHandle +get_current_context() +{ GILReleaseGuard gil; CUcontext ctx = nullptr; if (CUDA_SUCCESS != (err = DRIVER_CALL(cuCtxGetCurrent, &ctx))) { diff --git a/cuda_core/cuda/core/_cpp/rt/context_scope.hpp b/cuda_core/cuda/core/_cpp/rt/context_scope.hpp index 040a8eb9fa8..a60f15229b5 100644 --- a/cuda_core/cuda/core/_cpp/rt/context_scope.hpp +++ b/cuda_core/cuda/core/_cpp/rt/context_scope.hpp @@ -15,11 +15,14 @@ namespace cuda_core::rt::detail { // Implemented in context.cpp -CUresult enter_context(const ContextHandle& h_context, CUcontext* previous, int* changed) noexcept; +CUresult +enter_context(const ContextHandle &h_context, CUcontext *previous, int *changed) noexcept; // Implemented in context.cpp -CUresult restore_context(CUcontext previous) noexcept; +CUresult +restore_context(CUcontext previous) noexcept; // Implemented in context.cpp -CUresult exit_context(CUcontext previous, int changed, CUresult operation_status) noexcept; +CUresult +exit_context(CUcontext previous, int changed, CUresult operation_status) noexcept; // Require a callable to be invocable without throwing. #define ASSERT_NOTHROW_INVOCABLE(...) \ @@ -27,8 +30,10 @@ CUresult exit_context(CUcontext previous, int changed, CUresult operation_status // Run an operation with the requested context current. template -CUresult invoke_in_context(const ContextHandle& h_context, Fn&& operation, Args&&... args) noexcept { - ASSERT_NOTHROW_INVOCABLE(Fn&&, Args&&...); +CUresult +invoke_in_context(const ContextHandle &h_context, Fn &&operation, Args &&...args) noexcept +{ + ASSERT_NOTHROW_INVOCABLE(Fn &&, Args && ...); if (!h_context) { return CUDA_ERROR_INVALID_CONTEXT; } @@ -46,10 +51,16 @@ CUresult invoke_in_context(const ContextHandle& h_context, Fn&& operation, Args& // after verifying that the target context remains current; otherwise the // resource leaks rather than risking cleanup in the wrong context. template -CUresult invoke_in_context_or_undo(const ContextHandle& h_context, Fn&& operation, - Undo&& undo, bool undo_requires_target_context) noexcept { - ASSERT_NOTHROW_INVOCABLE(Fn&&); - ASSERT_NOTHROW_INVOCABLE(Undo&&); +CUresult +invoke_in_context_or_undo( + const ContextHandle &h_context, + Fn &&operation, + Undo &&undo, + bool undo_requires_target_context +) noexcept +{ + ASSERT_NOTHROW_INVOCABLE(Fn &&); + ASSERT_NOTHROW_INVOCABLE(Undo &&); if (!h_context) { return CUDA_ERROR_INVALID_CONTEXT; } @@ -65,16 +76,18 @@ CUresult invoke_in_context_or_undo(const ContextHandle& h_context, Fn&& operatio bool undo_ok = true; if (undo_requires_target_context) { CUcontext current = nullptr; - undo_ok = DRIVER_CALL(cuCtxGetCurrent, ¤t) == CUDA_SUCCESS - && current == as_cu(h_context); + undo_ok = DRIVER_CALL(cuCtxGetCurrent, ¤t) == CUDA_SUCCESS && current == as_cu(h_context); } if (undo_ok) { std::invoke(std::forward(undo)); - } else { + } + else { report_cuda_error( - "cuCtxSetCurrent (restoring the caller's context)", composite, + "cuCtxSetCurrent (restoring the caller's context)", + composite, "failed; cleanup of the new resource skipped because its context " - "is no longer current (resource leaked)"); + "is no longer current (resource leaked)" + ); } } return composite; @@ -90,18 +103,25 @@ CUresult invoke_in_context_or_undo(const ContextHandle& h_context, Fn&& operatio // outcome; pass a hook that unlocks it. Returns the operation or activation // status; restoration never changes it. template -CUresult cleanup_in_context(const ContextHandle& h_context, const char* name, - unsigned long long handle, Fn&& operation, - AfterCleanup&& after_cleanup) noexcept { - ASSERT_NOTHROW_INVOCABLE(Fn&&); - ASSERT_NOTHROW_INVOCABLE(AfterCleanup&&); +CUresult +cleanup_in_context( + const ContextHandle &h_context, + const char *name, + unsigned long long handle, + Fn &&operation, + AfterCleanup &&after_cleanup +) noexcept +{ + ASSERT_NOTHROW_INVOCABLE(Fn &&); + ASSERT_NOTHROW_INVOCABLE(AfterCleanup &&); CUcontext previous = nullptr; int changed = 0; - const char* detail = nullptr; + const char *detail = nullptr; CUresult status = enter_context(h_context, &previous, &changed); if (status != CUDA_SUCCESS) { detail = "skipped (context activation failed; resource leaked)"; - } else { + } + else { status = std::invoke(std::forward(operation)); } CUresult restore = exit_context(previous, changed, CUDA_SUCCESS); @@ -126,10 +146,11 @@ CUresult cleanup_in_context(const ContextHandle& h_context, const char* name, // Same, with nothing to do after the cleanup. template -CUresult cleanup_in_context(const ContextHandle& h_context, const char* name, - unsigned long long handle, Fn&& operation) noexcept { - return cleanup_in_context(h_context, name, handle, std::forward(operation), - []() noexcept {}); +CUresult +cleanup_in_context(const ContextHandle &h_context, const char *name, unsigned long long handle, Fn &&operation) noexcept +{ + return cleanup_in_context(h_context, name, handle, std::forward(operation), []() noexcept { + }); } #undef ASSERT_NOTHROW_INVOCABLE diff --git a/cuda_core/cuda/core/_cpp/rt/driver_api.cpp b/cuda_core/cuda/core/_cpp/rt/driver_api.cpp index 1d222573617..752076af6d8 100644 --- a/cuda_core/cuda/core/_cpp/rt/driver_api.cpp +++ b/cuda_core/cuda/core/_cpp/rt/driver_api.cpp @@ -25,24 +25,26 @@ namespace { // "__" + the macro-expanded symbol that cuda.h maps the public name to, which // is how cuda-bindings keys its table. #name is the public name, unexpanded. #define CUDA_CORE_DRIVER_FN_ENTRY(name, introduced) \ - {"__" CUDA_CORE_XSTR(name), #name, reinterpret_cast(&p_##name), introduced}, + {"__" CUDA_CORE_XSTR(name), #name, reinterpret_cast(&p_##name), introduced}, const FnEntry driver_entries[] = {CUDA_CORE_DRIVER_FUNCTIONS(CUDA_CORE_DRIVER_FN_ENTRY)}; #undef CUDA_CORE_DRIVER_FN_ENTRY const FnEntry nvrtc_entries[] = { - {"__nvrtcDestroyProgram", "nvrtcDestroyProgram", reinterpret_cast(&p_nvrtcDestroyProgram), 0}, + {"__nvrtcDestroyProgram", "nvrtcDestroyProgram", reinterpret_cast(&p_nvrtcDestroyProgram), 0}, }; const FnEntry nvvm_entries[] = { - {"__nvvmDestroyProgram", "nvvmDestroyProgram", reinterpret_cast(&p_nvvmDestroyProgram), 0}, + {"__nvvmDestroyProgram", "nvvmDestroyProgram", reinterpret_cast(&p_nvvmDestroyProgram), 0}, }; const FnEntry nvjitlink_entries[] = { - {"__nvJitLinkDestroy", "nvJitLinkDestroy", reinterpret_cast(&p_nvJitLinkDestroy), 0}, + {"__nvJitLinkDestroy", "nvJitLinkDestroy", reinterpret_cast(&p_nvJitLinkDestroy), 0}, }; } // namespace -const FnEntry* fn_table_entries(FnTable table, std::size_t* count) noexcept { +const FnEntry * +fn_table_entries(FnTable table, std::size_t *count) noexcept +{ switch (table) { case FnTable::driver: *count = sizeof(driver_entries) / sizeof(driver_entries[0]); diff --git a/cuda_core/cuda/core/_cpp/rt/driver_api.hpp b/cuda_core/cuda/core/_cpp/rt/driver_api.hpp index 9237e0ae723..713e57a0fe3 100644 --- a/cuda_core/cuda/core/_cpp/rt/driver_api.hpp +++ b/cuda_core/cuda/core/_cpp/rt/driver_api.hpp @@ -52,93 +52,93 @@ namespace cuda_core::rt { // cuStreamDestroy -> cuStreamDestroy_v2, and the table key follows that // mapping. The header cuda.core compiles against must therefore be the one // cuda-bindings was generated from. build_hooks.py enforces this. -#define CUDA_CORE_DRIVER_FUNCTIONS(X) \ - /* Error formatting */ \ - X(cuGetErrorName, 6000) \ - X(cuGetErrorString, 6000) \ - /* Context */ \ - X(cuDevicePrimaryCtxRetain, 7000) \ - X(cuDevicePrimaryCtxRelease, 11000) \ - X(cuCtxGetCurrent, 4000) \ - X(cuCtxSetCurrent, 4000) \ - X(cuCtxSynchronize, 2000) \ - X(cuCtxGetStreamPriorityRange, 5050) \ - X(cuCtxGetDevice, 2000) \ - X(cuGraphNodeSetParams, 12020) \ - X(cuGreenCtxCreate, 12040) \ - X(cuGreenCtxDestroy, 12040) \ - X(cuCtxFromGreenCtx, 12040) \ - X(cuDevResourceGenerateDesc, 12040) \ - X(cuGreenCtxStreamCreate, 12050) \ - /* Stream */ \ - X(cuStreamCreateWithPriority, 5050) \ - X(cuStreamDestroy, 4000) \ - X(cuStreamGetCtx, 9020) \ - /* Event */ \ - X(cuEventCreate, 2000) \ - X(cuEventDestroy, 4000) \ - X(cuIpcOpenEventHandle, 4010) \ - /* Device */ \ - X(cuDeviceGetCount, 2000) \ - /* Memory pool */ \ - X(cuMemPoolSetAccess, 11020) \ - X(cuMemPoolDestroy, 11020) \ - X(cuMemPoolCreate, 11020) \ - X(cuDeviceGetMemPool, 11020) \ - X(cuMemPoolImportFromShareableHandle, 11020) \ - /* Memory allocation */ \ - X(cuMemAllocFromPoolAsync, 11020) \ - X(cuMemAllocAsync, 11020) \ - X(cuMemAlloc, 3020) \ - X(cuMemAllocHost, 3020) \ - X(cuMemFreeAsync, 11020) \ - X(cuMemFree, 3020) \ - X(cuMemFreeHost, 2000) \ - /* IPC */ \ - X(cuMemPoolImportPointer, 11020) \ - /* Library */ \ - X(cuLibraryLoadFromFile, 12000) \ - X(cuLibraryLoadData, 12000) \ - X(cuLibraryUnload, 12000) \ - X(cuLibraryGetKernel, 12000) \ - /* Graph */ \ - X(cuGraphDestroy, 10000) \ - X(cuGraphInstantiateWithParams, 12000) \ - X(cuGraphExecUpdate, 12000) \ - X(cuGraphExecDestroy, 10000) \ - X(cuUserObjectCreate, 11030) \ - X(cuUserObjectRelease, 11030) \ - X(cuGraphRetainUserObject, 11030) \ - X(cuGraphReleaseUserObject, 11030) \ - X(cuGraphNodeFindInClone, 10000) \ - X(cuGraphChildGraphNodeGetGraph, 10000) \ - /* Linker */ \ - X(cuLinkDestroy, 5050) \ +#define CUDA_CORE_DRIVER_FUNCTIONS(X) \ + /* Error formatting */ \ + X(cuGetErrorName, 6000) \ + X(cuGetErrorString, 6000) \ + /* Context */ \ + X(cuDevicePrimaryCtxRetain, 7000) \ + X(cuDevicePrimaryCtxRelease, 11000) \ + X(cuCtxGetCurrent, 4000) \ + X(cuCtxSetCurrent, 4000) \ + X(cuCtxSynchronize, 2000) \ + X(cuCtxGetStreamPriorityRange, 5050) \ + X(cuCtxGetDevice, 2000) \ + X(cuGraphNodeSetParams, 12020) \ + X(cuGreenCtxCreate, 12040) \ + X(cuGreenCtxDestroy, 12040) \ + X(cuCtxFromGreenCtx, 12040) \ + X(cuDevResourceGenerateDesc, 12040) \ + X(cuGreenCtxStreamCreate, 12050) \ + /* Stream */ \ + X(cuStreamCreateWithPriority, 5050) \ + X(cuStreamDestroy, 4000) \ + X(cuStreamGetCtx, 9020) \ + /* Event */ \ + X(cuEventCreate, 2000) \ + X(cuEventDestroy, 4000) \ + X(cuIpcOpenEventHandle, 4010) \ + /* Device */ \ + X(cuDeviceGetCount, 2000) \ + /* Memory pool */ \ + X(cuMemPoolSetAccess, 11020) \ + X(cuMemPoolDestroy, 11020) \ + X(cuMemPoolCreate, 11020) \ + X(cuDeviceGetMemPool, 11020) \ + X(cuMemPoolImportFromShareableHandle, 11020) \ + /* Memory allocation */ \ + X(cuMemAllocFromPoolAsync, 11020) \ + X(cuMemAllocAsync, 11020) \ + X(cuMemAlloc, 3020) \ + X(cuMemAllocHost, 3020) \ + X(cuMemFreeAsync, 11020) \ + X(cuMemFree, 3020) \ + X(cuMemFreeHost, 2000) \ + /* IPC */ \ + X(cuMemPoolImportPointer, 11020) \ + /* Library */ \ + X(cuLibraryLoadFromFile, 12000) \ + X(cuLibraryLoadData, 12000) \ + X(cuLibraryUnload, 12000) \ + X(cuLibraryGetKernel, 12000) \ + /* Graph */ \ + X(cuGraphDestroy, 10000) \ + X(cuGraphInstantiateWithParams, 12000) \ + X(cuGraphExecUpdate, 12000) \ + X(cuGraphExecDestroy, 10000) \ + X(cuUserObjectCreate, 11030) \ + X(cuUserObjectRelease, 11030) \ + X(cuGraphRetainUserObject, 11030) \ + X(cuGraphReleaseUserObject, 11030) \ + X(cuGraphNodeFindInClone, 10000) \ + X(cuGraphChildGraphNodeGetGraph, 10000) \ + /* Linker */ \ + X(cuLinkDestroy, 5050) \ /* Graphics interop. cuda-bindings requests 7000 (PTDS) or 3000 (legacy) */ \ - X(cuGraphicsUnmapResources, 7000) \ - X(cuGraphicsUnregisterResource, 3000) \ - /* Texture / surface / array (PR #467) */ \ - X(cuArray3DCreate, 3020) \ - X(cuArrayDestroy, 2000) \ - X(cuMipmappedArrayCreate, 5000) \ - X(cuMipmappedArrayDestroy, 5000) \ - X(cuMipmappedArrayGetLevel, 5000) \ - X(cuTexObjectCreate, 5000) \ - X(cuTexObjectDestroy, 5000) \ - X(cuSurfObjectCreate, 5000) \ - X(cuSurfObjectDestroy, 5000) \ - /* Virtual memory management (VMM_DESIGN.md) */ \ - X(cuMemCreate, 10020) \ - X(cuMemRelease, 10020) \ - X(cuMemAddressReserve, 10020) \ - X(cuMemAddressFree, 10020) \ - X(cuMemMap, 10020) \ - X(cuMemUnmap, 10020) \ - X(cuMemSetAccess, 10020) \ - X(cuMemImportFromShareableHandle, 10020) \ - /* cuda-bindings requests 7000 (PTDS) or 2000 (legacy) */ \ - X(cuStreamSynchronize, 7000) \ - X(cuStreamIsCapturing, 10000) \ + X(cuGraphicsUnmapResources, 7000) \ + X(cuGraphicsUnregisterResource, 3000) \ + /* Texture / surface / array (PR #467) */ \ + X(cuArray3DCreate, 3020) \ + X(cuArrayDestroy, 2000) \ + X(cuMipmappedArrayCreate, 5000) \ + X(cuMipmappedArrayDestroy, 5000) \ + X(cuMipmappedArrayGetLevel, 5000) \ + X(cuTexObjectCreate, 5000) \ + X(cuTexObjectDestroy, 5000) \ + X(cuSurfObjectCreate, 5000) \ + X(cuSurfObjectDestroy, 5000) \ + /* Virtual memory management (VMM_DESIGN.md) */ \ + X(cuMemCreate, 10020) \ + X(cuMemRelease, 10020) \ + X(cuMemAddressReserve, 10020) \ + X(cuMemAddressFree, 10020) \ + X(cuMemMap, 10020) \ + X(cuMemUnmap, 10020) \ + X(cuMemSetAccess, 10020) \ + X(cuMemImportFromShareableHandle, 10020) \ + /* cuda-bindings requests 7000 (PTDS) or 2000 (legacy) */ \ + X(cuStreamSynchronize, 7000) \ + X(cuStreamIsCapturing, 10000) \ X(cuThreadExchangeStreamCaptureMode, 10010) #define CUDA_CORE_DECLARE_DRIVER_FN(name, introduced) extern decltype(&name) p_##name; @@ -152,9 +152,9 @@ CUDA_CORE_DRIVER_FUNCTIONS(CUDA_CORE_DECLARE_DRIVER_FN) // `nvvmResult nvvmDestroyProgram(nvvmProgram*)` and // `nvJitLinkResult nvJitLinkDestroy(nvJitLinkHandle*)` as int-sized enums. extern decltype(&nvrtcDestroyProgram) p_nvrtcDestroyProgram; -using NvvmDestroyProgramFn = int (*)(nvvmProgram*); +using NvvmDestroyProgramFn = int (*)(nvvmProgram *); extern NvvmDestroyProgramFn p_nvvmDestroyProgram; -using NvJitLinkDestroyFn = int (*)(nvJitLink_t*); +using NvJitLinkDestroyFn = int (*)(nvJitLink_t *); extern NvJitLinkDestroyFn p_nvJitLinkDestroy; // ============================================================================ @@ -164,17 +164,19 @@ extern NvJitLinkDestroyFn p_nvJitLinkDestroy; enum class FnTable { driver, nvrtc, nvvm, nvjitlink }; struct FnEntry { - const char* key; // cuda-bindings' name for the slot, e.g. "__cuStreamDestroy_v2" - const char* name; // public name, for messages, e.g. "cuStreamDestroy" - void** slot; // the p_ pointer, as storage + const char *key; // cuda-bindings' name for the slot, e.g. "__cuStreamDestroy_v2" + const char *name; // public name, for messages, e.g. "cuStreamDestroy" + void **slot; // the p_ pointer, as storage int introduced; // CUDA version cuda-bindings requests the symbol at (0 = n/a) }; // The entries of a table. Implemented in driver_api.cpp. -const FnEntry* fn_table_entries(FnTable table, std::size_t* count) noexcept; +const FnEntry * +fn_table_entries(FnTable table, std::size_t *count) noexcept; // Whether a table is filled. An acquire load: a true result orders the slots. -bool fn_table_ready(FnTable table) noexcept; +bool +fn_table_ready(FnTable table) noexcept; // Fill a table from cuda-bindings if it is not ready. The fill acquires the // GIL, imports cuda.bindings._internal., calls @@ -189,18 +191,21 @@ bool fn_table_ready(FnTable table) noexcept; // A failed fill is latched: later calls return false at once and do not // retry. The fill preserves a pending Python exception and never leaves one // set. Implemented in py_driver_fns.cpp. -bool ensure_fn_table(FnTable table) noexcept; +bool +ensure_fn_table(FnTable table) noexcept; // Copy the reason the fill of `table` failed into `buffer`. Return false if // it did not fail. Implemented in py_driver_fns.cpp. -bool fn_table_error(FnTable table, char* buffer, std::size_t size) noexcept; +bool +fn_table_error(FnTable table, char *buffer, std::size_t size) noexcept; // Record a call to `name` while it is unavailable. After a failed fill, this // function attaches the fill's reason to the error the caller raises. The fill // reported that reason when it failed. A null entry in a filled table is a gate // bug, which this function reports once per table. Implemented in // py_driver_fns.cpp. -void report_unavailable_fn(FnTable table, const char* name) noexcept; +void +report_unavailable_fn(FnTable table, const char *name) noexcept; namespace detail { @@ -228,13 +233,19 @@ template struct Unavailable; template struct Unavailable { - static R call(A...) noexcept { return UnavailableStatus::value; } + static R + call(A...) noexcept + { + return UnavailableStatus::value; + } }; // Fills the table on first use and returns the resolved pointer, or the // trampoline when the function is unavailable after the fill. Never throws. template -inline F fn_or_unavailable(F& slot, FnTable table, const char* name) noexcept { +inline F +fn_or_unavailable(F &slot, FnTable table, const char *name) noexcept +{ if (!fn_table_ready(table)) { ensure_fn_table(table); } @@ -256,11 +267,18 @@ inline F fn_or_unavailable(F& slot, FnTable table, const char* name) noexcept { // cuMemFree -> cuMemFree_v2. #define DRIVER_FN(name) \ (::cuda_core::rt::detail::fn_or_unavailable(::cuda_core::rt::p_##name, ::cuda_core::rt::FnTable::driver, #name)) -#define DRIVER_CALL(name, ...) \ - (::cuda_core::rt::detail::fn_or_unavailable(::cuda_core::rt::p_##name, ::cuda_core::rt::FnTable::driver, #name)(__VA_ARGS__)) -#define NVRTC_CALL(name, ...) \ - (::cuda_core::rt::detail::fn_or_unavailable(::cuda_core::rt::p_##name, ::cuda_core::rt::FnTable::nvrtc, #name)(__VA_ARGS__)) -#define NVVM_CALL(name, ...) \ - (::cuda_core::rt::detail::fn_or_unavailable(::cuda_core::rt::p_##name, ::cuda_core::rt::FnTable::nvvm, #name)(__VA_ARGS__)) +#define DRIVER_CALL(name, ...) \ + (::cuda_core::rt::detail::fn_or_unavailable(::cuda_core::rt::p_##name, ::cuda_core::rt::FnTable::driver, #name)( \ + __VA_ARGS__ \ + )) +#define NVRTC_CALL(name, ...) \ + (::cuda_core::rt::detail::fn_or_unavailable(::cuda_core::rt::p_##name, ::cuda_core::rt::FnTable::nvrtc, #name)( \ + __VA_ARGS__ \ + )) +#define NVVM_CALL(name, ...) \ + (::cuda_core::rt::detail::fn_or_unavailable(::cuda_core::rt::p_##name, ::cuda_core::rt::FnTable::nvvm, #name)( \ + __VA_ARGS__ \ + )) #define NVJITLINK_CALL(name, ...) \ - (::cuda_core::rt::detail::fn_or_unavailable(::cuda_core::rt::p_##name, ::cuda_core::rt::FnTable::nvjitlink, #name)(__VA_ARGS__)) + (::cuda_core::rt::detail:: \ + fn_or_unavailable(::cuda_core::rt::p_##name, ::cuda_core::rt::FnTable::nvjitlink, #name)(__VA_ARGS__)) diff --git a/cuda_core/cuda/core/_cpp/rt/error.cpp b/cuda_core/cuda/core/_cpp/rt/error.cpp index 0761c5eb7dc..72402122a09 100644 --- a/cuda_core/cuda/core/_cpp/rt/error.cpp +++ b/cuda_core/cuda/core/_cpp/rt/error.cpp @@ -34,17 +34,19 @@ thread_local CUresult last_error_detail_status = CUDA_SUCCESS; namespace detail { // Format " : : " for a failed CUDA call. -void format_cuda_error(char* buffer, size_t size, const char* operation, CUresult status, - const char* detail) noexcept { - const char* error_name = nullptr; - const char* error_description = nullptr; +void +format_cuda_error(char *buffer, size_t size, const char *operation, CUresult status, const char *detail) noexcept +{ + const char *error_name = nullptr; + const char *error_description = nullptr; // If the table is unavailable, the trampolines fail and the numeric branch below runs. - bool decoded = DRIVER_CALL(cuGetErrorName, status, &error_name) == CUDA_SUCCESS - && DRIVER_CALL(cuGetErrorString, status, &error_description) == CUDA_SUCCESS; - const char* outcome = detail ? detail : "failed"; + bool decoded = DRIVER_CALL(cuGetErrorName, status, &error_name) == CUDA_SUCCESS && + DRIVER_CALL(cuGetErrorString, status, &error_description) == CUDA_SUCCESS; + const char *outcome = detail ? detail : "failed"; if (decoded) { std::snprintf(buffer, size, "%s %s: %s: %s", operation, outcome, error_name, error_description); - } else { + } + else { std::snprintf(buffer, size, "%s %s (CUDA error %d)", operation, outcome, static_cast(status)); } } @@ -52,15 +54,18 @@ void format_cuda_error(char* buffer, size_t size, const char* operation, CUresul // the same call distinct. Python's warning registry collapses repeated warnings // with identical text from one call site, so without the handle only the first // of several leaked resources would be reported. -void format_operation(char* buffer, size_t size, const char* operation, - unsigned long long handle) noexcept { +void +format_operation(char *buffer, size_t size, const char *operation, unsigned long long handle) noexcept +{ std::snprintf(buffer, size, "%s(%#llx)", operation, handle); } } // namespace detail // Report a failed non-CUDA call (NVRTC, NVVM, nvJitLink) from a path that // cannot raise. -void report_status_code(const char* operation, long code) noexcept { +void +report_status_code(const char *operation, long code) noexcept +{ char message[256]; std::snprintf(message, sizeof(message), "%s failed (status %ld)", operation, code); report_message(message); @@ -69,7 +74,9 @@ void report_status_code(const char* operation, long code) noexcept { // Report a failed CUDA call from a path that cannot raise. CUDA_ERROR_DEINITIALIZED // is not reported: it means the driver is shutting down, which makes cleanup // failures expected and uninteresting. -void report_cuda_error(const char* operation, CUresult status, const char* detail) noexcept { +void +report_cuda_error(const char *operation, CUresult status, const char *detail) noexcept +{ if (status == CUDA_SUCCESS || status == CUDA_ERROR_DEINITIALIZED) { return; } @@ -78,7 +85,9 @@ void report_cuda_error(const char* operation, CUresult status, const char* detai report_message(message); } -const char* take_last_error_detail(CUresult status) noexcept { +const char * +take_last_error_detail(CUresult status) noexcept +{ if (!last_error_detail[0] || status != last_error_detail_status) { return nullptr; } @@ -87,12 +96,16 @@ const char* take_last_error_detail(CUresult status) noexcept { return taken_error_detail; } -void clear_last_error_detail() noexcept { +void +clear_last_error_detail() noexcept +{ last_error_detail[0] = 0; last_error_detail_status = CUDA_SUCCESS; } -void note_driver_table_failure(const char* reason) noexcept { +void +note_driver_table_failure(const char *reason) noexcept +{ std::snprintf(last_error_detail, sizeof(last_error_detail), "%s", reason); last_error_detail_status = CUDA_ERROR_NOT_INITIALIZED; } @@ -102,29 +115,38 @@ namespace detail { // CUresult about to be returned and raised: the operation status if the // operation failed too, else the restoration status. For a double failure the // detail also names the restoration error, which the raised error does not. -void note_context_not_restored(CUcontext previous, CUresult operation_status, - CUresult restore_status) noexcept { +void +note_context_not_restored(CUcontext previous, CUresult operation_status, CUresult restore_status) noexcept +{ CUcontext current = nullptr; if (DRIVER_CALL(cuCtxGetCurrent, ¤t) != CUDA_SUCCESS) { current = nullptr; } char cause[128] = {0}; if (operation_status != CUDA_SUCCESS) { - const char* error_name = nullptr; + const char *error_name = nullptr; if (DRIVER_CALL(cuGetErrorName, restore_status, &error_name) == CUDA_SUCCESS) { std::snprintf(cause, sizeof(cause), " after this failure (cuCtxSetCurrent: %s)", error_name); - } else { - std::snprintf(cause, sizeof(cause), " after this failure (cuCtxSetCurrent: CUDA error %d)", - static_cast(restore_status)); + } + else { + std::snprintf( + cause, + sizeof(cause), + " after this failure (cuCtxSetCurrent: CUDA error %d)", + static_cast(restore_status) + ); } } - std::snprintf(last_error_detail, sizeof(last_error_detail), - "the calling thread's CUDA context (%#llx) could not be restored%s; " - "context %#llx is now current. Call Device.set_current() before issuing " - "further CUDA work on this thread", - static_cast(reinterpret_cast(previous)), - cause, - static_cast(reinterpret_cast(current))); + std::snprintf( + last_error_detail, + sizeof(last_error_detail), + "the calling thread's CUDA context (%#llx) could not be restored%s; " + "context %#llx is now current. Call Device.set_current() before issuing " + "further CUDA work on this thread", + static_cast(reinterpret_cast(previous)), + cause, + static_cast(reinterpret_cast(current)) + ); last_error_detail_status = operation_status != CUDA_SUCCESS ? operation_status : restore_status; } } // namespace detail @@ -137,18 +159,24 @@ void note_context_not_restored(CUcontext previous, CUresult operation_status, thread_local CUresult err = CUDA_SUCCESS; // Return and clear the calling thread's most recent CUDA error. -CUresult get_last_error() noexcept { +CUresult +get_last_error() noexcept +{ CUresult e = err; err = CUDA_SUCCESS; return e; } // Return the calling thread's most recent CUDA error without clearing it. -CUresult peek_last_error() noexcept { +CUresult +peek_last_error() noexcept +{ return err; } -void clear_last_error() noexcept { +void +clear_last_error() noexcept +{ err = CUDA_SUCCESS; } diff --git a/cuda_core/cuda/core/_cpp/rt/error.hpp b/cuda_core/cuda/core/_cpp/rt/error.hpp index 4c4905411c7..7e3663bcd89 100644 --- a/cuda_core/cuda/core/_cpp/rt/error.hpp +++ b/cuda_core/cuda/core/_cpp/rt/error.hpp @@ -13,13 +13,16 @@ namespace cuda_core::rt { // ============================================================================ // Get and clear the last CUDA error (like cudaGetLastError) -CUresult get_last_error() noexcept; +CUresult +get_last_error() noexcept; // Get the last CUDA error without clearing it (like cudaPeekAtLastError) -CUresult peek_last_error() noexcept; +CUresult +peek_last_error() noexcept; // Explicitly clear the last error -void clear_last_error() noexcept; +void +clear_last_error() noexcept; // Thread-local status of the most recent CUDA API call in this module. Defined // in error.cpp; every family source writes it. @@ -40,21 +43,25 @@ extern thread_local CUresult err; // Report a failed CUDA call. `detail` replaces the default "failed" wording, // e.g. "skipped (context activation failed; resource leaked)". // CUDA_ERROR_DEINITIALIZED (driver shutting down) is never reported. -void report_cuda_error(const char* operation, CUresult status, const char* detail = nullptr) noexcept; +void +report_cuda_error(const char *operation, CUresult status, const char *detail = nullptr) noexcept; // Report a message that is not tied to a CUresult. // Implemented in py_report.cpp -void report_message(const char* message) noexcept; +void +report_message(const char *message) noexcept; // Report a failed NVRTC/NVVM/nvJitLink call by raw status code. -void report_status_code(const char* operation, long code) noexcept; +void +report_status_code(const char *operation, long code) noexcept; // Attach a failed CUDA call to the Python exception currently being handled // (PEP 678 note, Python 3.11+): for rollback failures inside `except` blocks // whose original exception is about to be re-raised. When no exception is // being handled or notes are unavailable, falls back to report_cuda_error(). // Implemented in py_report.cpp -void attach_rollback_failure(const char* operation, CUresult status, const char* detail = nullptr) noexcept; +void +attach_rollback_failure(const char *operation, CUresult status, const char *detail = nullptr) noexcept; // Detail recorded by a context-scoped helper for the CUresult it is about to // return, e.g. that the caller's context could not be restored. The Cython @@ -63,18 +70,22 @@ void attach_rollback_failure(const char* operation, CUresult status, const char* // thread) and clears it when `status` is the CUresult it was recorded for, and // returns nullptr otherwise, so a detail whose status was never raised cannot // attach to an unrelated error. -const char* take_last_error_detail(CUresult status) noexcept; -void clear_last_error_detail() noexcept; +const char * +take_last_error_detail(CUresult status) noexcept; +void +clear_last_error_detail() noexcept; // Record why the driver function table is unavailable. The reason becomes the // detail of the CUDA_ERROR_NOT_INITIALIZED that the trampoline is about to // return (see driver_api.hpp), so the raised CUDAError explains the failed // fill and does not suggest a missed cuInit() call. Implemented in error.cpp. -void note_driver_table_failure(const char* reason) noexcept; +void +note_driver_table_failure(const char *reason) noexcept; // Tests only: make the next context restoration on this thread fail with // `status`, leaving the target context current as a real failure would. // Implemented in context.cpp -void set_context_restore_fault_for_testing(CUresult status) noexcept; +void +set_context_restore_fault_for_testing(CUresult status) noexcept; } // namespace cuda_core::rt diff --git a/cuda_core/cuda/core/_cpp/rt/event.cpp b/cuda_core/cuda/core/_cpp/rt/event.cpp index 73b6e600a0e..7659ae2d020 100644 --- a/cuda_core/cuda/core/_cpp/rt/event.cpp +++ b/cuda_core/cuda/core/_cpp/rt/event.cpp @@ -31,54 +31,76 @@ struct EventBox { }; } // namespace -static const EventBox* get_box(const EventHandle& h) { - const CUevent* p = h.get(); - return reinterpret_cast( - reinterpret_cast(p) - offsetof(EventBox, resource) - ); +static const EventBox * +get_box(const EventHandle &h) +{ + const CUevent *p = h.get(); + return reinterpret_cast(reinterpret_cast(p) - offsetof(EventBox, resource)); } -bool get_event_timing_enabled(const EventHandle& h) noexcept { +bool +get_event_timing_enabled(const EventHandle &h) noexcept +{ return h ? get_box(h)->timing_enabled : false; } -bool get_event_is_blocking_sync(const EventHandle& h) noexcept { +bool +get_event_is_blocking_sync(const EventHandle &h) noexcept +{ return h ? get_box(h)->is_blocking_sync : false; } -bool get_event_ipc_enabled(const EventHandle& h) noexcept { +bool +get_event_ipc_enabled(const EventHandle &h) noexcept +{ return h ? get_box(h)->ipc_enabled : false; } -int get_event_device_id(const EventHandle& h) noexcept { +int +get_event_device_id(const EventHandle &h) noexcept +{ return h ? get_box(h)->device_id : -1; } // Return the context retained by an event handle. -ContextHandle get_event_context(const EventHandle& h) noexcept { +ContextHandle +get_event_context(const EventHandle &h) noexcept +{ return h ? get_box(h)->h_context : ContextHandle{}; } // See REGISTRY_DESIGN.md (Level 1: Driver Handle -> Resource Handle) static HandleRegistry event_registry; -EventHandle create_event_handle(const ContextHandle& h_ctx, unsigned int flags, - bool timing_enabled, bool is_blocking_sync, - bool ipc_enabled, int device_id) { +EventHandle +create_event_handle( + const ContextHandle &h_ctx, + unsigned int flags, + bool timing_enabled, + bool is_blocking_sync, + bool ipc_enabled, + int device_id +) +{ GILReleaseGuard gil; CUevent event = nullptr; err = invoke_in_context_or_undo( h_ctx, - [&]() noexcept { return DRIVER_CALL(cuEventCreate, &event, flags); }, - [&]() noexcept { pw_cuEventDestroy(event); }, - /*undo_requires_target_context=*/false); + [&]() noexcept { + return DRIVER_CALL(cuEventCreate, &event, flags); + }, + [&]() noexcept { + pw_cuEventDestroy(event); + }, + /*undo_requires_target_context=*/false + ); if (err != CUDA_SUCCESS) { return {}; } auto box = std::shared_ptr( new EventBox{event, timing_enabled, is_blocking_sync, ipc_enabled, device_id, h_ctx}, - [](const EventBox* b) { + [](const EventBox *b) { event_registry.unregister_handle(b->resource); GILReleaseGuard gil; pw_cuEventDestroy(b->resource); @@ -90,7 +112,9 @@ EventHandle create_event_handle(const ContextHandle& h_ctx, unsigned int flags, return h; } -EventHandle create_event_handle_for_stream(CUstream stream, unsigned int flags) { +EventHandle +create_event_handle_for_stream(CUstream stream, unsigned int flags) +{ // Resolve the stream's owning context (for default-stream tokens this is // the current context, per cuStreamGetCtx) and create the event there, so // it can be recorded on `stream` no matter which context is current. @@ -109,7 +133,9 @@ EventHandle create_event_handle_for_stream(CUstream stream, unsigned int flags) return create_event_handle(create_context_handle_ref(ctx), flags, false, false, false, -1); } -EventHandle create_event_handle_ref(CUevent event) { +EventHandle +create_event_handle_ref(CUevent event) +{ if (auto h = event_registry.lookup(event)) { return h; } @@ -117,8 +143,9 @@ EventHandle create_event_handle_ref(CUevent event) { return EventHandle(box, &box->resource); } -EventHandle create_event_handle_ipc(const CUipcEventHandle& ipc_handle, - bool is_blocking_sync) { +EventHandle +create_event_handle_ipc(const CUipcEventHandle &ipc_handle, bool is_blocking_sync) +{ GILReleaseGuard gil; CUevent event; if (CUDA_SUCCESS != (err = DRIVER_CALL(cuIpcOpenEventHandle, &event, ipc_handle))) { @@ -127,7 +154,7 @@ EventHandle create_event_handle_ipc(const CUipcEventHandle& ipc_handle, auto box = std::shared_ptr( new EventBox{event, false, is_blocking_sync, true, -1, {}}, - [](const EventBox* b) { + [](const EventBox *b) { event_registry.unregister_handle(b->resource); GILReleaseGuard gil; pw_cuEventDestroy(b->resource); diff --git a/cuda_core/cuda/core/_cpp/rt/graph.cpp b/cuda_core/cuda/core/_cpp/rt/graph.cpp index 47ffc80b941..5fb3b892522 100644 --- a/cuda_core/cuda/core/_cpp/rt/graph.cpp +++ b/cuda_core/cuda/core/_cpp/rt/graph.cpp @@ -28,9 +28,14 @@ using namespace detail; // caller can publish the metadata that depends on the successful update before // raising it; if the update itself failed, its status is returned with the // restoration failure recorded as its detail and *restore_status is CUDA_SUCCESS. -CUresult graph_node_set_params(CUgraphNode node, CUgraphNodeParams* params, - const ContextHandle& h_context, - CUresult* restore_status) noexcept { +CUresult +graph_node_set_params( + CUgraphNode node, + CUgraphNodeParams *params, + const ContextHandle &h_context, + CUresult *restore_status +) noexcept +{ *restore_status = CUDA_SUCCESS; CUcontext previous = nullptr; int changed = 0; @@ -63,7 +68,7 @@ CUresult graph_node_set_params(CUgraphNode node, CUgraphNodeParams* params, namespace { struct NodeAttachment; -using GraphAttachmentMap = std::map; +using GraphAttachmentMap = std::map; struct GraphHierarchy; @@ -75,30 +80,32 @@ struct GraphBoxBase { // Canonical state for one CUgraph. Its GraphHandle aliases resource, whose // address remains stable for the lifetime of the hierarchy. struct GraphBox : GraphBoxBase { - GraphHierarchy* hierarchy = nullptr; // Non-owning back-reference. - GraphBox* parent = nullptr; // Null for the root graph. + GraphHierarchy *hierarchy = nullptr; // Non-owning back-reference. + GraphBox *parent = nullptr; // Null for the root graph. CUgraphNode owner_node = nullptr; // Node in parent that owns this graph. GraphAttachmentMap attachments; // Non-owning attachment index. HandleRegistry node_handles; GraphBox( - CUgraph resource_, - GraphHierarchy* hierarchy_, - GraphBox* parent_ = nullptr, - CUgraphNode owner_node_ = nullptr) noexcept - : GraphBoxBase{resource_}, - hierarchy(hierarchy_), - parent(parent_), - owner_node(owner_node_) {} + CUgraph resource_, + GraphHierarchy *hierarchy_, + GraphBox *parent_ = nullptr, + CUgraphNode owner_node_ = nullptr + ) noexcept + : GraphBoxBase{resource_}, hierarchy(hierarchy_), parent(parent_), owner_node(owner_node_) + { + } }; // Shared owner of stable GraphBox storage. Every GraphHandle aliases the same // control block, so any graph handle keeps the entire hierarchy alive. struct GraphHierarchy { - std::list graphs; // Parent boxes precede their descendants. + std::list graphs; // Parent boxes precede their descendants. std::list graveyard; // Retired child graph tombstones. - GraphBox* root() noexcept { + GraphBox * + root() noexcept + { return graphs.empty() ? nullptr : &graphs.front(); } }; @@ -114,35 +121,40 @@ struct NodeAttachment : DeferredCleanupItem { CUuserObject object = nullptr; std::array owners; - NodeAttachment(OpaqueHandle owner0, OpaqueHandle owner1) - : owners{std::move(owner0), std::move(owner1)} {} + NodeAttachment(OpaqueHandle owner0, OpaqueHandle owner1) : owners{std::move(owner0), std::move(owner1)} {} }; // shared_ptr deleters for the payloads that need one. Typed handles convert to // OpaqueHandle by assignment and reuse their own control block, so they need no // deleter here. The Python deleter follows the owner-release pattern used by // the stream/deviceptr handles above. -void py_deleter(const void* p) noexcept { +void +py_deleter(const void *p) noexcept +{ GILAcquireGuard gil; if (gil.acquired()) { - Py_DECREF(const_cast(static_cast(p))); + Py_DECREF(const_cast(static_cast(p))); } } -void free_deleter(const void* p) noexcept { - std::free(const_cast(p)); +void +free_deleter(const void *p) noexcept +{ + std::free(const_cast(p)); } -GraphBox* get_box(const GraphHandle& h) noexcept { - auto* value = reinterpret_cast(h.get()); - return const_cast( - static_cast(value)); +GraphBox * +get_box(const GraphHandle &h) noexcept +{ + auto *value = reinterpret_cast(h.get()); + return const_cast(static_cast(value)); } // Rekey a staged attachment map from source nodes to their cloned nodes. // The caller must release the GIL before calling this function. -CUresult rekey_attachments( - GraphAttachmentMap& attachments, CUgraph cloned_graph) { +CUresult +rekey_attachments(GraphAttachmentMap &attachments, CUgraph cloned_graph) +{ if (!cloned_graph) { return CUDA_ERROR_INVALID_VALUE; } @@ -150,8 +162,7 @@ CUresult rekey_attachments( while (!attachments.empty()) { auto attachment = attachments.extract(attachments.begin()); CUgraphNode cloned_node = nullptr; - CUresult status = DRIVER_CALL(cuGraphNodeFindInClone, - &cloned_node, attachment.key(), cloned_graph); + CUresult status = DRIVER_CALL(cuGraphNodeFindInClone, &cloned_node, attachment.key(), cloned_graph); if (status != CUDA_SUCCESS) { return status; } @@ -165,57 +176,48 @@ CUresult rekey_attachments( } struct StagedGraphMetadata { - const GraphBox* source; - GraphBox* clone; - GraphAttachmentMap* attachments; + const GraphBox *source; + GraphBox *clone; + GraphAttachmentMap *attachments; }; using StagedGraphMetadataList = std::vector; // Copy a source hierarchy into detached metadata before CUDA mutation. -void stage_graph_metadata( - const GraphBox& source, - GraphBox& clone, - GraphAttachmentMap& attachments, - std::list& subgraphs, - StagedGraphMetadataList& staged) { +void +stage_graph_metadata( + const GraphBox &source, + GraphBox &clone, + GraphAttachmentMap &attachments, + std::list &subgraphs, + StagedGraphMetadataList &staged +) +{ attachments = source.attachments; staged.push_back({&source, &clone, &attachments}); - for (const GraphBox& source_child : source.hierarchy->graphs) { + for (const GraphBox &source_child : source.hierarchy->graphs) { if (source_child.parent != &source || !source_child.resource) { continue; } - GraphBox& cloned_child = subgraphs.emplace_back( - nullptr, - clone.hierarchy, - &clone, - nullptr); - stage_graph_metadata( - source_child, - cloned_child, - cloned_child.attachments, - subgraphs, - staged); + GraphBox &cloned_child = subgraphs.emplace_back(nullptr, clone.hierarchy, &clone, nullptr); + stage_graph_metadata(source_child, cloned_child, cloned_child.attachments, subgraphs, staged); } } // Bind staged metadata to a CUDA-cloned hierarchy. The root clone resource // must be populated before entry. The caller must release the GIL. -CUresult rekey_graph_metadata( - StagedGraphMetadataList& staged) { +CUresult +rekey_graph_metadata(StagedGraphMetadataList &staged) +{ CUresult status; for (size_t i = 0; i < staged.size(); ++i) { - const GraphBox& source = *staged[i].source; - GraphBox& clone = *staged[i].clone; + const GraphBox &source = *staged[i].source; + GraphBox &clone = *staged[i].clone; if (i != 0) { CUgraphNode cloned_owner = nullptr; - status = DRIVER_CALL(cuGraphNodeFindInClone, - &cloned_owner, - source.owner_node, - clone.parent->resource); + status = DRIVER_CALL(cuGraphNodeFindInClone, &cloned_owner, source.owner_node, clone.parent->resource); if (status == CUDA_SUCCESS) { - status = DRIVER_CALL(cuGraphChildGraphNodeGetGraph, - cloned_owner, &clone.resource); + status = DRIVER_CALL(cuGraphChildGraphNodeGetGraph, cloned_owner, &clone.resource); } if (status != CUDA_SUCCESS) { return status; @@ -223,8 +225,7 @@ CUresult rekey_graph_metadata( clone.owner_node = cloned_owner; } - status = rekey_attachments( - *staged[i].attachments, clone.resource); + status = rekey_attachments(*staged[i].attachments, clone.resource); if (status != CUDA_SUCCESS) { return status; } @@ -234,13 +235,17 @@ CUresult rekey_graph_metadata( } // namespace -OpaqueHandle make_opaque_py(PyObject* obj) { +OpaqueHandle +make_opaque_py(PyObject *obj) +{ Py_INCREF(obj); - return OpaqueHandle(static_cast(obj), py_deleter); + return OpaqueHandle(static_cast(obj), py_deleter); } -OpaqueHandle make_opaque_malloc(void* buf) { - return OpaqueHandle(static_cast(buf), free_deleter); +OpaqueHandle +make_opaque_malloc(void *buf) +{ + return OpaqueHandle(static_cast(buf), free_deleter); } // State held by PreparedAttachment between preparation and commit. It keeps the @@ -249,24 +254,23 @@ OpaqueHandle make_opaque_malloc(void* buf) { // the staged user-object retain unless graph_commit_attachment publishes it. struct PreparedAttachmentState { GraphHandle h_graph; - NodeAttachment* replacement = nullptr; + NodeAttachment *replacement = nullptr; GraphAttachmentMap::node_type replacement_entry; - explicit PreparedAttachmentState(GraphHandle h_graph_) - : h_graph(std::move(h_graph_)) {} + explicit PreparedAttachmentState(GraphHandle h_graph_) : h_graph(std::move(h_graph_)) {} }; -void rollback_prepared_attachment( - PreparedAttachmentState* state) noexcept { +void +rollback_prepared_attachment(PreparedAttachmentState *state) noexcept +{ if (!state) { return; } if (state->replacement) { - GraphBox* box = get_box(state->h_graph); + GraphBox *box = get_box(state->h_graph); if (box->resource) { GILReleaseGuard gil; - pw_cuGraphReleaseUserObject( - box->resource, state->replacement->object, 1); + pw_cuGraphReleaseUserObject(box->resource, state->replacement->object, 1); } } delete state; @@ -278,57 +282,55 @@ void rollback_prepared_attachment( struct PreparedChildGraphUpdateState { GraphHandle h_parent; GraphHandle h_source; - GraphBox* old_root = nullptr; + GraphBox *old_root = nullptr; CUgraphNode owner_node = nullptr; std::list replacement; StagedGraphMetadataList staged; std::vector handles; PreparedChildGraphUpdateState( - GraphHandle h_parent_, - GraphHandle h_source_, - GraphBox* old_root_, - CUgraphNode owner_node_) - : h_parent(std::move(h_parent_)), - h_source(std::move(h_source_)), - old_root(old_root_), - owner_node(owner_node_) {} + GraphHandle h_parent_, + GraphHandle h_source_, + GraphBox *old_root_, + CUgraphNode owner_node_ + ) + : h_parent(std::move(h_parent_)), h_source(std::move(h_source_)), old_root(old_root_), owner_node(owner_node_) + { + } }; -GraphHandle create_graph_handle(CUgraph graph) { +GraphHandle +create_graph_handle(CUgraph graph) +{ ensure_fn_table(FnTable::driver); // the deleter calls the driver: resolve the table before it can run if (!graph) { return {}; } - auto hierarchy = std::shared_ptr( - new GraphHierarchy{}, - [](GraphHierarchy* hierarchy) { - for (const GraphBox& box : hierarchy->graphs) { - if (box.resource) { - graph_registry.unregister_handle(box.resource); - } + auto hierarchy = std::shared_ptr(new GraphHierarchy{}, [](GraphHierarchy *hierarchy) { + for (const GraphBox &box : hierarchy->graphs) { + if (box.resource) { + graph_registry.unregister_handle(box.resource); } - GraphBox* root = hierarchy->root(); - if (root && root->resource) { - GILReleaseGuard gil; - pw_cuGraphDestroy(root->resource); - } - retry_deferred_cleanup(); - delete hierarchy; } - ); - GraphBox& root = hierarchy->graphs.emplace_back( - graph, hierarchy.get()); + GraphBox *root = hierarchy->root(); + if (root && root->resource) { + GILReleaseGuard gil; + pw_cuGraphDestroy(root->resource); + } + retry_deferred_cleanup(); + delete hierarchy; + }); + GraphBox &root = hierarchy->graphs.emplace_back(graph, hierarchy.get()); GraphHandle h_graph(hierarchy, &root.resource); graph_registry.register_handle(graph, h_graph); return h_graph; } -GraphHandle create_child_graph_handle( - CUgraph child_graph, const GraphHandle& h_parent, - CUgraphNode owner_node) { +GraphHandle +create_child_graph_handle(CUgraph child_graph, const GraphHandle &h_parent, CUgraphNode owner_node) +{ if (!child_graph || !h_parent || !owner_node) { return {}; } @@ -336,85 +338,80 @@ GraphHandle create_child_graph_handle( return h_graph; } - GraphBox* parent = get_box(h_parent); - GraphHierarchy* hierarchy = parent->hierarchy; - GraphBox& child = hierarchy->graphs.emplace_back( - child_graph, hierarchy, parent, owner_node); + GraphBox *parent = get_box(h_parent); + GraphHierarchy *hierarchy = parent->hierarchy; + GraphBox &child = hierarchy->graphs.emplace_back(child_graph, hierarchy, parent, owner_node); GraphHandle h_child(h_parent, &child.resource); graph_registry.register_handle(child_graph, h_child); return h_child; } -CUresult graph_prepare_child_graph_update( - const GraphHandle& h_parent, - const GraphHandle& h_old_child, - CUgraphNode owner_node, - const GraphHandle& h_source, - PreparedChildGraphUpdate* out_prepared) { - if (!h_parent || !h_old_child || !owner_node || - !h_source || !out_prepared) { +CUresult +graph_prepare_child_graph_update( + const GraphHandle &h_parent, + const GraphHandle &h_old_child, + CUgraphNode owner_node, + const GraphHandle &h_source, + PreparedChildGraphUpdate *out_prepared +) +{ + if (!h_parent || !h_old_child || !owner_node || !h_source || !out_prepared) { return CUDA_ERROR_INVALID_VALUE; } out_prepared->reset(); - GraphBox* parent = get_box(h_parent); - GraphBox* old_root = get_box(h_old_child); - GraphBox* source = get_box(h_source); + GraphBox *parent = get_box(h_parent); + GraphBox *old_root = get_box(h_old_child); + GraphBox *source = get_box(h_source); // A source from the destination hierarchy can include the old embedded // subtree whose raw node keys CUDA destroys during replacement. - if (!parent->resource || !old_root->resource || !source->resource || - old_root->parent != parent || - old_root->owner_node != owner_node || - source->hierarchy == parent->hierarchy) { + if (!parent->resource || !old_root->resource || !source->resource || old_root->parent != parent || + old_root->owner_node != owner_node || source->hierarchy == parent->hierarchy) { return CUDA_ERROR_INVALID_VALUE; } PreparedChildGraphUpdate prepared = - std::make_shared( - h_parent, h_source, old_root, owner_node); + std::make_shared(h_parent, h_source, old_root, owner_node); - GraphBox& replacement_root = - prepared->replacement.emplace_back( - nullptr, parent->hierarchy, parent, owner_node); + GraphBox &replacement_root = prepared->replacement.emplace_back(nullptr, parent->hierarchy, parent, owner_node); stage_graph_metadata( *source, replacement_root, replacement_root.attachments, prepared->replacement, - prepared->staged); + prepared->staged + ); const size_t graph_count = prepared->staged.size(); prepared->handles.reserve(graph_count); - for (const StagedGraphMetadata& graph : prepared->staged) { - prepared->handles.emplace_back( - h_parent, &graph.clone->resource); + for (const StagedGraphMetadata &graph : prepared->staged) { + prepared->handles.emplace_back(h_parent, &graph.clone->resource); } *out_prepared = std::move(prepared); return CUDA_SUCCESS; } -void publish_child_graph_update( - PreparedChildGraphUpdateState& state, - GraphHandle* out_child) { - GraphBox* parent = get_box(state.h_parent); - parent->hierarchy->graphs.splice( - parent->hierarchy->graphs.end(), state.replacement); +void +publish_child_graph_update(PreparedChildGraphUpdateState &state, GraphHandle *out_child) +{ + GraphBox *parent = get_box(state.h_parent); + parent->hierarchy->graphs.splice(parent->hierarchy->graphs.end(), state.replacement); *out_child = state.handles.front(); graph_registry.register_handles(state.handles); } -CUresult graph_commit_child_graph_update( - PreparedChildGraphUpdate& prepared, - GraphHandle* out_child) { +CUresult +graph_commit_child_graph_update(PreparedChildGraphUpdate &prepared, GraphHandle *out_child) +{ if (!prepared || !out_child) { return CUDA_ERROR_INVALID_VALUE; } out_child->reset(); - PreparedChildGraphUpdateState& state = *prepared; - GraphBox* parent = get_box(state.h_parent); + PreparedChildGraphUpdateState &state = *prepared; + GraphBox *parent = get_box(state.h_parent); if (!parent->resource || !state.old_root->resource) { prepared.reset(); return CUDA_ERROR_INVALID_VALUE; @@ -424,8 +421,7 @@ CUresult graph_commit_child_graph_update( CUgraph cloned_root = nullptr; { GILReleaseGuard gil; - status = DRIVER_CALL(cuGraphChildGraphNodeGetGraph, - state.owner_node, &cloned_root); + status = DRIVER_CALL(cuGraphChildGraphNodeGetGraph, state.owner_node, &cloned_root); if (status == CUDA_SUCCESS) { state.staged.front().clone->resource = cloned_root; status = rekey_graph_metadata(state.staged); @@ -434,13 +430,11 @@ CUresult graph_commit_child_graph_update( // CUDA has already destroyed the old embedded graph. No replacement // metadata is visible yet, so this selects only the old generation. - invalidate_child_graph_state( - state.h_parent, state.owner_node); + invalidate_child_graph_state(state.h_parent, state.owner_node); if (status != CUDA_SUCCESS) { prepared.reset(); - throw std::runtime_error( - "failed to update graph metadata after child graph replacement"); + throw std::runtime_error("failed to update graph metadata after child graph replacement"); } publish_child_graph_update(state, out_child); @@ -448,9 +442,9 @@ CUresult graph_commit_child_graph_update( return status; } -CUresult graph_get_attachment( - const GraphHandle& h_graph, CUgraphNode node, - OpaqueHandle* owner0, OpaqueHandle* owner1) { +CUresult +graph_get_attachment(const GraphHandle &h_graph, CUgraphNode node, OpaqueHandle *owner0, OpaqueHandle *owner1) +{ if (!h_graph || !node || (!owner0 && !owner1)) { return CUDA_ERROR_INVALID_VALUE; } @@ -461,7 +455,7 @@ CUresult graph_get_attachment( owner1->reset(); } - GraphBox* box = get_box(h_graph); + GraphBox *box = get_box(h_graph); if (!box->resource) { return CUDA_ERROR_INVALID_VALUE; } @@ -477,11 +471,14 @@ CUresult graph_get_attachment( return CUDA_SUCCESS; } -CUresult graph_prepare_attachment( - const GraphHandle& h_graph, - OpaqueHandle owner0, - OpaqueHandle owner1, - PreparedAttachment* out_prepared) { +CUresult +graph_prepare_attachment( + const GraphHandle &h_graph, + OpaqueHandle owner0, + OpaqueHandle owner1, + PreparedAttachment *out_prepared +) +{ if (!out_prepared) { return CUDA_ERROR_INVALID_VALUE; } @@ -490,39 +487,41 @@ CUresult graph_prepare_attachment( return CUDA_ERROR_INVALID_VALUE; } - GraphBox* box = get_box(h_graph); + GraphBox *box = get_box(h_graph); if (!box->resource) { return CUDA_ERROR_INVALID_VALUE; } PreparedAttachment prepared( new PreparedAttachmentState(h_graph), - PreparedAttachmentDeleter{rollback_prepared_attachment}); + PreparedAttachmentDeleter{rollback_prepared_attachment} + ); if (owner0 || owner1) { ensure_deferred_cleanup_ready(); - prepared->replacement = new NodeAttachment( - std::move(owner0), std::move(owner1)); + prepared->replacement = new NodeAttachment(std::move(owner0), std::move(owner1)); GraphAttachmentMap staged; try { staged.emplace(nullptr, prepared->replacement); - prepared->replacement_entry = - staged.extract(staged.begin()); - } catch (...) { + prepared->replacement_entry = staged.extract(staged.begin()); + } + catch (...) { delete prepared->replacement; prepared->replacement = nullptr; throw; } - auto* cleanup_item = - static_cast( - prepared->replacement); + auto *cleanup_item = static_cast(prepared->replacement); CUuserObject object = nullptr; CUresult status; { GILReleaseGuard gil; - status = DRIVER_CALL(cuUserObjectCreate, - &object, cleanup_item, + status = DRIVER_CALL( + cuUserObjectCreate, + &object, + cleanup_item, reinterpret_cast(enqueue_cleanup), - 1, CU_USER_OBJECT_NO_DESTRUCTOR_SYNC); + 1, + CU_USER_OBJECT_NO_DESTRUCTOR_SYNC + ); if (status != CUDA_SUCCESS) { prepared->replacement_entry.mapped() = nullptr; delete prepared->replacement; @@ -530,8 +529,7 @@ CUresult graph_prepare_attachment( return status; } prepared->replacement->object = object; - status = DRIVER_CALL(cuGraphRetainUserObject, - box->resource, object, 1, CU_GRAPH_USER_OBJECT_MOVE); + status = DRIVER_CALL(cuGraphRetainUserObject, box->resource, object, 1, CU_GRAPH_USER_OBJECT_MOVE); if (status != CUDA_SUCCESS) { prepared->replacement_entry.mapped() = nullptr; prepared->replacement = nullptr; @@ -545,15 +543,15 @@ CUresult graph_prepare_attachment( return CUDA_SUCCESS; } -CUresult graph_commit_attachment( - PreparedAttachment& prepared, - CUgraphNode node) { +CUresult +graph_commit_attachment(PreparedAttachment &prepared, CUgraphNode node) +{ if (!prepared) { return CUDA_ERROR_INVALID_VALUE; } GraphHandle h_graph = prepared->h_graph; - GraphBox* box = get_box(h_graph); + GraphBox *box = get_box(h_graph); if (!box->resource || (!node && !prepared->replacement)) { delete prepared.release(); return CUDA_ERROR_INVALID_VALUE; @@ -566,25 +564,25 @@ CUresult graph_commit_attachment( // Publish the replacement or removal before releasing the previous graph // reference; that release can make the previous payload eligible for // destruction. - NodeAttachment* previous = nullptr; + NodeAttachment *previous = nullptr; auto it = box->attachments.find(node); if (it == box->attachments.end()) { if (prepared->replacement) { prepared->replacement_entry.key() = node; - auto result = box->attachments.insert( - std::move(prepared->replacement_entry)); + auto result = box->attachments.insert(std::move(prepared->replacement_entry)); if (!result.inserted) { - prepared->replacement_entry = - std::move(result.node); + prepared->replacement_entry = std::move(result.node); delete prepared.release(); return CUDA_ERROR_INVALID_VALUE; } } - } else { + } + else { previous = it->second; if (prepared->replacement) { it->second = prepared->replacement; - } else { + } + else { box->attachments.erase(it); } } @@ -594,21 +592,19 @@ CUresult graph_commit_attachment( return CUDA_SUCCESS; } GILReleaseGuard gil; - return DRIVER_CALL(cuGraphReleaseUserObject, - box->resource, previous->object, 1); + return DRIVER_CALL(cuGraphReleaseUserObject, box->resource, previous->object, 1); } -CUresult graph_clone_attachments( - const GraphHandle& h_clone, - const GraphHandle& h_source) { +CUresult +graph_clone_attachments(const GraphHandle &h_clone, const GraphHandle &h_source) +{ if (!h_clone || !h_source) { return CUDA_ERROR_INVALID_VALUE; } - GraphBox* clone = get_box(h_clone); - GraphBox* source = get_box(h_source); - if (!clone->resource || !source->resource || - !clone->attachments.empty()) { + GraphBox *clone = get_box(h_clone); + GraphBox *source = get_box(h_source); + if (!clone->resource || !source->resource || !clone->attachments.empty()) { return CUDA_ERROR_INVALID_VALUE; } @@ -617,12 +613,11 @@ CUresult graph_clone_attachments( GraphAttachmentMap attachments; std::list subgraphs; StagedGraphMetadataList staged; - stage_graph_metadata( - *source, *clone, attachments, subgraphs, staged); + stage_graph_metadata(*source, *clone, attachments, subgraphs, staged); std::vector handles; handles.reserve(subgraphs.size()); - for (GraphBox& graph : subgraphs) { + for (GraphBox &graph : subgraphs) { handles.emplace_back(h_clone, &graph.resource); } @@ -640,8 +635,7 @@ CUresult graph_clone_attachments( return CUDA_SUCCESS; } - clone->hierarchy->graphs.splice( - clone->hierarchy->graphs.end(), subgraphs); + clone->hierarchy->graphs.splice(clone->hierarchy->graphs.end(), subgraphs); graph_registry.register_handles(handles); return CUDA_SUCCESS; } @@ -653,20 +647,20 @@ struct GraphNodeBox { }; } // namespace -static const GraphNodeBox* get_box(const GraphNodeHandle& h) { - const CUgraphNode* p = h.get(); - return reinterpret_cast( - reinterpret_cast(p) - offsetof(GraphNodeBox, resource) - ); +static const GraphNodeBox * +get_box(const GraphNodeHandle &h) +{ + const CUgraphNode *p = h.get(); + return reinterpret_cast(reinterpret_cast(p) - offsetof(GraphNodeBox, resource)); } // Retire one box of a graph that CUDA has destroyed: detach its node handles, // drop its registry entry and attachments, and move it to the graveyard. // `graph` must be an iterator into hierarchy.graphs; the splice invalidates it. -static void retire_graph( - GraphHierarchy& hierarchy, - std::list::iterator graph) noexcept { - for (auto& entry : graph->node_handles.drain()) { +static void +retire_graph(GraphHierarchy &hierarchy, std::list::iterator graph) noexcept +{ + for (auto &entry : graph->node_handles.drain()) { if (GraphNodeHandle h_node = entry.second.lock()) { get_box(h_node)->resource = nullptr; } @@ -676,32 +670,28 @@ static void retire_graph( graph->resource = nullptr; } graph->attachments.clear(); - hierarchy.graveyard.splice( - hierarchy.graveyard.end(), hierarchy.graphs, graph); + hierarchy.graveyard.splice(hierarchy.graveyard.end(), hierarchy.graphs, graph); } // graphs is ordered parent-before-child. Nulling a selected box marks its // later descendants, whose parent pointers remain valid after list splicing. // This permits one allocation-free sweep of the hierarchy. -void invalidate_child_graph_state( - const GraphHandle& h_parent, - CUgraphNode owner_node) noexcept { +void +invalidate_child_graph_state(const GraphHandle &h_parent, CUgraphNode owner_node) noexcept +{ if (!h_parent || !owner_node) { return; } - GraphBox* parent = get_box(h_parent); + GraphBox *parent = get_box(h_parent); if (!parent->resource) { return; } - GraphHierarchy& hierarchy = *parent->hierarchy; - for (auto it = hierarchy.graphs.begin(); - it != hierarchy.graphs.end();) { + GraphHierarchy &hierarchy = *parent->hierarchy; + for (auto it = hierarchy.graphs.begin(); it != hierarchy.graphs.end();) { auto graph = it++; - bool is_owned_root = graph->parent == parent && - graph->owner_node == owner_node; - bool is_descendant = graph->parent && - !graph->parent->resource; + bool is_owned_root = graph->parent == parent && graph->owner_node == owner_node; + bool is_descendant = graph->parent && !graph->parent->resource; if (is_owned_root || is_descendant) { retire_graph(hierarchy, graph); } @@ -711,54 +701,57 @@ void invalidate_child_graph_state( // CUDA destroyed the root graph and, with it, every child. Retire every box // so that no registry entry resolves to the dead graphs and the hierarchy's // deleter finds no root to destroy. -void invalidate_root_graph_state(const GraphHandle& h_root) noexcept { +void +invalidate_root_graph_state(const GraphHandle &h_root) noexcept +{ if (!h_root) { return; } - GraphBox* root = get_box(h_root); + GraphBox *root = get_box(h_root); if (!root->resource || root->parent) { return; } - GraphHierarchy& hierarchy = *root->hierarchy; - for (auto it = hierarchy.graphs.begin(); - it != hierarchy.graphs.end();) { + GraphHierarchy &hierarchy = *root->hierarchy; + for (auto it = hierarchy.graphs.begin(); it != hierarchy.graphs.end();) { retire_graph(hierarchy, it++); } } -GraphNodeHandle create_graph_node_handle(CUgraphNode node, const GraphHandle& h_graph) { +GraphNodeHandle +create_graph_node_handle(CUgraphNode node, const GraphHandle &h_graph) +{ if (!node) { - auto box = std::make_shared( - GraphNodeBox{nullptr, h_graph}); + auto box = std::make_shared(GraphNodeBox{nullptr, h_graph}); return GraphNodeHandle(box, &box->resource); } - GraphBox* graph = get_box(h_graph); - return graph->node_handles.get_or_create( - node, - [node, &h_graph] { - auto box = std::make_shared( - GraphNodeBox{node, h_graph}); - return GraphNodeHandle(box, &box->resource); - }); + GraphBox *graph = get_box(h_graph); + return graph->node_handles.get_or_create(node, [node, &h_graph] { + auto box = std::make_shared(GraphNodeBox{node, h_graph}); + return GraphNodeHandle(box, &box->resource); + }); } -GraphHandle graph_node_get_graph(const GraphNodeHandle& h) noexcept { +GraphHandle +graph_node_get_graph(const GraphNodeHandle &h) noexcept +{ return h ? get_box(h)->h_graph : GraphHandle{}; } -void invalidate_graph_node(const GraphNodeHandle& h) noexcept { +void +invalidate_graph_node(const GraphNodeHandle &h) noexcept +{ if (!h) { return; } - const GraphNodeBox* node_box = get_box(h); + const GraphNodeBox *node_box = get_box(h); CUgraphNode node = node_box->resource; if (!node) { return; } - GraphBox* graph = get_box(node_box->h_graph); + GraphBox *graph = get_box(node_box->h_graph); graph->node_handles.unregister_handle(node); node_box->resource = nullptr; } diff --git a/cuda_core/cuda/core/_cpp/rt/graph_exec.cpp b/cuda_core/cuda/core/_cpp/rt/graph_exec.cpp index 4927274fb79..0dfccd889dc 100644 --- a/cuda_core/cuda/core/_cpp/rt/graph_exec.cpp +++ b/cuda_core/cuda/core/_cpp/rt/graph_exec.cpp @@ -31,9 +31,10 @@ struct ExecAttachments : DeferredCleanupItem { struct GraphExecBox { CUgraphExec resource = nullptr; - ExecAttachments* attachments = nullptr; // Non-owning. + ExecAttachments *attachments = nullptr; // Non-owning. - ~GraphExecBox() noexcept { + ~GraphExecBox() noexcept + { if (resource) { GILReleaseGuard gil; pw_cuGraphExecDestroy(resource); @@ -43,17 +44,20 @@ struct GraphExecBox { } }; -GraphExecBox* get_exec_box(const GraphExecHandle& h) noexcept { - return const_cast( - reinterpret_cast(h.get())); +GraphExecBox * +get_exec_box(const GraphExecHandle &h) noexcept +{ + return const_cast(reinterpret_cast(h.get())); } -GraphExecHandle make_graph_exec_handle( - CUgraphExec graph_exec, ExecAttachments* attachments) { +GraphExecHandle +make_graph_exec_handle(CUgraphExec graph_exec, ExecAttachments *attachments) +{ struct RawGraphExecGuard { CUgraphExec resource; - ~RawGraphExecGuard() noexcept { + ~RawGraphExecGuard() noexcept + { if (resource) { GILReleaseGuard gil; pw_cuGraphExecDestroy(resource); @@ -75,14 +79,16 @@ GraphExecHandle make_graph_exec_handle( // alive, and otherwise this drops its last reference. struct ExecAttachmentStaging { GraphHandle h_source; - ExecAttachments* accumulator = nullptr; + ExecAttachments *accumulator = nullptr; - ~ExecAttachmentStaging() noexcept { - report_cuda_error("cuGraphReleaseUserObject", release(), - "failed while dropping a staged graph attachment"); + ~ExecAttachmentStaging() noexcept + { + report_cuda_error("cuGraphReleaseUserObject", release(), "failed while dropping a staged graph attachment"); } - CUresult release() noexcept { + CUresult + release() noexcept + { if (!h_source || !accumulator) { return CUDA_SUCCESS; } @@ -96,28 +102,30 @@ struct ExecAttachmentStaging { // Create an accumulator and retain it on h_source, so that a following // instantiation or whole-graph update propagates a reference into the exec. -CUresult stage_exec_attachments( - const GraphHandle& h_source, ExecAttachmentStaging* out_staging) { +CUresult +stage_exec_attachments(const GraphHandle &h_source, ExecAttachmentStaging *out_staging) +{ ensure_deferred_cleanup_ready(); - auto* accumulator = new ExecAttachments; + auto *accumulator = new ExecAttachments; CUuserObject object = nullptr; CUresult status; { GILReleaseGuard gil; - status = DRIVER_CALL(cuUserObjectCreate, + status = DRIVER_CALL( + cuUserObjectCreate, &object, - static_cast(accumulator), + static_cast(accumulator), reinterpret_cast(enqueue_cleanup), 1, - CU_USER_OBJECT_NO_DESTRUCTOR_SYNC); + CU_USER_OBJECT_NO_DESTRUCTOR_SYNC + ); if (status != CUDA_SUCCESS) { delete accumulator; return status; } accumulator->object = object; - status = DRIVER_CALL(cuGraphRetainUserObject, - *h_source, object, 1, CU_GRAPH_USER_OBJECT_MOVE); + status = DRIVER_CALL(cuGraphRetainUserObject, *h_source, object, 1, CU_GRAPH_USER_OBJECT_MOVE); if (status != CUDA_SUCCESS) { // Dropping the last reference retires the accumulator. pw_cuUserObjectRelease(object, 1); @@ -137,20 +145,18 @@ CUresult stage_exec_attachments( // rollback can drop owners staged for a mutation that CUDA rejected. struct PreparedExecAttachmentState { GraphExecHandle h_exec; - ExecAttachments* attachments = nullptr; + ExecAttachments *attachments = nullptr; size_t original_size = 0; - PreparedExecAttachmentState( - GraphExecHandle h_exec_, - ExecAttachments* attachments_, - size_t original_size_) - : h_exec(std::move(h_exec_)), - attachments(attachments_), - original_size(original_size_) {} + PreparedExecAttachmentState(GraphExecHandle h_exec_, ExecAttachments *attachments_, size_t original_size_) + : h_exec(std::move(h_exec_)), attachments(attachments_), original_size(original_size_) + { + } }; -void rollback_prepared_exec_attachment( - PreparedExecAttachmentState* state) noexcept { +void +rollback_prepared_exec_attachment(PreparedExecAttachmentState *state) noexcept +{ if (!state) { return; } @@ -162,9 +168,9 @@ void rollback_prepared_exec_attachment( delete state; } -GraphExecHandle create_graph_exec_handle( - const GraphHandle& h_source, - CUDA_GRAPH_INSTANTIATE_PARAMS* params) { +GraphExecHandle +create_graph_exec_handle(const GraphHandle &h_source, CUDA_GRAPH_INSTANTIATE_PARAMS *params) +{ if (!h_source || !*h_source || !params) { err = CUDA_ERROR_INVALID_VALUE; return {}; @@ -193,22 +199,20 @@ GraphExecHandle create_graph_exec_handle( return {}; } - GraphExecHandle h_exec = make_graph_exec_handle( - graph_exec, staging.accumulator); + GraphExecHandle h_exec = make_graph_exec_handle(graph_exec, staging.accumulator); if (CUDA_SUCCESS != (err = staging.release())) { return {}; } return h_exec; } -CUresult graph_exec_update( - const GraphExecHandle& h_exec, - const GraphHandle& h_source, - CUgraphExecUpdateResultInfo* result_info) { +CUresult +graph_exec_update(const GraphExecHandle &h_exec, const GraphHandle &h_source, CUgraphExecUpdateResultInfo *result_info) +{ if (!h_exec || !h_source || !*h_source || !result_info) { return CUDA_ERROR_INVALID_VALUE; } - GraphExecBox* box = get_exec_box(h_exec); + GraphExecBox *box = get_exec_box(h_exec); if (!box->resource) { return CUDA_ERROR_INVALID_VALUE; } @@ -233,11 +237,14 @@ CUresult graph_exec_update( return staging.release(); } -CUresult graph_prepare_exec_attachment( - const GraphExecHandle& h_exec, - OpaqueHandle owner0, - OpaqueHandle owner1, - PreparedExecAttachment* out_prepared) { +CUresult +graph_prepare_exec_attachment( + const GraphExecHandle &h_exec, + OpaqueHandle owner0, + OpaqueHandle owner1, + PreparedExecAttachment *out_prepared +) +{ if (!out_prepared) { return CUDA_ERROR_INVALID_VALUE; } @@ -246,21 +253,21 @@ CUresult graph_prepare_exec_attachment( return CUDA_ERROR_INVALID_VALUE; } - GraphExecBox* box = get_exec_box(h_exec); + GraphExecBox *box = get_exec_box(h_exec); if (!box->resource || !box->attachments) { return CUDA_ERROR_INVALID_VALUE; } - ExecAttachments* attachments = box->attachments; + ExecAttachments *attachments = box->attachments; const size_t original_size = attachments->owners.size(); const size_t additions = - static_cast(static_cast(owner0)) + - static_cast(static_cast(owner1)); + static_cast(static_cast(owner0)) + static_cast(static_cast(owner1)); // Reserve before staging so that rollback and commit cannot allocate. attachments->owners.reserve(original_size + additions); PreparedExecAttachment prepared( new PreparedExecAttachmentState(h_exec, attachments, original_size), - PreparedExecAttachmentDeleter{rollback_prepared_exec_attachment}); + PreparedExecAttachmentDeleter{rollback_prepared_exec_attachment} + ); if (owner0) { attachments->owners.emplace_back(std::move(owner0)); } @@ -271,8 +278,9 @@ CUresult graph_prepare_exec_attachment( return CUDA_SUCCESS; } -void graph_commit_exec_attachment( - PreparedExecAttachment& prepared) noexcept { +void +graph_commit_exec_attachment(PreparedExecAttachment &prepared) noexcept +{ delete prepared.release(); } diff --git a/cuda_core/cuda/core/_cpp/rt/internal.hpp b/cuda_core/cuda/core/_cpp/rt/internal.hpp index 39b3dae0886..2378f98d2e6 100644 --- a/cuda_core/cuda/core/_cpp/rt/internal.hpp +++ b/cuda_core/cuda/core/_cpp/rt/internal.hpp @@ -21,32 +21,37 @@ namespace cuda_core::rt::detail { // Implemented in error.cpp -void format_cuda_error(char* buffer, size_t size, const char* operation, CUresult status, - const char* detail) noexcept; +void +format_cuda_error(char *buffer, size_t size, const char *operation, CUresult status, const char *detail) noexcept; // Implemented in error.cpp -void note_context_not_restored(CUcontext previous, CUresult operation_status, - CUresult restore_status) noexcept; +void +note_context_not_restored(CUcontext previous, CUresult operation_status, CUresult restore_status) noexcept; // Implemented in error.cpp -void format_operation(char* buffer, size_t size, const char* operation, - unsigned long long handle) noexcept; +void +format_operation(char *buffer, size_t size, const char *operation, unsigned long long handle) noexcept; // Bits of a resource handle for diagnostics: a pointer as its address, an // integer handle (CUdeviceptr, CUtexObject, ...) as its value, and a pointer // to a handle (nvrtcDestroyProgram(&prog) and friends) as the handle it points // to. template -unsigned long long handle_bits(const T& value) noexcept { +unsigned long long +handle_bits(const T &value) noexcept +{ using U = std::remove_cv_t>; if constexpr (std::is_pointer_v) { using P = std::remove_cv_t>; if constexpr (std::is_pointer_v

) { return value ? handle_bits(*value) : 0ull; - } else { + } + else { return static_cast(reinterpret_cast(value)); } - } else if constexpr (std::is_integral_v || std::is_enum_v) { + } + else if constexpr (std::is_integral_v || std::is_enum_v) { return static_cast(value); - } else { + } + else { return 0ull; } } @@ -58,15 +63,18 @@ struct DeallocationStream { }; // Implemented in stream.cpp -ContextHandle deallocation_context(const DeallocationStream& stream) noexcept; +ContextHandle +deallocation_context(const DeallocationStream &stream) noexcept; // Implemented in stream.cpp -bool make_deallocation_stream(const StreamHandle& h, DeallocationStream& out) noexcept; +bool +make_deallocation_stream(const StreamHandle &h, DeallocationStream &out) noexcept; // Implemented in virtual_memory.cpp. Synchronize the stream a VMM buffer // recorded before its mappings are released. Skips the sync, with a report // when a capture is the reason, if no stream was recorded, the interpreter is // finalizing, or the sync would disturb a graph capture (VMM_DESIGN.md). -void vmm_sync_before_release(const DeallocationStream& stream) noexcept; +void +vmm_sync_before_release(const DeallocationStream &stream) noexcept; // Decorate a status-returning cleanup call to report whenever it fails. CUDA // calls (CUresult) are reported with the error name and description; NVRTC, @@ -78,10 +86,10 @@ void vmm_sync_before_release(const DeallocationStream& stream) noexcept; // one while holding a C++ lock; the GIL must be the outermost lock. Where a // lock must stay held, call the p_ pointer, keep the status, and report after // the lock is released (see deviceptr_import_ipc and DESIGN.md). -template +template class WarnOnFailure { -public: - explicit WarnOnFailure(const char* operation) noexcept : operation_(operation) {} + public: + explicit WarnOnFailure(const char *operation) noexcept : operation_(operation) {} // The first argument is the resource being released; it is named in the // report so that independent failures are not collapsed by the warning @@ -89,16 +97,22 @@ class WarnOnFailure { // table like DRIVER_CALL: it reports an unavailable entry and yields an // error status, never a null dereference. template - auto operator()(First&& first, Rest&&... rest) const noexcept { + auto + operator()(First &&first, Rest &&...rest) const noexcept + { const unsigned long long handle = handle_bits(first); auto status = detail::fn_or_unavailable(Function, Table, operation_)( - std::forward(first), std::forward(rest)...); + std::forward(first), + std::forward(rest)... + ); report(status, handle); return status; } -private: - void report(CUresult status, unsigned long long handle) const noexcept { + private: + void + report(CUresult status, unsigned long long handle) const noexcept + { if (status == CUDA_SUCCESS || status == CUDA_ERROR_DEINITIALIZED) { return; } @@ -108,7 +122,9 @@ class WarnOnFailure { } template - void report(Status status, unsigned long long handle) const noexcept { + void + report(Status status, unsigned long long handle) const noexcept + { if (static_cast(status) != 0) { char operation[160]; format_operation(operation, sizeof(operation), operation_, handle); @@ -116,7 +132,7 @@ class WarnOnFailure { } } - const char* operation_; + const char *operation_; }; // Warning-decorated CUDA operations for deleters and cleanup paths. Each one @@ -147,14 +163,16 @@ const WarnOnFailure pw_nvJitLinkDestroy{ // Intrusive base for payloads transferred out of CUDA's callback. struct DeferredCleanupItem { - DeferredCleanupItem* next = nullptr; + DeferredCleanupItem *next = nullptr; virtual ~DeferredCleanupItem() noexcept = default; }; // Implemented in py_deferred_cleanup.cpp -void ensure_deferred_cleanup_ready(); +void +ensure_deferred_cleanup_ready(); // Implemented in py_deferred_cleanup.cpp -void enqueue_cleanup(void* item) noexcept; +void +enqueue_cleanup(void *item) noexcept; // ============================================================================ // Handle reverse-lookup registry @@ -164,31 +182,39 @@ void enqueue_cleanup(void* item) noexcept; // Uses weak_ptr to avoid preventing destruction. // ============================================================================ -template> +template > class HandleRegistry { -public: + public: using MapType = std::unordered_map, Hash>; - void register_handle(const Key& key, const Handle& h) { + void + register_handle(const Key &key, const Handle &h) + { std::lock_guard lock(mutex_); map_[key] = h; } - void unregister_handle(const Key& key) noexcept { + void + unregister_handle(const Key &key) noexcept + { std::lock_guard lock(mutex_); map_.erase(key); } - void register_handles(const std::vector& handles) { + void + register_handles(const std::vector &handles) + { std::lock_guard lock(mutex_); - for (const Handle& h : handles) { + for (const Handle &h : handles) { if (h) { map_[*h] = h; } } } - Handle lookup(const Key& key) { + Handle + lookup(const Key &key) + { std::lock_guard lock(mutex_); auto it = map_.find(key); if (it != map_.end()) { @@ -200,8 +226,10 @@ class HandleRegistry { return {}; } - template - Handle get_or_create(const Key& key, Factory&& create) { + template + Handle + get_or_create(const Key &key, Factory &&create) + { std::lock_guard lock(mutex_); auto it = map_.find(key); if (it != map_.end()) { @@ -218,14 +246,16 @@ class HandleRegistry { return h; } - MapType drain() noexcept { + MapType + drain() noexcept + { std::lock_guard lock(mutex_); MapType extracted; extracted.swap(map_); return extracted; } -private: + private: std::mutex mutex_; MapType map_; }; diff --git a/cuda_core/cuda/core/_cpp/rt/memory.cpp b/cuda_core/cuda/core/_cpp/rt/memory.cpp index 47bdd89e387..6cd6f6b82d8 100644 --- a/cuda_core/cuda/core/_cpp/rt/memory.cpp +++ b/cuda_core/cuda/core/_cpp/rt/memory.cpp @@ -48,7 +48,9 @@ std::atomic g_live_owned_mempools{0}; // revoked with its own request, and the owning device (-1 if the pool has // none) is not requested at all. // Must be noexcept since it's called from a shared_ptr deleter. -static void clear_mempool_peer_access(CUmemoryPool pool, int owner_device) noexcept { +static void +clear_mempool_peer_access(CUmemoryPool pool, int owner_device) noexcept +{ try { int device_count = 0; if (DRIVER_CALL(cuDeviceGetCount, &device_count) != CUDA_SUCCESS || device_count <= 0) { @@ -65,37 +67,41 @@ static void clear_mempool_peer_access(CUmemoryPool pool, int owner_device) noexc revoke.location.id = i; DRIVER_CALL(cuMemPoolSetAccess, pool, &revoke, 1); // Best effort } - } catch (...) { + } + catch (...) { // Swallow exceptions - this is best-effort cleanup in destructor context } } // owner_device is the ordinal of a device-located pool, or -1 when there is no // device to skip (host pools) or the location is unknown (imported pools). -static MemoryPoolHandle wrap_mempool_owned(CUmemoryPool pool, int owner_device) { - auto* raw = new MemoryPoolBox{pool}; +static MemoryPoolHandle +wrap_mempool_owned(CUmemoryPool pool, int owner_device) +{ + auto *raw = new MemoryPoolBox{pool}; // Count before constructing the shared_ptr: if its control-block allocation // throws, it runs the deleter on raw, whose decrement then balances this. g_live_owned_mempools.fetch_add(1, std::memory_order_relaxed); - auto box = std::shared_ptr( - raw, - [owner_device](const MemoryPoolBox* b) { - GILReleaseGuard gil; - clear_mempool_peer_access(b->resource, owner_device); - pw_cuMemPoolDestroy(b->resource); - g_live_owned_mempools.fetch_sub(1, std::memory_order_relaxed); - delete b; - } - ); + auto box = std::shared_ptr(raw, [owner_device](const MemoryPoolBox *b) { + GILReleaseGuard gil; + clear_mempool_peer_access(b->resource, owner_device); + pw_cuMemPoolDestroy(b->resource); + g_live_owned_mempools.fetch_sub(1, std::memory_order_relaxed); + delete b; + }); return MemoryPoolHandle(box, &box->resource); } // Test-suite instrumentation; see g_live_owned_mempools. -long long live_owned_mempool_count() noexcept { +long long +live_owned_mempool_count() noexcept +{ return g_live_owned_mempools.load(std::memory_order_relaxed); } -MemoryPoolHandle create_mempool_handle(const CUmemPoolProps& props) { +MemoryPoolHandle +create_mempool_handle(const CUmemPoolProps &props) +{ GILReleaseGuard gil; CUmemoryPool pool; if (CUDA_SUCCESS != (err = DRIVER_CALL(cuMemPoolCreate, &pool, &props))) { @@ -105,12 +111,16 @@ MemoryPoolHandle create_mempool_handle(const CUmemPoolProps& props) { return wrap_mempool_owned(pool, owner_device); } -MemoryPoolHandle create_mempool_handle_ref(CUmemoryPool pool) { +MemoryPoolHandle +create_mempool_handle_ref(CUmemoryPool pool) +{ auto box = std::make_shared(MemoryPoolBox{pool}); return MemoryPoolHandle(box, &box->resource); } -MemoryPoolHandle get_device_mempool(int device_id) { +MemoryPoolHandle +get_device_mempool(int device_id) +{ GILReleaseGuard gil; CUmemoryPool pool; if (CUDA_SUCCESS != (err = DRIVER_CALL(cuDeviceGetMemPool, &pool, device_id))) { @@ -119,10 +129,12 @@ MemoryPoolHandle get_device_mempool(int device_id) { return create_mempool_handle_ref(pool); } -MemoryPoolHandle create_mempool_handle_ipc(int fd, CUmemAllocationHandleType handle_type) { +MemoryPoolHandle +create_mempool_handle_ipc(int fd, CUmemAllocationHandleType handle_type) +{ GILReleaseGuard gil; CUmemoryPool pool; - auto handle_ptr = reinterpret_cast(static_cast(fd)); + auto handle_ptr = reinterpret_cast(static_cast(fd)); if (CUDA_SUCCESS != (err = DRIVER_CALL(cuMemPoolImportFromShareableHandle, &pool, handle_ptr, handle_type, 0))) { return {}; } @@ -157,21 +169,27 @@ struct VmmDevicePtrBox : DevicePtrBox { // &box->resource, so we can compute the containing struct using offsetof. // The const_cast is safe because we only use this to access the mutable // deallocation member or in the deleter (where the box is being destroyed). -static DevicePtrBox* get_box(const DevicePtrHandle& h) { - const CUdeviceptr* p = h.get(); - return reinterpret_cast( - reinterpret_cast(const_cast(p)) - offsetof(DevicePtrBox, resource) +static DevicePtrBox * +get_box(const DevicePtrHandle &h) +{ + const CUdeviceptr *p = h.get(); + return reinterpret_cast( + reinterpret_cast(const_cast(p)) - offsetof(DevicePtrBox, resource) ); } // Return the stream that orders a device pointer's deallocation; empty for an // empty handle, as set_deallocation_stream rejects one. -StreamHandle deallocation_stream(const DevicePtrHandle& h) noexcept { +StreamHandle +deallocation_stream(const DevicePtrHandle &h) noexcept +{ return h ? get_box(h)->deallocation.h_stream : StreamHandle{}; } // Replace the stream that orders a device pointer's deallocation. -CUresult set_deallocation_stream(const DevicePtrHandle& h, const StreamHandle& h_stream) noexcept { +CUresult +set_deallocation_stream(const DevicePtrHandle &h, const StreamHandle &h_stream) noexcept +{ if (!h) { return CUDA_ERROR_INVALID_VALUE; } @@ -183,7 +201,9 @@ CUresult set_deallocation_stream(const DevicePtrHandle& h, const StreamHandle& h return CUDA_SUCCESS; } -DevicePtrHandle deviceptr_alloc_from_pool(size_t size, const MemoryPoolHandle& h_pool, const StreamHandle& h_stream) { +DevicePtrHandle +deviceptr_alloc_from_pool(size_t size, const MemoryPoolHandle &h_pool, const StreamHandle &h_stream) +{ GILReleaseGuard gil; CUdeviceptr ptr; if (CUDA_SUCCESS != (err = DRIVER_CALL(cuMemAllocFromPoolAsync, &ptr, size, *h_pool, as_cu(h_stream)))) { @@ -196,24 +216,20 @@ DevicePtrHandle deviceptr_alloc_from_pool(size_t size, const MemoryPoolHandle& h return {}; } - auto box = std::shared_ptr( - new DevicePtrBox{ptr, std::move(ds)}, - [h_pool](DevicePtrBox* b) { - GILReleaseGuard gil; - const DeallocationStream& stream = b->deallocation; - cleanup_in_context( - deallocation_context(stream), "cuMemFreeAsync", handle_bits(b->resource), - [&]() noexcept { - return DRIVER_CALL(cuMemFreeAsync, - b->resource, as_cu(stream.h_stream)); - }); - delete b; - } - ); + auto box = std::shared_ptr(new DevicePtrBox{ptr, std::move(ds)}, [h_pool](DevicePtrBox *b) { + GILReleaseGuard gil; + const DeallocationStream &stream = b->deallocation; + cleanup_in_context(deallocation_context(stream), "cuMemFreeAsync", handle_bits(b->resource), [&]() noexcept { + return DRIVER_CALL(cuMemFreeAsync, b->resource, as_cu(stream.h_stream)); + }); + delete b; + }); return DevicePtrHandle(box, &box->resource); } -DevicePtrHandle deviceptr_alloc_async(size_t size, const StreamHandle& h_stream) { +DevicePtrHandle +deviceptr_alloc_async(size_t size, const StreamHandle &h_stream) +{ GILReleaseGuard gil; CUdeviceptr ptr; if (CUDA_SUCCESS != (err = DRIVER_CALL(cuMemAllocAsync, &ptr, size, as_cu(h_stream)))) { @@ -226,58 +242,64 @@ DevicePtrHandle deviceptr_alloc_async(size_t size, const StreamHandle& h_stream) return {}; } - auto box = std::shared_ptr( - new DevicePtrBox{ptr, std::move(ds)}, - [](DevicePtrBox* b) { - GILReleaseGuard gil; - const DeallocationStream& stream = b->deallocation; - cleanup_in_context( - deallocation_context(stream), "cuMemFreeAsync", handle_bits(b->resource), - [&]() noexcept { - return DRIVER_CALL(cuMemFreeAsync, - b->resource, as_cu(stream.h_stream)); - }); - delete b; - } - ); + auto box = std::shared_ptr(new DevicePtrBox{ptr, std::move(ds)}, [](DevicePtrBox *b) { + GILReleaseGuard gil; + const DeallocationStream &stream = b->deallocation; + cleanup_in_context(deallocation_context(stream), "cuMemFreeAsync", handle_bits(b->resource), [&]() noexcept { + return DRIVER_CALL(cuMemFreeAsync, b->resource, as_cu(stream.h_stream)); + }); + delete b; + }); return DevicePtrHandle(box, &box->resource); } // Allocate device memory synchronously with the provided context current. -CUresult deviceptr_alloc_raw(CUdeviceptr* ptr, size_t size, - const ContextHandle& h_context) noexcept { +CUresult +deviceptr_alloc_raw(CUdeviceptr *ptr, size_t size, const ContextHandle &h_context) noexcept +{ GILReleaseGuard gil; return invoke_in_context_or_undo( h_context, - [&]() noexcept { return DRIVER_CALL(cuMemAlloc, ptr, size); }, - [&]() noexcept { pw_cuMemFree(*ptr); }, - /*undo_requires_target_context=*/false); + [&]() noexcept { + return DRIVER_CALL(cuMemAlloc, ptr, size); + }, + [&]() noexcept { + pw_cuMemFree(*ptr); + }, + /*undo_requires_target_context=*/false + ); } -DevicePtrHandle deviceptr_alloc_host(size_t size) { +DevicePtrHandle +deviceptr_alloc_host(size_t size) +{ GILReleaseGuard gil; - void* ptr; + void *ptr; if (CUDA_SUCCESS != (err = DRIVER_CALL(cuMemAllocHost, &ptr, size))) { return {}; } auto box = std::shared_ptr( new DevicePtrBox{reinterpret_cast(ptr), DeallocationStream{}}, - [](DevicePtrBox* b) { + [](DevicePtrBox *b) { GILReleaseGuard gil; - pw_cuMemFreeHost(reinterpret_cast(b->resource)); + pw_cuMemFreeHost(reinterpret_cast(b->resource)); delete b; } ); return DevicePtrHandle(box, &box->resource); } -DevicePtrHandle deviceptr_create_ref(CUdeviceptr ptr) { +DevicePtrHandle +deviceptr_create_ref(CUdeviceptr ptr) +{ auto box = std::make_shared(DevicePtrBox{ptr, DeallocationStream{}}); return DevicePtrHandle(box, &box->resource); } -DevicePtrHandle deviceptr_create_with_owner(CUdeviceptr ptr, PyObject* owner) { +DevicePtrHandle +deviceptr_create_with_owner(CUdeviceptr ptr, PyObject *owner) +{ if (!owner) { return deviceptr_create_ref(ptr); } @@ -288,43 +310,41 @@ DevicePtrHandle deviceptr_create_with_owner(CUdeviceptr ptr, PyObject* owner) { return deviceptr_create_ref(ptr); } Py_INCREF(owner); - auto box = std::shared_ptr( - new DevicePtrBox{ptr, DeallocationStream{}}, - [owner](DevicePtrBox* b) { - GILAcquireGuard gil; - if (gil.acquired()) { - Py_DECREF(owner); - } - delete b; + auto box = std::shared_ptr(new DevicePtrBox{ptr, DeallocationStream{}}, [owner](DevicePtrBox *b) { + GILAcquireGuard gil; + if (gil.acquired()) { + Py_DECREF(owner); } - ); + delete b; + }); return DevicePtrHandle(box, &box->resource); } -DevicePtrHandle deviceptr_create_mapped_graphics( +DevicePtrHandle +deviceptr_create_mapped_graphics( CUdeviceptr ptr, - const GraphicsResourceHandle& h_resource, - const StreamHandle& h_stream -) { + const GraphicsResourceHandle &h_resource, + const StreamHandle &h_stream +) +{ DeallocationStream ds; if (!make_deallocation_stream(h_stream, ds)) { return {}; } - auto box = std::shared_ptr( - new DevicePtrBox{ptr, std::move(ds)}, - [h_resource](DevicePtrBox* b) { - GILReleaseGuard gil; - CUgraphicsResource resource = as_cu(h_resource); - const DeallocationStream& stream = b->deallocation; - cleanup_in_context( - deallocation_context(stream), "cuGraphicsUnmapResources", handle_bits(resource), - [&]() noexcept { - return DRIVER_CALL(cuGraphicsUnmapResources, - 1, &resource, as_cu(stream.h_stream)); - }); - delete b; - } - ); + auto box = std::shared_ptr(new DevicePtrBox{ptr, std::move(ds)}, [h_resource](DevicePtrBox *b) { + GILReleaseGuard gil; + CUgraphicsResource resource = as_cu(h_resource); + const DeallocationStream &stream = b->deallocation; + cleanup_in_context( + deallocation_context(stream), + "cuGraphicsUnmapResources", + handle_bits(resource), + [&]() noexcept { + return DRIVER_CALL(cuGraphicsUnmapResources, 1, &resource, as_cu(stream.h_stream)); + } + ); + delete b; + }); return DevicePtrHandle(box, &box->resource); } @@ -332,14 +352,16 @@ DevicePtrHandle deviceptr_create_mapped_graphics( // Virtual memory ranges (VMM_DESIGN.md) // ============================================================================ -DevicePtrHandle deviceptr_create_vmm(CUdeviceptr base, const VmmRangeHandle& range) { +DevicePtrHandle +deviceptr_create_vmm(CUdeviceptr base, const VmmRangeHandle &range) +{ if (!range) { // a null handle; an empty range is valid err = CUDA_ERROR_INVALID_VALUE; return {}; } auto box = std::shared_ptr( new VmmDevicePtrBox{{base, DeallocationStream{}}, range}, - [](VmmDevicePtrBox* b) { + [](VmmDevicePtrBox *b) { GILReleaseGuard gil; // Order the release on this buffer's recorded stream, then free // the box, which drops the range: every mapping this buffer was @@ -352,10 +374,12 @@ DevicePtrHandle deviceptr_create_vmm(CUdeviceptr base, const VmmRangeHandle& ran return DevicePtrHandle(box, &box->resource); } -VmmRangeHandle vmm_range(const DevicePtrHandle& h) noexcept { +VmmRangeHandle +vmm_range(const DevicePtrHandle &h) noexcept +{ // Only for handles from deviceptr_create_vmm; the VirtualMemoryBuffer // class guarantees that for its callers. - return h ? static_cast(get_box(h))->range : VmmRangeHandle{}; + return h ? static_cast(get_box(h))->range : VmmRangeHandle{}; } // ============================================================================ @@ -364,11 +388,15 @@ VmmRangeHandle vmm_range(const DevicePtrHandle& h) noexcept { static MRDeallocCallback mr_dealloc_cb = nullptr; -void register_mr_dealloc_callback(MRDeallocCallback cb) { +void +register_mr_dealloc_callback(MRDeallocCallback cb) +{ mr_dealloc_cb = cb; } -DevicePtrHandle deviceptr_create_with_mr(CUdeviceptr ptr, size_t size, PyObject* mr) { +DevicePtrHandle +deviceptr_create_with_mr(CUdeviceptr ptr, size_t size, PyObject *mr) +{ if (!mr) { return deviceptr_create_ref(ptr); } @@ -378,29 +406,29 @@ DevicePtrHandle deviceptr_create_with_mr(CUdeviceptr ptr, size_t size, PyObject* return deviceptr_create_ref(ptr); } Py_INCREF(mr); - auto box = std::shared_ptr( - new DevicePtrBox{ptr, DeallocationStream{}}, - [mr, size](DevicePtrBox* b) { - GILAcquireGuard gil; - if (gil.acquired()) { - // The last reference may go while an exception propagates - // through the releasing caller; deallocate() must run with a - // clean error state and leave that exception in place. - PendingExceptionGuard pending; - if (mr_dealloc_cb) { - const DeallocationStream& stream = b->deallocation; - cleanup_in_context( - deallocation_context(stream), "MemoryResource.deallocate", handle_bits(b->resource), - [&]() noexcept { - mr_dealloc_cb(mr, b->resource, size, stream.h_stream); - return CUDA_SUCCESS; - }); - } - Py_DECREF(mr); + auto box = std::shared_ptr(new DevicePtrBox{ptr, DeallocationStream{}}, [mr, size](DevicePtrBox *b) { + GILAcquireGuard gil; + if (gil.acquired()) { + // The last reference may go while an exception propagates + // through the releasing caller; deallocate() must run with a + // clean error state and leave that exception in place. + PendingExceptionGuard pending; + if (mr_dealloc_cb) { + const DeallocationStream &stream = b->deallocation; + cleanup_in_context( + deallocation_context(stream), + "MemoryResource.deallocate", + handle_bits(b->resource), + [&]() noexcept { + mr_dealloc_cb(mr, b->resource, size, stream.h_stream); + return CUDA_SUCCESS; + } + ); } - delete b; + Py_DECREF(mr); } - ); + delete b; + }); return DevicePtrHandle(box, &box->resource); } @@ -427,11 +455,12 @@ DevicePtrHandle deviceptr_create_with_mr(CUdeviceptr ptr, size_t size, PyObject* // The cache key is the export_data bytes (CUmemPoolPtrExportData), not the // returned pointer, because we must check before calling the driver API. - // TODO: When driver fix for nvbug 5570902 is available, consider whether // the cache is still needed for memory pool allocations (it will still be // needed for pinned memory). -static bool use_ipc_ptr_cache() { +static bool +use_ipc_ptr_cache() +{ return true; } @@ -440,16 +469,20 @@ namespace { struct ExportDataKey { CUmemPoolPtrExportData data; - bool operator==(const ExportDataKey& other) const { + bool + operator==(const ExportDataKey &other) const + { return std::memcmp(&data, &other.data, sizeof(data)) == 0; } }; struct ExportDataKeyHash { - std::size_t operator()(const ExportDataKey& key) const { + std::size_t + operator()(const ExportDataKey &key) const + { // Simple hash of the bytes std::size_t h = 0; - const auto* bytes = reinterpret_cast(&key.data); + const auto *bytes = reinterpret_cast(&key.data); for (std::size_t i = 0; i < sizeof(key.data); ++i) { h = h * 31 + bytes[i]; } @@ -457,14 +490,15 @@ struct ExportDataKeyHash { } }; -} +} // namespace static HandleRegistry ipc_ptr_cache; static std::mutex ipc_import_mutex; -DevicePtrHandle deviceptr_import_ipc(const MemoryPoolHandle& h_pool, const void* export_data, const StreamHandle& h_stream) { - auto data = const_cast( - reinterpret_cast(export_data)); +DevicePtrHandle +deviceptr_import_ipc(const MemoryPoolHandle &h_pool, const void *export_data, const StreamHandle &h_stream) +{ + auto data = const_cast(reinterpret_cast(export_data)); // Resolve the table before you take any lock: a fill acquires the GIL, and // nothing under ipc_import_mutex may acquire it (#2840). The raw p_ calls @@ -502,9 +536,8 @@ DevicePtrHandle deviceptr_import_ipc(const MemoryPoolHandle& h_pool, const void* DeallocationStream ds; if (make_deallocation_stream(h_stream, ds)) { - auto box = std::shared_ptr( - new DevicePtrBox{ptr, std::move(ds)}, - [h_pool, key](DevicePtrBox* b) { + auto box = + std::shared_ptr(new DevicePtrBox{ptr, std::move(ds)}, [h_pool, key](DevicePtrBox *b) { // Release the GIL first (the GIL is the outermost lock), then hold the // mutex across unregister + free: a concurrent import that finds this // entry expired must wait until the mapping is gone, or it re-imports @@ -513,19 +546,26 @@ DevicePtrHandle deviceptr_import_ipc(const MemoryPoolHandle& h_pool, const void* // deallocation context is resolved before the lock (it may report) and // the lock is released as soon as the cleanup is done. GILReleaseGuard gil; - const DeallocationStream& stream = b->deallocation; + const DeallocationStream &stream = b->deallocation; ContextHandle h_dealloc = deallocation_context(stream); std::unique_lock lock(ipc_import_mutex); ipc_ptr_cache.unregister_handle(key); cleanup_in_context( - h_dealloc, "cuMemFreeAsync", handle_bits(b->resource), + h_dealloc, + "cuMemFreeAsync", + handle_bits(b->resource), [&]() noexcept { - return p_cuMemFreeAsync(b->resource, as_cu(stream.h_stream)); // raw: under ipc_import_mutex + return p_cuMemFreeAsync( // raw: under ipc_import_mutex + b->resource, + as_cu(stream.h_stream) + ); }, - [&]() noexcept { lock.unlock(); }); + [&]() noexcept { + lock.unlock(); + } + ); delete b; - } - ); + }); DevicePtrHandle h(box, &box->resource); ipc_ptr_cache.register_handle(key, h); return h; @@ -539,13 +579,16 @@ DevicePtrHandle deviceptr_import_ipc(const MemoryPoolHandle& h_pool, const void* if (discard_status != CUDA_SUCCESS) { char operation[160]; format_operation(operation, sizeof(operation), "cuMemFreeAsync", discarded); - report_cuda_error(operation, discard_status, - "failed while discarding an IPC import that could not record a " - "deallocation stream; the mapping leaked"); + report_cuda_error( + operation, + discard_status, + "failed while discarding an IPC import that could not record a " + "deallocation stream; the mapping leaked" + ); } return {}; - - } else { + } + else { GILReleaseGuard gil; CUdeviceptr ptr; if (CUDA_SUCCESS != (err = DRIVER_CALL(cuMemPoolImportPointer, &ptr, *h_pool, data))) { @@ -558,20 +601,19 @@ DevicePtrHandle deviceptr_import_ipc(const MemoryPoolHandle& h_pool, const void* return {}; } - auto box = std::shared_ptr( - new DevicePtrBox{ptr, std::move(ds)}, - [h_pool](DevicePtrBox* b) { - GILReleaseGuard gil; - const DeallocationStream& stream = b->deallocation; - cleanup_in_context( - deallocation_context(stream), "cuMemFreeAsync", handle_bits(b->resource), - [&]() noexcept { - return DRIVER_CALL(cuMemFreeAsync, - b->resource, as_cu(stream.h_stream)); - }); - delete b; - } - ); + auto box = std::shared_ptr(new DevicePtrBox{ptr, std::move(ds)}, [h_pool](DevicePtrBox *b) { + GILReleaseGuard gil; + const DeallocationStream &stream = b->deallocation; + cleanup_in_context( + deallocation_context(stream), + "cuMemFreeAsync", + handle_bits(b->resource), + [&]() noexcept { + return DRIVER_CALL(cuMemFreeAsync, b->resource, as_cu(stream.h_stream)); + } + ); + delete b; + }); return DevicePtrHandle(box, &box->resource); } } @@ -580,23 +622,24 @@ DevicePtrHandle deviceptr_import_ipc(const MemoryPoolHandle& h_pool, const void* // File Descriptor Handles // ============================================================================ -FileDescriptorHandle create_fd_handle(int fd) { +FileDescriptorHandle +create_fd_handle(int fd) +{ #ifdef _WIN32 throw std::runtime_error("create_fd_handle is not supported on Windows"); #else - return FileDescriptorHandle( - new int(fd), - [](const int* p) { - if (::close(*p) != 0) { - report_message("close() failed for an IPC file descriptor; the descriptor may have leaked"); - } - delete p; + return FileDescriptorHandle(new int(fd), [](const int *p) { + if (::close(*p) != 0) { + report_message("close() failed for an IPC file descriptor; the descriptor may have leaked"); } - ); + delete p; + }); #endif } -FileDescriptorHandle create_fd_handle_ref(int fd) { +FileDescriptorHandle +create_fd_handle_ref(int fd) +{ #ifdef _WIN32 throw std::runtime_error("create_fd_handle_ref is not supported on Windows"); #else diff --git a/cuda_core/cuda/core/_cpp/rt/program.cpp b/cuda_core/cuda/core/_cpp/rt/program.cpp index 1f32c07bf71..bff0102d453 100644 --- a/cuda_core/cuda/core/_cpp/rt/program.cpp +++ b/cuda_core/cuda/core/_cpp/rt/program.cpp @@ -25,45 +25,47 @@ struct LibraryBox { }; } // namespace -LibraryHandle create_library_handle_from_file(const char* path) { +LibraryHandle +create_library_handle_from_file(const char *path) +{ GILReleaseGuard gil; CUlibrary library; - if (CUDA_SUCCESS != (err = DRIVER_CALL(cuLibraryLoadFromFile, &library, path, nullptr, nullptr, 0, nullptr, nullptr, 0))) { + if (CUDA_SUCCESS != + (err = DRIVER_CALL(cuLibraryLoadFromFile, &library, path, nullptr, nullptr, 0, nullptr, nullptr, 0))) { return {}; } - auto box = std::shared_ptr( - new LibraryBox{library}, - [](const LibraryBox* b) { - GILReleaseGuard gil; - // TODO: re-enable once LibraryBox tracks its owning context - // p_cuLibraryUnload(b->resource); - delete b; - } - ); + auto box = std::shared_ptr(new LibraryBox{library}, [](const LibraryBox *b) { + GILReleaseGuard gil; + // TODO: re-enable once LibraryBox tracks its owning context + // p_cuLibraryUnload(b->resource); + delete b; + }); return LibraryHandle(box, &box->resource); } -LibraryHandle create_library_handle_from_data(const void* data) { +LibraryHandle +create_library_handle_from_data(const void *data) +{ GILReleaseGuard gil; CUlibrary library; - if (CUDA_SUCCESS != (err = DRIVER_CALL(cuLibraryLoadData, &library, data, nullptr, nullptr, 0, nullptr, nullptr, 0))) { + if (CUDA_SUCCESS != + (err = DRIVER_CALL(cuLibraryLoadData, &library, data, nullptr, nullptr, 0, nullptr, nullptr, 0))) { return {}; } - auto box = std::shared_ptr( - new LibraryBox{library}, - [](const LibraryBox* b) { - GILReleaseGuard gil; - // TODO: re-enable once LibraryBox tracks its owning context - // p_cuLibraryUnload(b->resource); - delete b; - } - ); + auto box = std::shared_ptr(new LibraryBox{library}, [](const LibraryBox *b) { + GILReleaseGuard gil; + // TODO: re-enable once LibraryBox tracks its owning context + // p_cuLibraryUnload(b->resource); + delete b; + }); return LibraryHandle(box, &box->resource); } -LibraryHandle create_library_handle_ref(CUlibrary library) { +LibraryHandle +create_library_handle_ref(CUlibrary library) +{ auto box = std::make_shared(LibraryBox{library}); return LibraryHandle(box, &box->resource); } @@ -79,17 +81,19 @@ struct KernelBox { }; } // namespace -static const KernelBox* get_box(const KernelHandle& h) { - const CUkernel* p = h.get(); - return reinterpret_cast( - reinterpret_cast(p) - offsetof(KernelBox, resource) - ); +static const KernelBox * +get_box(const KernelHandle &h) +{ + const CUkernel *p = h.get(); + return reinterpret_cast(reinterpret_cast(p) - offsetof(KernelBox, resource)); } // See REGISTRY_DESIGN.md (Level 1: Driver Handle -> Resource Handle) static HandleRegistry kernel_registry; -KernelHandle create_kernel_handle(const LibraryHandle& h_library, const char* name) { +KernelHandle +create_kernel_handle(const LibraryHandle &h_library, const char *name) +{ GILReleaseGuard gil; CUkernel kernel; if (CUDA_SUCCESS != (err = DRIVER_CALL(cuLibraryGetKernel, &kernel, *h_library, name))) { @@ -102,7 +106,9 @@ KernelHandle create_kernel_handle(const LibraryHandle& h_library, const char* na return h; } -KernelHandle create_kernel_handle_ref(CUkernel kernel) { +KernelHandle +create_kernel_handle_ref(CUkernel kernel) +{ if (auto h = kernel_registry.lookup(kernel)) { return h; } @@ -110,8 +116,11 @@ KernelHandle create_kernel_handle_ref(CUkernel kernel) { return KernelHandle(box, &box->resource); } -LibraryHandle get_kernel_library(const KernelHandle& h) noexcept { - if (!h) return {}; +LibraryHandle +get_kernel_library(const KernelHandle &h) noexcept +{ + if (!h) + return {}; return get_box(h)->h_library; } @@ -125,24 +134,25 @@ struct NvrtcProgramBox { }; } // namespace -NvrtcProgramHandle create_nvrtc_program_handle(nvrtcProgram prog) { +NvrtcProgramHandle +create_nvrtc_program_handle(nvrtcProgram prog) +{ // Resolve the table now, while the library that created `prog` is loaded, // so that the deleter never triggers a fill. ensure_fn_table(FnTable::nvrtc); - auto box = std::shared_ptr( - new NvrtcProgramBox{prog}, - [](NvrtcProgramBox* b) { - // Note: nvrtcDestroyProgram takes nvrtcProgram* and nulls it, - // but we're deleting the box anyway so nulling is harmless. - GILReleaseGuard gil; - pw_nvrtcDestroyProgram(&b->resource); - delete b; - } - ); + auto box = std::shared_ptr(new NvrtcProgramBox{prog}, [](NvrtcProgramBox *b) { + // Note: nvrtcDestroyProgram takes nvrtcProgram* and nulls it, + // but we're deleting the box anyway so nulling is harmless. + GILReleaseGuard gil; + pw_nvrtcDestroyProgram(&b->resource); + delete b; + }); return NvrtcProgramHandle(box, &box->resource); } -NvrtcProgramHandle create_nvrtc_program_handle_ref(nvrtcProgram prog) { +NvrtcProgramHandle +create_nvrtc_program_handle_ref(nvrtcProgram prog) +{ auto box = std::make_shared(NvrtcProgramBox{prog}); return NvrtcProgramHandle(box, &box->resource); } @@ -157,22 +167,23 @@ struct NvvmProgramBox { }; } // namespace -NvvmProgramHandle create_nvvm_program_handle(nvvmProgram prog) { +NvvmProgramHandle +create_nvvm_program_handle(nvvmProgram prog) +{ ensure_fn_table(FnTable::nvvm); - auto box = std::shared_ptr( - new NvvmProgramBox{{prog}}, - [](NvvmProgramBox* b) { - // Note: nvvmDestroyProgram takes nvvmProgram* and nulls it, - // but we're deleting the box anyway so nulling is harmless. - GILReleaseGuard gil; - pw_nvvmDestroyProgram(&b->resource.raw); - delete b; - } - ); + auto box = std::shared_ptr(new NvvmProgramBox{{prog}}, [](NvvmProgramBox *b) { + // Note: nvvmDestroyProgram takes nvvmProgram* and nulls it, + // but we're deleting the box anyway so nulling is harmless. + GILReleaseGuard gil; + pw_nvvmDestroyProgram(&b->resource.raw); + delete b; + }); return NvvmProgramHandle(box, &box->resource); } -NvvmProgramHandle create_nvvm_program_handle_ref(nvvmProgram prog) { +NvvmProgramHandle +create_nvvm_program_handle_ref(nvvmProgram prog) +{ auto box = std::make_shared(NvvmProgramBox{{prog}}); return NvvmProgramHandle(box, &box->resource); } @@ -187,22 +198,23 @@ struct NvJitLinkBox { }; } // namespace -NvJitLinkHandle create_nvjitlink_handle(nvJitLink_t handle) { +NvJitLinkHandle +create_nvjitlink_handle(nvJitLink_t handle) +{ ensure_fn_table(FnTable::nvjitlink); - auto box = std::shared_ptr( - new NvJitLinkBox{{handle}}, - [](NvJitLinkBox* b) { - // Note: nvJitLinkDestroy takes nvJitLinkHandle* and nulls it, - // but we're deleting the box anyway so nulling is harmless. - GILReleaseGuard gil; - pw_nvJitLinkDestroy(&b->resource.raw); - delete b; - } - ); + auto box = std::shared_ptr(new NvJitLinkBox{{handle}}, [](NvJitLinkBox *b) { + // Note: nvJitLinkDestroy takes nvJitLinkHandle* and nulls it, + // but we're deleting the box anyway so nulling is harmless. + GILReleaseGuard gil; + pw_nvJitLinkDestroy(&b->resource.raw); + delete b; + }); return NvJitLinkHandle(box, &box->resource); } -NvJitLinkHandle create_nvjitlink_handle_ref(nvJitLink_t handle) { +NvJitLinkHandle +create_nvjitlink_handle_ref(nvJitLink_t handle) +{ auto box = std::make_shared(NvJitLinkBox{{handle}}); return NvJitLinkHandle(box, &box->resource); } @@ -217,21 +229,22 @@ struct CuLinkBox { }; } // namespace -CuLinkHandle create_culink_handle(CUlinkState state) { +CuLinkHandle +create_culink_handle(CUlinkState state) +{ ensure_fn_table(FnTable::driver); - auto box = std::shared_ptr( - new CuLinkBox{state}, - [](CuLinkBox* b) { - // cuLinkDestroy takes CUlinkState by value (not pointer). - GILReleaseGuard gil; - pw_cuLinkDestroy(b->resource); - delete b; - } - ); + auto box = std::shared_ptr(new CuLinkBox{state}, [](CuLinkBox *b) { + // cuLinkDestroy takes CUlinkState by value (not pointer). + GILReleaseGuard gil; + pw_cuLinkDestroy(b->resource); + delete b; + }); return CuLinkHandle(box, &box->resource); } -CuLinkHandle create_culink_handle_ref(CUlinkState state) { +CuLinkHandle +create_culink_handle_ref(CUlinkState state) +{ auto box = std::make_shared(CuLinkBox{state}); return CuLinkHandle(box, &box->resource); } diff --git a/cuda_core/cuda/core/_cpp/rt/py.hpp b/cuda_core/cuda/core/_cpp/rt/py.hpp index 04c50be4a4e..94e64c5b1f6 100644 --- a/cuda_core/cuda/core/_cpp/rt/py.hpp +++ b/cuda_core/cuda/core/_cpp/rt/py.hpp @@ -12,7 +12,8 @@ namespace cuda_core::rt { #if PY_VERSION_HEX < 0x030D0000 -extern "C" int _Py_IsFinalizing(void); +extern "C" int +_Py_IsFinalizing(void); #endif // Best-effort probe for interpreter shutdown. @@ -31,7 +32,9 @@ extern "C" int _Py_IsFinalizing(void); // work once shutdown is underway and accept an intentional leak or skipped // Python conversion in that edge case rather than add more complex deferral // machinery. -inline bool py_is_finalizing() noexcept { +inline bool +py_is_finalizing() noexcept +{ #if PY_VERSION_HEX >= 0x030D0000 return Py_IsFinalizing(); #else @@ -41,8 +44,9 @@ inline bool py_is_finalizing() noexcept { // Conditionally release the GIL while calling into the CUDA driver. class GILReleaseGuard { -public: - GILReleaseGuard() noexcept { + public: + GILReleaseGuard() noexcept + { if (!Py_IsInitialized() || py_is_finalizing()) { return; } @@ -51,24 +55,27 @@ class GILReleaseGuard { } } - ~GILReleaseGuard() { + ~GILReleaseGuard() + { if (tstate_) { PyEval_RestoreThread(tstate_); } } - GILReleaseGuard(const GILReleaseGuard&) = delete; - GILReleaseGuard& operator=(const GILReleaseGuard&) = delete; + GILReleaseGuard(const GILReleaseGuard &) = delete; + GILReleaseGuard & + operator=(const GILReleaseGuard &) = delete; -private: - PyThreadState* tstate_ = nullptr; + private: + PyThreadState *tstate_ = nullptr; }; // Helper to acquire the GIL when we might not hold it. // Use in C++ destructors that need to manipulate Python objects. class GILAcquireGuard { -public: - GILAcquireGuard() : acquired_(false) { + public: + GILAcquireGuard() : acquired_(false) + { // Don't try to acquire GIL if Python is finalizing if (!Py_IsInitialized() || py_is_finalizing()) { return; @@ -77,19 +84,25 @@ class GILAcquireGuard { acquired_ = true; } - ~GILAcquireGuard() { + ~GILAcquireGuard() + { if (acquired_) { PyGILState_Release(gstate_); } } - bool acquired() const { return acquired_; } + bool + acquired() const + { + return acquired_; + } // Non-copyable, non-movable - GILAcquireGuard(const GILAcquireGuard&) = delete; - GILAcquireGuard& operator=(const GILAcquireGuard&) = delete; + GILAcquireGuard(const GILAcquireGuard &) = delete; + GILAcquireGuard & + operator=(const GILAcquireGuard &) = delete; -private: + private: PyGILState_STATE gstate_; bool acquired_; }; @@ -102,8 +115,9 @@ class GILAcquireGuard { // or the caller returns an error with no exception set. Allocates nothing. // Construct with the GIL held and destroy before releasing it. class PendingExceptionGuard { -public: - PendingExceptionGuard() noexcept { + public: + PendingExceptionGuard() noexcept + { #if PY_VERSION_HEX >= 0x030C0000 pending_ = PyErr_GetRaisedException(); #else @@ -111,7 +125,8 @@ class PendingExceptionGuard { #endif } - ~PendingExceptionGuard() { + ~PendingExceptionGuard() + { PyErr_Clear(); #if PY_VERSION_HEX >= 0x030C0000 PyErr_SetRaisedException(pending_); @@ -120,139 +135,194 @@ class PendingExceptionGuard { #endif } - PendingExceptionGuard(const PendingExceptionGuard&) = delete; - PendingExceptionGuard& operator=(const PendingExceptionGuard&) = delete; + PendingExceptionGuard(const PendingExceptionGuard &) = delete; + PendingExceptionGuard & + operator=(const PendingExceptionGuard &) = delete; -private: + private: #if PY_VERSION_HEX >= 0x030C0000 - PyObject* pending_ = nullptr; + PyObject *pending_ = nullptr; #else - PyObject* type_ = nullptr; - PyObject* value_ = nullptr; - PyObject* tb_ = nullptr; + PyObject *type_ = nullptr; + PyObject *value_ = nullptr; + PyObject *tb_ = nullptr; #endif }; // as_py() - convert handle to Python wrapper object (returns new reference) namespace detail { // n.b. class lookup is not cached to avoid deadlock hazard, see DESIGN.md -inline PyObject* make_py(const char* module_name, const char* class_name, std::intptr_t value) noexcept { +inline PyObject * +make_py(const char *module_name, const char *class_name, std::intptr_t value) noexcept +{ if (py_is_finalizing()) { Py_RETURN_NONE; } - PyObject* mod = PyImport_ImportModule(module_name); - if (!mod) return nullptr; - PyObject* cls = PyObject_GetAttrString(mod, class_name); + PyObject *mod = PyImport_ImportModule(module_name); + if (!mod) + return nullptr; + PyObject *cls = PyObject_GetAttrString(mod, class_name); Py_DECREF(mod); - if (!cls) return nullptr; - PyObject* result = PyObject_CallFunction(cls, "L", value); + if (!cls) + return nullptr; + PyObject *result = PyObject_CallFunction(cls, "L", value); Py_DECREF(cls); return result; } } // namespace detail -inline PyObject* as_py(const ContextHandle& h) noexcept { +inline PyObject * +as_py(const ContextHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUcontext", as_intptr(h)); } -inline PyObject* as_py(const GreenCtxHandle& h) noexcept { +inline PyObject * +as_py(const GreenCtxHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUgreenCtx", as_intptr(h)); } -inline PyObject* as_py(const StreamHandle& h) noexcept { +inline PyObject * +as_py(const StreamHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUstream", as_intptr(h)); } -inline PyObject* as_py(const EventHandle& h) noexcept { +inline PyObject * +as_py(const EventHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUevent", as_intptr(h)); } -inline PyObject* as_py(const MemoryPoolHandle& h) noexcept { +inline PyObject * +as_py(const MemoryPoolHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUmemoryPool", as_intptr(h)); } -inline PyObject* as_py(const DevicePtrHandle& h) noexcept { +inline PyObject * +as_py(const DevicePtrHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUdeviceptr", as_intptr(h)); } -inline PyObject* as_py(const LibraryHandle& h) noexcept { +inline PyObject * +as_py(const LibraryHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUlibrary", as_intptr(h)); } -inline PyObject* as_py(const CUmodule& h) noexcept { +inline PyObject * +as_py(const CUmodule &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUmodule", as_intptr(h)); } -inline PyObject* as_py(const KernelHandle& h) noexcept { +inline PyObject * +as_py(const KernelHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUkernel", as_intptr(h)); } -inline PyObject* as_py(const GraphHandle& h) noexcept { +inline PyObject * +as_py(const GraphHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUgraph", as_intptr(h)); } -inline PyObject* as_py(const GraphExecHandle& h) noexcept { +inline PyObject * +as_py(const GraphExecHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUgraphExec", as_intptr(h)); } -inline PyObject* as_py(const GraphNodeHandle& h) noexcept { +inline PyObject * +as_py(const GraphNodeHandle &h) noexcept +{ if (!as_intptr(h)) { Py_RETURN_NONE; } return detail::make_py("cuda.bindings.driver", "CUgraphNode", as_intptr(h)); } -inline PyObject* as_py(const NvrtcProgramHandle& h) noexcept { +inline PyObject * +as_py(const NvrtcProgramHandle &h) noexcept +{ return detail::make_py("cuda.bindings.nvrtc", "nvrtcProgram", as_intptr(h)); } -inline PyObject* as_py(const NvvmProgramHandle& h) noexcept { +inline PyObject * +as_py(const NvvmProgramHandle &h) noexcept +{ // NVVM bindings use raw integers, not wrapper classes return PyLong_FromSsize_t(as_intptr(h)); } -inline PyObject* as_py(const NvJitLinkHandle& h) noexcept { +inline PyObject * +as_py(const NvJitLinkHandle &h) noexcept +{ // nvJitLink bindings use raw integers, not wrapper classes return PyLong_FromSsize_t(as_intptr(h)); } -inline PyObject* as_py(const CuLinkHandle& h) noexcept { +inline PyObject * +as_py(const CuLinkHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUlinkState", as_intptr(h)); } -inline PyObject* as_py(const GraphicsResourceHandle& h) noexcept { +inline PyObject * +as_py(const GraphicsResourceHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUgraphicsResource", as_intptr(h)); } -inline PyObject* as_py(const FileDescriptorHandle& h) noexcept { +inline PyObject * +as_py(const FileDescriptorHandle &h) noexcept +{ return PyLong_FromSsize_t(as_intptr(h)); } -inline PyObject* as_py(const OpaqueArrayHandle& h) noexcept { +inline PyObject * +as_py(const OpaqueArrayHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUarray", as_intptr(h)); } -inline PyObject* as_py(const MipmappedArrayHandle& h) noexcept { +inline PyObject * +as_py(const MipmappedArrayHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUmipmappedArray", as_intptr(h)); } -inline PyObject* as_py(const TexObjectHandle& h) noexcept { +inline PyObject * +as_py(const TexObjectHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUtexObject", as_intptr(h)); } -inline PyObject* as_py(const SurfObjectHandle& h) noexcept { +inline PyObject * +as_py(const SurfObjectHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUsurfObject", as_intptr(h)); } -inline PyObject* as_py(const MemAllocationHandle& h) noexcept { +inline PyObject * +as_py(const MemAllocationHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUmemGenericAllocationHandle", as_intptr(h)); } -inline PyObject* as_py(const VaReservationHandle& h) noexcept { +inline PyObject * +as_py(const VaReservationHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUdeviceptr", as_intptr(h)); } -inline PyObject* as_py(const VaMappingHandle& h) noexcept { +inline PyObject * +as_py(const VaMappingHandle &h) noexcept +{ return detail::make_py("cuda.bindings.driver", "CUdeviceptr", as_intptr(h)); } @@ -262,43 +332,48 @@ inline PyObject* as_py(const VaMappingHandle& h) noexcept { // Register the Python warning category used by report_* (cuda.core.CUDAWarning). // Implemented in py_report.cpp -void register_warning_category(PyObject* category) noexcept; +void +register_warning_category(PyObject *category) noexcept; // Create a non-owning stream handle that prevents a Python owner from being GC'd. // The owner's refcount is incremented; decremented when handle is released. // The owner is responsible for keeping the stream's context alive. // Implemented in stream.cpp -StreamHandle create_stream_handle_with_owner(CUstream stream, PyObject* owner); +StreamHandle +create_stream_handle_with_owner(CUstream stream, PyObject *owner); // Create a non-owning device pointer handle that prevents a Python owner from being GC'd. // The owner's refcount is incremented; decremented when handle is released. // The pointer will NOT be freed when the handle is released. // If owner is nullptr, equivalent to deviceptr_create_ref. // Implemented in memory.cpp -DevicePtrHandle deviceptr_create_with_owner(CUdeviceptr ptr, PyObject* owner); +DevicePtrHandle +deviceptr_create_with_owner(CUdeviceptr ptr, PyObject *owner); // Callback type for MemoryResource deallocation. // Called from the shared_ptr deleter when a handle created via // deviceptr_create_with_mr is destroyed. The implementation is responsible // for converting raw C types to Python objects and calling // mr.deallocate(ptr, size, stream). -using MRDeallocCallback = void (*)(PyObject* mr, CUdeviceptr ptr, - size_t size, const StreamHandle& stream); +using MRDeallocCallback = void (*)(PyObject *mr, CUdeviceptr ptr, size_t size, const StreamHandle &stream); // Register the MR deallocation callback. // Implemented in memory.cpp -void register_mr_dealloc_callback(MRDeallocCallback cb); +void +register_mr_dealloc_callback(MRDeallocCallback cb); // Create a device pointer handle whose destructor calls mr.deallocate() // via the registered callback. The mr's refcount is incremented and // decremented when the handle is released. // If mr is nullptr, equivalent to deviceptr_create_ref. // Implemented in memory.cpp -DevicePtrHandle deviceptr_create_with_mr(CUdeviceptr ptr, size_t size, PyObject* mr); +DevicePtrHandle +deviceptr_create_with_mr(CUdeviceptr ptr, size_t size, PyObject *mr); // Build an OpaqueHandle from a Python object: increments its refcount now and // decrements it (under the GIL) on release. The caller must hold the GIL. // Implemented in graph.cpp -OpaqueHandle make_opaque_py(PyObject* obj); +OpaqueHandle +make_opaque_py(PyObject *obj); } // namespace cuda_core::rt diff --git a/cuda_core/cuda/core/_cpp/rt/py_deferred_cleanup.cpp b/cuda_core/cuda/core/_cpp/rt/py_deferred_cleanup.cpp index 126a90879b1..4a3c75b995e 100644 --- a/cuda_core/cuda/core/_cpp/rt/py_deferred_cleanup.cpp +++ b/cuda_core/cuda/core/_cpp/rt/py_deferred_cleanup.cpp @@ -25,37 +25,46 @@ using namespace detail; namespace { // Process-lifetime MPSC queue that drains payloads from Python's main thread. class DeferredCleanupQueue { -public: + public: // Transfer one preallocated cleanup item from a producer to the queue. - void enqueue(DeferredCleanupItem* item) noexcept { - DeferredCleanupItem* head = head_.load(std::memory_order_relaxed); + void + enqueue(DeferredCleanupItem *item) noexcept + { + DeferredCleanupItem *head = head_.load(std::memory_order_relaxed); do { item->next = head; - } while (!head_.compare_exchange_weak( - head, item, std::memory_order_release, std::memory_order_relaxed)); + } while (!head_.compare_exchange_weak(head, item, std::memory_order_release, std::memory_order_relaxed)); schedule(); } // Permanently disable pending-call scheduling during interpreter shutdown. - void stop() noexcept { + void + stop() noexcept + { accepting_.store(false, std::memory_order_release); } // Reattempt scheduling after Py_AddPendingCall() found CPython's bounded // pending-call queue full and left payloads queued for a later safe entry. - void retry_schedule() noexcept { + void + retry_schedule() noexcept + { schedule(); } -private: + private: // Adapt queue draining to CPython's int (*)(void*) callback ABI. - static int pending_call(void* arg) noexcept { - static_cast(arg)->drain(); + static int + pending_call(void *arg) noexcept + { + static_cast(arg)->drain(); return 0; } // Coalesce all queued work behind at most one CPython pending call. - void schedule() noexcept { + void + schedule() noexcept + { if (!accepting_.load(std::memory_order_acquire)) { return; } @@ -67,9 +76,7 @@ class DeferredCleanupQueue { return; } bool expected = false; - if (!scheduled_.compare_exchange_strong( - expected, true, std::memory_order_acq_rel, - std::memory_order_relaxed)) { + if (!scheduled_.compare_exchange_strong(expected, true, std::memory_order_acq_rel, std::memory_order_relaxed)) { return; } if (Py_AddPendingCall(&DeferredCleanupQueue::pending_call, this) != 0) { @@ -80,17 +87,18 @@ class DeferredCleanupQueue { } // Detach and destroy all queued payloads from Python's main thread. - void drain() noexcept { + void + drain() noexcept + { if (!Py_IsInitialized() || py_is_finalizing()) { stop(); scheduled_.store(false, std::memory_order_release); return; // Intentionally leak intact payloads during shutdown. } - while (DeferredCleanupItem* list = - head_.exchange(nullptr, std::memory_order_acquire)) { + while (DeferredCleanupItem *list = head_.exchange(nullptr, std::memory_order_acquire)) { while (list) { - DeferredCleanupItem* next = list->next; + DeferredCleanupItem *next = list->next; delete list; list = next; } @@ -103,7 +111,7 @@ class DeferredCleanupQueue { } // Head of the intrusive multi-producer, single-consumer payload stack. - std::atomic head_{nullptr}; + std::atomic head_{nullptr}; // True while one cuda-core drain callback is pending or executing. std::atomic scheduled_{false}; // False once shutdown begins, causing later payloads to be leaked safely. @@ -111,13 +119,14 @@ class DeferredCleanupQueue { }; // Published once at module initialization and intentionally never freed. -std::atomic deferred_cleanup_queue{nullptr}; +std::atomic deferred_cleanup_queue{nullptr}; } // namespace namespace detail { -void ensure_deferred_cleanup_ready() { - DeferredCleanupQueue* queue = - deferred_cleanup_queue.load(std::memory_order_acquire); +void +ensure_deferred_cleanup_ready() +{ + DeferredCleanupQueue *queue = deferred_cleanup_queue.load(std::memory_order_acquire); if (!queue) { throw std::runtime_error("deferred cleanup is not initialized"); } @@ -125,10 +134,11 @@ void ensure_deferred_cleanup_ready() { } // CUDA's CUhostFn ABI is void (*)(void*); recover and enqueue the cleanup item. -void enqueue_cleanup(void* item) noexcept { - auto* cleanup = static_cast(item); - if (DeferredCleanupQueue* queue = - deferred_cleanup_queue.load(std::memory_order_acquire)) { +void +enqueue_cleanup(void *item) noexcept +{ + auto *cleanup = static_cast(item); + if (DeferredCleanupQueue *queue = deferred_cleanup_queue.load(std::memory_order_acquire)) { queue->enqueue(cleanup); } } @@ -136,20 +146,23 @@ void enqueue_cleanup(void* item) noexcept { // Module initialization calls this once with the GIL held, which serializes // the check, allocation, and publication below. -void initialize_deferred_cleanup() { +void +initialize_deferred_cleanup() +{ if (deferred_cleanup_queue.load(std::memory_order_acquire)) { return; } - auto* queue = new DeferredCleanupQueue(); + auto *queue = new DeferredCleanupQueue(); deferred_cleanup_queue.store(queue, std::memory_order_release); } -void retry_deferred_cleanup() noexcept { +void +retry_deferred_cleanup() noexcept +{ if (!Py_IsInitialized() || py_is_finalizing()) { return; } - if (DeferredCleanupQueue* queue = - deferred_cleanup_queue.load(std::memory_order_acquire)) { + if (DeferredCleanupQueue *queue = deferred_cleanup_queue.load(std::memory_order_acquire)) { queue->retry_schedule(); } } diff --git a/cuda_core/cuda/core/_cpp/rt/py_driver_fns.cpp b/cuda_core/cuda/core/_cpp/rt/py_driver_fns.cpp index bb969e289f5..fd36385def8 100644 --- a/cuda_core/cuda/core/_cpp/rt/py_driver_fns.cpp +++ b/cuda_core/cuda/core/_cpp/rt/py_driver_fns.cpp @@ -48,24 +48,40 @@ std::atomic unavailable_reported[kTables]; std::mutex fill_mutex; char fill_error[kTables][512] = {}; // guarded by fill_mutex -std::size_t index_of(FnTable table) noexcept { return static_cast(table); } +std::size_t +index_of(FnTable table) noexcept +{ + return static_cast(table); +} -const char* module_name(FnTable table) noexcept { +const char * +module_name(FnTable table) noexcept +{ switch (table) { - case FnTable::driver: return "cuda.bindings._internal.driver"; - case FnTable::nvrtc: return "cuda.bindings._internal.nvrtc"; - case FnTable::nvvm: return "cuda.bindings._internal.nvvm"; - case FnTable::nvjitlink: return "cuda.bindings._internal.nvjitlink"; + case FnTable::driver: + return "cuda.bindings._internal.driver"; + case FnTable::nvrtc: + return "cuda.bindings._internal.nvrtc"; + case FnTable::nvvm: + return "cuda.bindings._internal.nvvm"; + case FnTable::nvjitlink: + return "cuda.bindings._internal.nvjitlink"; } return "cuda.bindings._internal"; } -const char* library_name(FnTable table) noexcept { +const char * +library_name(FnTable table) noexcept +{ switch (table) { - case FnTable::driver: return "CUDA driver"; - case FnTable::nvrtc: return "NVRTC"; - case FnTable::nvvm: return "NVVM"; - case FnTable::nvjitlink: return "nvJitLink"; + case FnTable::driver: + return "CUDA driver"; + case FnTable::nvrtc: + return "NVRTC"; + case FnTable::nvvm: + return "NVVM"; + case FnTable::nvjitlink: + return "nvJitLink"; } return "library"; } @@ -82,23 +98,25 @@ const char* library_name(FnTable table) noexcept { // PyErr_SetInterrupt() once its own Python work (the warning) is done. Python // then raises KeyboardInterrupt at the next bytecode boundary, so the user // sees the interrupt, not the CUDAError from the failed call. -bool take_python_error(char* buf, std::size_t size, bool* interrupted) noexcept { - const bool transient = PyErr_Occurred() - && (!PyErr_ExceptionMatches(PyExc_Exception) || PyErr_ExceptionMatches(PyExc_MemoryError) - || PyErr_ExceptionMatches(PyExc_RecursionError)); +bool +take_python_error(char *buf, std::size_t size, bool *interrupted) noexcept +{ + const bool transient = + PyErr_Occurred() && (!PyErr_ExceptionMatches(PyExc_Exception) || PyErr_ExceptionMatches(PyExc_MemoryError) || + PyErr_ExceptionMatches(PyExc_RecursionError)); *interrupted = PyErr_Occurred() && PyErr_ExceptionMatches(PyExc_KeyboardInterrupt); #if PY_VERSION_HEX >= 0x030C0000 - PyObject* exc = PyErr_GetRaisedException(); + PyObject *exc = PyErr_GetRaisedException(); #else PyObject *type, *value, *traceback; PyErr_Fetch(&type, &value, &traceback); PyErr_NormalizeException(&type, &value, &traceback); - PyObject* exc = value; + PyObject *exc = value; Py_XDECREF(type); Py_XDECREF(traceback); #endif - PyObject* text = exc ? PyObject_Str(exc) : nullptr; - const char* utf8 = text ? PyUnicode_AsUTF8(text) : nullptr; + PyObject *text = exc ? PyObject_Str(exc) : nullptr; + const char *utf8 = text ? PyUnicode_AsUTF8(text) : nullptr; if (utf8 == nullptr || *utf8 == '\0') { utf8 = exc ? Py_TYPE(exc)->tp_name : "unknown error"; // str(KeyboardInterrupt()) is empty } @@ -109,7 +127,9 @@ bool take_python_error(char* buf, std::size_t size, bool* interrupted) noexcept return transient; } -void record_failure(FnTable table, const char* message, bool latch = true) noexcept { +void +record_failure(FnTable table, const char *message, bool latch = true) noexcept +{ { std::lock_guard lock(fill_mutex); std::snprintf(fill_error[index_of(table)], sizeof(fill_error[0]), "%s", message); @@ -122,13 +142,17 @@ void record_failure(FnTable table, const char* message, bool latch = true) noexc } // namespace -bool fn_table_ready(FnTable table) noexcept { +bool +fn_table_ready(FnTable table) noexcept +{ return table_ready[index_of(table)].load(std::memory_order_acquire); } -bool fn_table_error(FnTable table, char* buffer, std::size_t size) noexcept { +bool +fn_table_error(FnTable table, char *buffer, std::size_t size) noexcept +{ std::lock_guard lock(fill_mutex); - const char* text = fill_error[index_of(table)]; + const char *text = fill_error[index_of(table)]; if (!text[0]) { return false; } @@ -136,7 +160,9 @@ bool fn_table_error(FnTable table, char* buffer, std::size_t size) noexcept { return true; } -bool ensure_fn_table(FnTable table) noexcept { +bool +ensure_fn_table(FnTable table) noexcept +{ const std::size_t idx = index_of(table); if (table_ready[idx].load(std::memory_order_acquire)) { return true; @@ -145,7 +171,7 @@ bool ensure_fn_table(FnTable table) noexcept { return false; // latched: the fill reported its reason when it failed } std::size_t count = 0; - const FnEntry* entries = fn_table_entries(table, &count); + const FnEntry *entries = fn_table_entries(table, &count); if (entries == nullptr || count > kMaxEntries) { record_failure(table, "internal cuda.core error, please report: function table has an unexpected size"); return false; @@ -157,7 +183,7 @@ bool ensure_fn_table(FnTable table) noexcept { char message[640]; char cause[384]; - void* values[kMaxEntries] = {}; + void *values[kMaxEntries] = {}; GILAcquireGuard gil; if (!gil.acquired()) { @@ -167,18 +193,23 @@ bool ensure_fn_table(FnTable table) noexcept { PendingExceptionGuard pending; bool interrupted = false; - PyObject* module = PyImport_ImportModule(module_name(table)); + PyObject *module = PyImport_ImportModule(module_name(table)); if (module == nullptr) { const bool transient = take_python_error(cause, sizeof(cause), &interrupted); - std::snprintf(message, sizeof(message), - "cuda.core cannot import %s from the installed cuda-bindings: %s", module_name(table), cause); + std::snprintf( + message, + sizeof(message), + "cuda.core cannot import %s from the installed cuda-bindings: %s", + module_name(table), + cause + ); record_failure(table, message, !transient); if (interrupted) { PyErr_SetInterrupt(); // after the warning: the handler must not consume the interrupt } return false; } - PyObject* pointers = PyObject_CallMethod(module, "_inspect_function_pointers", nullptr); + PyObject *pointers = PyObject_CallMethod(module, "_inspect_function_pointers", nullptr); Py_DECREF(module); if (pointers == nullptr) { const bool transient = take_python_error(cause, sizeof(cause), &interrupted); @@ -191,31 +222,43 @@ bool ensure_fn_table(FnTable table) noexcept { } if (!PyDict_Check(pointers)) { Py_DECREF(pointers); - std::snprintf(message, sizeof(message), - "internal cuda.core error, please report: %s._inspect_function_pointers() did not return a dict", - module_name(table)); + std::snprintf( + message, + sizeof(message), + "internal cuda.core error, please report: %s._inspect_function_pointers() did not return a dict", + module_name(table) + ); record_failure(table, message); return false; } for (std::size_t i = 0; i < count; ++i) { - PyObject* item = PyDict_GetItemString(pointers, entries[i].key); // borrowed; no error on a miss + PyObject *item = PyDict_GetItemString(pointers, entries[i].key); // borrowed; no error on a miss if (item == nullptr) { Py_DECREF(pointers); - std::snprintf(message, sizeof(message), - "the installed cuda-bindings has no entry for %s (%s). This cuda.core build used a " - "cuda.h that names this symbol differently from the installed cuda-bindings. Install the " - "cuda-bindings this cuda.core requires", - entries[i].name, entries[i].key); + std::snprintf( + message, + sizeof(message), + "the installed cuda-bindings has no entry for %s (%s). This cuda.core build used a " + "cuda.h that names this symbol differently from the installed cuda-bindings. Install the " + "cuda-bindings this cuda.core requires", + entries[i].name, + entries[i].key + ); record_failure(table, message); return false; } - void* address = PyLong_AsVoidPtr(item); + void *address = PyLong_AsVoidPtr(item); if (address == nullptr && PyErr_Occurred()) { Py_DECREF(pointers); take_python_error(cause, sizeof(cause), &interrupted); - std::snprintf(message, sizeof(message), - "internal cuda.core error, please report: the entry for %s in %s is not an address: %s", - entries[i].key, module_name(table), cause); + std::snprintf( + message, + sizeof(message), + "internal cuda.core error, please report: the entry for %s in %s is not an address: %s", + entries[i].key, + module_name(table), + cause + ); record_failure(table, message); return false; } @@ -230,11 +273,16 @@ bool ensure_fn_table(FnTable table) noexcept { if (table == FnTable::driver) { for (std::size_t i = 0; i < count; ++i) { if (values[i] == nullptr && entries[i].introduced <= CUDA_CORE_BUILD_MAJOR * 1000) { - std::snprintf(message, sizeof(message), - "the installed CUDA driver lacks %s, which every CUDA %d driver provides. " - "CUDA %d.%d introduced this function. This cuda.core build needs a newer driver", - entries[i].name, CUDA_CORE_BUILD_MAJOR, entries[i].introduced / 1000, - entries[i].introduced / 10 % 100); + std::snprintf( + message, + sizeof(message), + "the installed CUDA driver lacks %s, which every CUDA %d driver provides. " + "CUDA %d.%d introduced this function. This cuda.core build needs a newer driver", + entries[i].name, + CUDA_CORE_BUILD_MAJOR, + entries[i].introduced / 1000, + entries[i].introduced / 10 % 100 + ); record_failure(table, message); return false; } @@ -254,17 +302,24 @@ bool ensure_fn_table(FnTable table) noexcept { return true; } -void report_unavailable_fn(FnTable table, const char* name) noexcept { +void +report_unavailable_fn(FnTable table, const char *name) noexcept +{ char reason[sizeof(fill_error[0])]; const bool failed_fill = fn_table_error(table, reason, sizeof(reason)); char message[768]; if (failed_fill) { std::snprintf(message, sizeof(message), "cuda.core could not call %s: %s", name, reason); - } else { - std::snprintf(message, sizeof(message), - "internal cuda.core error, please report: cuda.core called %s but the installed %s does not " - "provide it. A feature gate is missing or wrong. The call returned an error instead.", - name, library_name(table)); + } + else { + std::snprintf( + message, + sizeof(message), + "internal cuda.core error, please report: cuda.core called %s but the installed %s does not " + "provide it. A feature gate is missing or wrong. The call returned an error instead.", + name, + library_name(table) + ); } if (table == FnTable::driver) { // The trampoline returns CUDA_ERROR_NOT_INITIALIZED. The Cython error diff --git a/cuda_core/cuda/core/_cpp/rt/py_report.cpp b/cuda_core/cuda/core/_cpp/rt/py_report.cpp index 626f19b2cd7..6bc24bf3a5c 100644 --- a/cuda_core/cuda/core/_cpp/rt/py_report.cpp +++ b/cuda_core/cuda/core/_cpp/rt/py_report.cpp @@ -14,7 +14,7 @@ using namespace detail; namespace { // Warning category registered by _utils/cuda_utils.pyx (cuda.core.CUDAWarning). -std::atomic warning_category{nullptr}; +std::atomic warning_category{nullptr}; } // namespace // Report a message that could not be raised. Emits cuda.core.CUDAWarning via @@ -28,8 +28,10 @@ std::atomic warning_category{nullptr}; // table fill re-armed; see py_driver_fns.cpp) can fire inside it. The warning // is then emitted again and the interrupt re-armed with PyErr_SetInterrupt(), // so Python raises it at the next bytecode boundary in the caller's code. -void report_message(const char* message) noexcept { - PyObject* category = warning_category.load(std::memory_order_acquire); +void +report_message(const char *message) noexcept +{ + PyObject *category = warning_category.load(std::memory_order_acquire); if (category && Py_IsInitialized() && !py_is_finalizing()) { GILAcquireGuard gil; if (gil.acquired()) { @@ -42,7 +44,7 @@ void report_message(const char* message) noexcept { PyErr_Clear(); } if (!interrupted || PyErr_WarnEx(category, message, 1) != 0) { - PyObject* subject = PyUnicode_FromString(message); + PyObject *subject = PyUnicode_FromString(message); PyErr_WriteUnraisable(subject); Py_XDECREF(subject); } @@ -57,14 +59,18 @@ void report_message(const char* message) noexcept { std::fflush(stderr); } -void register_warning_category(PyObject* category) noexcept { +void +register_warning_category(PyObject *category) noexcept +{ warning_category.store(category, std::memory_order_release); } namespace { // Attach `message` as a PEP 678 note to the exception currently being handled. // Returns false when there is none or the interpreter cannot be used. -bool add_note_to_handled_exception(const char* message) noexcept { +bool +add_note_to_handled_exception(const char *message) noexcept +{ #if PY_VERSION_HEX >= 0x030B0000 if (!Py_IsInitialized() || py_is_finalizing()) { return false; @@ -73,11 +79,11 @@ bool add_note_to_handled_exception(const char* message) noexcept { if (!gil.acquired()) { return false; } - PyObject* exc = PyErr_GetHandledException(); + PyObject *exc = PyErr_GetHandledException(); if (!exc) { return false; } - PyObject* result = PyObject_CallMethod(exc, "add_note", "s", message); + PyObject *result = PyObject_CallMethod(exc, "add_note", "s", message); Py_DECREF(exc); if (!result) { PyErr_Clear(); @@ -92,7 +98,9 @@ bool add_note_to_handled_exception(const char* message) noexcept { } } // namespace -void attach_rollback_failure(const char* operation, CUresult status, const char* detail) noexcept { +void +attach_rollback_failure(const char *operation, CUresult status, const char *detail) noexcept +{ if (status == CUDA_SUCCESS || status == CUDA_ERROR_DEINITIALIZED) { return; } diff --git a/cuda_core/cuda/core/_cpp/rt/stream.cpp b/cuda_core/cuda/core/_cpp/rt/stream.cpp index 0497c0fabb5..9ef79d40212 100644 --- a/cuda_core/cuda/core/_cpp/rt/stream.cpp +++ b/cuda_core/cuda/core/_cpp/rt/stream.cpp @@ -19,7 +19,9 @@ using namespace detail; namespace { // Return whether a stream handle needs a current context to resolve it. -bool is_default_stream(CUstream stream) noexcept { +bool +is_default_stream(CUstream stream) noexcept +{ return stream == nullptr || stream == CU_STREAM_LEGACY || stream == CU_STREAM_PER_THREAD; } } // namespace @@ -28,17 +30,19 @@ namespace detail { // Return the context a deallocation-stream token must run under. Real streams // resolve their own context; default-stream tokens use the context bound at // allocation time. Warn when PTDS deallocation crosses host threads. -ContextHandle deallocation_context(const DeallocationStream& stream) noexcept { +ContextHandle +deallocation_context(const DeallocationStream &stream) noexcept +{ if (!is_default_stream(as_cu(stream.h_stream))) { return {}; } - if (stream.ptds_tid != std::thread::id{} - && stream.ptds_tid != std::this_thread::get_id()) { + if (stream.ptds_tid != std::thread::id{} && stream.ptds_tid != std::this_thread::get_id()) { report_message( "Buffer deallocation for a per-thread default stream " "is running on a different host thread than the one that recorded " "the deallocation stream; ordering relative to the allocating " - "thread's PTDS is not preserved"); + "thread's PTDS is not preserved" + ); } return get_stream_context(stream.h_stream); } @@ -54,66 +58,72 @@ struct StreamBox { ContextHandle h_context; }; -static const StreamBox* get_box(const StreamHandle& h) noexcept { - const CUstream* p = h.get(); - return reinterpret_cast( - reinterpret_cast(p) - offsetof(StreamBox, resource) - ); +static const StreamBox * +get_box(const StreamHandle &h) noexcept +{ + const CUstream *p = h.get(); + return reinterpret_cast(reinterpret_cast(p) - offsetof(StreamBox, resource)); } // See REGISTRY_DESIGN.md (Level 1: Driver Handle -> Resource Handle) static HandleRegistry stream_registry; } // namespace -StreamHandle create_stream_handle(const ContextHandle& h_ctx, unsigned int flags, int priority) { +StreamHandle +create_stream_handle(const ContextHandle &h_ctx, unsigned int flags, int priority) +{ GILReleaseGuard gil; CUstream stream = nullptr; GreenCtxHandle h_green = get_context_green_ctx(h_ctx); if (h_green) { // Cython gates this on driver >= 12.5, which introduced cuGreenCtxStreamCreate. err = DRIVER_CALL(cuGreenCtxStreamCreate, &stream, as_cu(h_green), flags, priority); - } else { + } + else { err = invoke_in_context_or_undo( h_ctx, - [&]() noexcept { return DRIVER_CALL(cuStreamCreateWithPriority, &stream, flags, priority); }, - [&]() noexcept { pw_cuStreamDestroy(stream); }, - /*undo_requires_target_context=*/false); + [&]() noexcept { + return DRIVER_CALL(cuStreamCreateWithPriority, &stream, flags, priority); + }, + [&]() noexcept { + pw_cuStreamDestroy(stream); + }, + /*undo_requires_target_context=*/false + ); } if (err != CUDA_SUCCESS) { return {}; } - auto box = std::shared_ptr( - new StreamBox{stream, h_ctx}, - [](const StreamBox* b) { - stream_registry.unregister_handle(b->resource); - GILReleaseGuard gil; - pw_cuStreamDestroy(b->resource); - delete b; - } - ); + auto box = std::shared_ptr(new StreamBox{stream, h_ctx}, [](const StreamBox *b) { + stream_registry.unregister_handle(b->resource); + GILReleaseGuard gil; + pw_cuStreamDestroy(b->resource); + delete b; + }); StreamHandle h(box, &box->resource); stream_registry.register_handle(stream, h); return h; } -StreamHandle create_stream_handle_ref(CUstream stream) { +StreamHandle +create_stream_handle_ref(CUstream stream) +{ if (auto h = stream_registry.lookup(stream)) { return h; } - auto box = std::shared_ptr( - new StreamBox{stream, {}}, - [](const StreamBox* b) { - stream_registry.unregister_handle(b->resource); - delete b; - } - ); + auto box = std::shared_ptr(new StreamBox{stream, {}}, [](const StreamBox *b) { + stream_registry.unregister_handle(b->resource); + delete b; + }); StreamHandle h(box, &box->resource); stream_registry.register_handle(stream, h); return h; } -StreamHandle create_stream_handle_with_owner(CUstream stream, PyObject* owner) { +StreamHandle +create_stream_handle_with_owner(CUstream stream, PyObject *owner) +{ if (auto h = stream_registry.lookup(stream)) { // Reuse handles that already carry structural context metadata, e.g. // cuda-core-owned streams. @@ -134,35 +144,40 @@ StreamHandle create_stream_handle_with_owner(CUstream stream, PyObject* owner) { // Owner-backed handles are NOT registered in the stream registry to avoid // corruption when multiple owners wrap the same CUstream (each stacks its // own Py_INCREF/Py_DECREF independently). - auto box = std::shared_ptr( - new StreamBox{stream, {}}, - [owner](const StreamBox* b) { - GILAcquireGuard gil; - if (gil.acquired()) { - Py_DECREF(owner); - } - delete b; + auto box = std::shared_ptr(new StreamBox{stream, {}}, [owner](const StreamBox *b) { + GILAcquireGuard gil; + if (gil.acquired()) { + Py_DECREF(owner); } - ); + delete b; + }); return StreamHandle(box, &box->resource); } // Return the context retained by a stream handle. -ContextHandle get_stream_context(const StreamHandle& h) noexcept { +ContextHandle +get_stream_context(const StreamHandle &h) noexcept +{ return h ? get_box(h)->h_context : ContextHandle{}; } -StreamHandle get_legacy_stream() { +StreamHandle +get_legacy_stream() +{ static StreamHandle handle = create_stream_handle_ref(CU_STREAM_LEGACY); return handle; } -StreamHandle get_per_thread_stream() { +StreamHandle +get_per_thread_stream() +{ static StreamHandle handle = create_stream_handle_ref(CU_STREAM_PER_THREAD); return handle; } -StreamHandle create_context_bound_legacy_stream(const ContextHandle& h_context) { +StreamHandle +create_context_bound_legacy_stream(const ContextHandle &h_context) +{ if (!h_context) { return {}; } @@ -190,8 +205,9 @@ namespace detail { // Real streams are copied unchanged. Default-stream tokens without an embedded // context are bound to the current context. Returns false (and sets err) when a // default-stream token cannot be bound because no context is current. -bool make_deallocation_stream( - const StreamHandle& h, DeallocationStream& out) noexcept { +bool +make_deallocation_stream(const StreamHandle &h, DeallocationStream &out) noexcept +{ out = {}; if (!h) { return true; @@ -214,8 +230,7 @@ bool make_deallocation_stream( } // Do not register in stream_registry: the token value alone is not // a unique stream identity (context is part of the meaning). - auto box = std::shared_ptr( - new StreamBox{stream, h_ctx}); + auto box = std::shared_ptr(new StreamBox{stream, h_ctx}); h_bound = StreamHandle(box, &box->resource); } diff --git a/cuda_core/cuda/core/_cpp/rt/texture.cpp b/cuda_core/cuda/core/_cpp/rt/texture.cpp index 309e34aa4c9..bad8be58fad 100644 --- a/cuda_core/cuda/core/_cpp/rt/texture.cpp +++ b/cuda_core/cuda/core/_cpp/rt/texture.cpp @@ -26,16 +26,16 @@ struct GraphicsResourceBox { }; } // namespace -GraphicsResourceHandle create_graphics_resource_handle(CUgraphicsResource resource) { +GraphicsResourceHandle +create_graphics_resource_handle(CUgraphicsResource resource) +{ ensure_fn_table(FnTable::driver); // the deleter calls the driver: resolve the table before it can run - auto box = std::shared_ptr( - new GraphicsResourceBox{resource}, - [](const GraphicsResourceBox* b) { + auto box = + std::shared_ptr(new GraphicsResourceBox{resource}, [](const GraphicsResourceBox *b) { GILReleaseGuard gil; pw_cuGraphicsUnregisterResource(b->resource); delete b; - } - ); + }); return GraphicsResourceHandle(box, &box->resource); } @@ -79,50 +79,61 @@ struct SurfObjectBox { }; // Recover an array's owning box from its aliased resource pointer. -const ArrayBox* get_box(const OpaqueArrayHandle& h) noexcept { - const CUarray* p = h.get(); - return reinterpret_cast( - reinterpret_cast(p) - offsetof(ArrayBox, resource)); +const ArrayBox * +get_box(const OpaqueArrayHandle &h) noexcept +{ + const CUarray *p = h.get(); + return reinterpret_cast(reinterpret_cast(p) - offsetof(ArrayBox, resource)); } // Recover a mipmapped array's owning box from its aliased resource pointer. -const MipmappedArrayBox* get_box(const MipmappedArrayHandle& h) noexcept { - const CUmipmappedArray* p = h.get(); - return reinterpret_cast( - reinterpret_cast(p) - - offsetof(MipmappedArrayBox, resource)); +const MipmappedArrayBox * +get_box(const MipmappedArrayHandle &h) noexcept +{ + const CUmipmappedArray *p = h.get(); + return reinterpret_cast( + reinterpret_cast(p) - offsetof(MipmappedArrayBox, resource) + ); } // Wrap an array with shared owning-destruction behavior. -static OpaqueArrayHandle wrap_array_owned(CUarray arr, ContextHandle h_context) { - auto box = std::shared_ptr( - new ArrayBox{arr, {}, std::move(h_context)}, - [](const ArrayBox* b) { - GILReleaseGuard gil; - pw_cuArrayDestroy(b->resource); - delete b; - } - ); +static OpaqueArrayHandle +wrap_array_owned(CUarray arr, ContextHandle h_context) +{ + auto box = std::shared_ptr(new ArrayBox{arr, {}, std::move(h_context)}, [](const ArrayBox *b) { + GILReleaseGuard gil; + pw_cuArrayDestroy(b->resource); + delete b; + }); return OpaqueArrayHandle(box, &box->resource); } } // namespace -OpaqueArrayHandle create_array_handle(const ContextHandle& h_context, const CUDA_ARRAY3D_DESCRIPTOR& desc) { +OpaqueArrayHandle +create_array_handle(const ContextHandle &h_context, const CUDA_ARRAY3D_DESCRIPTOR &desc) +{ GILReleaseGuard gil; CUarray arr = nullptr; err = invoke_in_context_or_undo( h_context, - [&]() noexcept { return DRIVER_CALL(cuArray3DCreate, &arr, &desc); }, - [&]() noexcept { pw_cuArrayDestroy(arr); }, - /*undo_requires_target_context=*/false); + [&]() noexcept { + return DRIVER_CALL(cuArray3DCreate, &arr, &desc); + }, + [&]() noexcept { + pw_cuArrayDestroy(arr); + }, + /*undo_requires_target_context=*/false + ); if (err != CUDA_SUCCESS) { return {}; } return wrap_array_owned(arr, h_context); } -OpaqueArrayHandle create_array_handle_ref(CUarray arr) { +OpaqueArrayHandle +create_array_handle_ref(CUarray arr) +{ if (!arr) { return {}; } @@ -130,7 +141,9 @@ OpaqueArrayHandle create_array_handle_ref(CUarray arr) { return OpaqueArrayHandle(box, &box->resource); } -OpaqueArrayHandle create_array_handle_owning(CUarray arr) { +OpaqueArrayHandle +create_array_handle_owning(CUarray arr) +{ ensure_fn_table(FnTable::driver); // the deleter calls the driver: resolve the table before it can run if (!arr) { return {}; @@ -139,11 +152,15 @@ OpaqueArrayHandle create_array_handle_owning(CUarray arr) { } // Return the context retained by an array handle. -ContextHandle get_array_context(const OpaqueArrayHandle& h) noexcept { +ContextHandle +get_array_context(const OpaqueArrayHandle &h) noexcept +{ return h ? get_box(h)->h_context : ContextHandle{}; } -OpaqueArrayHandle create_array_level_handle(const MipmappedArrayHandle& h_mip, unsigned int level) { +OpaqueArrayHandle +create_array_level_handle(const MipmappedArrayHandle &h_mip, unsigned int level) +{ GILReleaseGuard gil; CUarray arr; ContextHandle h_context = h_mip ? get_box(h_mip)->h_context : ContextHandle{}; @@ -152,60 +169,77 @@ OpaqueArrayHandle create_array_level_handle(const MipmappedArrayHandle& h_mip, u } // Non-owning level view: storage belongs to the mipmap. Embed the mipmap // handle so the parent outlives this level; the deleter does not destroy. - auto box = std::shared_ptr( - new ArrayBox{arr, h_mip, h_context}, - [](const ArrayBox* b) { delete b; } - ); + auto box = std::shared_ptr(new ArrayBox{arr, h_mip, h_context}, [](const ArrayBox *b) { + delete b; + }); return OpaqueArrayHandle(box, &box->resource); } -MipmappedArrayHandle create_mipmapped_array_handle(const ContextHandle& h_context, - const CUDA_ARRAY3D_DESCRIPTOR& desc, - unsigned int num_levels) { +MipmappedArrayHandle +create_mipmapped_array_handle( + const ContextHandle &h_context, + const CUDA_ARRAY3D_DESCRIPTOR &desc, + unsigned int num_levels +) +{ GILReleaseGuard gil; CUmipmappedArray mip = nullptr; err = invoke_in_context_or_undo( h_context, - [&]() noexcept { return DRIVER_CALL(cuMipmappedArrayCreate, &mip, &desc, num_levels); }, - [&]() noexcept { pw_cuMipmappedArrayDestroy(mip); }, - /*undo_requires_target_context=*/false); + [&]() noexcept { + return DRIVER_CALL(cuMipmappedArrayCreate, &mip, &desc, num_levels); + }, + [&]() noexcept { + pw_cuMipmappedArrayDestroy(mip); + }, + /*undo_requires_target_context=*/false + ); if (err != CUDA_SUCCESS) { return {}; } - auto box = std::shared_ptr( - new MipmappedArrayBox{mip, h_context}, - [](const MipmappedArrayBox* b) { + auto box = + std::shared_ptr(new MipmappedArrayBox{mip, h_context}, [](const MipmappedArrayBox *b) { GILReleaseGuard gil; pw_cuMipmappedArrayDestroy(b->resource); delete b; - } - ); + }); return MipmappedArrayHandle(box, &box->resource); } // Return the context retained by a mipmapped array handle. -ContextHandle get_mipmapped_array_context(const MipmappedArrayHandle& h) noexcept { +ContextHandle +get_mipmapped_array_context(const MipmappedArrayHandle &h) noexcept +{ return h ? get_box(h)->h_context : ContextHandle{}; } namespace { -TexObjectHandle make_tex_object_handle(const CUDA_RESOURCE_DESC& res, - const CUDA_TEXTURE_DESC& tex, - std::shared_ptr h_backing, - const ContextHandle& h_context) { +TexObjectHandle +make_tex_object_handle( + const CUDA_RESOURCE_DESC &res, + const CUDA_TEXTURE_DESC &tex, + std::shared_ptr h_backing, + const ContextHandle &h_context +) +{ GILReleaseGuard gil; CUtexObject obj = 0; err = invoke_in_context_or_undo( h_context, - [&]() noexcept { return DRIVER_CALL(cuTexObjectCreate, &obj, &res, &tex, nullptr); }, - [&]() noexcept { pw_cuTexObjectDestroy(obj); }, - /*undo_requires_target_context=*/true); + [&]() noexcept { + return DRIVER_CALL(cuTexObjectCreate, &obj, &res, &tex, nullptr); + }, + [&]() noexcept { + pw_cuTexObjectDestroy(obj); + }, + /*undo_requires_target_context=*/true + ); if (err != CUDA_SUCCESS) { return {}; } auto box = std::shared_ptr( new TexObjectBox{TexObjectValue{obj}, std::move(h_backing), h_context}, - [](const TexObjectBox* b) { + [](const TexObjectBox *b) { GILReleaseGuard gil; cleanup_in_context(b->h_context, "cuTexObjectDestroy", handle_bits(b->resource.raw), [&]() noexcept { return DRIVER_CALL(cuTexObjectDestroy, b->resource.raw); @@ -217,43 +251,64 @@ TexObjectHandle make_tex_object_handle(const CUDA_RESOURCE_DESC& res, } } // namespace -TexObjectHandle create_tex_object_handle_array(const ContextHandle& h_context, - const CUDA_RESOURCE_DESC& res, - const CUDA_TEXTURE_DESC& tex, - const OpaqueArrayHandle& h_backing) { +TexObjectHandle +create_tex_object_handle_array( + const ContextHandle &h_context, + const CUDA_RESOURCE_DESC &res, + const CUDA_TEXTURE_DESC &tex, + const OpaqueArrayHandle &h_backing +) +{ return make_tex_object_handle(res, tex, h_backing, h_context); } -TexObjectHandle create_tex_object_handle_mipmap(const ContextHandle& h_context, - const CUDA_RESOURCE_DESC& res, - const CUDA_TEXTURE_DESC& tex, - const MipmappedArrayHandle& h_backing) { +TexObjectHandle +create_tex_object_handle_mipmap( + const ContextHandle &h_context, + const CUDA_RESOURCE_DESC &res, + const CUDA_TEXTURE_DESC &tex, + const MipmappedArrayHandle &h_backing +) +{ return make_tex_object_handle(res, tex, h_backing, h_context); } -TexObjectHandle create_tex_object_handle_linear(const ContextHandle& h_context, - const CUDA_RESOURCE_DESC& res, - const CUDA_TEXTURE_DESC& tex, - const DevicePtrHandle& h_backing) { +TexObjectHandle +create_tex_object_handle_linear( + const ContextHandle &h_context, + const CUDA_RESOURCE_DESC &res, + const CUDA_TEXTURE_DESC &tex, + const DevicePtrHandle &h_backing +) +{ return make_tex_object_handle(res, tex, h_backing, h_context); } -SurfObjectHandle create_surf_object_handle(const ContextHandle& h_context, - const CUDA_RESOURCE_DESC& res, - const OpaqueArrayHandle& h_backing) { +SurfObjectHandle +create_surf_object_handle( + const ContextHandle &h_context, + const CUDA_RESOURCE_DESC &res, + const OpaqueArrayHandle &h_backing +) +{ GILReleaseGuard gil; CUsurfObject obj = 0; err = invoke_in_context_or_undo( h_context, - [&]() noexcept { return DRIVER_CALL(cuSurfObjectCreate, &obj, &res); }, - [&]() noexcept { pw_cuSurfObjectDestroy(obj); }, - /*undo_requires_target_context=*/true); + [&]() noexcept { + return DRIVER_CALL(cuSurfObjectCreate, &obj, &res); + }, + [&]() noexcept { + pw_cuSurfObjectDestroy(obj); + }, + /*undo_requires_target_context=*/true + ); if (err != CUDA_SUCCESS) { return {}; } auto box = std::shared_ptr( new SurfObjectBox{SurfObjectValue{obj}, h_backing, h_context}, - [](const SurfObjectBox* b) { + [](const SurfObjectBox *b) { GILReleaseGuard gil; cleanup_in_context(b->h_context, "cuSurfObjectDestroy", handle_bits(b->resource.raw), [&]() noexcept { return DRIVER_CALL(cuSurfObjectDestroy, b->resource.raw); diff --git a/cuda_core/cuda/core/_cpp/rt/types.hpp b/cuda_core/cuda/core/_cpp/rt/types.hpp index d0a831d2105..08e5bb65e49 100644 --- a/cuda_core/cuda/core/_cpp/rt/types.hpp +++ b/cuda_core/cuda/core/_cpp/rt/types.hpp @@ -13,11 +13,11 @@ // Forward declaration for NVVM - avoids nvvm.h dependency // Use void* to match cuda.bindings.cynvvm's typedef -using nvvmProgram = void*; +using nvvmProgram = void *; // Forward declaration for nvJitLink - avoids nvJitLink.h dependency // Use void* to match cuda.bindings.cynvjitlink's typedef -using nvJitLink_t = void*; +using nvJitLink_t = void *; namespace cuda_core::rt { @@ -29,7 +29,7 @@ namespace cuda_core::rt { // value with a unique tag type, making each shared_ptr type distinct. // ============================================================================ -template +template struct TaggedHandle { T raw; }; @@ -97,241 +97,338 @@ using DevicePtrHandle = std::shared_ptr; using OpaqueHandle = std::shared_ptr; struct PreparedAttachmentState; -using PreparedAttachmentRollback = - void (*)(PreparedAttachmentState*) noexcept; +using PreparedAttachmentRollback = void (*)(PreparedAttachmentState *) noexcept; struct PreparedAttachmentDeleter { PreparedAttachmentRollback rollback = nullptr; - void operator()(PreparedAttachmentState* state) const noexcept { + void + operator()(PreparedAttachmentState *state) const noexcept + { rollback(state); } }; -using PreparedAttachment = - std::unique_ptr; +using PreparedAttachment = std::unique_ptr; struct PreparedChildGraphUpdateState; // Opaque unpublished hierarchy transaction; releasing it discards staged // metadata unless graph_commit_child_graph_update publishes the replacement. -using PreparedChildGraphUpdate = - std::shared_ptr; +using PreparedChildGraphUpdate = std::shared_ptr; struct PreparedExecAttachmentState; -using PreparedExecAttachmentRollback = - void (*)(PreparedExecAttachmentState*) noexcept; +using PreparedExecAttachmentRollback = void (*)(PreparedExecAttachmentState *) noexcept; struct PreparedExecAttachmentDeleter { PreparedExecAttachmentRollback rollback = nullptr; - void operator()(PreparedExecAttachmentState* state) const noexcept { + void + operator()(PreparedExecAttachmentState *state) const noexcept + { rollback(state); } }; // Opaque append transaction. Releasing it rolls back newly appended owners // unless graph_commit_exec_attachment has kept them. -using PreparedExecAttachment = - std::unique_ptr; +using PreparedExecAttachment = std::unique_ptr; // ============================================================================ // Overloaded helper functions to extract raw resources from handles // ============================================================================ // as_cu() - extract the raw CUDA handle -inline CUcontext as_cu(const ContextHandle& h) noexcept { +inline CUcontext +as_cu(const ContextHandle &h) noexcept +{ return h ? *h : nullptr; } -inline CUgreenCtx as_cu(const GreenCtxHandle& h) noexcept { +inline CUgreenCtx +as_cu(const GreenCtxHandle &h) noexcept +{ return h ? *h : nullptr; } -inline CUstream as_cu(const StreamHandle& h) noexcept { +inline CUstream +as_cu(const StreamHandle &h) noexcept +{ return h ? *h : nullptr; } -inline CUevent as_cu(const EventHandle& h) noexcept { +inline CUevent +as_cu(const EventHandle &h) noexcept +{ return h ? *h : nullptr; } -inline CUmemoryPool as_cu(const MemoryPoolHandle& h) noexcept { +inline CUmemoryPool +as_cu(const MemoryPoolHandle &h) noexcept +{ return h ? *h : nullptr; } -inline CUdeviceptr as_cu(const DevicePtrHandle& h) noexcept { +inline CUdeviceptr +as_cu(const DevicePtrHandle &h) noexcept +{ return h ? *h : 0; } -inline CUlibrary as_cu(const LibraryHandle& h) noexcept { +inline CUlibrary +as_cu(const LibraryHandle &h) noexcept +{ return h ? *h : nullptr; } -inline CUmodule as_cu(const CUmodule& h) noexcept { +inline CUmodule +as_cu(const CUmodule &h) noexcept +{ return h; } -inline CUkernel as_cu(const KernelHandle& h) noexcept { +inline CUkernel +as_cu(const KernelHandle &h) noexcept +{ return h ? *h : nullptr; } -inline CUgraph as_cu(const GraphHandle& h) noexcept { +inline CUgraph +as_cu(const GraphHandle &h) noexcept +{ return h ? *h : nullptr; } -inline CUgraphExec as_cu(const GraphExecHandle& h) noexcept { +inline CUgraphExec +as_cu(const GraphExecHandle &h) noexcept +{ return h ? *h : nullptr; } -inline CUgraphNode as_cu(const GraphNodeHandle& h) noexcept { +inline CUgraphNode +as_cu(const GraphNodeHandle &h) noexcept +{ return h ? *h : nullptr; } -inline CUgraphicsResource as_cu(const GraphicsResourceHandle& h) noexcept { +inline CUgraphicsResource +as_cu(const GraphicsResourceHandle &h) noexcept +{ return h ? *h : nullptr; } -inline nvrtcProgram as_cu(const NvrtcProgramHandle& h) noexcept { +inline nvrtcProgram +as_cu(const NvrtcProgramHandle &h) noexcept +{ return h ? *h : nullptr; } -inline nvvmProgram as_cu(const NvvmProgramHandle& h) noexcept { +inline nvvmProgram +as_cu(const NvvmProgramHandle &h) noexcept +{ return h ? h->raw : nullptr; } -inline nvJitLink_t as_cu(const NvJitLinkHandle& h) noexcept { +inline nvJitLink_t +as_cu(const NvJitLinkHandle &h) noexcept +{ return h ? h->raw : nullptr; } -inline CUlinkState as_cu(const CuLinkHandle& h) noexcept { +inline CUlinkState +as_cu(const CuLinkHandle &h) noexcept +{ return h ? *h : nullptr; } -inline CUarray as_cu(const OpaqueArrayHandle& h) noexcept { +inline CUarray +as_cu(const OpaqueArrayHandle &h) noexcept +{ return h ? *h : nullptr; } -inline CUmipmappedArray as_cu(const MipmappedArrayHandle& h) noexcept { +inline CUmipmappedArray +as_cu(const MipmappedArrayHandle &h) noexcept +{ return h ? *h : nullptr; } // CUtexObject / CUsurfObject are integer-valued (like CUdeviceptr); null is 0. // The raw value lives in the tagged wrapper's `raw` field. -inline CUtexObject as_cu(const TexObjectHandle& h) noexcept { +inline CUtexObject +as_cu(const TexObjectHandle &h) noexcept +{ return h ? h->raw : 0; } -inline CUsurfObject as_cu(const SurfObjectHandle& h) noexcept { +inline CUsurfObject +as_cu(const SurfObjectHandle &h) noexcept +{ return h ? h->raw : 0; } -inline CUmemGenericAllocationHandle as_cu(const MemAllocationHandle& h) noexcept { +inline CUmemGenericAllocationHandle +as_cu(const MemAllocationHandle &h) noexcept +{ return h ? h->raw : 0; } -inline CUdeviceptr as_cu(const VaReservationHandle& h) noexcept { +inline CUdeviceptr +as_cu(const VaReservationHandle &h) noexcept +{ return h ? h->raw : 0; } -inline CUdeviceptr as_cu(const VaMappingHandle& h) noexcept { +inline CUdeviceptr +as_cu(const VaMappingHandle &h) noexcept +{ return h ? h->raw : 0; } // as_intptr() - extract handle as intptr_t for Python interop // Using signed intptr_t per C standard convention and issue #1342 -inline std::intptr_t as_intptr(const ContextHandle& h) noexcept { +inline std::intptr_t +as_intptr(const ContextHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const GreenCtxHandle& h) noexcept { +inline std::intptr_t +as_intptr(const GreenCtxHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const StreamHandle& h) noexcept { +inline std::intptr_t +as_intptr(const StreamHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const EventHandle& h) noexcept { +inline std::intptr_t +as_intptr(const EventHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const MemoryPoolHandle& h) noexcept { +inline std::intptr_t +as_intptr(const MemoryPoolHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const DevicePtrHandle& h) noexcept { +inline std::intptr_t +as_intptr(const DevicePtrHandle &h) noexcept +{ return static_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const LibraryHandle& h) noexcept { +inline std::intptr_t +as_intptr(const LibraryHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const CUmodule& h) noexcept { +inline std::intptr_t +as_intptr(const CUmodule &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const KernelHandle& h) noexcept { +inline std::intptr_t +as_intptr(const KernelHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const GraphHandle& h) noexcept { +inline std::intptr_t +as_intptr(const GraphHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const GraphExecHandle& h) noexcept { +inline std::intptr_t +as_intptr(const GraphExecHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const GraphNodeHandle& h) noexcept { +inline std::intptr_t +as_intptr(const GraphNodeHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const GraphicsResourceHandle& h) noexcept { +inline std::intptr_t +as_intptr(const GraphicsResourceHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const NvrtcProgramHandle& h) noexcept { +inline std::intptr_t +as_intptr(const NvrtcProgramHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const NvvmProgramHandle& h) noexcept { +inline std::intptr_t +as_intptr(const NvvmProgramHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const NvJitLinkHandle& h) noexcept { +inline std::intptr_t +as_intptr(const NvJitLinkHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const CuLinkHandle& h) noexcept { +inline std::intptr_t +as_intptr(const CuLinkHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const FileDescriptorHandle& h) noexcept { +inline std::intptr_t +as_intptr(const FileDescriptorHandle &h) noexcept +{ return h ? static_cast(*h) : -1; } -inline std::intptr_t as_intptr(const OpaqueArrayHandle& h) noexcept { +inline std::intptr_t +as_intptr(const OpaqueArrayHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const MipmappedArrayHandle& h) noexcept { +inline std::intptr_t +as_intptr(const MipmappedArrayHandle &h) noexcept +{ return reinterpret_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const TexObjectHandle& h) noexcept { +inline std::intptr_t +as_intptr(const TexObjectHandle &h) noexcept +{ return static_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const SurfObjectHandle& h) noexcept { +inline std::intptr_t +as_intptr(const SurfObjectHandle &h) noexcept +{ return static_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const MemAllocationHandle& h) noexcept { +inline std::intptr_t +as_intptr(const MemAllocationHandle &h) noexcept +{ return static_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const VaReservationHandle& h) noexcept { +inline std::intptr_t +as_intptr(const VaReservationHandle &h) noexcept +{ return static_cast(as_cu(h)); } -inline std::intptr_t as_intptr(const VaMappingHandle& h) noexcept { +inline std::intptr_t +as_intptr(const VaMappingHandle &h) noexcept +{ return static_cast(as_cu(h)); } diff --git a/cuda_core/cuda/core/_cpp/rt/versions.hpp b/cuda_core/cuda/core/_cpp/rt/versions.hpp index 88dbd8f4e16..eb19e5bf224 100644 --- a/cuda_core/cuda/core/_cpp/rt/versions.hpp +++ b/cuda_core/cuda/core/_cpp/rt/versions.hpp @@ -49,12 +49,14 @@ #ifdef CUDA_CORE_MIN_CUDA_VERSION #if CUDA_VERSION < CUDA_CORE_MIN_CUDA_VERSION -#error "cuda.h is older than the floor of this cuda.core release for its CUDA major series (see the cuda.core support policy)" +#error \ + "cuda.h is older than the floor of this cuda.core release for its CUDA major series (see the cuda.core support policy)" #endif #endif #ifdef CUDA_CORE_BINDINGS_CUDA_VERSION #if (CUDA_VERSION / 10) != (CUDA_CORE_BINDINGS_CUDA_VERSION / 10) -#error "the cuda.h the compiler resolved is not of the major.minor that the installed cuda-bindings was generated from (check CUDA_PATH, CUDA_HOME and the compiler's include path)" +#error \ + "the cuda.h the compiler resolved is not of the major.minor that the installed cuda-bindings was generated from (check CUDA_PATH, CUDA_HOME and the compiler's include path)" #endif #endif diff --git a/cuda_core/cuda/core/_cpp/rt/virtual_memory.cpp b/cuda_core/cuda/core/_cpp/rt/virtual_memory.cpp index 0c9db6355af..538f07aecd7 100644 --- a/cuda_core/cuda/core/_cpp/rt/virtual_memory.cpp +++ b/cuda_core/cuda/core/_cpp/rt/virtual_memory.cpp @@ -36,8 +36,8 @@ struct MemAllocationBox { }; struct VaReservationBox { - VaReservationValue resource; // base address from cuMemAddressReserve - size_t size; // exact reserved size; cuMemAddressFree needs both + VaReservationValue resource; // base address from cuMemAddressReserve + size_t size; // exact reserved size; cuMemAddressFree needs both }; // Process-wide count of live VA reservations created here with @@ -49,27 +49,34 @@ struct VaReservationBox { std::atomic g_live_va_reservations{0}; struct VaMappingBox { - VaMappingValue resource; // mapped address - size_t size; // == the allocation's size - MemAllocationHandle h_alloc; // released after the unmap - VaReservationHandle h_reservation; // freed after the unmap + VaMappingValue resource; // mapped address + size_t size; // == the allocation's size + MemAllocationHandle h_alloc; // released after the unmap + VaReservationHandle h_reservation; // freed after the unmap }; // Recover a box from the aliased handle; the resource is the first member. template -const Box* box_of(const Handle& h) noexcept { - return reinterpret_cast( - reinterpret_cast(h.get()) - offsetof(Box, resource)); +const Box * +box_of(const Handle &h) noexcept +{ + return reinterpret_cast(reinterpret_cast(h.get()) - offsetof(Box, resource)); } // Wrap an allocation handle this process holds one reference to, whether // cuMemCreate or cuMemImportFromShareableHandle produced it. The last // reference releases it; the memory is freed once no mapping remains. -MemAllocationHandle wrap_mem_allocation(CUmemGenericAllocationHandle handle, size_t size, - std::vector&& access, bool imported) { +MemAllocationHandle +wrap_mem_allocation( + CUmemGenericAllocationHandle handle, + size_t size, + std::vector &&access, + bool imported +) +{ auto box = std::shared_ptr( new MemAllocationBox{{handle}, size, std::move(access), imported}, - [](const MemAllocationBox* b) { + [](const MemAllocationBox *b) { GILReleaseGuard gil; pw_cuMemRelease(b->resource.raw); delete b; @@ -84,8 +91,9 @@ MemAllocationHandle wrap_mem_allocation(CUmemGenericAllocationHandle handle, siz // Physical allocations // ============================================================================ -MemAllocationHandle create_mem_allocation_handle(size_t size, const CUmemAllocationProp& prop, - const CUmemAccessDesc* descs, size_t count) { +MemAllocationHandle +create_mem_allocation_handle(size_t size, const CUmemAllocationProp &prop, const CUmemAccessDesc *descs, size_t count) +{ // Copy the descriptors before the driver call so a failed copy leaves // nothing to undo. std::vector access(descs, descs + count); @@ -97,8 +105,15 @@ MemAllocationHandle create_mem_allocation_handle(size_t size, const CUmemAllocat return wrap_mem_allocation(handle, size, std::move(access), false); } -MemAllocationHandle import_mem_allocation_handle(void* os_handle, CUmemAllocationHandleType handle_type, - size_t size, const CUmemAccessDesc* descs, size_t count) { +MemAllocationHandle +import_mem_allocation_handle( + void *os_handle, + CUmemAllocationHandleType handle_type, + size_t size, + const CUmemAccessDesc *descs, + size_t count +) +{ // Copy the descriptors before the driver call so a failed copy leaves // nothing to undo. std::vector access(descs, descs + count); @@ -110,11 +125,15 @@ MemAllocationHandle import_mem_allocation_handle(void* os_handle, CUmemAllocatio return wrap_mem_allocation(handle, size, std::move(access), true); } -size_t mem_allocation_size(const MemAllocationHandle& h) noexcept { +size_t +mem_allocation_size(const MemAllocationHandle &h) noexcept +{ return h ? box_of(h)->size : 0; } -bool mem_allocation_is_imported(const MemAllocationHandle& h) noexcept { +bool +mem_allocation_is_imported(const MemAllocationHandle &h) noexcept +{ return h ? box_of(h)->imported : false; } @@ -122,34 +141,37 @@ bool mem_allocation_is_imported(const MemAllocationHandle& h) noexcept { // Address reservations // ============================================================================ -VaReservationHandle create_va_reservation_handle(size_t size, size_t alignment, CUdeviceptr hint) { +VaReservationHandle +create_va_reservation_handle(size_t size, size_t alignment, CUdeviceptr hint) +{ GILReleaseGuard gil; CUdeviceptr ptr = 0; if (CUDA_SUCCESS != (err = DRIVER_CALL(cuMemAddressReserve, &ptr, size, alignment, hint, 0))) { return {}; } - auto* raw = new VaReservationBox{{ptr}, size}; + auto *raw = new VaReservationBox{{ptr}, size}; // Count before constructing the shared_ptr: if its control-block allocation // throws, it runs the deleter on raw, whose decrement then balances this. g_live_va_reservations.fetch_add(1, std::memory_order_relaxed); - auto box = std::shared_ptr( - raw, - [](const VaReservationBox* b) { - GILReleaseGuard gil; - pw_cuMemAddressFree(b->resource.raw, b->size); - g_live_va_reservations.fetch_sub(1, std::memory_order_relaxed); - delete b; - } - ); + auto box = std::shared_ptr(raw, [](const VaReservationBox *b) { + GILReleaseGuard gil; + pw_cuMemAddressFree(b->resource.raw, b->size); + g_live_va_reservations.fetch_sub(1, std::memory_order_relaxed); + delete b; + }); return VaReservationHandle(box, &box->resource); } // Test-suite instrumentation; see g_live_va_reservations. -long long live_va_reservation_count() noexcept { +long long +live_va_reservation_count() noexcept +{ return g_live_va_reservations.load(std::memory_order_relaxed); } -size_t va_reservation_size(const VaReservationHandle& h) noexcept { +size_t +va_reservation_size(const VaReservationHandle &h) noexcept +{ return h ? box_of(h)->size : 0; } @@ -157,14 +179,15 @@ size_t va_reservation_size(const VaReservationHandle& h) noexcept { // Mappings // ============================================================================ -VaMappingHandle create_va_mapping_handle(CUdeviceptr ptr, const MemAllocationHandle& h_alloc, - const VaReservationHandle& h_res) { +VaMappingHandle +create_va_mapping_handle(CUdeviceptr ptr, const MemAllocationHandle &h_alloc, const VaReservationHandle &h_res) +{ if (!h_alloc || !h_res) { err = CUDA_ERROR_INVALID_VALUE; return {}; } - const MemAllocationBox* alloc = box_of(h_alloc); - const VaReservationBox* res = box_of(h_res); + const MemAllocationBox *alloc = box_of(h_alloc); + const VaReservationBox *res = box_of(h_res); const CUdeviceptr base = res->resource.raw; if (ptr < base || ptr - base > res->size || alloc->size > res->size - (ptr - base)) { err = CUDA_ERROR_INVALID_VALUE; @@ -178,7 +201,8 @@ VaMappingHandle create_va_mapping_handle(CUdeviceptr ptr, const MemAllocationHan // cuMemSetAccess rejects an empty descriptor list; a mapping with no // descriptors is mapped but not accessible, which is what the caller asked for. if (!alloc->access.empty()) { - const CUresult status = DRIVER_CALL(cuMemSetAccess, ptr, alloc->size, alloc->access.data(), alloc->access.size()); + const CUresult status = + DRIVER_CALL(cuMemSetAccess, ptr, alloc->size, alloc->access.data(), alloc->access.size()); if (status != CUDA_SUCCESS) { pw_cuMemUnmap(ptr, alloc->size); err = status; @@ -187,7 +211,7 @@ VaMappingHandle create_va_mapping_handle(CUdeviceptr ptr, const MemAllocationHan } auto box = std::shared_ptr( new VaMappingBox{{ptr}, alloc->size, h_alloc, h_res}, - [](const VaMappingBox* b) { + [](const VaMappingBox *b) { GILReleaseGuard gil; pw_cuMemUnmap(b->resource.raw, b->size); delete b; // then the allocation and the reservation release @@ -196,11 +220,15 @@ VaMappingHandle create_va_mapping_handle(CUdeviceptr ptr, const MemAllocationHan return VaMappingHandle(box, &box->resource); } -size_t va_mapping_size(const VaMappingHandle& h) noexcept { +size_t +va_mapping_size(const VaMappingHandle &h) noexcept +{ return h ? box_of(h)->size : 0; } -MemAllocationHandle va_mapping_allocation(const VaMappingHandle& h) noexcept { +MemAllocationHandle +va_mapping_allocation(const VaMappingHandle &h) noexcept +{ return h ? box_of(h)->h_alloc : MemAllocationHandle{}; } @@ -214,7 +242,9 @@ MemAllocationHandle va_mapping_allocation(const VaMappingHandle& h) noexcept { // CUDA_ERROR_STREAM_CAPTURE_IMPLICIT). cuStreamSynchronize would invalidate // such a capture; cuStreamIsCapturing does not. It is the query // cuStreamGetCaptureInfo makes first, with one signature on every CUDA major. -static bool sync_would_disturb_capture(CUstream stream) noexcept { +static bool +sync_would_disturb_capture(CUstream stream) noexcept +{ CUstreamCaptureStatus status = CU_STREAM_CAPTURE_STATUS_NONE; const CUresult result = DRIVER_CALL(cuStreamIsCapturing, stream, &status); if (result == CUDA_ERROR_STREAM_CAPTURE_IMPLICIT) { @@ -235,19 +265,24 @@ static bool sync_would_disturb_capture(CUstream stream) noexcept { // every global-mode capture in the process, and any non-relaxed capture this // thread began, on streams unrelated to `s`. Relaxed mode disables that // interaction; the stream's own capture state is still checked above. -static void sync_recorded_stream(const DeallocationStream& ds, bool& capture_skipped) noexcept { +static void +sync_recorded_stream(const DeallocationStream &ds, bool &capture_skipped) noexcept +{ const CUstream s = as_cu(ds.h_stream); CUcontext previous = nullptr; int changed = 0; - const char* detail = nullptr; + const char *detail = nullptr; CUresult status = enter_context(deallocation_context(ds), &previous, &changed); if (status != CUDA_SUCCESS) { detail = "skipped (context activation failed); the range was unmapped without synchronizing it"; - } else if (sync_would_disturb_capture(s)) { + } + else if (sync_would_disturb_capture(s)) { capture_skipped = true; - } else { + } + else { CUstreamCaptureMode mode = CU_STREAM_CAPTURE_MODE_RELAXED; - const CUresult swapped = DRIVER_CALL(cuThreadExchangeStreamCaptureMode, &mode); // `mode` now holds the previous mode + const CUresult swapped = + DRIVER_CALL(cuThreadExchangeStreamCaptureMode, &mode); // `mode` now holds the previous mode status = DRIVER_CALL(cuStreamSynchronize, s); if (swapped == CUDA_SUCCESS) { DRIVER_CALL(cuThreadExchangeStreamCaptureMode, &mode); // restore the previous mode @@ -272,7 +307,9 @@ static void sync_recorded_stream(const DeallocationStream& ds, bool& capture_ski } namespace detail { -void vmm_sync_before_release(const DeallocationStream& stream) noexcept { +void +vmm_sync_before_release(const DeallocationStream &stream) noexcept +{ // Order the release on this buffer's stream, as the other DevicePtrBox // deleters do, unless the interpreter is finalizing or the sync would // disturb a capture. A host-located buffer recorded no stream. @@ -285,23 +322,30 @@ void vmm_sync_before_release(const DeallocationStream& stream) noexcept { report_message( "a VirtualMemoryResource buffer was released while its deallocation stream " "is capturing, or is the legacy stream while another stream in its context is " - "capturing; the buffer was unmapped without synchronizing that stream"); + "capturing; the buffer was unmapped without synchronizing that stream" + ); } } } // namespace detail -VmmRangeHandle create_vmm_range(const std::vector& mappings) { +VmmRangeHandle +create_vmm_range(const std::vector &mappings) +{ return std::make_shared(mappings); } -std::vector vmm_range_mappings(const VmmRangeHandle& range) { +std::vector +vmm_range_mappings(const VmmRangeHandle &range) +{ return range ? *range : VmmRange{}; } -size_t vmm_range_total(const VmmRangeHandle& range) noexcept { +size_t +vmm_range_total(const VmmRangeHandle &range) noexcept +{ size_t total = 0; if (range) { - for (const VaMappingHandle& m : *range) { + for (const VaMappingHandle &m : *range) { total += va_mapping_size(m); } } diff --git a/cuda_core/cuda/core/_include/layout.hpp b/cuda_core/cuda/core/_include/layout.hpp index f92401a0205..52385dc8c73 100644 --- a/cuda_core/cuda/core/_include/layout.hpp +++ b/cuda_core/cuda/core/_include/layout.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. // // SPDX-License-Identifier: Apache-2.0 @@ -13,33 +13,40 @@ #define STRIDED_LAYOUT_MAX_NDIM 64 #define AXES_MASK_ALL 0xFFFFFFFFFFFFFFFEULL -inline int64_t _c_abs(int64_t x) { return std::abs(x); } +inline int64_t +_c_abs(int64_t x) +{ + return std::abs(x); +} template -void _swap(T &a, T &b) noexcept { std::swap(a, b); } +void +_swap(T &a, T &b) noexcept +{ + std::swap(a, b); +} -inline void _order_from_strides(std::vector &indices, const int64_t *shape, - const int64_t *strides, int ndim) +inline void +_order_from_strides(std::vector &indices, const int64_t *shape, const int64_t *strides, int ndim) { indices.resize(ndim); std::iota(indices.begin(), indices.end(), 0); - if (!strides) - { + if (!strides) { return; } - std::sort(indices.begin(), indices.end(), [&strides, &shape](int i, int j) - { - int64_t stride_i = _c_abs(strides[i]); - int64_t stride_j = _c_abs(strides[j]); - if (stride_i != stride_j) { - return stride_i > stride_j; - } - int64_t shape_i = shape[i]; - int64_t shape_j = shape[j]; - if (shape_i != shape_j) { - return shape_i > shape_j; - } - return i < j; }); + std::sort(indices.begin(), indices.end(), [&strides, &shape](int i, int j) { + int64_t stride_i = _c_abs(strides[i]); + int64_t stride_j = _c_abs(strides[j]); + if (stride_i != stride_j) { + return stride_i > stride_j; + } + int64_t shape_i = shape[i]; + int64_t shape_j = shape[j]; + if (shape_i != shape_j) { + return shape_i > shape_j; + } + return i < j; + }); } -#endif // CUDA_CORE_LAYOUT_HPP +#endif // CUDA_CORE_LAYOUT_HPP diff --git a/cuda_core/cuda/core/_include/utility.hpp b/cuda_core/cuda/core/_include/utility.hpp index 64b357ac324..f3ef396f268 100644 --- a/cuda_core/cuda/core/_include/utility.hpp +++ b/cuda_core/cuda/core/_include/utility.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. // // SPDX-License-Identifier: Apache-2.0 @@ -10,14 +10,16 @@ // We want to have the C++ layer supporting all past 12.x versions, so some tricks are needed. // Since there's no std::has_member so we use SFINAE to create the same effect. -template ::_pvt_ptr)>, int> = 0> -inline auto& get_cuda_native_handle(const T& obj) { +template ::_pvt_ptr)>, int> = 0> +inline auto & +get_cuda_native_handle(const T &obj) +{ return *(obj->_pvt_ptr); } -template ::_ptr)>, int> = 0> -inline auto& get_cuda_native_handle(const T& obj) { +template ::_ptr)>, int> = 0> +inline auto & +get_cuda_native_handle(const T &obj) +{ return *(obj->_ptr); } diff --git a/cuda_core/tests/test_binaries/saxpy.cu b/cuda_core/tests/test_binaries/saxpy.cu index f0afc58637c..38890a648e6 100644 --- a/cuda_core/tests/test_binaries/saxpy.cu +++ b/cuda_core/tests/test_binaries/saxpy.cu @@ -3,12 +3,16 @@ #include -__device__ float saxpy_step(float a, float x, float y) { +__device__ float +saxpy_step(float a, float x, float y) +{ return a * x + y; } -template -__global__ void saxpy(const T a, const T* x, const T* y, T* out, size_t N) { +template +__global__ void +saxpy(const T a, const T *x, const T *y, T *out, size_t N) +{ const unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x; for (size_t i = tid; i < N; i += gridDim.x * blockDim.x) { out[tid] = a * x[tid] + y[tid];