Add CUDA error checks after every kernel launch

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);
 }
 

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));

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());
   }
 }
 

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));

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

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

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

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;

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());
   }
 }
 

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());
   }
 }
 

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

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;
 }

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

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

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

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

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

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

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));
     }
   }

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