feat[cuda]: dict decimal take

vortex-array/src/arrays/decimal/array.rs

@@ -316,6 +316,7 @@ impl DecimalArray {
         }
     }
 
+    /// Return the `DecimalType` used to represent the values in the array.
     pub fn values_type(&self) -> DecimalType {
         self.values_type
     }

vortex-cuda/Cargo.toml

@@ -30,7 +30,7 @@ vortex-alp = { workspace = true }
 vortex-array = { workspace = true }
 vortex-buffer = { workspace = true }
 vortex-decimal-byte-parts = { workspace = true }
-vortex-dtype = { workspace = true }
+vortex-dtype = { workspace = true, features = ["cudarc"] }
 vortex-error = { workspace = true }
 vortex-fastlanes = { workspace = true }
 vortex-nvcomp = { path = "nvcomp" }
@@ -45,6 +45,7 @@ rstest = { workspace = true }
 tokio = { workspace = true, features = ["rt", "macros"] }
 vortex-array = { workspace = true, features = ["_test-harness"] }
 vortex-cuda = { path = ".", features = ["_test-harness"] }
+vortex-dtype = { workspace = true, features = ["cudarc"] }
 
 [build-dependencies]
 

vortex-cuda/benches/dict_cuda.rs

@@ -102,7 +102,7 @@ mod cuda_benchmarks {
         let events = vortex_cuda::launch_cuda_kernel!(
             execution_ctx: cuda_ctx,
             module: "dict",
-            ptypes: &[value_ptype, code_ptype],
+            ptypes: &[value_ptype.to_string().as_str(), code_ptype.to_string().as_str()],
             launch_args: [codes_view, codes_len_u64, values_view, output_view],
             event_recording: CU_EVENT_BLOCKING_SYNC,
             array_len: codes_len

vortex-cuda/benches/for_cuda.rs

@@ -91,7 +91,7 @@ mod cuda_benchmarks {
         let events = vortex_cuda::launch_cuda_kernel!(
             execution_ctx: cuda_ctx,
             module: "for",
-            ptypes: &[for_array.ptype()],
+            ptypes: &[for_array.ptype().to_string().as_str()],
             launch_args: [device_data, reference, array_len_u64],
             event_recording: CU_EVENT_BLOCKING_SYNC,
             array_len: for_array.len()
@@ -112,7 +112,7 @@ mod cuda_benchmarks {
         let events = vortex_cuda::launch_cuda_kernel!(
             execution_ctx: cuda_ctx,
             module: "for",
-            ptypes: &[for_array.ptype()],
+            ptypes: &[for_array.ptype().to_string().as_str()],
             launch_args: [device_data, reference, array_len_u64],
             event_recording: CU_EVENT_BLOCKING_SYNC,
             array_len: for_array.len()
@@ -133,7 +133,7 @@ mod cuda_benchmarks {
         let events = vortex_cuda::launch_cuda_kernel!(
             execution_ctx: cuda_ctx,
             module: "for",
-            ptypes: &[for_array.ptype()],
+            ptypes: &[for_array.ptype().to_string().as_str()],
             launch_args: [device_data, reference, array_len_u64],
             event_recording: CU_EVENT_BLOCKING_SYNC,
             array_len: for_array.len()
@@ -154,7 +154,7 @@ mod cuda_benchmarks {
         let events = vortex_cuda::launch_cuda_kernel!(
             execution_ctx: cuda_ctx,
             module: "for",
-            ptypes: &[for_array.ptype()],
+            ptypes: &[for_array.ptype().to_string().as_str()],
             launch_args: [device_data, reference, array_len_u64],
             event_recording: CU_EVENT_BLOCKING_SYNC,
             array_len: for_array.len()

vortex-cuda/kernels/dict.cu

@@ -6,6 +6,7 @@
 #include <stdint.h>
 
 #include "config.cuh"
+#include "types.cuh"
 
 template<typename ValueT, typename IndexT>
 __device__ void dict_kernel(
@@ -55,3 +56,7 @@ GENERATE_DICT_KERNELS_FOR_VALUE(u32, uint32_t)
 GENERATE_DICT_KERNELS_FOR_VALUE(i32, int32_t)
 GENERATE_DICT_KERNELS_FOR_VALUE(u64, uint64_t)
 GENERATE_DICT_KERNELS_FOR_VALUE(i64, int64_t)
+
+// Decimal types (128-bit and 256-bit)
+GENERATE_DICT_KERNELS_FOR_VALUE(i128, int128_t)
+GENERATE_DICT_KERNELS_FOR_VALUE(i256, int256_t)

vortex-cuda/kernels/types.cuh

@@ -0,0 +1,20 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright the Vortex contributors
+
+#ifndef VORTEX_CUDA_TYPES_CUH
+#define VORTEX_CUDA_TYPES_CUH
+
+#include <stdint.h>
+
+// 128-bit signed integer type for decimal values
+struct __align__(16) int128_t {
+    int64_t lo;
+    int64_t hi;
+};
+
+// 256-bit signed integer type for decimal values
+struct __align__(32) int256_t {
+    int64_t parts[4];
+};
+
+#endif // VORTEX_CUDA_TYPES_CUH

vortex-cuda/src/executor.rs

@@ -98,8 +98,42 @@ impl CudaExecutionCtx {
     /// # Errors
     ///
     /// Returns an error if kernel loading fails.
-    pub fn load_function(&self, module_name: &str, ptypes: &[PType]) -> VortexResult<CudaFunction> {
-        self.cuda_session.load_function(module_name, ptypes)
+    pub fn load_function_ptype(
+        &self,
+        module_name: &str,
+        ptypes: &[PType],
+    ) -> VortexResult<CudaFunction> {
+        let type_suffixes: Vec<String> = ptypes.iter().map(|ptype| ptype.to_string()).collect();
+        self.load_function(
+            module_name,
+            type_suffixes
+                .iter()
+                .map(|t| t.as_str())
+                .collect::<Vec<_>>()
+                .as_slice(),
+        )
+    }
+
+    /// Loads a CUDA kernel function by module name and type suffixes.
+    ///
+    /// This is a lower-level version of `load_function` that accepts string suffixes
+    /// directly, useful for types that don't have a `PType` (e.g., i128, i256).
+    ///
+    /// # Arguments
+    ///
+    /// * `module_name` - Name of the module (`kernels/{module_name}.ptx`)
+    /// * `type_suffixes` - List of type suffix strings for the kernel name
+    ///
+    /// # Errors
+    ///
+    /// Returns an error if kernel loading fails.
+    pub fn load_function(
+        &self,
+        module_name: &str,
+        type_suffixes: &[&str],
+    ) -> VortexResult<CudaFunction> {
+        self.cuda_session
+            .load_function_with_suffixes(module_name, type_suffixes)
     }
 
     /// Returns a launch builder for a CUDA kernel function.