diff --git a/.github/workflows/breeze.yml b/.github/workflows/breeze.yml
new file mode 100644
index 000000000000..0d778f6cac86
--- /dev/null
+++ b/.github/workflows/breeze.yml
@@ -0,0 +1,114 @@
+# Copyright (c) Facebook, Inc. and its affiliates.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+name: Breeze Linux Build
+
+on:
+  push:
+    branches:
+      - "main"
+    paths:
+      - "velox/experimental/breeze/**"
+      - "velox/external/perfetto/**"
+      - "CMake/**"
+      - "scripts/setup-ubuntu.sh"
+      - "scripts/setup-helper-functions.sh"
+      - ".github/workflows/breeze.yml"
+
+  pull_request:
+    paths:
+      - "velox/experimental/breeze/**"
+      - "velox/external/perfetto/**"
+      - "CMake/**"
+      - "scripts/setup-ubuntu.sh"
+      - "scripts/setup-helper-functions.sh"
+      - ".github/workflows/breeze.yml"
+
+permissions:
+  contents: read
+
+concurrency:
+  group: ${{ github.workflow }}-${{ github.repository }}-${{ github.head_ref || github.sha }}
+  cancel-in-progress: true
+
+jobs:
+  ubuntu-debug:
+    runs-on: ubuntu-22.04
+    # prevent errors when forks ff their main branch
+    if: ${{ github.repository == 'facebookincubator/velox' }}
+    name: "Ubuntu debug"
+    defaults:
+      run:
+        shell: bash
+        working-directory: velox
+    steps:
+
+      - uses: actions/checkout@v4
+        with:
+          path: velox
+
+      - name: Install Dependencies
+        run: |
+          source scripts/setup-ubuntu.sh && install_apt_deps
+
+      - name: Make Debug Build
+        env:
+          VELOX_DEPENDENCY_SOURCE: BUNDLED
+        # OpenMP build with asan+ubsan enabled
+        run: |
+          cmake -S velox/experimental/breeze -B _build-breeze/debug \
+                -DCMAKE_BUILD_TYPE=Asan \
+                -DCMAKE_CXX_FLAGS="-fsanitize=undefined" \
+                -DBUILD_GENERATE_TEST_FIXTURES=OFF \
+                -DBUILD_OPENMP=ON
+          cmake --build _build-breeze/debug -j 8
+
+      - name: Run Tests
+        run: |
+          cd _build-breeze/debug && ctest -j 8 --output-on-failure --no-tests=error
+
+  ubuntu-gpu-relwithdebinfo:
+    runs-on: 4-core-ubuntu-gpu-t4
+    # prevent errors when forks ff their main branch
+    if: ${{ github.repository == 'facebookincubator/velox' }}
+    name: "Ubuntu GPU debug"
+    env:
+      CUDA_VERSION: "12.2"
+    defaults:
+      run:
+        shell: bash
+        working-directory: velox
+    steps:
+
+      - uses: actions/checkout@v4
+        with:
+          path: velox
+
+      - name: Install Dependencies
+        run: |
+          source scripts/setup-ubuntu.sh && install_apt_deps && install_cuda ${CUDA_VERSION}
+          sudo chmod 755 -R /usr/local/lib/python3.10/dist-packages
+
+      - name: Make RelWithDebInfo Build
+        run: |
+          cmake -S velox/experimental/breeze -B _build-breeze/relwithdebinfo \
+                -DCMAKE_BUILD_TYPE=RelWithDebInfo \
+                -DBUILD_GENERATE_TEST_FIXTURES=OFF \
+                -DBUILD_CUDA=ON \
+                -DCMAKE_NVCC_FLAGS="-arch=native"
+          cmake --build _build-breeze/relwithdebinfo -j 8
+
+      - name: Run Tests
+        run: |
+          cd _build-breeze/relwithdebinfo && ctest -j 8 --output-on-failure --no-tests=error
diff --git a/velox/experimental/breeze/CMakeLists.txt b/velox/experimental/breeze/CMakeLists.txt
index b3beedf31a64..122f0c995b04 100644
--- a/velox/experimental/breeze/CMakeLists.txt
+++ b/velox/experimental/breeze/CMakeLists.txt
@@ -39,6 +39,7 @@ option(BUILD_OPENCL "Build OpenCL tests." OFF)
 option(BUILD_OPENMP "Build OpenMP tests." OFF)
 option(BUILD_METAL "Build Metal tests." OFF)
 option(BUILD_TRACING "Build tracing." ON)
+option(BUILD_GENERATE_TEST_FIXTURES "Generate test fixtures at build time." ON)
 
 if(NOT DEFINED PERFTEST_EXT_TYPES)
   set(PERFTEST_EXT_TYPES
diff --git a/velox/experimental/breeze/cmake/cuda.cmake b/velox/experimental/breeze/cmake/cuda.cmake
index c4761b10dcd8..732a8a23aef2 100644
--- a/velox/experimental/breeze/cmake/cuda.cmake
+++ b/velox/experimental/breeze/cmake/cuda.cmake
@@ -96,7 +96,7 @@ function(breeze_add_cuda_test target source)
     OUTPUT ${target}
     COMMAND
       ${NVCC_EXECUTABLE} -o ${target} ${target}.o ${arg_LIBS}
-      $<TARGET_FILE_DIR:gtest>/libgtest.a
+      $<TARGET_FILE_DIR:GTest::gtest>/libgtest.a
       $<TARGET_FILE_DIR:test_main>/libtest_main.a
       $<$<BOOL:${BUILD_TRACING}>:$<TARGET_FILE_DIR:perfetto>/libperfetto.a>
       ${ARCH_LINK_FLAGS}
@@ -107,8 +107,6 @@ function(breeze_add_cuda_test target source)
                PROPERTY IMPORTED_LOCATION ${CMAKE_CURRENT_BINARY_DIR}/${target})
   gtest_discover_tests(${target}_TESTS TEST_PREFIX cuda: DISCOVERY_MODE
                                                          PRE_TEST)
