diff --git a/cpp/src/bitmask/null_mask.cu b/cpp/src/bitmask/null_mask.cu
index cdbbb64660cf..4add660bd18f 100644
--- a/cpp/src/bitmask/null_mask.cu
+++ b/cpp/src/bitmask/null_mask.cu
@@ -495,6 +495,7 @@ std::vector<size_type> batch_count_set_bits(host_span<bitmask_type const* const>
     dim3{static_cast<unsigned int>(grid.num_blocks), static_cast<unsigned int>(num_bitmasks), 1};
   count_set_bits_kernel<block_size><<<kernel_grid, block_size, 0, stream.value()>>>(
     d_bitmasks, start, stop - 1, d_non_zero_count.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   // Use pinned memory to copy the result back to the host, then copy again to the output vector.
   auto h_non_zero_count = cudf::detail::make_pinned_vector<size_type>(num_bitmasks, stream);
@@ -807,6 +808,7 @@ size_type index_of_first_set_bit(bitmask_type const* bitmask,
   find_first_set_bit_kernel<block_size>
     <<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
       bitmask, start, stop, bit_count, d_index.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
   return d_index.value(stream);
 }
 
diff --git a/cpp/src/copying/concatenate.cu b/cpp/src/copying/concatenate.cu
index 8af2b12249b6..6aebe179ec31 100644
--- a/cpp/src/copying/concatenate.cu
+++ b/cpp/src/copying/concatenate.cu
@@ -165,6 +165,7 @@ size_type concatenate_masks(device_span<column_device_view const> d_views,
       dest_mask,
       output_size,
       d_valid_count.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
   return output_size - d_valid_count.value(stream);
 }
 
@@ -271,6 +272,7 @@ std::unique_ptr<column> fused_concatenate(host_span<column_view const> views,
     static_cast<size_type>(d_views.size()),
     *d_out_view,
     d_valid_count.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   if (has_nulls) {
     out_col->set_null_count(output_size - d_valid_count.value(stream));
diff --git a/cpp/src/copying/contiguous_split.cu b/cpp/src/copying/contiguous_split.cu
index 4795460c4fed..2c3a9a191514 100644
--- a/cpp/src/copying/contiguous_split.cu
+++ b/cpp/src/copying/contiguous_split.cu
@@ -1759,6 +1759,7 @@ void copy_data(int num_batches_to_copy,
     auto index_to_buffer = [user_buffer] __device__(unsigned int) { return user_buffer; };
     copy_partitions<block_size><<<num_batches_to_copy, block_size, 0, stream.value()>>>(
       index_to_buffer, d_src_bufs, d_dst_buf_info.data() + starting_batch);
+    CUDF_CUDA_TRY(cudaGetLastError());
   } else {
     auto index_to_buffer = [d_dst_bufs,
                             dst_buf_info = d_dst_buf_info.data(),
@@ -1768,6 +1769,7 @@ void copy_data(int num_batches_to_copy,
     };
     copy_partitions<block_size><<<num_batches_to_copy, block_size, 0, stream.value()>>>(
       index_to_buffer, d_src_bufs, d_dst_buf_info.data() + starting_batch);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 }
 
diff --git a/cpp/src/copying/scatter.cu b/cpp/src/copying/scatter.cu
index 364347e58e63..23396b48bc16 100644
--- a/cpp/src/copying/scatter.cu
+++ b/cpp/src/copying/scatter.cu
@@ -88,6 +88,7 @@ void scatter_scalar_bitmask_inplace(std::reference_wrapper<scalar const> const&
                                           : marking_bitmask_kernel<false, decltype(scatter_map)>;
     bitmask_kernel<<<grid_size, block_size, 0, stream.value()>>>(
       *target_view, scatter_map, num_scatter_rows);
+    CUDF_CUDA_TRY(cudaGetLastError());
 
     target.set_null_count(
       cudf::detail::null_count(target.view().null_mask(), 0, target.size(), stream));
diff --git a/cpp/src/groupby/hash/compute_mapping_indices.cuh b/cpp/src/groupby/hash/compute_mapping_indices.cuh
index 43b66e70f58f..bdf42b1c72fa 100644
--- a/cpp/src/groupby/hash/compute_mapping_indices.cuh
+++ b/cpp/src/groupby/hash/compute_mapping_indices.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -179,5 +179,6 @@ void compute_mapping_indices(size_type grid_size,
     global_mapping_indices,
     block_cardinality,
     needs_global_memory_fallback);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 }  // namespace cudf::groupby::detail::hash
diff --git a/cpp/src/groupby/hash/compute_shared_memory_aggs.cu b/cpp/src/groupby/hash/compute_shared_memory_aggs.cu
index 7c8d5389d78f..810eb38c06a5 100644
--- a/cpp/src/groupby/hash/compute_shared_memory_aggs.cu
+++ b/cpp/src/groupby/hash/compute_shared_memory_aggs.cu
@@ -315,5 +315,6 @@ void compute_shared_memory_aggs(cudf::size_type grid_size,
     d_agg_kinds,
     shmem_agg_size,
     offsets_size);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 }  // namespace cudf::groupby::detail::hash
diff --git a/cpp/src/io/avro/avro_gpu.cu b/cpp/src/io/avro/avro_gpu.cu
index 1ad78947633c..4a0929c0484e 100644
--- a/cpp/src/io/avro/avro_gpu.cu
+++ b/cpp/src/io/avro/avro_gpu.cu
@@ -424,6 +424,7 @@ void DecodeAvroColumnData(device_span<block_desc_s const> blocks,
 
   gpuDecodeAvroColumnData<<<dim_grid, dim_block, 0, stream.value()>>>(
     blocks, schema, global_dictionary, avro_data, schema_len, min_row_size);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 }  // namespace gpu
diff --git a/cpp/src/io/comp/debrotli.cu b/cpp/src/io/comp/debrotli.cu
index e31c1083dc60..1e6903007498 100644
--- a/cpp/src/io/comp/debrotli.cu
+++ b/cpp/src/io/comp/debrotli.cu
@@ -2102,6 +2102,7 @@ void gpu_debrotli(device_span<device_span<uint8_t const> const> inputs,
     scratch.data() + fb_heap_size, get_brotli_dictionary(), sizeof(brotli_dictionary_s), stream));
   gpu_debrotli_kernel<<<dim_grid, dim_block, 0, stream.value()>>>(
     inputs, outputs, results, scratch.data(), fb_heap_size);
+  CUDF_CUDA_TRY(cudaGetLastError());
 #if DUMP_FB_HEAP
   uint32_t dump[2];
   uint32_t cur = 0;
diff --git a/cpp/src/io/comp/gpuinflate.cu b/cpp/src/io/comp/gpuinflate.cu
index 2462734da949..a0cbfcf63ccb 100644
--- a/cpp/src/io/comp/gpuinflate.cu
+++ b/cpp/src/io/comp/gpuinflate.cu
@@ -1370,6 +1370,7 @@ void gpuinflate(device_span<device_span<uint8_t const> const> inputs,
   if (inputs.size() > 0) {
     inflate_kernel_no_racecheck<block_size>
       <<<inputs.size(), block_size, 0, stream.value()>>>(inputs, outputs, results, parse_hdr);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 }
 
@@ -1380,6 +1381,7 @@ void gpu_copy_uncompressed_blocks(device_span<device_span<uint8_t const> const>
   constexpr auto block_size = 1024;
   if (inputs.size() > 0) {
     copy_uncompressed_kernel<<<inputs.size(), block_size, 0, stream.value()>>>(inputs, outputs);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 }
 
diff --git a/cpp/src/io/comp/snap.cu b/cpp/src/io/comp/snap.cu
index 3a11c67b38a1..db27c7a0b734 100644
--- a/cpp/src/io/comp/snap.cu
+++ b/cpp/src/io/comp/snap.cu
@@ -319,6 +319,7 @@ void gpu_snap(device_span<device_span<uint8_t const> const> inputs,
   dim3 dim_grid(inputs.size(), 1);
   if (inputs.size() > 0) {
     snap_kernel_no_racecheck<<<dim_grid, dim_block, 0, stream.value()>>>(inputs, outputs, results);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 }
 
diff --git a/cpp/src/io/comp/unsnap.cu b/cpp/src/io/comp/unsnap.cu
index b078618f6254..d21667221eac 100644
--- a/cpp/src/io/comp/unsnap.cu
+++ b/cpp/src/io/comp/unsnap.cu
@@ -713,14 +713,17 @@ void gpu_unsnap(device_span<device_span<uint8_t const> const> inputs,
                 device_span<codec_exec_result> results,
                 rmm::cuda_stream_view stream)
 {
+  if (inputs.empty()) { return; }
+
   dim3 dim_block(128, 1);           // 4 warps per stream, 1 stream per block
   dim3 dim_grid(inputs.size(), 1);  // TODO: Check max grid dimensions vs max expected count
 
   unsnap_kernel_no_racecheck<128>
     <<<dim_grid, dim_block, 0, stream.value()>>>(inputs, outputs, results);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
-__global__ void get_snappy_uncompressed_size_kernel(
+CUDF_KERNEL void get_snappy_uncompressed_size_kernel(
   device_span<device_span<uint8_t const> const> inputs, device_span<size_t> uncompressed_sizes)
 {
   auto const idx = cudf::detail::grid_1d::global_thread_id();
@@ -750,12 +753,15 @@ void get_snappy_uncompressed_size(device_span<device_span<uint8_t const> const>
                                   device_span<size_t> uncompressed_sizes,
                                   rmm::cuda_stream_view stream)
 {
+  if (inputs.empty()) { return; }
+
   int threads_per_block = 128;
   auto const num_blocks =
     cudf::util::div_rounding_up_safe<size_t>(inputs.size(), threads_per_block);
 
   get_snappy_uncompressed_size_kernel<<<num_blocks, threads_per_block, 0, stream.value()>>>(
     inputs, uncompressed_sizes);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 }  // namespace cudf::io::detail
diff --git a/cpp/src/io/csv/csv_gpu.cu b/cpp/src/io/csv/csv_gpu.cu
index 90a112a9a436..ef24422400b9 100644
--- a/cpp/src/io/csv/csv_gpu.cu
+++ b/cpp/src/io/csv/csv_gpu.cu
@@ -854,6 +854,7 @@ cudf::detail::host_vector<column_type_histogram> detect_column_types(
 
   data_type_detection<<<grid_size, block_size, 0, stream.value()>>>(
     options, data, column_flags, row_starts, d_stats);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   return cudf::detail::make_host_vector(d_stats, stream);
 }
@@ -883,6 +884,7 @@ void decode_row_column_data(cudf::io::parse_options_view const& options,
                                                                     valids,
                                                                     valid_counts,
                                                                     is_quoted_flags);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 uint32_t __host__ gather_row_offsets(parse_options_view const& options,
@@ -918,6 +920,7 @@ uint32_t __host__ gather_row_offsets(parse_options_view const& options,
     (options.quotechar) ? options.quotechar : 0x100,
     /*(options.escapechar) ? options.escapechar :*/ 0x100,
     (options.comment) ? options.comment : 0x100);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   return dim_grid;
 }
diff --git a/cpp/src/io/fst/dispatch_dfa.cuh b/cpp/src/io/fst/dispatch_dfa.cuh
index 5841813a51a0..6d1d98963de6 100644
--- a/cpp/src/io/fst/dispatch_dfa.cuh
+++ b/cpp/src/io/fst/dispatch_dfa.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -462,6 +462,7 @@ struct DispatchFSM : DeviceFSMPolicy {
       uint32_t num_fst_init_blocks    = cuda::ceil_div(num_blocks, FST_INIT_TPB);
       initialization_pass_kernel<<<num_fst_init_blocks, FST_INIT_TPB, 0, stream>>>(
         fst_offset_tile_state, num_blocks);
+      CUDF_CUDA_TRY(cudaGetLastError());
     }
 
     //------------------------------------------------------------------------------
@@ -477,6 +478,7 @@ struct DispatchFSM : DeviceFSMPolicy {
       uint32_t num_stv_init_blocks    = cuda::ceil_div(num_blocks, STV_INIT_TPB);
       initialization_pass_kernel<<<num_stv_init_blocks, STV_INIT_TPB, 0, stream>>>(stv_tile_state,
                                                                                    num_blocks);
+      CUDF_CUDA_TRY(cudaGetLastError());
     } else {
       // Compute state-transition vectors
       // TODO tag dispatch or constexpr if depending on single-pass config to avoid superfluous
diff --git a/cpp/src/io/orc/dict_enc.cu b/cpp/src/io/orc/dict_enc.cu
index 0a45d9c5f7b6..2b8c474e0ee2 100644
--- a/cpp/src/io/orc/dict_enc.cu
+++ b/cpp/src/io/orc/dict_enc.cu
@@ -64,6 +64,7 @@ void rowgroup_char_counts(device_2dspan<size_type> counts,
 
   rowgroup_char_counts_kernel<<<num_blocks, block_size, 0, stream.value()>>>(
     counts, orc_columns, rowgroup_bounds, str_col_indexes);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 struct equality_functor {
@@ -231,6 +232,7 @@ void populate_dictionary_hash_maps(device_2dspan<stripe_dictionary> dictionaries
   constexpr int block_size = 256;
   populate_dictionary_hash_maps_kernel<block_size>
     <<<dictionaries.count(), block_size, 0, stream.value()>>>(dictionaries, columns);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void collect_map_entries(device_2dspan<stripe_dictionary> dictionaries,
@@ -240,6 +242,7 @@ void collect_map_entries(device_2dspan<stripe_dictionary> dictionaries,
   constexpr int block_size = 1024;
   collect_map_entries_kernel<block_size>
     <<<dictionaries.count(), block_size, 0, stream.value()>>>(dictionaries);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void get_dictionary_indices(device_2dspan<stripe_dictionary> dictionaries,
@@ -250,6 +253,7 @@ void get_dictionary_indices(device_2dspan<stripe_dictionary> dictionaries,
   constexpr int block_size = 1024;
   get_dictionary_indices_kernel<block_size>
     <<<dictionaries.count(), block_size, 0, stream.value()>>>(dictionaries, columns);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 }  // namespace cudf::io::orc::detail
diff --git a/cpp/src/io/orc/stats_enc.cu b/cpp/src/io/orc/stats_enc.cu
index c2423032c682..5b8fdff7adc5 100644
--- a/cpp/src/io/orc/stats_enc.cu
+++ b/cpp/src/io/orc/stats_enc.cu
@@ -451,6 +451,7 @@ void orc_init_statistics_groups(statistics_group* groups,
   dim3 dim_block(init_threads_per_group, init_groups_per_block);
   gpu_init_statistics_groups<<<num_blocks, dim_block, 0, stream.value()>>>(
     groups, cols, rowgroup_bounds);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 /**
@@ -468,6 +469,7 @@ void orc_init_statistics_buffersize(statistics_merge_group* groups,
 {
   gpu_init_statistics_buffersize<block_size>
     <<<1, block_size, 0, stream.value()>>>(groups, chunks, statistics_count);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 /**
@@ -490,6 +492,7 @@ void orc_encode_statistics(uint8_t* blob_bfr,
   dim3 dim_block(encode_threads_per_chunk, encode_chunks_per_block);
   gpu_encode_statistics<<<num_blocks, dim_block, 0, stream.value()>>>(
     blob_bfr, groups, chunks, statistics_count);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 }  // namespace cudf::io::orc::detail
diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu
index f4e8c5b5781a..8c7789ae9c6c 100644
--- a/cpp/src/io/orc/stripe_data.cu
+++ b/cpp/src/io/orc/stripe_data.cu
@@ -2072,6 +2072,7 @@ void __host__ decode_nulls_and_string_dictionaries(column_desc* chunks,
   decode_nulls_and_string_dictionaries_kernel<block_size>
     <<<dim_grid, block_size, 0, stream.value()>>>(
       chunks, global_dictionary, num_columns, num_stripes, first_row);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 /**
@@ -2106,6 +2107,7 @@ void __host__ decode_column_data(column_desc* chunks,
   auto const num_blocks = num_columns * (num_rowgroups > 0 ? num_rowgroups : num_stripes);
   decode_column_data_kernel<block_size><<<num_blocks, block_size, 0, stream.value()>>>(
     chunks, global_dictionary, tz_table, row_groups, first_row, rowidx_stride, level, error_count);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 }  // namespace cudf::io::orc::detail
diff --git a/cpp/src/io/orc/stripe_enc.cu b/cpp/src/io/orc/stripe_enc.cu
index ff143de6bfd5..1a18a2dcc1ea 100644
--- a/cpp/src/io/orc/stripe_enc.cu
+++ b/cpp/src/io/orc/stripe_enc.cu
@@ -1304,6 +1304,7 @@ void encode_orc_column_data(device_2dspan<encoder_chunk const> chunks,
   auto const num_blocks = chunks.size().first * chunks.size().second;
   encode_column_data_kernel<encode_block_size>
     <<<num_blocks, encode_block_size, 0, stream.value()>>>(chunks, streams);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void encode_stripe_dictionaries(stripe_dictionary const* stripes,
@@ -1318,6 +1319,7 @@ void encode_stripe_dictionaries(stripe_dictionary const* stripes,
   dim3 dim_grid(num_string_columns * num_stripes, 2);
   encode_string_dictionaries_kernel<block_size>
     <<<dim_grid, block_size, 0, stream.value()>>>(stripes, columns, chunks, enc_streams);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void compact_orc_data_streams(device_2dspan<stripe_stream> strm_desc,
@@ -1340,6 +1342,7 @@ void compact_orc_data_streams(device_2dspan<stripe_stream> strm_desc,
     strm_desc.size().second;
   init_batched_memcpy_kernel<<<num_blocks, compact_streams_block_size, 0, stream.value()>>>(
     strm_desc, enc_streams, srcs, dsts, lengths);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   // Copy streams in a batched manner.
   cudf::detail::batched_memcpy_async(
@@ -1372,11 +1375,13 @@ std::optional<writer_compression_statistics> compress_orc_data_streams(
                                                                          comp_blk_size,
                                                                          max_comp_blk_size,
                                                                          comp_block_align);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   cudf::io::detail::compress(compression, comp_in, comp_out, comp_res, stream);
 
   compact_compressed_blocks_kernel<<<num_blocks, 1024, 0, stream.value()>>>(
     strm_desc, comp_in, comp_out, comp_res, compressed_data, comp_blk_size, max_comp_blk_size);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   if (collect_statistics) {
     return cudf::io::detail::collect_compression_statistics(comp_in, comp_res, stream);
@@ -1407,6 +1412,7 @@ void decimal_sizes_to_offsets(device_2dspan<rowgroup_rows const> rg_bounds,
   auto const num_blocks = elem_sizes.size() * rg_bounds.size().first;
   decimal_sizes_to_offsets_kernel<block_size>
     <<<num_blocks, block_size, 0, stream.value()>>>(rg_bounds, d_sizes);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 }  // namespace cudf::io::orc::detail
diff --git a/cpp/src/io/orc/stripe_init.cu b/cpp/src/io/orc/stripe_init.cu
index 7f5492a5be7e..d5514ffb5a53 100644
--- a/cpp/src/io/orc/stripe_init.cu
+++ b/cpp/src/io/orc/stripe_init.cu
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -555,6 +555,7 @@ void __host__ parse_compressed_stripe_data(compressed_stream_info* strm_info,
   if (num_blocks > 0) {
     parse_compressed_stripe_data_kernel<<<num_blocks, 128, 0, stream.value()>>>(
       strm_info, num_streams, compression_block_size, log2maxcr);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 }
 
@@ -566,6 +567,7 @@ void __host__ post_decompression_reassemble(compressed_stream_info* strm_info,
   if (num_blocks > 0) {
     post_decompression_reassemble_kernel<<<num_blocks, 128, 0, stream.value()>>>(strm_info,
                                                                                  num_streams);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 }
 
@@ -581,6 +583,7 @@ void __host__ parse_row_group_index(row_group* row_groups,
   auto const num_blocks = num_columns * num_stripes;
   parse_row_group_index_kernel<<<num_blocks, 128, 0, stream.value()>>>(
     row_groups, strm_info, chunks, num_columns, num_stripes, rowidx_stride, use_base_stride);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void __host__ reduce_pushdown_masks(device_span<orc_column_device_view const> columns,
@@ -592,6 +595,7 @@ void __host__ reduce_pushdown_masks(device_span<orc_column_device_view const> co
   constexpr int block_size = 128;
   reduce_pushdown_masks_kernel<block_size>
     <<<num_blocks, block_size, 0, stream.value()>>>(columns, rowgroups, valid_counts);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 }  // namespace cudf::io::orc::detail
diff --git a/cpp/src/io/orc/writer_impl.cu b/cpp/src/io/orc/writer_impl.cu
index 815dddbd4f62..9768c578dd1e 100644
--- a/cpp/src/io/orc/writer_impl.cu
+++ b/cpp/src/io/orc/writer_impl.cu
@@ -428,6 +428,7 @@ void persisted_statistics::persist(int num_table_rows,
         offsets.data(),
         intermediate_stats.stripe_stat_chunks.data(),
         intermediate_stats.stripe_stat_merge.device_ptr());
+      CUDF_CUDA_TRY(cudaGetLastError());
       string_pools.emplace_back(std::move(string_pool));
     }
   }
diff --git a/cpp/src/io/parquet/chunk_dict.cu b/cpp/src/io/parquet/chunk_dict.cu
index e7800da5e5ca..dddd82c5b303 100644
--- a/cpp/src/io/parquet/chunk_dict.cu
+++ b/cpp/src/io/parquet/chunk_dict.cu
@@ -483,6 +483,7 @@ void populate_chunk_hash_maps(device_span<slot_type> const map_storage,
   dim3 const dim_grid(frags.size().second, frags.size().first);
   populate_chunk_hash_maps_kernel<DEFAULT_BLOCK_SIZE>
     <<<dim_grid, DEFAULT_BLOCK_SIZE, 0, stream.value()>>>(map_storage, frags);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void collect_map_entries(device_span<slot_type> const map_storage,
@@ -496,6 +497,7 @@ void collect_map_entries(device_span<slot_type> const map_storage,
                 "each histogram bucket.");
   collect_map_entries_kernel<block_size>
     <<<chunks.size(), block_size, 0, stream.value()>>>(map_storage, chunks, frags);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void get_dictionary_indices(device_span<slot_type> const map_storage,
@@ -505,6 +507,7 @@ void get_dictionary_indices(device_span<slot_type> const map_storage,
   dim3 const dim_grid(frags.size().second, frags.size().first);
   get_dictionary_indices_kernel<DEFAULT_BLOCK_SIZE>
     <<<dim_grid, DEFAULT_BLOCK_SIZE, 0, stream.value()>>>(map_storage, frags);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void compute_per_page_dict_bits(device_span<EncPage> pages, rmm::cuda_stream_view stream)
@@ -514,6 +517,7 @@ void compute_per_page_dict_bits(device_span<EncPage> pages, rmm::cuda_stream_vie
   auto const num_blocks =
     cudf::util::div_rounding_up_safe(static_cast<size_type>(pages.size()), warps_per_block);
   compute_page_dict_bits_kernel<<<num_blocks, DEFAULT_BLOCK_SIZE, 0, stream.value()>>>(pages);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 }  // namespace cudf::io::parquet::detail
diff --git a/cpp/src/io/parquet/decode_fixed.cu b/cpp/src/io/parquet/decode_fixed.cu
index 9b649cc32664..b002598c82ab 100644
--- a/cpp/src/io/parquet/decode_fixed.cu
+++ b/cpp/src/io/parquet/decode_fixed.cu
@@ -1296,6 +1296,7 @@ void decode_page_data(cudf::detail::hostdevice_span<PageInfo> pages,
                                                      initial_str_offsets,
                                                      page_string_offset_indices,
                                                      error_code);
+      CUDF_CUDA_TRY(cudaGetLastError());
     } else {
       decode_page_data_generic<uint16_t, decode_block_size, mask>
         <<<dim_grid, dim_block, 0, stream.value()>>>(pages.device_ptr(),
@@ -1306,6 +1307,7 @@ void decode_page_data(cudf::detail::hostdevice_span<PageInfo> pages,
                                                      initial_str_offsets,
                                                      page_string_offset_indices,
                                                      error_code);
+      CUDF_CUDA_TRY(cudaGetLastError());
     }
   };
 
diff --git a/cpp/src/io/parquet/decode_preprocess.cu b/cpp/src/io/parquet/decode_preprocess.cu
index cc9549166c9d..4aa40f3cd1ab 100644
--- a/cpp/src/io/parquet/decode_preprocess.cu
+++ b/cpp/src/io/parquet/decode_preprocess.cu
@@ -474,6 +474,8 @@ void compute_page_sizes(cudf::detail::hostdevice_span<PageInfo> pages,
 {
   CUDF_FUNC_RANGE();
 
+  if (pages.size() == 0) { return; }
+
   dim3 dim_block(preprocess_block_size, 1);
   dim3 dim_grid(pages.size(), 1);  // 1 threadblock per page
 
@@ -485,9 +487,11 @@ void compute_page_sizes(cudf::detail::hostdevice_span<PageInfo> pages,
   if (level_type_size == 1) {
     compute_page_sizes_kernel<uint8_t><<<dim_grid, dim_block, 0, stream.value()>>>(
       pages.device_ptr(), chunks, page_mask, min_row, num_rows, compute_num_rows);
+    CUDF_CUDA_TRY(cudaGetLastError());
   } else {
     compute_page_sizes_kernel<uint16_t><<<dim_grid, dim_block, 0, stream.value()>>>(
       pages.device_ptr(), chunks, page_mask, min_row, num_rows, compute_num_rows);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 }
 
@@ -513,10 +517,12 @@ void preprocess_levels(cudf::detail::hostdevice_span<PageInfo> pages,
     preprocess_levels_kernel<uint8_t, level_decode_block_size>
       <<<dim_grid, dim_block, 0, stream.value()>>>(
         pages.device_ptr(), chunks, page_mask, min_row, num_rows);
+    CUDF_CUDA_TRY(cudaGetLastError());
   } else {
     preprocess_levels_kernel<uint16_t, level_decode_block_size>
       <<<dim_grid, dim_block, 0, stream.value()>>>(
         pages.device_ptr(), chunks, page_mask, min_row, num_rows);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 }
 
diff --git a/cpp/src/io/parquet/experimental/dictionary_page_filter.cu b/cpp/src/io/parquet/experimental/dictionary_page_filter.cu
index 3dc0bb817cc7..d535a11273cb 100644
--- a/cpp/src/io/parquet/experimental/dictionary_page_filter.cu
+++ b/cpp/src/io/parquet/experimental/dictionary_page_filter.cu
@@ -1093,6 +1093,7 @@ struct dictionary_caster {
                                                                      num_dictionary_columns,
                                                                      dictionary_col_idx,
                                                                      error_code.data());
+      CUDF_CUDA_TRY(cudaGetLastError());
 
     } else {
       // Check if the decode block size is a multiple of the warp size
@@ -1119,6 +1120,7 @@ struct dictionary_caster {
         num_dictionary_columns,
         dictionary_col_idx,
         error_code.data());
+      CUDF_CUDA_TRY(cudaGetLastError());
     }
 
     // Check if there are any errors in data decoding
@@ -1158,6 +1160,7 @@ struct dictionary_caster {
                                                                     value_offsets.data(),
                                                                     total_row_groups,
                                                                     physical_type);
+    CUDF_CUDA_TRY(cudaGetLastError());
 
     // Build the BOOL8 columns from the results buffers
     return build_columns(results_buffers, stream, mr);
@@ -1233,6 +1236,7 @@ struct dictionary_caster {
                                                                      num_dictionary_columns,
                                                                      dictionary_col_idx,
                                                                      error_code.data());
+      CUDF_CUDA_TRY(cudaGetLastError());
     } else {
       static_assert(DECODE_BLOCK_SIZE % cudf::detail::warp_size == 0,
                     "decoder block size must be a multiple of warp_size");
@@ -1256,6 +1260,7 @@ struct dictionary_caster {
         num_dictionary_columns,
         dictionary_col_idx,
         error_code.data());
+      CUDF_CUDA_TRY(cudaGetLastError());
     }
 
     // Check if there are any errors in data decoding
diff --git a/cpp/src/io/parquet/page_data.cu b/cpp/src/io/parquet/page_data.cu
index 18d26c516e52..5b2f854fce08 100644
--- a/cpp/src/io/parquet/page_data.cu
+++ b/cpp/src/io/parquet/page_data.cu
@@ -564,9 +564,11 @@ void decode_page_data(cudf::detail::hostdevice_span<PageInfo> pages,
   if (level_type_size == 1) {
     decode_page_data<rolling_buf_size, uint8_t><<<dim_grid, dim_block, 0, stream.value()>>>(
       pages.device_ptr(), chunks, min_row, num_rows, page_mask, error_code);
+    CUDF_CUDA_TRY(cudaGetLastError());
   } else {
     decode_page_data<rolling_buf_size, uint16_t><<<dim_grid, dim_block, 0, stream.value()>>>(
       pages.device_ptr(), chunks, min_row, num_rows, page_mask, error_code);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 }
 
@@ -591,10 +593,12 @@ void decode_split_page_data(cudf::detail::hostdevice_span<PageInfo> pages,
     decode_split_page_data_kernel<rolling_buf_size, uint8_t>
       <<<dim_grid, dim_block, 0, stream.value()>>>(
         pages.device_ptr(), chunks, min_row, num_rows, page_mask, error_code);
+    CUDF_CUDA_TRY(cudaGetLastError());
   } else {
     decode_split_page_data_kernel<rolling_buf_size, uint16_t>
       <<<dim_grid, dim_block, 0, stream.value()>>>(
         pages.device_ptr(), chunks, min_row, num_rows, page_mask, error_code);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 }
 
diff --git a/cpp/src/io/parquet/page_delta_decode.cu b/cpp/src/io/parquet/page_delta_decode.cu
index 82ac47391811..acdb0840f01d 100644
--- a/cpp/src/io/parquet/page_delta_decode.cu
+++ b/cpp/src/io/parquet/page_delta_decode.cu
@@ -925,9 +925,11 @@ void decode_delta_binary(cudf::detail::hostdevice_span<PageInfo> pages,
   if (level_type_size == 1) {
     decode_delta_binary_kernel<uint8_t><<<dim_grid, dim_block, 0, stream.value()>>>(
       pages.device_ptr(), chunks, min_row, num_rows, page_mask, error_code);
+    CUDF_CUDA_TRY(cudaGetLastError());
   } else {
     decode_delta_binary_kernel<uint16_t><<<dim_grid, dim_block, 0, stream.value()>>>(
       pages.device_ptr(), chunks, min_row, num_rows, page_mask, error_code);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 }
 
@@ -952,9 +954,11 @@ void decode_delta_byte_array(cudf::detail::hostdevice_span<PageInfo> pages,
   if (level_type_size == 1) {
     decode_delta_byte_array_kernel<uint8_t><<<dim_grid, dim_block, 0, stream.value()>>>(
       pages.device_ptr(), chunks, min_row, num_rows, page_mask, initial_str_offsets, error_code);
+    CUDF_CUDA_TRY(cudaGetLastError());
   } else {
     decode_delta_byte_array_kernel<uint16_t><<<dim_grid, dim_block, 0, stream.value()>>>(
       pages.device_ptr(), chunks, min_row, num_rows, page_mask, initial_str_offsets, error_code);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 }
 
@@ -979,9 +983,11 @@ void decode_delta_length_byte_array(cudf::detail::hostdevice_span<PageInfo> page
   if (level_type_size == 1) {
     decode_delta_length_byte_array_kernel<uint8_t><<<dim_grid, dim_block, 0, stream.value()>>>(
       pages.device_ptr(), chunks, min_row, num_rows, page_mask, initial_str_offsets, error_code);
+    CUDF_CUDA_TRY(cudaGetLastError());
   } else {
     decode_delta_length_byte_array_kernel<uint16_t><<<dim_grid, dim_block, 0, stream.value()>>>(
       pages.device_ptr(), chunks, min_row, num_rows, page_mask, initial_str_offsets, error_code);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 }
 
diff --git a/cpp/src/io/parquet/page_enc.cu b/cpp/src/io/parquet/page_enc.cu
index bf0edd14dc10..0f635b8235fb 100644
--- a/cpp/src/io/parquet/page_enc.cu
+++ b/cpp/src/io/parquet/page_enc.cu
@@ -3435,6 +3435,7 @@ void InitRowGroupFragments(device_2dspan<PageFragment> frag,
   dim3 const dim_grid(num_columns, grid_y);  // 1 threadblock per fragment
   gpuInitRowGroupFragments<512><<<dim_grid, 512, 0, stream.value()>>>(
     frag, col_desc, partitions, part_frag_offset, fragment_size);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void CalculatePageFragments(device_span<PageFragment> frag,
@@ -3442,6 +3443,7 @@ void CalculatePageFragments(device_span<PageFragment> frag,
                             rmm::cuda_stream_view stream)
 {
   gpuCalculatePageFragments<512><<<frag.size(), 512, 0, stream.value()>>>(frag, column_frag_sizes);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void InitFragmentStatistics(device_span<statistics_group> groups,
@@ -3452,6 +3454,7 @@ void InitFragmentStatistics(device_span<statistics_group> groups,
   int const dim =
     util::div_rounding_up_safe(num_fragments, encode_block_size / cudf::detail::warp_size);
   gpuInitFragmentStats<<<dim, encode_block_size, 0, stream.value()>>>(groups, fragments);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void InitEncoderPages(device_2dspan<EncColumnChunk> chunks,
@@ -3482,6 +3485,7 @@ void InitEncoderPages(device_2dspan<EncColumnChunk> chunks,
                                                                    max_page_size_rows,
                                                                    page_align,
                                                                    write_v2_headers);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void EncodePages(device_span<EncPage> pages,
@@ -3509,43 +3513,55 @@ void EncodePages(device_span<EncPage> pages,
     auto const strm = streams[s_idx++];
     gpuEncodePageLevels<encode_block_size><<<num_pages, encode_block_size, 0, strm.value()>>>(
       pages, write_v2_headers, encode_kernel_mask::PLAIN);
+    CUDF_CUDA_TRY(cudaGetLastError());
     gpuEncodePages<encode_block_size><<<num_pages, encode_block_size, 0, strm.value()>>>(
       pages, comp_in, comp_out, comp_results, write_v2_headers, false);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
   if (BitAnd(kernel_mask, encode_kernel_mask::BYTE_STREAM_SPLIT) != 0) {
     auto const strm = streams[s_idx++];
     gpuEncodePageLevels<encode_block_size><<<num_pages, encode_block_size, 0, strm.value()>>>(
       pages, write_v2_headers, encode_kernel_mask::BYTE_STREAM_SPLIT);
+    CUDF_CUDA_TRY(cudaGetLastError());
     gpuEncodePages<encode_block_size><<<num_pages, encode_block_size, 0, strm.value()>>>(
       pages, comp_in, comp_out, comp_results, write_v2_headers, true);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
   if (BitAnd(kernel_mask, encode_kernel_mask::DELTA_BINARY) != 0) {
     auto const strm = streams[s_idx++];
     gpuEncodePageLevels<encode_block_size><<<num_pages, encode_block_size, 0, strm.value()>>>(
       pages, write_v2_headers, encode_kernel_mask::DELTA_BINARY);
+    CUDF_CUDA_TRY(cudaGetLastError());
     gpuEncodeDeltaBinaryPages<encode_block_size>
       <<<num_pages, encode_block_size, 0, strm.value()>>>(pages, comp_in, comp_out, comp_results);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
   if (BitAnd(kernel_mask, encode_kernel_mask::DELTA_LENGTH_BA) != 0) {
     auto const strm = streams[s_idx++];
     gpuEncodePageLevels<encode_block_size><<<num_pages, encode_block_size, 0, strm.value()>>>(
       pages, write_v2_headers, encode_kernel_mask::DELTA_LENGTH_BA);
+    CUDF_CUDA_TRY(cudaGetLastError());
     gpuEncodeDeltaLengthByteArrayPages<encode_block_size>
       <<<num_pages, encode_block_size, 0, strm.value()>>>(pages, comp_in, comp_out, comp_results);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
   if (BitAnd(kernel_mask, encode_kernel_mask::DELTA_BYTE_ARRAY) != 0) {
     auto const strm = streams[s_idx++];
     gpuEncodePageLevels<encode_block_size><<<num_pages, encode_block_size, 0, strm.value()>>>(
       pages, write_v2_headers, encode_kernel_mask::DELTA_BYTE_ARRAY);
+    CUDF_CUDA_TRY(cudaGetLastError());
     gpuEncodeDeltaByteArrayPages<encode_block_size>
       <<<num_pages, encode_block_size, 0, strm.value()>>>(pages, comp_in, comp_out, comp_results);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
   if (BitAnd(kernel_mask, encode_kernel_mask::DICTIONARY) != 0) {
     auto const strm = streams[s_idx++];
     gpuEncodePageLevels<encode_block_size><<<num_pages, encode_block_size, 0, strm.value()>>>(
       pages, write_v2_headers, encode_kernel_mask::DICTIONARY);
+    CUDF_CUDA_TRY(cudaGetLastError());
     gpuEncodeDictPages<encode_block_size><<<num_pages, encode_block_size, 0, strm.value()>>>(
       pages, comp_in, comp_out, comp_results, write_v2_headers);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   cudf::detail::join_streams(streams, stream);
@@ -3559,6 +3575,7 @@ void decide_compression(device_span<EncColumnChunk> chunks,
     util::div_rounding_up_safe<int>(chunks.size(), decide_compression_warps_in_block);
   decide_compression_kernel<<<num_blocks, decide_compression_block_size, 0, stream.value()>>>(
     chunks, page_level_compression);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void EncodePageHeaders(device_span<EncPage> pages,
@@ -3570,11 +3587,13 @@ void EncodePageHeaders(device_span<EncPage> pages,
   auto const num_blocks = util::div_rounding_up_safe<int>(pages.size(), encode_block_size);
   gpuEncodePageHeaders<<<num_blocks, encode_block_size, 0, stream.value()>>>(
     pages, comp_results, page_stats, chunk_stats);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void GatherPages(device_span<EncColumnChunk> chunks, rmm::cuda_stream_view stream)
 {
   gpuGatherPages<<<chunks.size(), 1024, 0, stream.value()>>>(chunks);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void EncodeColumnIndexes(device_span<EncColumnChunk> chunks,
@@ -3584,6 +3603,7 @@ void EncodeColumnIndexes(device_span<EncColumnChunk> chunks,
 {
   gpuEncodeColumnIndexes<<<chunks.size(), 1, 0, stream.value()>>>(
     chunks, column_stats, column_index_truncate_length);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 }  // namespace cudf::io::parquet::detail
diff --git a/cpp/src/io/parquet/page_hdr.cu b/cpp/src/io/parquet/page_hdr.cu
index c311f8bb5f25..f659363f0e4e 100644
--- a/cpp/src/io/parquet/page_hdr.cu
+++ b/cpp/src/io/parquet/page_hdr.cu
@@ -930,6 +930,7 @@ void count_page_headers(cudf::detail::hostdevice_span<ColumnChunkDesc> chunks,
   dim3 dim_grid(num_blocks, 1);
 
   count_page_headers_kernel<<<dim_grid, dim_block, 0, stream.value()>>>(chunks, error_code);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void decode_page_headers(cudf::device_span<ColumnChunkDesc const> chunks,
@@ -950,6 +951,7 @@ void decode_page_headers(cudf::device_span<ColumnChunkDesc const> chunks,
 
   decode_page_headers_kernel<<<dim_grid, dim_block, 0, stream.value()>>>(
     chunks, chunk_pages, error_code);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 void decode_page_headers_with_pgidx(cudf::device_span<ColumnChunkDesc const> chunks,
@@ -986,6 +988,7 @@ void build_string_dictionary_index(ColumnChunkDesc* chunks,
 
   build_string_dictionary_index_kernel<<<dim_grid, dim_block, 0, stream.value()>>>(
     chunks, num_chunks, error_code);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 }  // namespace cudf::io::parquet::detail
diff --git a/cpp/src/io/parquet/page_string_decode.cu b/cpp/src/io/parquet/page_string_decode.cu
index 4143389096f1..663a828e4a2f 100644
--- a/cpp/src/io/parquet/page_string_decode.cu
+++ b/cpp/src/io/parquet/page_string_decode.cu
@@ -948,9 +948,11 @@ void compute_page_string_sizes_pass1(cudf::detail::hostdevice_span<PageInfo> pag
   if (level_type_size == 1) {
     compute_string_page_bounds_kernel<uint8_t><<<dim_grid, dim_block, 0, stream.value()>>>(
       pages.device_ptr(), chunks, page_mask, min_row, num_rows, all_rows);
+    CUDF_CUDA_TRY(cudaGetLastError());
   } else {
     compute_string_page_bounds_kernel<uint16_t><<<dim_grid, dim_block, 0, stream.value()>>>(
       pages.device_ptr(), chunks, page_mask, min_row, num_rows, all_rows);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   // kernel mask may contain other kernels we don't need to count
@@ -963,6 +965,7 @@ void compute_page_string_sizes_pass1(cudf::detail::hostdevice_span<PageInfo> pag
     dim3 dim_delta(delta_preproc_block_size, 1);
     compute_delta_page_string_sizes_kernel<<<dim_grid, dim_delta, 0, streams[s_idx++].value()>>>(
       pages.device_ptr(), chunks, page_mask, min_row, num_rows);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
   if (BitAnd(kernel_mask, decode_kernel_mask::DELTA_LENGTH_BA) != 0) {
     dim3 dim_delta(delta_length_block_size, 1);
@@ -971,10 +974,12 @@ void compute_page_string_sizes_pass1(cudf::detail::hostdevice_span<PageInfo> pag
                                                     0,
                                                     streams[s_idx++].value()>>>(
       pages.device_ptr(), chunks, page_mask, min_row, num_rows);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
   if (BitAnd(kernel_mask, STRINGS_MASK_NON_DELTA) != 0) {
     compute_page_string_sizes_kernel<<<dim_grid, dim_block, 0, streams[s_idx++].value()>>>(
       pages.device_ptr(), chunks, page_mask, page_string_offset_indices, min_row, num_rows);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   // synchronize the streams
@@ -1344,6 +1349,7 @@ void preprocess_string_offsets(cudf::detail::hostdevice_span<PageInfo> pages,
   preprocess_string_offsets_kernel<preprocess_block_size, prefetch_size>
     <<<dim_grid, dim_block, 0, stream.value()>>>(
       pages.device_ptr(), chunks, page_string_offset_indices, page_mask, min_row, num_rows);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 }  // namespace cudf::io::parquet::detail
diff --git a/cpp/src/io/statistics/column_statistics.cuh b/cpp/src/io/statistics/column_statistics.cuh
index 8e791e96c524..81964cf2eabe 100644
--- a/cpp/src/io/statistics/column_statistics.cuh
+++ b/cpp/src/io/statistics/column_statistics.cuh
@@ -338,6 +338,7 @@ void calculate_group_statistics(statistics_chunk* chunks,
   constexpr int block_size = 256;
   gpu_calculate_group_statistics<block_size, IO>
     <<<num_chunks, block_size, 0, stream.value()>>>(chunks, groups, int96_timestamps);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 /**
@@ -392,6 +393,7 @@ void merge_group_statistics(statistics_chunk* chunks_out,
   constexpr int block_size = 256;
   gpu_merge_group_statistics<block_size, IO>
     <<<num_chunks, block_size, 0, stream.value()>>>(chunks_out, chunks_in, groups);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 }  // namespace detail
diff --git a/cpp/src/io/text/multibyte_split.cu b/cpp/src/io/text/multibyte_split.cu
index b638c10bdcde..8812077a1890 100644
--- a/cpp/src/io/text/multibyte_split.cu
+++ b/cpp/src/io/text/multibyte_split.cu
@@ -349,6 +349,7 @@ std::unique_ptr<cudf::column> multibyte_split(cudf::io::text::data_chunk_source
       tile_multistates,
       tile_offsets,
       cudf::io::text::detail::scan_tile_status::oob);
+    CUDF_CUDA_TRY(cudaGetLastError());
 
     auto multistate_seed = multistate();
     multistate_seed.enqueue(0, 0);  // this represents the first state in the pattern.
@@ -425,6 +426,7 @@ std::unique_ptr<cudf::column> multibyte_split(cudf::io::text::data_chunk_source
           delimiter[0],
           *chunk,
           row_offsets);
+        CUDF_CUDA_TRY(cudaGetLastError());
       } else {
         multibyte_split_kernel<<<tiles_in_launch,
                                  THREADS_PER_TILE,
@@ -438,6 +440,7 @@ std::unique_ptr<cudf::column> multibyte_split(cudf::io::text::data_chunk_source
           {device_delim.data(), static_cast<std::size_t>(device_delim.size())},
           *chunk,
           row_offsets);
+        CUDF_CUDA_TRY(cudaGetLastError());
       }
 
       // load the next chunk
diff --git a/cpp/src/io/utilities/data_casting.cu b/cpp/src/io/utilities/data_casting.cu
index d803ddecdb00..80de4860aa18 100644
--- a/cpp/src/io/utilities/data_casting.cu
+++ b/cpp/src/io/utilities/data_casting.cu
@@ -837,6 +837,7 @@ static std::unique_ptr<column> parse_string(string_view_pair_it str_tuples,
         d_sizes,
         cudf::detail::input_offsetalator{},
         nullptr);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   if (max_length > WARP_THRESHOLD) {
@@ -853,6 +854,7 @@ static std::unique_ptr<column> parse_string(string_view_pair_it str_tuples,
         d_sizes,
         cudf::detail::input_offsetalator{},
         nullptr);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   auto [offsets, bytes] =
@@ -884,6 +886,7 @@ static std::unique_ptr<column> parse_string(string_view_pair_it str_tuples,
         d_sizes,
         d_offsets,
         d_chars);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   if (max_length > WARP_THRESHOLD) {
@@ -900,6 +903,7 @@ static std::unique_ptr<column> parse_string(string_view_pair_it str_tuples,
         d_sizes,
         d_offsets,
         d_chars);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   return make_strings_column(col_size,
diff --git a/cpp/src/io/utilities/type_inference.cu b/cpp/src/io/utilities/type_inference.cu
index 5ecb46631449..a86383d57412 100644
--- a/cpp/src/io/utilities/type_inference.cu
+++ b/cpp/src/io/utilities/type_inference.cu
@@ -239,6 +239,7 @@ cudf::io::column_type_histogram infer_column_type(OptionsView const& options,
 
   infer_column_type_kernel<block_size><<<grid_size, block_size, 0, stream.value()>>>(
     options, data, offset_length_begin, size, d_column_info.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   return d_column_info.value(stream);
 }
diff --git a/cpp/src/join/conditional_join.cu b/cpp/src/join/conditional_join.cu
index d45912139c4c..8e2688b99f3b 100644
--- a/cpp/src/join/conditional_join.cu
+++ b/cpp/src/join/conditional_join.cu
@@ -81,10 +81,12 @@ std::unique_ptr<rmm::device_uvector<size_type>> conditional_join_anti_semi(
       compute_conditional_join_output_size<DEFAULT_JOIN_BLOCK_SIZE, true>
         <<<config.num_blocks, config.num_threads_per_block, shmem_size_per_block, stream.value()>>>(
           *left_table, *right_table, join_type, parser.device_expression_data, false, size.data());
+      CUDF_CUDA_TRY(cudaGetLastError());
     } else {
       compute_conditional_join_output_size<DEFAULT_JOIN_BLOCK_SIZE, false>
         <<<config.num_blocks, config.num_threads_per_block, shmem_size_per_block, stream.value()>>>(
           *left_table, *right_table, join_type, parser.device_expression_data, false, size.data());
+      CUDF_CUDA_TRY(cudaGetLastError());
     }
     join_size = size.value(stream);
   }
@@ -106,6 +108,7 @@ std::unique_ptr<rmm::device_uvector<size_type>> conditional_join_anti_semi(
         write_index.data(),
         parser.device_expression_data,
         join_size);
+    CUDF_CUDA_TRY(cudaGetLastError());
   } else {
     conditional_join_anti_semi<DEFAULT_JOIN_BLOCK_SIZE, DEFAULT_CACHE_SIZE, false>
       <<<config.num_blocks, config.num_threads_per_block, shmem_size_per_block, stream.value()>>>(
@@ -116,6 +119,7 @@ std::unique_ptr<rmm::device_uvector<size_type>> conditional_join_anti_semi(
         write_index.data(),
         parser.device_expression_data,
         join_size);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
   return left_indices;
 }
@@ -204,6 +208,7 @@ conditional_join(table_view const& left,
           parser.device_expression_data,
           swap_tables,
           size.data());
+      CUDF_CUDA_TRY(cudaGetLastError());
     } else {
       compute_conditional_join_output_size<DEFAULT_JOIN_BLOCK_SIZE, false>
         <<<config.num_blocks, config.num_threads_per_block, shmem_size_per_block, stream.value()>>>(
@@ -213,6 +218,7 @@ conditional_join(table_view const& left,
           parser.device_expression_data,
           swap_tables,
           size.data());
+      CUDF_CUDA_TRY(cudaGetLastError());
     }
     join_size = size.value(stream);
   }
@@ -249,6 +255,7 @@ conditional_join(table_view const& left,
         parser.device_expression_data,
         join_size,
         swap_tables);
+    CUDF_CUDA_TRY(cudaGetLastError());
   } else {
     conditional_join<DEFAULT_JOIN_BLOCK_SIZE, DEFAULT_CACHE_SIZE, false>
       <<<config.num_blocks, config.num_threads_per_block, shmem_size_per_block, stream.value()>>>(
@@ -261,6 +268,7 @@ conditional_join(table_view const& left,
         parser.device_expression_data,
         join_size,
         swap_tables);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   auto join_indices = std::pair(std::move(left_indices), std::move(right_indices));
@@ -353,6 +361,7 @@ std::size_t compute_conditional_join_output_size(table_view const& left,
         parser.device_expression_data,
         swap_tables,
         size.data());
+    CUDF_CUDA_TRY(cudaGetLastError());
   } else {
     compute_conditional_join_output_size<DEFAULT_JOIN_BLOCK_SIZE, false>
       <<<config.num_blocks, config.num_threads_per_block, shmem_size_per_block, stream.value()>>>(
@@ -362,6 +371,7 @@ std::size_t compute_conditional_join_output_size(table_view const& left,
         parser.device_expression_data,
         swap_tables,
         size.data());
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
   return size.value(stream);
 }
diff --git a/cpp/src/join/filtered_join.cu b/cpp/src/join/filtered_join.cu
index 88a8fc12600a..144c224c46e0 100644
--- a/cpp/src/join/filtered_join.cu
+++ b/cpp/src/join/filtered_join.cu
@@ -113,6 +113,7 @@ void filtered_join::insert_right_table(Ref const& insert_ref, rmm::cuda_stream_v
           cuda::counting_iterator<size_type>{0},
           row_is_valid{row_bitmask_ptr},
           insert_ref);
+      CUDF_CUDA_TRY(cudaGetLastError());
     } else {
       cuco::detail::open_addressing_ns::insert_if_n<CGSize, cuco::detail::default_block_size()>
         <<<grid_size, cuco::detail::default_block_size(), 0, stream.value()>>>(
@@ -121,6 +122,7 @@ void filtered_join::insert_right_table(Ref const& insert_ref, rmm::cuda_stream_v
           cuda::constant_iterator<bool>{true},
           cuda::std::identity{},
           insert_ref);
+      CUDF_CUDA_TRY(cudaGetLastError());
     }
   };
 
@@ -175,6 +177,7 @@ std::unique_ptr<rmm::device_uvector<cudf::size_type>> distinct_filtered_join::qu
           row_is_valid{row_bitmask_ptr},
           contains_iter,
           query_ref);
+      CUDF_CUDA_TRY(cudaGetLastError());
     } else {
       cuco::detail::open_addressing_ns::contains_if_n<CGSize, cuco::detail::default_block_size()>
         <<<grid_size, cuco::detail::default_block_size(), 0, stream.value()>>>(
@@ -184,6 +187,7 @@ std::unique_ptr<rmm::device_uvector<cudf::size_type>> distinct_filtered_join::qu
           cuda::std::identity{},
           contains_iter,
           query_ref);
+      CUDF_CUDA_TRY(cudaGetLastError());
     }
   };
 
diff --git a/cpp/src/join/hash_join/partitioned_count_kernels.cuh b/cpp/src/join/hash_join/partitioned_count_kernels.cuh
index f9d3809a0a4c..cf21086b68e9 100644
--- a/cpp/src/join/hash_join/partitioned_count_kernels.cuh
+++ b/cpp/src/join/hash_join/partitioned_count_kernels.cuh
@@ -89,6 +89,7 @@ void launch_partitioned_count(probe_key_type const* keys,
 
   partitioned_count_kernel<IsOuter>
     <<<config.num_blocks, config.num_threads_per_block, 0, stream.value()>>>(keys, n, output, ref);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 }  // namespace cudf::detail
diff --git a/cpp/src/join/hash_join/partitioned_retrieve_kernels.cuh b/cpp/src/join/hash_join/partitioned_retrieve_kernels.cuh
index 7e822fed12ba..f86e990bfe9a 100644
--- a/cpp/src/join/hash_join/partitioned_retrieve_kernels.cuh
+++ b/cpp/src/join/hash_join/partitioned_retrieve_kernels.cuh
@@ -250,6 +250,7 @@ launch_partitioned_retrieve(probe_key_type const* keys,
 
   partitioned_retrieve_kernel<IsOuter><<<num_blocks, DEFAULT_JOIN_BLOCK_SIZE, 0, stream.value()>>>(
     keys, n, left_offset, left_indices->data(), right_indices->data(), output_counter.data(), ref);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   return std::pair(std::move(left_indices), std::move(right_indices));
 }
diff --git a/cpp/src/join/key_remapping.cu b/cpp/src/join/key_remapping.cu
index 7ec7f0922d6e..514a833e9723 100644
--- a/cpp/src/join/key_remapping.cu
+++ b/cpp/src/join/key_remapping.cu
@@ -353,6 +353,7 @@ class key_remap_table : public key_remap_table_interface {
 
     insert_and_count_kernel<<<grid.num_blocks, KEY_REMAP_BLOCK_SIZE, 0, stream.value()>>>(
       build_num_rows, set_ref, key_iter, counts.data(), d_distinct_count.data(), bitmask_ptr);
+    CUDF_CUDA_TRY(cudaGetLastError());
 
     _distinct_count = d_distinct_count.value(stream);
 
diff --git a/cpp/src/join/mark_join.cu b/cpp/src/join/mark_join.cu
index 6d72f6b6ec84..6c11115251fe 100644
--- a/cpp/src/join/mark_join.cu
+++ b/cpp/src/join/mark_join.cu
@@ -391,6 +391,7 @@ void mark_join::clear_marks(rmm::cuda_stream_view stream)
   auto const grid_size = cudf::util::div_rounding_up_unsafe(num_buckets, mark_block_size);
   clear_marks_kernel<mark_block_size><<<grid_size, mark_block_size, 0, stream.value()>>>(
     storage_ref, static_cast<slot_type>(masked_empty_sentinel), num_buckets);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 template <typename Comparator>
@@ -421,6 +422,7 @@ cudf::size_type mark_join::mark_probe_without_prefilter(storage_ref_type storage
       num_right_rows,
       d_mark_counter.data(),
       right_row_bitmask);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   return d_mark_counter.value(stream);
@@ -459,6 +461,7 @@ cudf::size_type mark_join::mark_probe_with_prefilter(storage_ref_type storage_re
   compact_if_kernel<mark_block_size, right_key_type, right_key_type, prefilter_operator_type>
     <<<filter_grid_size, mark_block_size, 0, stream.value()>>>(
       filtered_right_rows.data(), d_filtered_count.data(), prefilter_op);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   auto const filtered_count = d_filtered_count.value(stream);
   return mark_probe_without_prefilter(
@@ -558,6 +561,7 @@ std::unique_ptr<rmm::device_uvector<cudf::size_type>> mark_join::mark_probe_and_
           result.data(),
           d_scan_offset.data(),
           num_buckets);
+      CUDF_CUDA_TRY(cudaGetLastError());
     } else {
       mark_retrieve_kernel<mark_block_size, true>
         <<<grid_size, mark_block_size, 0, stream.value()>>>(
@@ -566,6 +570,7 @@ std::unique_ptr<rmm::device_uvector<cudf::size_type>> mark_join::mark_probe_and_
           result.data(),
           d_scan_offset.data(),
           num_buckets);
+      CUDF_CUDA_TRY(cudaGetLastError());
     }
   }
 
@@ -650,6 +655,7 @@ mark_join::mark_join(cudf::table_view const& left,
           cuda::counting_iterator<size_type>{0},
           row_is_valid{row_bitmask_ptr},
           insert_ref);
+      CUDF_CUDA_TRY(cudaGetLastError());
     } else {
       cuco::detail::open_addressing_ns::insert_if_n<masked_probing_scheme::cg_size,
                                                     cuco::detail::default_block_size()>
@@ -659,6 +665,7 @@ mark_join::mark_join(cudf::table_view const& left,
           cuda::constant_iterator<bool>{true},
           cuda::std::identity{},
           insert_ref);
+      CUDF_CUDA_TRY(cudaGetLastError());
     }
   };
 
diff --git a/cpp/src/join/mixed_join_kernel.cuh b/cpp/src/join/mixed_join_kernel.cuh
index 8687c13f859d..4c865b44d3b4 100644
--- a/cpp/src/join/mixed_join_kernel.cuh
+++ b/cpp/src/join/mixed_join_kernel.cuh
@@ -174,6 +174,7 @@ void launch_mixed_join(
       join_output_l,
       join_output_r,
       join_result_offsets);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 }  // namespace detail
diff --git a/cpp/src/join/mixed_join_kernels_semi.cu b/cpp/src/join/mixed_join_kernels_semi.cu
index a387b4a6c40b..0302c25c679e 100644
--- a/cpp/src/join/mixed_join_kernels_semi.cu
+++ b/cpp/src/join/mixed_join_kernels_semi.cu
@@ -86,6 +86,7 @@ void launch_mixed_join_semi(bool has_nulls,
         set_ref,
         left_table_keep_mask,
         device_expression_data);
+    CUDF_CUDA_TRY(cudaGetLastError());
   } else {
     mixed_join_semi<DEFAULT_JOIN_BLOCK_SIZE, false>
       <<<config.num_blocks, config.num_threads_per_block, shmem_size_per_block, stream.value()>>>(
@@ -97,6 +98,7 @@ void launch_mixed_join_semi(bool has_nulls,
         set_ref,
         left_table_keep_mask,
         device_expression_data);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 }
 
diff --git a/cpp/src/join/mixed_join_size_kernel.cuh b/cpp/src/join/mixed_join_size_kernel.cuh
index 3526ca23ad36..9a719170d7a3 100644
--- a/cpp/src/join/mixed_join_size_kernel.cuh
+++ b/cpp/src/join/mixed_join_size_kernel.cuh
@@ -130,6 +130,7 @@ void launch_mixed_join_count(
       hash_indices,
       device_expression_data,
       matches_per_row);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 }  // namespace cudf::detail
diff --git a/cpp/src/json/json_path.cu b/cpp/src/json/json_path.cu
index 8b1098d2941e..861bf521ac0e 100644
--- a/cpp/src/json/json_path.cu
+++ b/cpp/src/json/json_path.cu
@@ -1015,6 +1015,7 @@ std::unique_ptr<cudf::column> get_json_object(cudf::strings_column_view const& c
       cuda::std::nullopt,
       cuda::std::nullopt,
       options);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   // convert sizes to offsets
   auto [offsets, output_size] =
@@ -1043,6 +1044,7 @@ std::unique_ptr<cudf::column> get_json_object(cudf::strings_column_view const& c
       static_cast<bitmask_type*>(validity.data()),
       d_valid_count.data(),
       options);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   auto result = make_strings_column(col.size(),
                                     std::move(offsets),
diff --git a/cpp/src/merge/merge.cu b/cpp/src/merge/merge.cu
index a007d5104c27..da696419d8ca 100644
--- a/cpp/src/merge/merge.cu
+++ b/cpp/src/merge/merge.cu
@@ -169,16 +169,19 @@ void materialize_bitmask(column_view const& left_col,
       materialize_merged_bitmask_kernel<true, true>
         <<<grid_config.num_blocks, grid_config.num_threads_per_block, 0, stream.value()>>>(
           left_valid, right_valid, out_validity, num_elements, merged_indices);
+      CUDF_CUDA_TRY(cudaGetLastError());
     } else {
       materialize_merged_bitmask_kernel<true, false>
         <<<grid_config.num_blocks, grid_config.num_threads_per_block, 0, stream.value()>>>(
           left_valid, right_valid, out_validity, num_elements, merged_indices);
+      CUDF_CUDA_TRY(cudaGetLastError());
     }
   } else {
     if (right_col.has_nulls()) {
       materialize_merged_bitmask_kernel<false, true>
         <<<grid_config.num_blocks, grid_config.num_threads_per_block, 0, stream.value()>>>(
           left_valid, right_valid, out_validity, num_elements, merged_indices);
+      CUDF_CUDA_TRY(cudaGetLastError());
     } else {
       CUDF_FAIL("materialize_merged_bitmask_kernel<false, false>() should never be called.");
     }
diff --git a/cpp/src/partitioning/partitioning.cu b/cpp/src/partitioning/partitioning.cu
index e2e627c3390d..bede1b88780f 100644
--- a/cpp/src/partitioning/partitioning.cu
+++ b/cpp/src/partitioning/partitioning.cu
@@ -364,6 +364,7 @@ void copy_block_partitions_impl(InputIter const input,
     row_partition_offset,
     block_partition_sizes,
     scanned_block_partition_sizes);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 rmm::device_uvector<size_type> compute_gather_map(size_type num_rows,
@@ -641,6 +642,7 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table(
                                                       row_partition_offset.data(),
                                                       block_partition_sizes.data(),
                                                       global_partition_sizes.data());
+    CUDF_CUDA_TRY(cudaGetLastError());
   } else {
     // Determines how the mapping between hash value and partition number is
     // computed
@@ -661,6 +663,7 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table(
                                                       row_partition_offset.data(),
                                                       block_partition_sizes.data(),
                                                       global_partition_sizes.data());
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   // Compute exclusive scan of all blocks' partition sizes in-place to determine
@@ -733,6 +736,7 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table(
                                    num_partitions * sizeof(size_type),
                                    stream.value()>>>(
       row_output_locations, num_rows, num_partitions, scanned_block_partition_sizes_ptr);
+    CUDF_CUDA_TRY(cudaGetLastError());
 
     // Use the resulting scatter map to materialize the output
     auto output = detail::scatter(input, row_partition_numbers, input, stream, mr);
diff --git a/cpp/src/quantiles/tdigest/tdigest.cu b/cpp/src/quantiles/tdigest/tdigest.cu
index 42d89b72bcf6..bf793aecb3bc 100644
--- a/cpp/src/quantiles/tdigest/tdigest.cu
+++ b/cpp/src/quantiles/tdigest/tdigest.cu
@@ -241,6 +241,7 @@ std::unique_ptr<column> compute_approx_percentiles(tdigest_column_view const& in
     tdv.max_begin(),
     cumulative_weights->view().begin<double>(),
     result->mutable_view().begin<double>());
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   return result;
 }
diff --git a/cpp/src/quantiles/tdigest/tdigest_aggregation.cu b/cpp/src/quantiles/tdigest/tdigest_aggregation.cu
index 79b0835c2f11..5ea133095b58 100644
--- a/cpp/src/quantiles/tdigest/tdigest_aggregation.cu
+++ b/cpp/src/quantiles/tdigest/tdigest_aggregation.cu
@@ -621,6 +621,7 @@ void generate_cluster_limits(int delta,
       group_num_clusters,
       group_cluster_start,
       has_nulls);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   // overlap CPU work
diff --git a/cpp/src/replace/nulls.cu b/cpp/src/replace/nulls.cu
index a39c88e49e26..90c91af54d23 100644
--- a/cpp/src/replace/nulls.cu
+++ b/cpp/src/replace/nulls.cu
@@ -129,6 +129,7 @@ struct replace_nulls_column_kernel_forwarder {
 
     replace<<<grid.num_blocks, BLOCK_SIZE, 0, stream.value()>>>(
       *device_in, *device_replacement, *device_out, valid_count);
+    CUDF_CUDA_TRY(cudaGetLastError());
 
     if (output_view.nullable()) {
       output->set_null_count(output->size() - valid_counter.value(stream));
diff --git a/cpp/src/replace/replace.cu b/cpp/src/replace/replace.cu
index 1cb7d38c3c56..af16841d482b 100644
--- a/cpp/src/replace/replace.cu
+++ b/cpp/src/replace/replace.cu
@@ -207,6 +207,7 @@ struct replace_kernel_forwarder {
                                                                 output_view.size(),
                                                                 *device_values_to_replace,
                                                                 *device_replacement_values);
+    CUDF_CUDA_TRY(cudaGetLastError());
 
     if (output_view.nullable()) {
       output->set_null_count(output->size() - valid_counter.value(stream));
diff --git a/cpp/src/rolling/detail/rolling.cuh b/cpp/src/rolling/detail/rolling.cuh
index 9eab3d598109..9b0109f7d4c5 100644
--- a/cpp/src/rolling/detail/rolling.cuh
+++ b/cpp/src/rolling/detail/rolling.cuh
@@ -464,6 +464,7 @@ struct rolling_window_launcher {
                                                              device_op,
                                                              preceding_window_begin,
                                                              following_window_begin);
+      CUDF_CUDA_TRY(cudaGetLastError());
 
       auto const valid_count = d_valid_count.value(stream);
       output->set_null_count(output->size() - valid_count);
diff --git a/cpp/src/sort/segmented_top_k.cu b/cpp/src/sort/segmented_top_k.cu
index d82e43b8ecc2..8e619c235342 100644
--- a/cpp/src/sort/segmented_top_k.cu
+++ b/cpp/src/sort/segmented_top_k.cu
@@ -98,6 +98,7 @@ std::unique_ptr<column> segmented_top_k_order(column_view const& col,
   auto const grid    = cudf::detail::grid_1d(indices->size(), 256);
   resolve_segment_indices<<<grid.num_blocks, grid.num_threads_per_block, 0, stream>>>(
     segment_offsets, k, span_indices, segment_sizes.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
   auto [offsets, total_elements] =
     cudf::detail::make_offsets_child_column(segment_sizes.begin(), segment_sizes.end(), stream, mr);
 
diff --git a/cpp/src/strings/attributes.cu b/cpp/src/strings/attributes.cu
index b4108c101af8..3bfc6dbfe62c 100644
--- a/cpp/src/strings/attributes.cu
+++ b/cpp/src/strings/attributes.cu
@@ -146,6 +146,7 @@ std::unique_ptr<column> count_characters_parallel(strings_column_view const& inp
   cudf::detail::grid_1d grid{input.size() * warp_size, block_size};
   count_characters_parallel_fn<<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
     *d_strings, d_lengths);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   // reset null count after call to mutable_view()
   results->set_null_count(input.null_count());
diff --git a/cpp/src/strings/case.cu b/cpp/src/strings/case.cu
index b69109e7e735..bf8deaf009f8 100644
--- a/cpp/src/strings/case.cu
+++ b/cpp/src/strings/case.cu
@@ -429,6 +429,7 @@ std::unique_ptr<column> convert_case(strings_column_view const& input,
   mismatch_multibytes_kernel<bytes_per_thread>
     <<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
       input_chars, first_offset, last_offset, mb_count.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
   if (mb_count.value(stream) == 0) {
     // optimization for the non-special case;
     // copying the input column automatically handles normalizing sliced inputs
@@ -438,6 +439,7 @@ std::unique_ptr<column> convert_case(strings_column_view const& input,
     multibyte_converter_kernel<bytes_per_thread>
       <<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
         ccfn, input_chars + first_offset, chars_size, d_chars);
+    CUDF_CUDA_TRY(cudaGetLastError());
     result->set_null_count(input.null_count());
     return result;
   }
@@ -451,6 +453,7 @@ std::unique_ptr<column> convert_case(strings_column_view const& input,
     count_bytes_kernel<bytes_per_thread>
       <<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
         ccfn, *d_strings, sizes.data());
+    CUDF_CUDA_TRY(cudaGetLastError());
     // convert sizes to offsets
     return cudf::strings::detail::make_offsets_child_column(sizes.begin(), sizes.end(), stream, mr);
   }();
diff --git a/cpp/src/strings/convert/convert_urls.cu b/cpp/src/strings/convert/convert_urls.cu
index f48a1d5ce4bf..5c962b433371 100644
--- a/cpp/src/strings/convert/convert_urls.cu
+++ b/cpp/src/strings/convert/convert_urls.cu
@@ -392,6 +392,7 @@ std::unique_ptr<column> url_decode(strings_column_view const& strings,
   auto row_sizes = rmm::device_uvector<size_type>(strings_count, stream);
   url_decode_char_counter<num_warps_per_threadblock, char_block_size>
     <<<num_threadblocks, threadblock_size, 0, stream.value()>>>(*d_strings, row_sizes.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
   // performs scan on the sizes and builds the appropriate offsets column
   auto [offsets_column, out_chars_bytes] = cudf::strings::detail::make_offsets_child_column(
     row_sizes.begin(), row_sizes.end(), stream, mr);
@@ -405,6 +406,7 @@ std::unique_ptr<column> url_decode(strings_column_view const& strings,
   // decode and copy the characters from the input column to the output column
   url_decode_char_replacer<num_warps_per_threadblock, char_block_size>
     <<<num_threadblocks, threadblock_size, 0, stream.value()>>>(*d_strings, d_out_chars, offsets);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   // copy null mask
   rmm::device_buffer null_mask = cudf::detail::copy_bitmask(strings.parent(), stream, mr);
diff --git a/cpp/src/strings/copying/concatenate.cu b/cpp/src/strings/copying/concatenate.cu
index 8d90d18c1515..bfde1f742698 100644
--- a/cpp/src/strings/copying/concatenate.cu
+++ b/cpp/src/strings/copying/concatenate.cu
@@ -251,6 +251,7 @@ std::unique_ptr<column> concatenate(host_span<column_view const> columns,
       itr_new_offsets,
       reinterpret_cast<bitmask_type*>(null_mask.data()),
       d_valid_count.data());
+    CUDF_CUDA_TRY(cudaGetLastError());
 
     if (has_nulls) { null_count = strings_count - d_valid_count.value(stream); }
   }
@@ -268,6 +269,7 @@ std::unique_ptr<column> concatenate(host_span<column_view const> columns,
                                                             static_cast<size_type>(columns.size()),
                                                             total_bytes,
                                                             d_new_chars);
+      CUDF_CUDA_TRY(cudaGetLastError());
     } else {
       // Memcpy each input chars column (more efficient for very large strings)
       for (auto column = columns.begin(); column != columns.end(); ++column) {
diff --git a/cpp/src/strings/like.cu b/cpp/src/strings/like.cu
index 588d99cd998b..3f7823950c1f 100644
--- a/cpp/src/strings/like.cu
+++ b/cpp/src/strings/like.cu
@@ -335,6 +335,7 @@ std::unique_ptr<column> like(strings_column_view const& input,
     auto const grid = cudf::detail::grid_1d(input.size() * warp_size, block_size);
     like_kernel<<<grid.num_blocks, grid.num_threads_per_block, 0, stream>>>(
       *d_strings, patterns_itr, d_escape, results->mutable_view().data<bool>());
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   results->set_null_count(input.null_count());
diff --git a/cpp/src/strings/regex/utilities.cuh b/cpp/src/strings/regex/utilities.cuh
index a65bac3844f7..c6bf18b1dc21 100644
--- a/cpp/src/strings/regex/utilities.cuh
+++ b/cpp/src/strings/regex/utilities.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -61,6 +61,7 @@ void launch_for_each_kernel(ForEachFunction fn,
   cudf::detail::grid_1d grid{thread_count, regex_launch_kernel_block_size};
   for_each_kernel<<<grid.num_blocks, grid.num_threads_per_block, shmem_size, stream.value()>>>(
     fn, d_prog, size);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 template <typename TransformFunction, typename OutputType>
@@ -99,6 +100,7 @@ void launch_transform_kernel(TransformFunction fn,
   cudf::detail::grid_1d grid{thread_count, regex_launch_kernel_block_size};
   transform_kernel<<<grid.num_blocks, grid.num_threads_per_block, shmem_size, stream.value()>>>(
     fn, d_prog, d_output, size);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 
 template <typename SizeAndExecuteFunction>
@@ -122,6 +124,7 @@ auto make_strings_children(SizeAndExecuteFunction size_and_exec_fn,
   if (strings_count > 0) {
     for_each_kernel<<<grid.num_blocks, grid.num_threads_per_block, shmem_size, stream.value()>>>(
       size_and_exec_fn, d_prog, strings_count);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
   // Convert the sizes to offsets
   auto [offsets, char_bytes] = cudf::strings::detail::make_offsets_child_column(
@@ -135,6 +138,7 @@ auto make_strings_children(SizeAndExecuteFunction size_and_exec_fn,
     size_and_exec_fn.d_chars = chars.data();
     for_each_kernel<<<grid.num_blocks, grid.num_threads_per_block, shmem_size, stream.value()>>>(
       size_and_exec_fn, d_prog, strings_count);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   return std::make_pair(std::move(offsets), std::move(chars));
diff --git a/cpp/src/strings/replace/multi.cu b/cpp/src/strings/replace/multi.cu
index ef0c2f569fb9..d9c53d9f8de9 100644
--- a/cpp/src/strings/replace/multi.cu
+++ b/cpp/src/strings/replace/multi.cu
@@ -334,6 +334,7 @@ std::unique_ptr<column> replace_character_parallel(strings_column_view const& in
   auto const num_blocks = util::div_rounding_up_safe(
     util::div_rounding_up_safe(chars_bytes, static_cast<int64_t>(bytes_per_thread)), block_size);
   count_targets<<<num_blocks, block_size, 0, stream.value()>>>(fn, chars_bytes, d_count.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
   auto target_count = d_count.value(stream);
   // Create a vector of every target position in the chars column.
   // These may also include overlapping targets which will be resolved later.
diff --git a/cpp/src/strings/replace/replace.cu b/cpp/src/strings/replace/replace.cu
index 77a9aea444f1..40298ba47739 100644
--- a/cpp/src/strings/replace/replace.cu
+++ b/cpp/src/strings/replace/replace.cu
@@ -288,6 +288,7 @@ std::unique_ptr<column> replace_character_parallel(strings_column_view const& in
     util::div_rounding_up_safe(chars_bytes, static_cast<int64_t>(bytes_per_thread)), block_size);
   count_targets_kernel<block_size, bytes_per_thread>
     <<<num_blocks, block_size, 0, stream.value()>>>(fn, chars_bytes, d_target_count.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
   auto target_count = d_target_count.value(stream);
 
   // Create a vector of every target position in the chars column.
diff --git a/cpp/src/strings/search/contains_multiple.cu b/cpp/src/strings/search/contains_multiple.cu
index 9edc64f53117..542aea6b4177 100644
--- a/cpp/src/strings/search/contains_multiple.cu
+++ b/cpp/src/strings/search/contains_multiple.cu
@@ -271,6 +271,7 @@ std::unique_ptr<table> contains_multiple(strings_column_view const& input,
                                                                            unique_count,
                                                                            nullptr,
                                                                            d_results);
+    CUDF_CUDA_TRY(cudaGetLastError());
   } else {
     constexpr cudf::thread_index_type tile_size = cudf::detail::warp_size;
 
@@ -292,6 +293,7 @@ std::unique_ptr<table> contains_multiple(strings_column_view const& input,
         unique_count,
         working_memory.data(),
         d_results);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   return std::make_unique<table>(std::move(results));
diff --git a/cpp/src/strings/search/find.cu b/cpp/src/strings/search/find.cu
index 4b72c21629c3..265b25321706 100644
--- a/cpp/src/strings/search/find.cu
+++ b/cpp/src/strings/search/find.cu
@@ -176,6 +176,7 @@ void find_utility(strings_column_view const& input,
     finder_warp_parallel_fn<TargetIterator, forward>
       <<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
         *d_strings, target_itr, start, stop, d_results);
+    CUDF_CUDA_TRY(cudaGetLastError());
   } else {
     // string-per-thread function
     thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
@@ -396,6 +397,7 @@ std::unique_ptr<column> contains_warp_parallel(strings_column_view const& input,
     cudf::detail::grid_1d grid{input.size() * warp_size, block_size};
     contains_warp_parallel_fn<<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
       *d_strings, d_target, results_view.data<bool>());
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
   results->set_null_count(input.null_count());
   return results;
diff --git a/cpp/src/strings/search/find_instance.cu b/cpp/src/strings/search/find_instance.cu
index 1e4267d4d4db..c12f6c1710d1 100644
--- a/cpp/src/strings/search/find_instance.cu
+++ b/cpp/src/strings/search/find_instance.cu
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -113,6 +113,7 @@ std::unique_ptr<column> find_instance(strings_column_view const& input,
                                    grid.num_threads_per_block,
                                    0,
                                    stream.value()>>>(*d_strings, d_target, instance, d_results);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   return results;
 }
diff --git a/cpp/src/strings/slice.cu b/cpp/src/strings/slice.cu
index 84d88436545e..47799478a2f0 100644
--- a/cpp/src/strings/slice.cu
+++ b/cpp/src/strings/slice.cu
@@ -260,6 +260,7 @@ std::unique_ptr<column> compute_substrings_from_fn(strings_column_view const& in
     auto const num_blocks = util::div_rounding_up_safe(threads, block_size);
     substring_from_kernel<IndexIterator>
       <<<num_blocks, block_size, 0, stream.value()>>>(*d_column, starts, stops, results.data());
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
   return make_strings_column(results.begin(), results.end(), stream, mr);
 }
diff --git a/cpp/src/strings/split/split.cuh b/cpp/src/strings/split/split.cuh
index b3fec39a2f86..bdfc0f06704b 100644
--- a/cpp/src/strings/split/split.cuh
+++ b/cpp/src/strings/split/split.cuh
@@ -525,6 +525,7 @@ std::pair<std::unique_ptr<column>, rmm::device_uvector<string_index_pair>> split
       util::div_rounding_up_safe(chars_bytes, static_cast<int64_t>(bytes_per_thread)), block_size);
     count_delimiters_kernel<DelimiterFn, block_size, bytes_per_thread>
       <<<num_blocks, block_size, 0, stream.value()>>>(delimiter_fn, chars_bytes, d_count.data());
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   // Create a vector of every delimiter position in the chars column.
diff --git a/cpp/src/strings/strings_column_factories.cu b/cpp/src/strings/strings_column_factories.cu
index 746f72353e89..3ad65d2c0945 100644
--- a/cpp/src/strings/strings_column_factories.cu
+++ b/cpp/src/strings/strings_column_factories.cu
@@ -106,6 +106,7 @@ std::vector<std::unique_ptr<column>> make_strings_column_batch(
         string_count,
         [] __device__(string_index_pair const pair) -> bool { return pair.first != nullptr; },
         d_valid_counts.data() + idx);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   auto const chars_sizes  = cudf::detail::make_std_vector_async(d_chars_sizes, stream);
diff --git a/cpp/src/text/bpe/byte_pair_encoding.cu b/cpp/src/text/bpe/byte_pair_encoding.cu
index 9bcdf5581708..47c120e7e93c 100644
--- a/cpp/src/text/bpe/byte_pair_encoding.cu
+++ b/cpp/src/text/bpe/byte_pair_encoding.cu
@@ -413,12 +413,14 @@ std::unique_ptr<cudf::column> byte_pair_encoding(cudf::strings_column_view const
     auto const pair_map = get_bpe_merge_pairs_impl(merge_pairs)->get_merge_pairs_ref();
     bpe_parallel_fn<decltype(pair_map)><<<tmp_size, block_size, 0, stream.value()>>>(
       *d_tmp_strings, d_input_chars, pair_map, d_spaces.data(), d_ranks.data(), d_rerank.data());
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   // compute the output sizes
   auto output_sizes = rmm::device_uvector<cudf::size_type>(input.size(), stream);
   bpe_finalize<<<input.size(), block_size, 0, stream.value()>>>(
     *d_strings, d_input_chars, d_spaces.data(), output_sizes.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   // convert sizes to offsets in-place
   auto [offsets, bytes] = cudf::strings::detail::make_offsets_child_column(
diff --git a/cpp/src/text/edit_distance.cu b/cpp/src/text/edit_distance.cu
index 0598355f646b..869ebb156ec9 100644
--- a/cpp/src/text/edit_distance.cu
+++ b/cpp/src/text/edit_distance.cu
@@ -274,6 +274,7 @@ std::unique_ptr<cudf::column> edit_distance(cudf::strings_column_view const& inp
   cudf::detail::grid_1d grid{input.size() * tile_size, block_size};
   levenshtein_kernel<tile_size><<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
     *d_strings, *d_targets, d_buffer, offsets.data(), d_results);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   return results;
 }
diff --git a/cpp/src/text/generate_ngrams.cu b/cpp/src/text/generate_ngrams.cu
index 124d98e749dc..f7a902656d3a 100644
--- a/cpp/src/text/generate_ngrams.cu
+++ b/cpp/src/text/generate_ngrams.cu
@@ -275,6 +275,7 @@ std::unique_ptr<cudf::column> generate_character_ngrams(cudf::strings_column_vie
       static_cast<cudf::thread_index_type>(input.size()) * tile_size, block_size);
     count_char_ngrams_kernel<<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
       *d_strings, ngrams, tile_size, counts.data());
+    CUDF_CUDA_TRY(cudaGetLastError());
     return cudf::detail::make_offsets_child_column(counts.begin(), counts.end(), stream, mr);
   }();
   auto d_offsets = offsets->view().data<cudf::size_type>();
@@ -375,6 +376,7 @@ std::unique_ptr<cudf::column> hash_character_ngrams(cudf::strings_column_view co
     auto counts = rmm::device_uvector<cudf::size_type>(input.size(), stream);
     count_char_ngrams_kernel<<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
       *d_strings, ngrams, cudf::detail::warp_size, counts.data());
+    CUDF_CUDA_TRY(cudaGetLastError());
     return cudf::detail::make_offsets_child_column(counts.begin(), counts.end(), stream, mr);
   }();
   auto d_offsets = offsets->view().data<cudf::size_type>();
@@ -389,6 +391,7 @@ std::unique_ptr<cudf::column> hash_character_ngrams(cudf::strings_column_view co
 
   character_ngram_hash_kernel<<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
     *d_strings, ngrams, seed, d_offsets, d_hashes);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   return make_lists_column(
     input.size(), std::move(offsets), std::move(hashes), 0, rmm::device_buffer{});
diff --git a/cpp/src/text/jaccard.cu b/cpp/src/text/jaccard.cu
index 74cafca0b3da..a4e7ff354a25 100644
--- a/cpp/src/text/jaccard.cu
+++ b/cpp/src/text/jaccard.cu
@@ -112,6 +112,7 @@ rmm::device_uvector<cudf::size_type> compute_unique_counts(uint32_t const* value
     static_cast<cudf::thread_index_type>(rows) * cudf::detail::warp_size, block_size);
   sorted_unique_fn<<<num_blocks, block_size, 0, stream.value()>>>(
     values, offsets, rows, d_results.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
   return d_results;
 }
 
@@ -186,6 +187,7 @@ rmm::device_uvector<cudf::size_type> compute_intersect_counts(uint32_t const* va
     static_cast<cudf::thread_index_type>(rows) * cudf::detail::warp_size, block_size);
   sorted_intersect_fn<<<num_blocks, block_size, 0, stream.value()>>>(
     values1, offsets1, values2, offsets2, rows, d_results.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
   return d_results;
 }
 
@@ -336,6 +338,7 @@ std::pair<rmm::device_uvector<uint32_t>, rmm::device_uvector<int64_t>> hash_subs
     static_cast<cudf::thread_index_type>(input.size()) * cudf::detail::warp_size, block_size);
   count_substrings_kernel<<<num_blocks, block_size, 0, stream.value()>>>(
     *d_strings, width, offsets.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
   auto const total_hashes =
     cudf::detail::sizes_to_offsets(offsets.begin(), offsets.end(), offsets.begin(), 0, stream);
 
@@ -343,6 +346,7 @@ std::pair<rmm::device_uvector<uint32_t>, rmm::device_uvector<int64_t>> hash_subs
   rmm::device_uvector<uint32_t> hashes(total_hashes, stream);
   substring_hash_kernel<<<num_blocks, block_size, 0, stream.value()>>>(
     *d_strings, width, offsets.data(), hashes.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   // sort hashes
   rmm::device_uvector<uint32_t> sorted(total_hashes, stream);
diff --git a/cpp/src/text/minhash.cu b/cpp/src/text/minhash.cu
index 772a1ca81554..8646b6e03ae2 100644
--- a/cpp/src/text/minhash.cu
+++ b/cpp/src/text/minhash.cu
@@ -476,6 +476,7 @@ std::unique_ptr<cudf::column> minhash_fn(cudf::strings_column_view const& input,
                                                                          d_threshold_count.data(),
                                                                          parameter_a.size(),
                                                                          d_results);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   auto transform_fn = [d_strings = *d_strings] __device__(auto idx) -> cudf::size_type {
     if (d_strings.is_null(idx)) { return 0; }
@@ -496,6 +497,7 @@ std::unique_ptr<cudf::column> minhash_fn(cudf::strings_column_view const& input,
     minhash_kernel<offsets_type, hash_value_type, 1>
       <<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
         input_offsets, d_indices, parameter_a, parameter_b, width, d_hashes.data(), d_results);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   // handle the strings above the threshold width
@@ -507,6 +509,7 @@ std::unique_ptr<cudf::column> minhash_fn(cudf::strings_column_view const& input,
     minhash_kernel<offsets_type, hash_value_type, blocks_per_row>
       <<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
         input_offsets, d_indices, parameter_a, parameter_b, width, d_hashes.data(), d_results);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   return results;
@@ -564,6 +567,7 @@ std::unique_ptr<cudf::column> minhash_ngrams_fn(
                                                                          d_threshold_count.data(),
                                                                          parameter_a.size(),
                                                                          d_results);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   auto sizes_fn = [d_list] __device__(auto idx) -> cudf::size_type {
     if (d_list.is_null(idx)) { return 0; }
@@ -583,6 +587,7 @@ std::unique_ptr<cudf::column> minhash_ngrams_fn(
     minhash_kernel<offset_type, hash_value_type, 1>
       <<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
         input_offsets, d_indices, parameter_a, parameter_b, ngrams, d_hashes.data(), d_results);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   // handle the strings above the threshold width
@@ -594,6 +599,7 @@ std::unique_ptr<cudf::column> minhash_ngrams_fn(
     minhash_kernel<offset_type, hash_value_type, blocks_per_row>
       <<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
         input_offsets, d_indices, parameter_a, parameter_b, ngrams, d_hashes.data(), d_results);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   return results;
diff --git a/cpp/src/text/normalize.cu b/cpp/src/text/normalize.cu
index ad272fe42412..fe89efcb6bb2 100644
--- a/cpp/src/text/normalize.cu
+++ b/cpp/src/text/normalize.cu
@@ -518,6 +518,7 @@ std::unique_ptr<cudf::column> normalize_characters(cudf::strings_column_view con
     parameters->aux_table.data(),
     parameters->do_lower_case,
     d_normalized.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   // This removes space added around any special tokens in the form of [ttt].
   // An alternate approach is to do a multi-replace of '[ ttt ]' with '[ttt]' right
@@ -526,6 +527,7 @@ std::unique_ptr<cudf::column> normalize_characters(cudf::strings_column_view con
   if (!special_tokens.empty()) {
     special_tokens_kernel<<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
       d_normalized.data(), chars_size, special_tokens, parameters->do_lower_case);
+    CUDF_CUDA_TRY(cudaGetLastError());
   }
 
   // Use segmented-reduce over the non-zero codepoints to get the size of the output rows
diff --git a/cpp/src/text/vocabulary_tokenize.cu b/cpp/src/text/vocabulary_tokenize.cu
index 0ea4c7219da3..99c543a0bd67 100644
--- a/cpp/src/text/vocabulary_tokenize.cu
+++ b/cpp/src/text/vocabulary_tokenize.cu
@@ -409,12 +409,14 @@ std::unique_ptr<cudf::column> tokenize_with_vocabulary(cudf::strings_column_view
                        grid_chars.num_threads_per_block,
                        0,
                        stream.value()>>>(d_input_chars, chars_size, d_delimiter, d_marks.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   // launch warp per string to compute token counts
   constexpr cudf::thread_index_type warp_size = cudf::detail::warp_size;
   cudf::detail::grid_1d grid{input.size() * warp_size, block_size};
   token_counts_fn<<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
     *d_strings, d_delimiter, d_token_counts.data(), d_marks.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
   auto [token_offsets, total_count] = cudf::detail::make_offsets_child_column(
     d_token_counts.begin(), d_token_counts.end(), stream, mr);
 
diff --git a/cpp/src/text/wordpiece_tokenize.cu b/cpp/src/text/wordpiece_tokenize.cu
index feb8dc1aafde..7673253ee3eb 100644
--- a/cpp/src/text/wordpiece_tokenize.cu
+++ b/cpp/src/text/wordpiece_tokenize.cu
@@ -561,6 +561,7 @@ rmm::device_uvector<cudf::size_type> compute_all_tokens(
   tokenize_all_kernel<decltype(map_ref), decltype(sub_map_ref)>
     <<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
       d_all_edges, d_input_chars, map_ref, sub_map_ref, unk_id, d_tokens.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   return d_tokens;
 }
@@ -791,6 +792,7 @@ rmm::device_uvector<cudf::size_type> compute_some_tokens(
   find_words_kernel<warp_size>
     <<<grid_find.num_blocks, grid_find.num_threads_per_block, 0, stream.value()>>>(
       *d_strings, d_input_chars, max_word_offsets.data(), start_words.data(), word_sizes.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   // remove the non-words
   auto const end =
@@ -826,6 +828,7 @@ rmm::device_uvector<cudf::size_type> compute_some_tokens(
   tokenize_kernel<decltype(map_ref), decltype(sub_map_ref)>
     <<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>(
       start_words, word_sizes, d_input_chars, map_ref, sub_map_ref, unk_id, d_tokens.data());
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   return d_tokens;
 }
diff --git a/cpp/src/transform/compute_column_kernel.cuh b/cpp/src/transform/compute_column_kernel.cuh
index ed18527238be..24bc4cbe2d3d 100644
--- a/cpp/src/transform/compute_column_kernel.cuh
+++ b/cpp/src/transform/compute_column_kernel.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -72,5 +72,6 @@ void launch_compute_column_kernel(table_device_view const& table_device,
   compute_column_kernel<MAX_BLOCK_SIZE, HasNull, HasComplexType>
     <<<config.num_blocks, config.num_threads_per_block, shmem_per_block, stream.value()>>>(
       table_device, device_expression_data, mutable_output_device);
+  CUDF_CUDA_TRY(cudaGetLastError());
 }
 }  // namespace cudf::detail
diff --git a/cpp/src/transform/row_bit_count.cu b/cpp/src/transform/row_bit_count.cu
index ce93c4777566..e604e47266f0 100644
--- a/cpp/src/transform/row_bit_count.cu
+++ b/cpp/src/transform/row_bit_count.cu
@@ -552,6 +552,7 @@ std::unique_ptr<column> segmented_row_bit_count(table_view const& t,
     {mcv.data<size_type>(), static_cast<std::size_t>(mcv.size())},
     segment_length,
     h_info.max_branch_depth);
+  CUDF_CUDA_TRY(cudaGetLastError());
 
   return output;
 }