[NVIDIA] Add support Thor, Spark and GB300 and CUDA 13 (#776)

* Update setup.py

* refactor: replace cub::CTA_SYNC() with __syncthreads() for better clarity

* Update setup.py

Author: johnnynunez

csrc/selective_scan/reverse_scan.cuh

@@ -129,7 +129,7 @@ struct WarpReverseScan {
     /// Constructor
     explicit __device__ __forceinline__
     WarpReverseScan()
-        : lane_id(cub::LaneId())
+        : lane_id(threadIdx.x & 0x1f)
         , warp_id(IS_ARCH_WARP ? 0 : (lane_id / LOGICAL_WARP_THREADS))
         , member_mask(cub::WarpMask<LOGICAL_WARP_THREADS>(warp_id))
     {
@@ -332,7 +332,7 @@ struct BlockReverseScan {
             // Place thread partial into shared memory raking grid
             T *placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
             detail::uninitialized_copy(placement_ptr, input);
-            cub::CTA_SYNC();
+            __syncthreads();
             // Reduce parallelism down to just raking threads
             if (linear_tid < RAKING_THREADS) {
                 WarpReverseScan warp_scan;
@@ -350,7 +350,7 @@ struct BlockReverseScan {
                 // Exclusive raking downsweep scan
                 ExclusiveDownsweep(scan_op, downsweep_postfix);
             }
-            cub::CTA_SYNC();
+            __syncthreads();
             // Grab thread postfix from shared memory
             exclusive_output = *placement_ptr;
 
@@ -382,7 +382,7 @@ struct BlockReverseScan {
             //     }
             // }
 
-            // cub::CTA_SYNC();
+            // __syncthreads();
 
             // // Incorporate thread block postfix into outputs
             // T block_postfix = temp_storage.block_postfix;

setup.py

@@ -195,6 +195,13 @@ def append_nvcc_threads(nvcc_extra_args):
             cc_flag.append("arch=compute_100,code=sm_100")
             cc_flag.append("-gencode")
             cc_flag.append("arch=compute_120,code=sm_120")
+        if bare_metal_version >= Version("13.0"):
+            cc_flag.append("-gencode")
+            cc_flag.append("arch=compute_103,code=sm_103")
+            cc_flag.append("-gencode")
+            cc_flag.append("arch=compute_110,code=sm_110")
+            cc_flag.append("-gencode")
+            cc_flag.append("arch=compute_121,code=sm_121")
 
 
     # HACK: The compiler flag -D_GLIBCXX_USE_CXX11_ABI is set to be the same as