-  install(PROGRAMS ${CMAKE_CURRENT_BINARY_DIR}/${target}
-          DESTINATION ${CMAKE_INSTALL_BINDIR})
   if(DEFINED CUDA_EXPECTED_RESOURCE_USAGE_DIR)
     if(EXISTS "${CUDA_EXPECTED_RESOURCE_USAGE_DIR}/${target}-expected.txt")
       set(GET_RESOURCE_USAGE_CMDLINE
diff --git a/velox/experimental/breeze/cmake/hip.cmake b/velox/experimental/breeze/cmake/hip.cmake
index 01c468110efc..f1462577138c 100644
--- a/velox/experimental/breeze/cmake/hip.cmake
+++ b/velox/experimental/breeze/cmake/hip.cmake
@@ -39,7 +39,7 @@ function(breeze_add_hip_test target source)
     OUTPUT ${target}
     COMMAND
       ${HIPCC_EXECUTABLE} -o ${target} ${target}.o
-      $<TARGET_FILE_DIR:gtest>/libgtest.a
+      $<TARGET_FILE_DIR:GTest::gtest>/libgtest.a
       $<TARGET_FILE_DIR:test_main>/libtest_main.a
       $<$<BOOL:${BUILD_TRACING}>:$<TARGET_FILE_DIR:perfetto>/libperfetto.a>
       ${ARCH_LINK_FLAGS}
@@ -49,6 +49,4 @@ function(breeze_add_hip_test target source)
   set_property(TARGET ${target}_TESTS
                PROPERTY IMPORTED_LOCATION ${CMAKE_CURRENT_BINARY_DIR}/${target})
   gtest_discover_tests(${target}_TESTS TEST_PREFIX hip: DISCOVERY_MODE PRE_TEST)
-  install(PROGRAMS ${CMAKE_CURRENT_BINARY_DIR}/${target}
-          DESTINATION ${CMAKE_INSTALL_BINDIR})
 endfunction()
diff --git a/velox/experimental/breeze/cmake/sycl.cmake b/velox/experimental/breeze/cmake/sycl.cmake
index 345b473b30a5..51c26acfeeb2 100644
--- a/velox/experimental/breeze/cmake/sycl.cmake
+++ b/velox/experimental/breeze/cmake/sycl.cmake
@@ -40,7 +40,7 @@ function(breeze_add_sycl_test target source)
     OUTPUT ${target}
     COMMAND
       ${SYCLCC_EXECUTABLE} -o ${target} ${target}.o
-      $<TARGET_FILE_DIR:gtest>/libgtest.a
+      $<TARGET_FILE_DIR:GTest::gtest>/libgtest.a
       $<TARGET_FILE_DIR:test_main>/libtest_main.a ${CMAKE_THREAD_LIBS_INIT}
       ${ARCH_LINK_FLAGS}
       $<$<BOOL:${BUILD_TRACING}>:$<TARGET_FILE_DIR:perfetto>/libperfetto.a>
@@ -51,6 +51,4 @@ function(breeze_add_sycl_test target source)
                PROPERTY IMPORTED_LOCATION ${CMAKE_CURRENT_BINARY_DIR}/${target})
   gtest_discover_tests(${target}_TESTS TEST_PREFIX sycl: DISCOVERY_MODE
                                                          PRE_TEST)
-  install(PROGRAMS ${CMAKE_CURRENT_BINARY_DIR}/${target}
-          DESTINATION ${CMAKE_INSTALL_BINDIR})
 endfunction()
diff --git a/velox/experimental/breeze/test/CMakeLists.txt b/velox/experimental/breeze/test/CMakeLists.txt
index 069f14e45bfd..0afdd04880c8 100644
--- a/velox/experimental/breeze/test/CMakeLists.txt
+++ b/velox/experimental/breeze/test/CMakeLists.txt
@@ -18,7 +18,7 @@
 
 include_directories(${CMAKE_SOURCE_DIR})
 
-cxx_library(test_main "${cxx_strict}" test_main.cpp)
+add_library(test_main test_main.cpp)
 target_link_libraries(
   test_main
   PUBLIC gtest $<$<BOOL:${BUILD_TRACING}>:perfetto>)
