                                      size_t N,
                                      cudaStream_t stream)
 {
-  std::vector<size_t> levels;
-  calcLevelSizes(levels, N);
-
-  std::vector<size_t> offsets;
-  offsets.push_back(0);
-  size_t alignment = 128 / sizeof(uint32_t);
-  for (size_t i = 0; i < levels.size(); i++) {
-    size_t levelSize = (levels[i] + alignment - 1) & ~alignment;
-    offsets.push_back(offsets[i] + levelSize);
-  }
-
   if (N <= std::numeric_limits<uint32_t>::max()) {
+    uint32_t n = static_cast<uint32_t>(N);
+
+    std::vector<uint32_t> levels;
+    calcLevelSizes(levels, n);
+
+    std::vector<uint32_t> offsets;
+    offsets.push_back(0);
+    size_t alignment = 128 / sizeof(uint32_t);
+    for (size_t i = 0; i < levels.size(); i++) {
+      size_t levelSize = (levels[i] + alignment - 1) & ~alignment;
+      offsets.push_back(offsets[i] + levelSize);
+    }
+
     uint32_t blockSize = SCAN_WARPS * WARPSIZE;
-    uint32_t n = static_cast<uint32_t>(N);
 
     if (levels.empty()) {
 

    
  }


  void calcLevelSizes(std::vector<size_t>& levels, size_t N)
  {
    size_t R = 4 * SCAN_WARPS*WARPSIZE; // Amount of reduction per level.
    levels.clear();
    while (1 < N) {
      N = (N + R - 1) / R;


    }
    else {
      uint32_t L = static_cast<uint32_t>(levels.size());
        
      // From input, populate level 0
      reduce << <levels[0], blockSize, 0, stream >> > (scratch_d + offsets[0],
                                                       counts_d,
                                                       n);
      // From level i-1, populate level i, up to including L-1.
      for (uint32_t i = 1; i < L; i++) {
        reduce << <levels[i], blockSize, 0, stream >> > (scratch_d + offsets[i],
                                                         scratch_d + offsets[i - 1],
                                                         levels[i - 1]);
      }

      // Run scan on last level L-1, and write off total sum to last element of output (offsets_d+N).
      auto l = levels.size() - 1;
      if (sum_d == nullptr) {
        scan<false, true, false><<<1, blockSize, 0, stream>>>(scratch_d + offsets[l],
