From 3c6c19531dd3274bbdc4ed289aae5cdd346d8797 Mon Sep 17 00:00:00 2001 From: Alexander Droste Date: Fri, 23 Oct 2026 12:19:10 +0100 Subject: [PATCH] feat(cuda): add minimal CUDA driver bindings Introduce vortex-cuda-driver with lazily loaded CUDA symbols, raw handle types, and fallible driver-call wrappers. Register the crate in the workspace while leaving existing cudarc consumers unchanged. Original-commit: 5fd3f54eba Signed-off-by: Alexander Droste --- Cargo.lock | 7 + Cargo.toml | 2 + vortex-cuda/driver/Cargo.toml | 23 +++ vortex-cuda/driver/src/lib.rs | 12 ++ vortex-cuda/driver/src/result.rs | 257 +++++++++++++++++++++++++++ vortex-cuda/driver/src/sys.rs | 290 +++++++++++++++++++++++++++++++ 6 files changed, 591 insertions(+) create mode 100644 vortex-cuda/driver/Cargo.toml create mode 100644 vortex-cuda/driver/src/lib.rs create mode 100644 vortex-cuda/driver/src/result.rs create mode 100644 vortex-cuda/driver/src/sys.rs diff --git a/Cargo.lock b/Cargo.lock index a471337bea4..e5b3a9f61ab 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -11762,6 +11762,13 @@ dependencies = [ "vortex-onpair", ] +[[package]] +name = "vortex-cuda-driver" +version = "0.1.0" +dependencies = [ + "libloading 0.8.9", +] + [[package]] name = "vortex-cuda-ffi" version = "0.1.0" diff --git a/Cargo.toml b/Cargo.toml index a551b0f397b..9fe805bfa3c 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -29,6 +29,7 @@ members = [ "vortex-duckdb", "vortex-cuda", "vortex-cuda/cub", + "vortex-cuda/driver", "vortex-cuda/ffi", "vortex-cuda/gpu-scan-cli", "vortex-cuda/macros", @@ -348,6 +349,7 @@ vortex-zstd = { version = "0.1.0", path = "./encodings/zstd", default-features = vortex-bench = { path = "./vortex-bench", default-features = false } vortex-bench-support = { path = "./benchmarks/bench-support" } vortex-cuda = { path = "./vortex-cuda", default-features = false } +vortex-cuda-driver = { path = "./vortex-cuda/driver" } vortex-cuda-macros = { path = "./vortex-cuda/macros" } vortex-duckdb = { path = "./vortex-duckdb", default-features = false } vortex-tui = { path = "./vortex-tui" } diff --git a/vortex-cuda/driver/Cargo.toml b/vortex-cuda/driver/Cargo.toml new file mode 100644 index 00000000000..1ae7ad19d96 --- /dev/null +++ b/vortex-cuda/driver/Cargo.toml @@ -0,0 +1,23 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright the Vortex contributors + +[package] +name = "vortex-cuda-driver" +authors.workspace = true +description = "Minimal CUDA driver API bindings for Vortex" +edition = { workspace = true } +categories = { workspace = true } +include = { workspace = true } +keywords = { workspace = true } +license = { workspace = true } +readme = { workspace = true } +repository = { workspace = true } +rust-version = { workspace = true } +version = { workspace = true } +publish = false + +[lints] +workspace = true + +[dependencies] +libloading = { workspace = true } diff --git a/vortex-cuda/driver/src/lib.rs b/vortex-cuda/driver/src/lib.rs new file mode 100644 index 00000000000..64cbe2ce72f --- /dev/null +++ b/vortex-cuda/driver/src/lib.rs @@ -0,0 +1,12 @@ +// SPDX-License-Identifier: Apache-2.0 +// SPDX-FileCopyrightText: Copyright the Vortex contributors + +//! Minimal CUDA driver API bindings for Vortex. +//! +//! The driver library is loaded lazily, so dependents build and run without CUDA until +//! they first use it. The raw bindings and result wrappers operate on CUDA driver handles. + +pub mod result; +pub mod sys; + +pub use result::DriverError; diff --git a/vortex-cuda/driver/src/result.rs b/vortex-cuda/driver/src/result.rs new file mode 100644 index 00000000000..9f5e27bdd34 --- /dev/null +++ b/vortex-cuda/driver/src/result.rs @@ -0,0 +1,257 @@ +// SPDX-License-Identifier: Apache-2.0 +// SPDX-FileCopyrightText: Copyright the Vortex contributors + +//! Driver calls returning [`Result`]s, for callers that need the raw handles. + +use std::ffi::CStr; +use std::ffi::c_uint; +use std::ffi::c_void; +use std::fmt; +use std::mem::MaybeUninit; + +use crate::sys; +use crate::sys::CUresult; + +/// A failed CUDA driver call. +#[derive(Clone, Copy, PartialEq, Eq, Hash)] +pub struct DriverError(pub CUresult); + +impl DriverError { + /// The driver's name for the error, such as `CUDA_ERROR_INVALID_VALUE`. + pub fn error_name(&self) -> Result<&'static CStr, DriverError> { + let mut name = std::ptr::null(); + // SAFETY: The driver writes a pointer to a static string. + unsafe { sys::cuGetErrorName(self.0, &raw mut name) }.result()?; + // SAFETY: Successful calls return a NUL-terminated static string. + Ok(unsafe { CStr::from_ptr(name) }) + } + + /// The driver's description of the error. + pub fn error_string(&self) -> Result<&'static CStr, DriverError> { + let mut description = std::ptr::null(); + // SAFETY: The driver writes a pointer to a static string. + unsafe { sys::cuGetErrorString(self.0, &raw mut description) }.result()?; + // SAFETY: Successful calls return a NUL-terminated static string. + Ok(unsafe { CStr::from_ptr(description) }) + } +} + +impl fmt::Debug for DriverError { + fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result { + match self.error_name() { + Ok(name) => write!(f, "DriverError({}, {})", self.0.0, name.to_string_lossy()), + Err(_) => write!(f, "DriverError({})", self.0.0), + } + } +} + +impl fmt::Display for DriverError { + fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result { + match (self.error_name(), self.error_string()) { + (Ok(name), Ok(description)) => write!( + f, + "{}: {}", + name.to_string_lossy(), + description.to_string_lossy() + ), + _ => write!(f, "CUDA driver error {}", self.0.0), + } + } +} + +impl std::error::Error for DriverError {} + +impl CUresult { + /// `Ok` for [`CUresult::CUDA_SUCCESS`], otherwise the error. + #[inline] + pub fn result(self) -> Result<(), DriverError> { + if self == CUresult::CUDA_SUCCESS { + Ok(()) + } else { + Err(DriverError(self)) + } + } +} + +/// Initializes the driver; repeated calls are cheap. +pub fn init() -> Result<(), DriverError> { + // SAFETY: cuInit has no preconditions. + unsafe { sys::cuInit(0) }.result() +} + +/// Allocates `bytes` of page-locked host memory with the given `cuMemHostAlloc` flags. +/// +/// # Safety +/// +/// A CUDA context must be current; the memory is uninitialized. +pub unsafe fn malloc_host(bytes: usize, flags: c_uint) -> Result<*mut c_void, DriverError> { + let mut pointer = MaybeUninit::uninit(); + // SAFETY: The driver writes the allocation into `pointer` on success. + unsafe { sys::cuMemHostAlloc(pointer.as_mut_ptr(), bytes, flags) }.result()?; + // SAFETY: Initialized by the successful call. + Ok(unsafe { pointer.assume_init() }) +} + +/// Frees memory allocated by [`malloc_host`]. +/// +/// # Safety +/// +/// `pointer` must come from [`malloc_host`] and no pending copy may still use it. +pub unsafe fn free_host(pointer: *mut c_void) -> Result<(), DriverError> { + // SAFETY: Under the caller's guarantees. + unsafe { sys::cuMemFreeHost(pointer) }.result() +} + +/// Synchronously copies `target.len()` elements from device memory at `source`. +/// +/// # Safety +/// +/// `source` must address at least `size_of_val(target)` readable device bytes. +pub unsafe fn memcpy_dtoh_sync( + target: &mut [T], + source: sys::CUdeviceptr, +) -> Result<(), DriverError> { + // SAFETY: Under the caller's guarantees; `target` is writable for its whole size. + unsafe { sys::cuMemcpyDtoH_v2(target.as_mut_ptr().cast(), source, size_of_val(target)) } + .result() +} + +/// Device queries. +pub mod device { + use super::*; + + /// The device with the given ordinal. + pub fn get(ordinal: i32) -> Result { + let mut device = MaybeUninit::uninit(); + // SAFETY: The driver writes the device into `device` on success. + unsafe { sys::cuDeviceGet(device.as_mut_ptr(), ordinal) }.result()?; + // SAFETY: Initialized by the successful call. + Ok(unsafe { device.assume_init() }) + } + + /// The number of devices. + pub fn get_count() -> Result { + let mut count = 0; + // SAFETY: The driver writes the count. + unsafe { sys::cuDeviceGetCount(&raw mut count) }.result()?; + Ok(count) + } + + /// A device attribute. + /// + /// # Safety + /// + /// `device` must be a valid device handle. + pub unsafe fn get_attribute( + device: sys::CUdevice, + attribute: sys::CUdevice_attribute, + ) -> Result { + let mut value = 0; + // SAFETY: Under the caller's guarantees. + unsafe { sys::cuDeviceGetAttribute(&raw mut value, attribute, device) }.result()?; + Ok(value) + } + + /// The device's current memory pool. + /// + /// # Safety + /// + /// `device` must be a valid device handle. + pub unsafe fn get_mem_pool(device: sys::CUdevice) -> Result { + let mut pool = MaybeUninit::uninit(); + // SAFETY: Under the caller's guarantees. + unsafe { sys::cuDeviceGetMemPool(pool.as_mut_ptr(), device) }.result()?; + // SAFETY: Initialized by the successful call. + Ok(unsafe { pool.assume_init() }) + } +} + +/// Memory pool attributes. +pub mod mem_pool { + use super::*; + + /// Reads a pool attribute into `value`. + /// + /// # Safety + /// + /// `value` must point to storage of the attribute's type. + pub unsafe fn get_attribute( + pool: sys::CUmemoryPool, + attribute: sys::CUmemPool_attribute, + value: *mut c_void, + ) -> Result<(), DriverError> { + // SAFETY: Under the caller's guarantees. + unsafe { sys::cuMemPoolGetAttribute(pool, attribute, value) }.result() + } + + /// Sets a pool attribute from `value`. + /// + /// # Safety + /// + /// `value` must point to a value of the attribute's type. + pub unsafe fn set_attribute( + pool: sys::CUmemoryPool, + attribute: sys::CUmemPool_attribute, + value: *mut c_void, + ) -> Result<(), DriverError> { + // SAFETY: Under the caller's guarantees. + unsafe { sys::cuMemPoolSetAttribute(pool, attribute, value) }.result() + } +} + +/// Stream operations on raw handles. +pub mod stream { + use super::*; + + /// The (least, greatest) stream priorities of the current context. + pub fn get_priority_range() -> Result<(i32, i32), DriverError> { + let (mut least, mut greatest) = (0, 0); + // SAFETY: The driver writes both values. + unsafe { sys::cuCtxGetStreamPriorityRange(&raw mut least, &raw mut greatest) }.result()?; + Ok((least, greatest)) + } + + /// Makes `stream` wait for the work recorded in `event`. + /// + /// # Safety + /// + /// Both handles must be valid and belong to the current context. + pub unsafe fn wait_event( + stream: sys::CUstream, + event: sys::CUevent, + flags: sys::CUevent_wait_flags, + ) -> Result<(), DriverError> { + // SAFETY: Under the caller's guarantees. + unsafe { sys::cuStreamWaitEvent(stream, event, flags.0) }.result() + } + + /// Enqueues a host function on `stream`. + /// + /// # Safety + /// + /// `stream` must be valid, and `function` must be safe to call once with `argument` from a + /// driver thread, without making CUDA calls. + pub unsafe fn launch_host_function( + stream: sys::CUstream, + function: unsafe extern "C" fn(*mut c_void), + argument: *mut c_void, + ) -> Result<(), DriverError> { + // SAFETY: Under the caller's guarantees. + unsafe { sys::cuLaunchHostFunc(stream, Some(function), argument) }.result() + } +} + +/// Event operations on raw handles. +pub mod event { + use super::*; + + /// Blocks until the work recorded in `event` completes. + /// + /// # Safety + /// + /// `event` must be a valid event handle. + pub unsafe fn synchronize(event: sys::CUevent) -> Result<(), DriverError> { + // SAFETY: Under the caller's guarantees. + unsafe { sys::cuEventSynchronize(event) }.result() + } +} diff --git a/vortex-cuda/driver/src/sys.rs b/vortex-cuda/driver/src/sys.rs new file mode 100644 index 00000000000..790810906f1 --- /dev/null +++ b/vortex-cuda/driver/src/sys.rs @@ -0,0 +1,290 @@ +// SPDX-License-Identifier: Apache-2.0 +// SPDX-FileCopyrightText: Copyright the Vortex contributors + +//! Raw CUDA driver API: the types, constants and functions Vortex uses. +//! +//! The driver library is opened lazily on first use, so binaries build and start on hosts +//! without CUDA; [`is_culib_present`] tests for it without failing. Each function forwards to +//! the driver symbol of the same name and returns the raw [`CUresult`]. + +#![allow( + non_camel_case_types, + non_snake_case, + missing_docs, + clippy::too_many_arguments, + reason = "mirrors the driver header's names and signatures" +)] + +use std::ffi::c_char; +use std::ffi::c_int; +use std::ffi::c_uint; +use std::ffi::c_void; +use std::sync::OnceLock; + +use libloading::Library; + +#[repr(C)] +pub struct CUctx_st { + _private: [u8; 0], +} +#[repr(C)] +pub struct CUstream_st { + _private: [u8; 0], +} +#[repr(C)] +pub struct CUevent_st { + _private: [u8; 0], +} +#[repr(C)] +pub struct CUmod_st { + _private: [u8; 0], +} +#[repr(C)] +pub struct CUfunc_st { + _private: [u8; 0], +} +#[repr(C)] +pub struct CUmemPoolHandle_st { + _private: [u8; 0], +} + +pub type CUdevice = c_int; +pub type CUdeviceptr = u64; +pub type CUcontext = *mut CUctx_st; +pub type CUstream = *mut CUstream_st; +pub type CUevent = *mut CUevent_st; +pub type CUmodule = *mut CUmod_st; +pub type CUfunction = *mut CUfunc_st; +pub type CUmemoryPool = *mut CUmemPoolHandle_st; +pub type CUhostFn = Option; + +/// A CUDA driver API status code. +#[repr(transparent)] +#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)] +pub struct CUresult(pub c_uint); + +impl CUresult { + pub const CUDA_SUCCESS: Self = Self(0); + pub const CUDA_ERROR_INVALID_VALUE: Self = Self(1); + pub const CUDA_ERROR_OUT_OF_MEMORY: Self = Self(2); + pub const CUDA_ERROR_NOT_INITIALIZED: Self = Self(3); + pub const CUDA_ERROR_NO_DEVICE: Self = Self(100); + pub const CUDA_ERROR_INVALID_CONTEXT: Self = Self(201); + pub const CUDA_ERROR_INVALID_HANDLE: Self = Self(400); + pub const CUDA_ERROR_NOT_READY: Self = Self(600); +} + +/// The status type under the name of the driver header's enum. +pub type cudaError_enum = CUresult; + +#[repr(transparent)] +#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)] +pub struct CUevent_flags(pub c_uint); + +impl CUevent_flags { + pub const CU_EVENT_DEFAULT: Self = Self(0x0); + pub const CU_EVENT_BLOCKING_SYNC: Self = Self(0x1); + pub const CU_EVENT_DISABLE_TIMING: Self = Self(0x2); +} + +#[repr(transparent)] +#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)] +pub struct CUevent_wait_flags(pub c_uint); + +impl CUevent_wait_flags { + pub const CU_EVENT_WAIT_DEFAULT: Self = Self(0x0); +} + +#[repr(transparent)] +#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)] +pub struct CUdevice_attribute(pub c_int); + +impl CUdevice_attribute { + pub const CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT: Self = Self(16); + pub const CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR: Self = Self(75); + pub const CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR: Self = Self(76); + pub const CU_DEVICE_ATTRIBUTE_PAGEABLE_MEMORY_ACCESS_USES_HOST_PAGE_TABLES: Self = Self(100); + pub const CU_DEVICE_ATTRIBUTE_MEMORY_POOLS_SUPPORTED: Self = Self(115); +} + +#[repr(transparent)] +#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)] +pub struct CUmemPool_attribute(pub c_int); + +impl CUmemPool_attribute { + pub const CU_MEMPOOL_ATTR_RELEASE_THRESHOLD: Self = Self(4); +} + +/// The memory pool attribute type under the name of the driver header's enum. +pub type CUmemPool_attribute_enum = CUmemPool_attribute; + +pub const CU_STREAM_NON_BLOCKING: c_uint = 0x1; +pub const CU_MEMHOSTALLOC_PORTABLE: c_uint = 0x1; +pub const CU_MEMHOSTALLOC_DEVICEMAP: c_uint = 0x2; +pub const CU_MEMHOSTALLOC_WRITECOMBINED: c_uint = 0x4; + +/// Shared-library names tried, in order, when loading the driver. +const DRIVER_LIBRARIES: &[&str] = &["libcuda.so.1", "libcuda.so", "nvcuda.dll"]; + +/// Opens the driver library, if any of its names resolve. +fn open_library() -> Option { + DRIVER_LIBRARIES.iter().find_map(|name| { + // SAFETY: Loading the CUDA driver runs only its own initializers, which have no + // preconditions on the calling process. + unsafe { Library::new(name) }.ok() + }) +} + +/// Whether the CUDA driver library can be loaded on this host. +/// +/// # Safety +/// +/// Loads a shared library, running its initializers. +pub unsafe fn is_culib_present() -> bool { + driver().is_ok() +} + +macro_rules! driver_api { + ($( fn $name:ident ( $( $arg:ident : $ty:ty ),* $(,)? ); )*) => { + /// Resolved driver entry points, kept alive with their library. + struct Driver { + $( $name: unsafe extern "C" fn( $( $ty ),* ) -> CUresult, )* + _library: Library, + } + + impl Driver { + fn load(library: Library) -> Result { + // SAFETY: Each symbol is declared with the signature of the driver header. + unsafe { + Ok(Self { + $( $name: *library + .get:: CUresult>( + concat!(stringify!($name), "\0").as_bytes(), + ) + .map_err(|error| format!("{}: {error}", stringify!($name)))?, )* + _library: library, + }) + } + } + } + + $( + /// Calls the driver function of the same name. + /// + /// # Safety + /// + /// The arguments must satisfy the driver API's requirements for this function. + /// + /// # Panics + /// + /// Panics if the CUDA driver library cannot be loaded. + #[inline] + pub unsafe fn $name( $( $arg: $ty ),* ) -> CUresult { + let driver = match driver() { + Ok(driver) => driver, + Err(error) => panic!("CUDA driver unavailable: {error}"), + }; + // SAFETY: Forwarded under the caller's guarantees. + unsafe { (driver.$name)( $( $arg ),* ) } + } + )* + }; +} + +driver_api! { + fn cuInit(flags: c_uint); + fn cuDeviceGet(device: *mut CUdevice, ordinal: c_int); + fn cuDeviceGetCount(count: *mut c_int); + fn cuDeviceGetAttribute(value: *mut c_int, attribute: CUdevice_attribute, device: CUdevice); + fn cuDeviceGetName(name: *mut c_char, len: c_int, device: CUdevice); + fn cuDeviceGetMemPool(pool: *mut CUmemoryPool, device: CUdevice); + fn cuDevicePrimaryCtxRetain(context: *mut CUcontext, device: CUdevice); + fn cuDevicePrimaryCtxRelease_v2(device: CUdevice); + fn cuCtxCreate_v2(context: *mut CUcontext, flags: c_uint, device: CUdevice); + fn cuCtxDestroy_v2(context: CUcontext); + fn cuCtxGetCurrent(context: *mut CUcontext); + fn cuCtxSetCurrent(context: CUcontext); + fn cuCtxSynchronize(); + fn cuCtxGetStreamPriorityRange(least: *mut c_int, greatest: *mut c_int); + fn cuMemGetInfo_v2(free: *mut usize, total: *mut usize); + fn cuStreamCreate(stream: *mut CUstream, flags: c_uint); + fn cuStreamCreateWithPriority(stream: *mut CUstream, flags: c_uint, priority: c_int); + fn cuStreamDestroy_v2(stream: CUstream); + fn cuStreamSynchronize(stream: CUstream); + fn cuStreamWaitEvent(stream: CUstream, event: CUevent, flags: c_uint); + fn cuLaunchHostFunc(stream: CUstream, function: CUhostFn, user_data: *mut c_void); + fn cuEventCreate(event: *mut CUevent, flags: c_uint); + fn cuEventDestroy_v2(event: CUevent); + fn cuEventRecord(event: CUevent, stream: CUstream); + fn cuEventSynchronize(event: CUevent); + fn cuEventQuery(event: CUevent); + fn cuEventElapsedTime(milliseconds: *mut f32, start: CUevent, end: CUevent); + fn cuMemAllocAsync(pointer: *mut CUdeviceptr, bytes: usize, stream: CUstream); + fn cuMemFreeAsync(pointer: CUdeviceptr, stream: CUstream); + fn cuMemAlloc_v2(pointer: *mut CUdeviceptr, bytes: usize); + fn cuMemFree_v2(pointer: CUdeviceptr); + fn cuMemsetD8Async(target: CUdeviceptr, value: u8, count: usize, stream: CUstream); + fn cuMemsetD32Async(target: CUdeviceptr, value: c_uint, count: usize, stream: CUstream); + fn cuMemcpyHtoDAsync_v2( + target: CUdeviceptr, + source: *const c_void, + bytes: usize, + stream: CUstream, + ); + fn cuMemcpyDtoHAsync_v2( + target: *mut c_void, + source: CUdeviceptr, + bytes: usize, + stream: CUstream, + ); + fn cuMemcpyDtoDAsync_v2( + target: CUdeviceptr, + source: CUdeviceptr, + bytes: usize, + stream: CUstream, + ); + fn cuMemcpyDtoH_v2(target: *mut c_void, source: CUdeviceptr, bytes: usize); + fn cuMemHostAlloc(pointer: *mut *mut c_void, bytes: usize, flags: c_uint); + fn cuMemFreeHost(pointer: *mut c_void); + fn cuMemPoolSetAttribute( + pool: CUmemoryPool, + attribute: CUmemPool_attribute, + value: *mut c_void, + ); + fn cuMemPoolGetAttribute( + pool: CUmemoryPool, + attribute: CUmemPool_attribute, + value: *mut c_void, + ); + fn cuModuleLoadData(module: *mut CUmodule, image: *const c_void); + fn cuModuleUnload(module: CUmodule); + fn cuModuleGetFunction(function: *mut CUfunction, module: CUmodule, name: *const c_char); + fn cuLaunchKernel( + function: CUfunction, + grid_x: c_uint, + grid_y: c_uint, + grid_z: c_uint, + block_x: c_uint, + block_y: c_uint, + block_z: c_uint, + shared_mem_bytes: c_uint, + stream: CUstream, + params: *mut *mut c_void, + extra: *mut *mut c_void, + ); + fn cuGetErrorName(error: CUresult, name: *mut *const c_char); + fn cuGetErrorString(error: CUresult, description: *mut *const c_char); +} + +/// The process-wide driver entry points, loaded on first use. +fn driver() -> Result<&'static Driver, &'static str> { + static DRIVER: OnceLock> = OnceLock::new(); + DRIVER + .get_or_init(|| { + let library = open_library() + .ok_or_else(|| format!("none of {DRIVER_LIBRARIES:?} could be loaded"))?; + Driver::load(library) + }) + .as_ref() + .map_err(String::as_str) +}