@@ -160,101 +160,103 @@ function(
 endfunction()
 
 function(breeze_add_unittests dir type)
-  if(BUILD_OPENCL)
-    generate_kernel(
-      ${type}_opencl_kernels_src
-      ${dir}
-      ${type}
-      "opencl"
-      "generated/${dir}/kernels-opencl.h")
-    generate_test_fixture(
-      ${type}_opencl_test_fixture_src
-      ${dir}
-      ${type}
-      "opencl"
-      "generated/${dir}/${type}_test-opencl.h")
-    set(_opencl_kernels_src ${type}_opencl_kernels_src)
-    set(_opencl_test_fixture_src ${type}_opencl_test_fixture_src)
-  endif()
-  if(BUILD_METAL)
-    generate_kernel(
-      ${type}_metal_kernels_src
-      ${dir}
-      ${type}
-      "metal"
-      "generated/${dir}/kernels-metal.h")
-    generate_test_fixture(
-      ${type}_metal_test_fixture_src
-      ${dir}
-      ${type}
-      "metal"
-      "generated/${dir}/${type}_test-metal.h")
-    set(_metal_kernels_src ${type}_metal_kernels_src)
-    set(_metal_test_fixture_src ${type}_metal_test_fixture_src)
-  endif()
-  if(BUILD_CUDA)
-    generate_kernel(
-      ${type}_cuda_kernels_src
-      ${dir}
-      ${type}
-      "cuda"
-      "generated/${dir}/kernels.cuh")
-    generate_test_fixture(
-      ${type}_cuda_test_fixture_src
-      ${dir}
-      ${type}
-      "cuda"
-      "generated/${dir}/${type}_test-cuda.cuh")
-    set(_cuda_kernels_src ${type}_cuda_kernels_src)
-    set(_cuda_test_fixture_src ${type}_cuda_test_fixture_src)
-  endif()
-  if(BUILD_HIP)
-    generate_kernel(
-      ${type}_hip_kernels_src
-      ${dir}
-      ${type}
-      "hip"
-      "generated/${dir}/kernels-hip.hpp")
-    generate_test_fixture(
-      ${type}_hip_test_fixture_src
-      ${dir}
-      ${type}
-      "hip"
-      "generated/${dir}/${type}_test-hip.hpp")
-    set(_hip_kernels_src ${type}_hip_kernels_src)
-    set(_hip_test_fixture_src ${type}_hip_test_fixture_src)
-  endif()
-  if(BUILD_SYCL)
-    generate_kernel(
-      ${type}_sycl_kernels_src
-      ${dir}
-      ${type}
-      "sycl"
-      "generated/${dir}/kernels-sycl.hpp")
-    generate_test_fixture(
-      ${type}_sycl_test_fixture_src
-      ${dir}
-      ${type}
-      "sycl"
-      "generated/${dir}/${type}_test-sycl.hpp")
-    set(_sycl_kernels_src ${type}_sycl_kernels_src)
-    set(_sycl_test_fixture_src ${type}_sycl_test_fixture_src)
-  endif()
-  if(BUILD_OPENMP)
-    generate_kernel(
-      ${type}_openmp_kernels_src
-      ${dir}
-      ${type}
-      "openmp"
-      "generated/${dir}/kernels-openmp.h")
-    generate_test_fixture(
-      ${type}_openmp_test_fixture_src
-      ${dir}
-      ${type}
-      "openmp"
-      "generated/${dir}/${type}_test-openmp.h")
-    set(_openmp_kernels_src ${type}_openmp_kernels_src)
-    set(_openmp_test_fixture_src ${type}_openmp_test_fixture_src)
+  if(BUILD_GENERATE_TEST_FIXTURES)
+    if(BUILD_OPENCL)
+      generate_kernel(
+        ${type}_opencl_kernels_src
+        ${dir}
+        ${type}
+        "opencl"
+        "generated/${dir}/kernels-opencl.h")
+      generate_test_fixture(
+        ${type}_opencl_test_fixture_src
+        ${dir}
+        ${type}
+        "opencl"
+        "generated/${dir}/${type}_test-opencl.h")
+      set(_opencl_kernels_src ${type}_opencl_kernels_src)
+      set(_opencl_test_fixture_src ${type}_opencl_test_fixture_src)
+    endif()
+    if(BUILD_METAL)
+      generate_kernel(
+        ${type}_metal_kernels_src
+        ${dir}
+        ${type}
+        "metal"
+        "generated/${dir}/kernels-metal.h")
+      generate_test_fixture(
+        ${type}_metal_test_fixture_src
+        ${dir}
+        ${type}
+        "metal"
+        "generated/${dir}/${type}_test-metal.h")
+      set(_metal_kernels_src ${type}_metal_kernels_src)
+      set(_metal_test_fixture_src ${type}_metal_test_fixture_src)
+    endif()
+    if(BUILD_CUDA)
+      generate_kernel(
+        ${type}_cuda_kernels_src
+        ${dir}
+        ${type}
+        "cuda"
+        "generated/${dir}/kernels-cuda.cuh")
+      generate_test_fixture(
+        ${type}_cuda_test_fixture_src
+        ${dir}
+        ${type}
+        "cuda"
+        "generated/${dir}/${type}_test-cuda.cuh")
+      set(_cuda_kernels_src ${type}_cuda_kernels_src)
+      set(_cuda_test_fixture_src ${type}_cuda_test_fixture_src)
+    endif()
+    if(BUILD_HIP)
+      generate_kernel(
+        ${type}_hip_kernels_src
+        ${dir}
+        ${type}
+        "hip"
+        "generated/${dir}/kernels-hip.hpp")
+      generate_test_fixture(
+        ${type}_hip_test_fixture_src
+        ${dir}
+        ${type}
+        "hip"
+        "generated/${dir}/${type}_test-hip.hpp")
+      set(_hip_kernels_src ${type}_hip_kernels_src)
+      set(_hip_test_fixture_src ${type}_hip_test_fixture_src)
+    endif()
+    if(BUILD_SYCL)
+      generate_kernel(
+        ${type}_sycl_kernels_src
+        ${dir}
+        ${type}
+        "sycl"
+        "generated/${dir}/kernels-sycl.hpp")
+      generate_test_fixture(
+        ${type}_sycl_test_fixture_src
+        ${dir}
+        ${type}
+        "sycl"
+        "generated/${dir}/${type}_test-sycl.hpp")
+      set(_sycl_kernels_src ${type}_sycl_kernels_src)
+      set(_sycl_test_fixture_src ${type}_sycl_test_fixture_src)
+    endif()
+    if(BUILD_OPENMP)
+      generate_kernel(
+        ${type}_openmp_kernels_src
+        ${dir}
+        ${type}
+        "openmp"
+        "generated/${dir}/kernels-openmp.h")
+      generate_test_fixture(
+        ${type}_openmp_test_fixture_src
+        ${dir}
+        ${type}
+        "openmp"
+        "generated/${dir}/${type}_test-openmp.h")
+      set(_openmp_kernels_src ${type}_openmp_kernels_src)
+      set(_openmp_test_fixture_src ${type}_openmp_test_fixture_src)
+    endif()
   endif()
   breeze_add_tests(
     ${dir}
diff --git a/velox/experimental/breeze/test/generate.sh b/velox/experimental/breeze/test/generate.sh
index fc41e4d5edad..1d65793bd99b 100755
--- a/velox/experimental/breeze/test/generate.sh
+++ b/velox/experimental/breeze/test/generate.sh
@@ -30,3 +30,5 @@ function generate {
 
 generate openmp "algorithm" h
 generate openmp "function" h
+generate cuda "algorithm" cuh
+generate cuda "function" cuh
diff --git a/velox/experimental/breeze/test/generated/algorithms/algorithm_test-cuda.cuh b/velox/experimental/breeze/test/generated/algorithms/algorithm_test-cuda.cuh
new file mode 100644
index 000000000000..2f96445832d4
--- /dev/null
+++ b/velox/experimental/breeze/test/generated/algorithms/algorithm_test-cuda.cuh
@@ -0,0 +1,88 @@
+/*
+ * Copyright (c) Facebook, Inc. and its affiliates.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+/*
+ * Copyright (c) 2024 by Rivos Inc.
+ * Licensed under the Apache License, Version 2.0, see LICENSE for details.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+/*
+ * This file is auto-generated from test_fixture_generator.py
+ * DO NOT EDIT!
+ */
+
+#include <gtest/gtest.h>
+
+#include <vector>
+
+#include "breeze/platforms/cuda.cuh"
+#include "test/generated/algorithms/kernels-cuda.cuh"
+#include "test/platforms/cuda_test.cuh"
+
+template <typename T>
+class AlgorithmTest : public ::testing::Test {
+ protected:
+  template <typename ReduceOp, int BLOCK_THREADS, int ITEMS_PER_THREAD,
+            typename U>
+  void Reduce(const std::vector<T>& in, U* out, int num_blocks) {
+    std::vector<U> vec_out(1, *out);
+    CudaTestLaunch<BLOCK_THREADS>(
+        num_blocks,
+        &kernels::Reduce<ReduceOp, BLOCK_THREADS, ITEMS_PER_THREAD, T, U>, in,
+        vec_out, in.size());
+    *out = vec_out[0];
+  }
+
+  template <typename ScanOp, int BLOCK_THREADS, int ITEMS_PER_THREAD,
+            int LOOKBACK_DISTANCE, typename U, typename V>
+  void Scan(const std::vector<T>& in, std::vector<U>& out, int* next_blocks_idx,
+            std::vector<V>& blocks, int num_blocks) {
+    std::vector<int> vec_next_blocks_idx(1, *next_blocks_idx);
+    CudaTestLaunch<BLOCK_THREADS>(
+        num_blocks,
+        &kernels::Scan<ScanOp, BLOCK_THREADS, ITEMS_PER_THREAD,
+                       LOOKBACK_DISTANCE, T, U, V>,
+        in, out, vec_next_blocks_idx, blocks, in.size());
+    *next_blocks_idx = vec_next_blocks_idx[0];
+  }
+
+  template <int BLOCK_THREADS, int ITEMS_PER_THREAD, int TILE_SIZE,
+            int RADIX_BITS>
+  void RadixSortHistogram(const std::vector<T>& in, std::vector<unsigned>& out,
+                          int num_blocks) {
+    CudaTestLaunch<BLOCK_THREADS>(
+        num_blocks,
+        &kernels::RadixSortHistogram<BLOCK_THREADS, ITEMS_PER_THREAD, TILE_SIZE,
+                                     RADIX_BITS, T>,
+        in, out, in.size());
+  }
+
+  template <int BLOCK_THREADS, int ITEMS_PER_THREAD, int RADIX_BITS>
+  void RadixSort(const std::vector<T>& in,
+                 const std::vector<unsigned>& in_offsets, int start_bit,
+                 int num_pass_bits, std::vector<T>& out,
+                 std::vector<int>& next_block_idx,
+                 std::vector<unsigned>& blocks, int num_blocks) {
+    const std::vector<int> vec_start_bit(1, start_bit);
+    const std::vector<int> vec_num_pass_bits(1, num_pass_bits);
+    CudaTestLaunch<BLOCK_THREADS>(
+        num_blocks,
+        &kernels::RadixSort<BLOCK_THREADS, ITEMS_PER_THREAD, RADIX_BITS, T>, in,
+        in_offsets, vec_start_bit, vec_num_pass_bits, out, next_block_idx,
+        blocks, in.size());
+  }
+};
diff --git a/velox/experimental/breeze/test/generated/algorithms/kernels-cuda.cuh b/velox/experimental/breeze/test/generated/algorithms/kernels-cuda.cuh
new file mode 100644
index 000000000000..e3e864b90698
--- /dev/null
+++ b/velox/experimental/breeze/test/generated/algorithms/kernels-cuda.cuh
@@ -0,0 +1,120 @@
+/*
+ * Copyright (c) Facebook, Inc. and its affiliates.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+/*
+ * Copyright (c) 2024 by Rivos Inc.
+ * Licensed under the Apache License, Version 2.0, see LICENSE for details.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+/*
+ * This file is auto-generated from kernel_generator.py
+ * DO NOT EDIT!
+ */
+
+#include "breeze/algorithms/reduce.h"
+#include "breeze/algorithms/scan.h"
+#include "breeze/algorithms/sort.h"
+#include "breeze/platforms/cuda.cuh"
+#include "breeze/platforms/platform.h"
+#include "breeze/utils/types.h"
+
+namespace kernels {
+
+enum { WARP_THREADS = 32 };
+template <typename Op, int BLOCK_THREADS, int ITEMS_PER_THREAD, typename T,
+          typename U>
+__global__ void Reduce(const T* in, U* out, int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  __shared__
+      typename breeze::algorithms::DeviceReduce<PlatformT, U>::Scratch scratch_;
+  auto scratch =
+      (typename breeze::algorithms::DeviceReduce<PlatformT,
+                                                 U>::Scratch*)&scratch_;
+
+  breeze::algorithms::DeviceReduce<PlatformT, U>::template Reduce<
+      Op, ITEMS_PER_THREAD>(
+      p, breeze::utils::make_slice<breeze::utils::GLOBAL>(in),
+      breeze::utils::make_slice<breeze::utils::GLOBAL>(out),
+      breeze::utils::make_slice<breeze::utils::SHARED>(scratch), num_items);
+}
+
+template <typename Op, int BLOCK_THREADS, int ITEMS_PER_THREAD,
+          int LOOKBACK_DISTANCE, typename T, typename U, typename V>
+__global__ void Scan(const T* in, U* out, int* next_block_idx, V* blocks,
+                     int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  __shared__ typename breeze::algorithms::DeviceScan<
+      PlatformT, U, ITEMS_PER_THREAD, LOOKBACK_DISTANCE>::Scratch scratch_;
+  auto scratch = (typename breeze::algorithms::DeviceScan<
+                  PlatformT, U, ITEMS_PER_THREAD,
+                  LOOKBACK_DISTANCE>::Scratch*)&scratch_;
+
+  breeze::algorithms::DeviceScan<PlatformT, U, ITEMS_PER_THREAD,
+                                 LOOKBACK_DISTANCE>::
+      template Scan<Op>(
+          p, breeze::utils::make_slice<breeze::utils::GLOBAL>(in),
+          breeze::utils::make_slice<breeze::utils::GLOBAL>(out),
+          breeze::utils::make_slice<breeze::utils::GLOBAL>(next_block_idx),
+          breeze::utils::make_slice<breeze::utils::GLOBAL>(blocks),
+          breeze::utils::make_slice<breeze::utils::SHARED>(scratch), num_items);
+}
+
+template <int BLOCK_THREADS, int ITEMS_PER_THREAD, int TILE_SIZE,
+          int RADIX_BITS, typename T>
+__global__ void RadixSortHistogram(const T* in, unsigned* out, int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  __shared__ typename breeze::algorithms::DeviceRadixSortHistogram<
+      RADIX_BITS, T>::Scratch scratch_;
+  auto scratch = (typename breeze::algorithms::DeviceRadixSortHistogram<
+                  RADIX_BITS, T>::Scratch*)&scratch_;
+
+  breeze::algorithms::DeviceRadixSortHistogram<RADIX_BITS, T>::template Build<
+      ITEMS_PER_THREAD, TILE_SIZE>(
+      p, breeze::utils::make_slice<breeze::utils::GLOBAL>(in),
+      breeze::utils::make_slice<breeze::utils::GLOBAL>(out),
+      breeze::utils::make_slice<breeze::utils::SHARED>(scratch), num_items);
+}
+
+template <int BLOCK_THREADS, int ITEMS_PER_THREAD, int RADIX_BITS, typename T>
+__global__ void RadixSort(const T* in, const unsigned* in_offsets,
+                          const int* start_bit, const int* num_pass_bits,
+                          T* out, int* next_block_idx, unsigned* blocks,
+                          int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  __shared__ typename breeze::algorithms::DeviceRadixSort<
+      PlatformT, ITEMS_PER_THREAD, RADIX_BITS, T>::Scratch scratch_;
+  auto scratch =
+      (typename breeze::algorithms::DeviceRadixSort<
+          PlatformT, ITEMS_PER_THREAD, RADIX_BITS, T>::Scratch*)&scratch_;
+
+  breeze::algorithms::DeviceRadixSort<PlatformT, ITEMS_PER_THREAD, RADIX_BITS,
+                                      T>::
+      template Sort<unsigned>(
+          p, breeze::utils::make_slice<breeze::utils::GLOBAL>(in),
+          breeze::utils::make_slice<breeze::utils::GLOBAL>(in_offsets),
+          *start_bit, *num_pass_bits,
+          breeze::utils::make_slice<breeze::utils::GLOBAL>(out),
+          breeze::utils::make_slice<breeze::utils::GLOBAL>(next_block_idx),
+          breeze::utils::make_slice<breeze::utils::GLOBAL>(blocks),
+          breeze::utils::make_slice<breeze::utils::SHARED>(scratch), num_items);
+}
+
+}  // namespace kernels
diff --git a/velox/experimental/breeze/test/generated/functions/function_test-cuda.cuh b/velox/experimental/breeze/test/generated/functions/function_test-cuda.cuh
new file mode 100644
index 000000000000..ce451a90c8b7
--- /dev/null
+++ b/velox/experimental/breeze/test/generated/functions/function_test-cuda.cuh
@@ -0,0 +1,160 @@
+/*
+ * Copyright (c) Facebook, Inc. and its affiliates.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+/*
+ * Copyright (c) 2024 by Rivos Inc.
+ * Licensed under the Apache License, Version 2.0, see LICENSE for details.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+/*
+ * This file is auto-generated from test_fixture_generator.py
+ * DO NOT EDIT!
+ */
+
+#include <gtest/gtest.h>
+
+#include <vector>
+
+#include "breeze/platforms/cuda.cuh"
+#include "test/generated/functions/kernels-cuda.cuh"
+#include "test/platforms/cuda_test.cuh"
+
+template <typename T>
+class FunctionTest : public ::testing::Test {
+ protected:
+  template <int BLOCK_THREADS, int ITEMS_PER_THREAD>
+  void BlockLoad(const std::vector<T>& in, std::vector<T>& out) {
+    CudaTestLaunch<BLOCK_THREADS>(
+        /*num_blocks=*/1,
+        &kernels::BlockLoad<BLOCK_THREADS, ITEMS_PER_THREAD, T>, in, out,
+        in.size());
+  }
+
+  template <int BLOCK_THREADS, int ITEMS_PER_THREAD>
+  void BlockLoadIf(const std::vector<T>& in,
+                   const std::vector<int>& selection_flags,
+                   std::vector<T>& out) {
+    CudaTestLaunch<BLOCK_THREADS>(
+        /*num_blocks=*/1,
+        &kernels::BlockLoadIf<BLOCK_THREADS, ITEMS_PER_THREAD, T>, in,
+        selection_flags, out, in.size());
+  }
+
+  template <int BLOCK_THREADS, int ITEMS_PER_THREAD>
+  void BlockLoadFrom(const std::vector<T>& in, const std::vector<int>& offsets,
+                     std::vector<T>& out) {
+    CudaTestLaunch<BLOCK_THREADS>(
+        /*num_blocks=*/1,
+        &kernels::BlockLoadFrom<BLOCK_THREADS, ITEMS_PER_THREAD, T>, in,
+        offsets, out, in.size());
+  }
+
+  template <int BLOCK_THREADS, int ITEMS_PER_THREAD>
+  void BlockStore(const std::vector<T>& in, std::vector<T>& out) {
+    CudaTestLaunch<BLOCK_THREADS>(
+        /*num_blocks=*/1,
+        &kernels::BlockStore<BLOCK_THREADS, ITEMS_PER_THREAD, T>, in, out,
+        out.size());
+  }
+
+  template <int BLOCK_THREADS, int ITEMS_PER_THREAD>
+  void BlockStoreIf(const std::vector<T>& in,
+                    const std::vector<int>& selection_flags,
+                    std::vector<T>& out) {
+    CudaTestLaunch<BLOCK_THREADS>(
+        /*num_blocks=*/1,
+        &kernels::BlockStoreIf<BLOCK_THREADS, ITEMS_PER_THREAD, T>, in,
+        selection_flags, out, out.size());
+  }
+
+  template <int BLOCK_THREADS, int ITEMS_PER_THREAD>
+  void BlockStoreAt(const std::vector<T>& in, const std::vector<int>& offsets,
+                    std::vector<T>& out) {
+    CudaTestLaunch<BLOCK_THREADS>(
+        /*num_blocks=*/1,
+        &kernels::BlockStoreAt<BLOCK_THREADS, ITEMS_PER_THREAD, T>, in, offsets,
+        out, out.size());
+  }
+
+  template <int BLOCK_THREADS, int ITEMS_PER_THREAD>
+  void BlockStoreAtIf(const std::vector<T>& in, const std::vector<int>& offsets,
+                      const std::vector<int>& selection_flags,
+                      std::vector<T>& out) {
+    CudaTestLaunch<BLOCK_THREADS>(
+        /*num_blocks=*/1,
+        &kernels::BlockStoreAtIf<BLOCK_THREADS, ITEMS_PER_THREAD, T>, in,
+        offsets, selection_flags, out, out.size());
+  }
+
+  template <int BLOCK_THREADS, int ITEMS_PER_THREAD>
+  void BlockFill(T value, std::vector<T>& out) {
+    const std::vector<T> vec_value(1, value);
+    CudaTestLaunch<BLOCK_THREADS>(
+        /*num_blocks=*/1,
+        &kernels::BlockFill<BLOCK_THREADS, ITEMS_PER_THREAD, T>, vec_value, out,
+        out.size());
+  }
+
+  template <int BLOCK_THREADS, int ITEMS_PER_THREAD>
+  void BlockFillAtIf(T value, const std::vector<int>& offsets,
+                     const std::vector<int>& selection_flags,
+                     std::vector<T>& out) {
+    const std::vector<T> vec_value(1, value);
+    CudaTestLaunch<BLOCK_THREADS>(
+        /*num_blocks=*/1,
+        &kernels::BlockFillAtIf<BLOCK_THREADS, ITEMS_PER_THREAD, T>, vec_value,
+        offsets, selection_flags, out, out.size());
+  }
+
+  template <typename ReduceOp, int BLOCK_THREADS, int ITEMS_PER_THREAD,
+            typename U>
+  void BlockReduce(const std::vector<T>& in, U* out) {
+    std::vector<U> vec_out(1, *out);
+    CudaTestLaunch<BLOCK_THREADS>(
+        /*num_blocks=*/1,
+        &kernels::BlockReduce<ReduceOp, BLOCK_THREADS, ITEMS_PER_THREAD, T, U>,
+        in, vec_out, in.size());
+    *out = vec_out[0];
+  }
+
+  template <typename ScanOp, int BLOCK_THREADS, int ITEMS_PER_THREAD,
+            typename U>
+  void BlockScan(const std::vector<T>& in, std::vector<U>& out) {
+    CudaTestLaunch<BLOCK_THREADS>(
+        /*num_blocks=*/1,
+        &kernels::BlockScan<ScanOp, BLOCK_THREADS, ITEMS_PER_THREAD, T, U>, in,
+        out, in.size());
+  }
+
+  template <int BLOCK_THREADS, int ITEMS_PER_THREAD, int RADIX_BITS>
+  void BlockRadixRank(const std::vector<T>& in, std::vector<int>& out) {
+    CudaTestLaunch<BLOCK_THREADS>(
+        /*num_blocks=*/1,
+        &kernels::BlockRadixRank<BLOCK_THREADS, ITEMS_PER_THREAD, RADIX_BITS,
+                                 T>,
+        in, out, in.size());
+  }
+
+  template <int BLOCK_THREADS, int ITEMS_PER_THREAD, int RADIX_BITS>
+  void BlockRadixSort(const std::vector<T>& in, std::vector<T>& out) {
+    CudaTestLaunch<BLOCK_THREADS>(
+        /*num_blocks=*/1,
+        &kernels::BlockRadixSort<BLOCK_THREADS, ITEMS_PER_THREAD, RADIX_BITS,
+                                 T>,
+        in, out, in.size());
+  }
+};
diff --git a/velox/experimental/breeze/test/generated/functions/kernels-cuda.cuh b/velox/experimental/breeze/test/generated/functions/kernels-cuda.cuh
new file mode 100644
index 000000000000..3aa089397bbb
--- /dev/null
+++ b/velox/experimental/breeze/test/generated/functions/kernels-cuda.cuh
@@ -0,0 +1,284 @@
+/*
+ * Copyright (c) Facebook, Inc. and its affiliates.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+/*
+ * Copyright (c) 2024 by Rivos Inc.
+ * Licensed under the Apache License, Version 2.0, see LICENSE for details.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+/*
+ * This file is auto-generated from kernel_generator.py
+ * DO NOT EDIT!
+ */
+
+#include "breeze/functions/load.h"
+#include "breeze/functions/reduce.h"
+#include "breeze/functions/scan.h"
+#include "breeze/functions/sort.h"
+#include "breeze/functions/store.h"
+#include "breeze/platforms/cuda.cuh"
+#include "breeze/platforms/platform.h"
+
+namespace kernels {
+
+enum { WARP_THREADS = 32 };
+template <int BLOCK_THREADS, int ITEMS_PER_THREAD, typename T>
+__global__ void BlockLoad(const T* in, T* out, int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  T items[ITEMS_PER_THREAD];
+  breeze::functions::BlockLoad<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p, breeze::utils::make_slice<breeze::utils::GLOBAL>(in),
+      breeze::utils::make_slice(items), num_items);
+  breeze::functions::BlockStore<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p, breeze::utils::make_slice(items),
+      breeze::utils::make_slice<breeze::utils::GLOBAL>(out), num_items);
+}
+
+template <int BLOCK_THREADS, int ITEMS_PER_THREAD, typename T>
+__global__ void BlockLoadIf(const T* in, const int* in_selection_flags, T* out,
+                            int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  int selection_flags[ITEMS_PER_THREAD];
+  breeze::functions::BlockLoad<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p, breeze::utils::make_slice<breeze::utils::GLOBAL>(in_selection_flags),
+      breeze::utils::make_slice(selection_flags), num_items);
+  T items[ITEMS_PER_THREAD];
+  breeze::functions::BlockLoadIf<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p, breeze::utils::make_slice<breeze::utils::GLOBAL>(in),
+      breeze::utils::make_slice(selection_flags),
+      breeze::utils::make_slice(items), num_items);
+  breeze::functions::BlockStoreIf<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p, breeze::utils::make_slice(items),
+      breeze::utils::make_slice(selection_flags),
+      breeze::utils::make_slice<breeze::utils::GLOBAL>(out), num_items);
+}
+
+template <int BLOCK_THREADS, int ITEMS_PER_THREAD, typename T>
+__global__ void BlockLoadFrom(const T* in, const int* in_offsets, T* out,
+                              int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  int offsets[ITEMS_PER_THREAD];
+  breeze::functions::BlockLoad<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p, breeze::utils::make_slice<breeze::utils::GLOBAL>(in_offsets),
+      breeze::utils::make_slice(offsets), num_items);
+  T items[ITEMS_PER_THREAD];
+  breeze::functions::BlockLoadFrom<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p, breeze::utils::make_slice<breeze::utils::GLOBAL>(in),
+      breeze::utils::make_slice(offsets), breeze::utils::make_slice(items),
+      num_items);
+  breeze::functions::BlockStore<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p, breeze::utils::make_slice(items),
+      breeze::utils::make_slice<breeze::utils::GLOBAL>(out), num_items);
+}
+
+template <int BLOCK_THREADS, int ITEMS_PER_THREAD, typename T>
+__global__ void BlockStore(const T* in, T* out, int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  breeze::functions::BlockStore<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p,
+      breeze::utils::make_slice<breeze::utils::GLOBAL, breeze::utils::STRIPED>(
+          in),
+      breeze::utils::make_slice<breeze::utils::GLOBAL>(out), num_items);
+}
+
+template <int BLOCK_THREADS, int ITEMS_PER_THREAD, typename T>
+__global__ void BlockStoreIf(const T* in, const int* selection_flags, T* out,
+                             int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  breeze::functions::BlockStoreIf<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p,
+      breeze::utils::make_slice<breeze::utils::GLOBAL, breeze::utils::STRIPED>(
+          in),
+      breeze::utils::make_slice<breeze::utils::GLOBAL, breeze::utils::STRIPED>(
+          selection_flags),
+      breeze::utils::make_slice<breeze::utils::GLOBAL>(out), num_items);
+}
+
+template <int BLOCK_THREADS, int ITEMS_PER_THREAD, typename T>
+__global__ void BlockStoreAt(const T* in, const int* offsets, T* out,
+                             int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  breeze::functions::BlockStoreAt<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p,
+      breeze::utils::make_slice<breeze::utils::GLOBAL, breeze::utils::STRIPED>(
+          in),
+      breeze::utils::make_slice<breeze::utils::GLOBAL, breeze::utils::STRIPED>(
+          offsets),
+      breeze::utils::make_slice<breeze::utils::GLOBAL>(out), num_items);
+}
+
+template <int BLOCK_THREADS, int ITEMS_PER_THREAD, typename T>
+__global__ void BlockStoreAtIf(const T* in, const int* offsets,
+                               const int* selection_flags, T* out,
+                               int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  breeze::functions::BlockStoreAtIf<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p,
+      breeze::utils::make_slice<breeze::utils::GLOBAL, breeze::utils::STRIPED>(
+          in),
+      breeze::utils::make_slice<breeze::utils::GLOBAL, breeze::utils::STRIPED>(
+          offsets),
+      breeze::utils::make_slice<breeze::utils::GLOBAL, breeze::utils::STRIPED>(
+          selection_flags),
+      breeze::utils::make_slice<breeze::utils::GLOBAL>(out), num_items);
+}
+
+template <int BLOCK_THREADS, int ITEMS_PER_THREAD, typename T>
+__global__ void BlockFill(const T* value, T* out, int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  breeze::functions::BlockFill<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p, *value, breeze::utils::make_slice<breeze::utils::GLOBAL>(out),
+      num_items);
+}
+
+template <int BLOCK_THREADS, int ITEMS_PER_THREAD, typename T>
+__global__ void BlockFillAtIf(const T* value, const int* offsets,
+                              const int* selection_flags, T* out,
+                              int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  breeze::functions::BlockFillAtIf<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p, *value,
+      breeze::utils::make_slice<breeze::utils::GLOBAL, breeze::utils::STRIPED>(
+          offsets),
+      breeze::utils::make_slice<breeze::utils::GLOBAL, breeze::utils::STRIPED>(
+          selection_flags),
+      breeze::utils::make_slice<breeze::utils::GLOBAL>(out), num_items);
+}
+
+template <typename Op, int BLOCK_THREADS, int ITEMS_PER_THREAD, typename T,
+          typename U>
+__global__ void BlockReduce(const T* in, U* out, int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  __shared__
+      typename breeze::functions::BlockReduce<PlatformT, U>::Scratch scratch_;
+  auto scratch =
+      (typename breeze::functions::BlockReduce<PlatformT,
+                                               U>::Scratch*)&scratch_;
+
+  T items[ITEMS_PER_THREAD];
+  breeze::functions::BlockLoad<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p, breeze::utils::make_slice<breeze::utils::GLOBAL>(in),
+      breeze::utils::make_slice(items), num_items);
+  U aggregate = breeze::functions::BlockReduce<PlatformT, U>::template Reduce<
+      Op, ITEMS_PER_THREAD>(
+      p, breeze::utils::make_slice(items),
+      breeze::utils::make_slice<breeze::utils::SHARED>(scratch), num_items);
+  p.syncthreads();
+  if (p.thread_idx() == 0) {
+    *out = aggregate;
+  }
+}
+
+template <typename Op, int BLOCK_THREADS, int ITEMS_PER_THREAD, typename T,
+          typename U>
+__global__ void BlockScan(const T* in, U* out, int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  __shared__
+      typename breeze::functions::BlockScan<PlatformT, U,
+                                            ITEMS_PER_THREAD>::Scratch scratch_;
+  auto scratch = (typename breeze::functions::BlockScan<
+                  PlatformT, U, ITEMS_PER_THREAD>::Scratch*)&scratch_;
+
+  T items[ITEMS_PER_THREAD];
+  breeze::functions::BlockLoad<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p, breeze::utils::make_slice<breeze::utils::GLOBAL>(in),
+      breeze::utils::make_slice(items), num_items);
+  U sums[ITEMS_PER_THREAD];
+  breeze::functions::BlockScan<PlatformT, U, ITEMS_PER_THREAD>::template Scan<
+      Op>(p, breeze::utils::make_slice(items), breeze::utils::make_slice(sums),
+          breeze::utils::make_slice<breeze::utils::SHARED>(scratch), num_items);
+  breeze::functions::BlockStore<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p, breeze::utils::make_slice(sums),
+      breeze::utils::make_slice<breeze::utils::GLOBAL>(out), num_items);
+}
+
+template <int BLOCK_THREADS, int ITEMS_PER_THREAD, int RADIX_BITS, typename T>
+__global__ void BlockRadixRank(const T* in, int* out, int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  __shared__
+      typename breeze::functions::BlockRadixRank<PlatformT, ITEMS_PER_THREAD,
+                                                 RADIX_BITS>::Scratch scratch_;
+  auto scratch = (typename breeze::functions::BlockRadixRank<
+                  PlatformT, ITEMS_PER_THREAD, RADIX_BITS>::Scratch*)&scratch_;
+
+  T items[ITEMS_PER_THREAD];
+  // initialize invalid items to max value
+  for (int i = 0; i < ITEMS_PER_THREAD; ++i) {
+    items[i] = static_cast<T>((1 << RADIX_BITS) - 1);
+  }
+  breeze::functions::BlockLoad<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p, breeze::utils::make_slice<breeze::utils::GLOBAL>(in),
+      breeze::utils::make_slice<breeze::utils::THREAD,
+                                breeze::utils::WARP_STRIPED>(items),
+      num_items);
+  int ranks[ITEMS_PER_THREAD];
+  breeze::functions::BlockRadixRank<PlatformT, ITEMS_PER_THREAD, RADIX_BITS>::
+      Rank(p,
+           breeze::utils::make_slice<breeze::utils::THREAD,
+                                     breeze::utils::WARP_STRIPED>(items),
+           breeze::utils::make_slice<breeze::utils::THREAD,
+                                     breeze::utils::WARP_STRIPED>(ranks),
+           breeze::utils::make_slice<breeze::utils::SHARED>(scratch));
+  breeze::functions::BlockStore<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p,
+      breeze::utils::make_slice<breeze::utils::THREAD,
+                                breeze::utils::WARP_STRIPED>(ranks),
+      breeze::utils::make_slice<breeze::utils::GLOBAL>(out), num_items);
+}
+
+template <int BLOCK_THREADS, int ITEMS_PER_THREAD, int RADIX_BITS, typename T>
+__global__ void BlockRadixSort(const T* in, T* out, int num_items) {
+  using PlatformT = CudaPlatform<BLOCK_THREADS, WARP_THREADS>;
+  PlatformT p;
+  __shared__ typename breeze::functions::BlockRadixSort<
+      PlatformT, ITEMS_PER_THREAD, RADIX_BITS, T>::Scratch scratch_;
+  auto scratch =
+      (typename breeze::functions::BlockRadixSort<
+          PlatformT, ITEMS_PER_THREAD, RADIX_BITS, T>::Scratch*)&scratch_;
+
+  T items[ITEMS_PER_THREAD];
+  breeze::functions::BlockLoad<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p, breeze::utils::make_slice<breeze::utils::GLOBAL>(in),
+      breeze::utils::make_slice<breeze::utils::THREAD,
+                                breeze::utils::WARP_STRIPED>(items),
+      num_items);
+  breeze::functions::
+      BlockRadixSort<PlatformT, ITEMS_PER_THREAD, RADIX_BITS, T>::Sort(
+          p,
+          breeze::utils::make_slice<breeze::utils::THREAD,
+                                    breeze::utils::WARP_STRIPED>(items),
+          breeze::utils::make_slice<breeze::utils::SHARED>(scratch), num_items);
+  breeze::functions::BlockStore<BLOCK_THREADS, ITEMS_PER_THREAD>(
+      p,
+      breeze::utils::make_slice<breeze::utils::THREAD,
+                                breeze::utils::WARP_STRIPED>(items),
+      breeze::utils::make_slice<breeze::utils::GLOBAL>(out), num_items);
+}
+
+}  // namespace kernels
diff --git a/velox/experimental/breeze/test/test_fixture_generator.py b/velox/experimental/breeze/test/test_fixture_generator.py
index 36425fac93d9..cbb6d7d19e81 100755
--- a/velox/experimental/breeze/test/test_fixture_generator.py
+++ b/velox/experimental/breeze/test/test_fixture_generator.py
@@ -275,7 +275,7 @@ def __init__(self):
     def includes(self, fixture_name):
         test_type = fixture_name.replace("Test", "").lower()
         return f"""
-            #include "test/generated/{test_type}s/kernels.cuh"
+            #include "test/generated/{test_type}s/kernels-cuda.cuh"
             #include "breeze/platforms/cuda.cuh"
             #include "test/platforms/cuda_test.cuh"
         """