diff --git a/cpp/src/strings/copying/concatenate.cu b/cpp/src/strings/copying/concatenate.cu
index 287910c9a6f..26cd4fff09b 100644
--- a/cpp/src/strings/copying/concatenate.cu
+++ b/cpp/src/strings/copying/concatenate.cu
@@ -121,8 +121,8 @@ __global__ void fused_concatenate_string_offset_kernel(column_device_view const*
                                                        bitmask_type* output_mask,
                                                        size_type* out_valid_count)
 {
-  size_type output_index     = threadIdx.x + blockIdx.x * blockDim.x;
-  size_type warp_valid_count = 0;
+  cudf::thread_index_type output_index = threadIdx.x + blockIdx.x * blockDim.x;
+  size_type warp_valid_count           = 0;
 
   unsigned active_mask;
   if (Nullable) { active_mask = __ballot_sync(0xFFFF'FFFFu, output_index < output_size); }
@@ -175,7 +175,7 @@ __global__ void fused_concatenate_string_chars_kernel(column_device_view const*
                                                       size_type const output_size,
                                                       char* output_data)
 {
-  size_type output_index = threadIdx.x + blockIdx.x * blockDim.x;
+  cudf::thread_index_type output_index = threadIdx.x + blockIdx.x * blockDim.x;
 
   while (output_index < output_size) {
     // Lookup input index by searching for output index in offsets
diff --git a/cpp/tests/copying/concatenate_tests.cpp b/cpp/tests/copying/concatenate_tests.cpp
index 0c6394637fc..c81f1772d10 100644
--- a/cpp/tests/copying/concatenate_tests.cpp
+++ b/cpp/tests/copying/concatenate_tests.cpp
@@ -162,6 +162,37 @@ TEST_F(StringColumnTest, ConcatenateColumnView)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
 }
 
+TEST_F(StringColumnTest, ConcatenateColumnViewLarge)
+{
+  // Test large concatenate, causes out of bound device memory errors if kernel
+  // indexing is not int64_t.
+  // 1.5GB bytes, 5k columns
+  constexpr size_t num_strings        = 10000;
+  constexpr size_t string_length      = 150000;
+  constexpr size_t strings_per_column = 2;
+  constexpr size_t num_columns        = num_strings / strings_per_column;
+
+  std::vector<std::string> strings;
+  std::vector<char const*> h_strings;
+  std::vector<cudf::test::strings_column_wrapper> strings_column_wrappers;
+  std::vector<cudf::column_view> strings_columns;
+
+  std::string s(string_length, 'a');
+  for (size_t i = 0; i < num_strings; ++i)
+    h_strings.push_back(s.data());
+
+  for (size_t i = 0; i < num_columns; ++i)
+    strings_column_wrappers.push_back(cudf::test::strings_column_wrapper(
+      h_strings.data() + i * strings_per_column, h_strings.data() + (i + 1) * strings_per_column));
+  for (auto& wrapper : strings_column_wrappers)
+    strings_columns.push_back(wrapper);
+
+  auto results = cudf::concatenate(strings_columns);
+
+  cudf::test::strings_column_wrapper expected(h_strings.begin(), h_strings.end());
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
+}
+
 TEST_F(StringColumnTest, ConcatenateTooManyColumns)
 {
   std::vector<char const*> h_strings{"aaa",