cust: use usize in BlockSize and GridSize.

Instead of `u32`, because using `usize` for dimensions and indices is
more natural in Rust and avoids lots of casts.

Author: nnethercote

crates/cust/src/function.rs

@@ -18,43 +18,43 @@ use crate::module::Module;
 #[derive(Debug, Clone, Copy, PartialEq, Eq)]
 pub struct GridSize {
     /// Width of grid in blocks
-    pub x: u32,
+    pub x: usize,
     /// Height of grid in blocks
-    pub y: u32,
+    pub y: usize,
     /// Depth of grid in blocks
-    pub z: u32,
+    pub z: usize,
 }
 impl GridSize {
     /// Create a one-dimensional grid of `x` blocks
     #[inline]
-    pub fn x(x: u32) -> GridSize {
+    pub fn x(x: usize) -> GridSize {
         GridSize { x, y: 1, z: 1 }
     }
 
     /// Create a two-dimensional grid of `x * y` blocks
     #[inline]
-    pub fn xy(x: u32, y: u32) -> GridSize {
+    pub fn xy(x: usize, y: usize) -> GridSize {
         GridSize { x, y, z: 1 }
     }
 
     /// Create a three-dimensional grid of `x * y * z` blocks
     #[inline]
-    pub fn xyz(x: u32, y: u32, z: u32) -> GridSize {
+    pub fn xyz(x: usize, y: usize, z: usize) -> GridSize {
         GridSize { x, y, z }
     }
 }
-impl From<u32> for GridSize {
-    fn from(x: u32) -> GridSize {
+impl From<usize> for GridSize {
+    fn from(x: usize) -> GridSize {
         GridSize::x(x)
     }
 }
-impl From<(u32, u32)> for GridSize {
-    fn from((x, y): (u32, u32)) -> GridSize {
+impl From<(usize, usize)> for GridSize {
+    fn from((x, y): (usize, usize)) -> GridSize {
         GridSize::xy(x, y)
     }
 }
-impl From<(u32, u32, u32)> for GridSize {
-    fn from((x, y, z): (u32, u32, u32)) -> GridSize {
+impl From<(usize, usize, usize)> for GridSize {
+    fn from((x, y, z): (usize, usize, usize)) -> GridSize {
         GridSize::xyz(x, y, z)
     }
 }
@@ -64,52 +64,28 @@ impl From<&GridSize> for GridSize {
     }
 }
 #[cfg(feature = "vek")]
-impl From<vek::Vec2<u32>> for GridSize {
-    fn from(vec: vek::Vec2<u32>) -> Self {
-        GridSize::xy(vec.x, vec.y)
-    }
-}
-#[cfg(feature = "vek")]
-impl From<vek::Vec3<u32>> for GridSize {
-    fn from(vec: vek::Vec3<u32>) -> Self {
-        GridSize::xyz(vec.x, vec.y, vec.z)
-    }
-}
-#[cfg(feature = "vek")]
 impl From<vek::Vec2<usize>> for GridSize {
     fn from(vec: vek::Vec2<usize>) -> Self {
-        GridSize::xy(vec.x as u32, vec.y as u32)
+        GridSize::xy(vec.x, vec.y)
     }
 }
 #[cfg(feature = "vek")]
 impl From<vek::Vec3<usize>> for GridSize {
     fn from(vec: vek::Vec3<usize>) -> Self {
-        GridSize::xyz(vec.x as u32, vec.y as u32, vec.z as u32)
-    }
-}
-
-#[cfg(feature = "glam")]
-impl From<glam::UVec2> for GridSize {
-    fn from(vec: glam::UVec2) -> Self {
-        GridSize::xy(vec.x, vec.y)
-    }
-}
-#[cfg(feature = "glam")]
-impl From<glam::UVec3> for GridSize {
-    fn from(vec: glam::UVec3) -> Self {
         GridSize::xyz(vec.x, vec.y, vec.z)
     }
 }
+
 #[cfg(feature = "glam")]
 impl From<glam::USizeVec2> for GridSize {
     fn from(vec: glam::USizeVec2) -> Self {
-        GridSize::xy(vec.x as u32, vec.y as u32)
+        GridSize::xy(vec.x, vec.y)
     }
 }
 #[cfg(feature = "glam")]
 impl From<glam::USizeVec3> for GridSize {
     fn from(vec: glam::USizeVec3) -> Self {
-        GridSize::xyz(vec.x as u32, vec.y as u32, vec.z as u32)
+        GridSize::xyz(vec.x, vec.y, vec.z)
     }
 }
 
@@ -123,43 +99,43 @@ impl From<glam::USizeVec3> for GridSize {
 #[derive(Debug, Clone, Copy, PartialEq, Eq)]
 pub struct BlockSize {
     /// X dimension of each thread block
-    pub x: u32,
+    pub x: usize,
     /// Y dimension of each thread block
-    pub y: u32,
+    pub y: usize,
     /// Z dimension of each thread block
-    pub z: u32,
+    pub z: usize,
 }
 impl BlockSize {
     /// Create a one-dimensional block of `x` threads
     #[inline]
-    pub fn x(x: u32) -> BlockSize {
+    pub fn x(x: usize) -> BlockSize {
         BlockSize { x, y: 1, z: 1 }
     }
 
     /// Create a two-dimensional block of `x * y` threads
     #[inline]
-    pub fn xy(x: u32, y: u32) -> BlockSize {
+    pub fn xy(x: usize, y: usize) -> BlockSize {
         BlockSize { x, y, z: 1 }
     }
 
     /// Create a three-dimensional block of `x * y * z` threads
     #[inline]
-    pub fn xyz(x: u32, y: u32, z: u32) -> BlockSize {
+    pub fn xyz(x: usize, y: usize, z: usize) -> BlockSize {
         BlockSize { x, y, z }
     }
 }
-impl From<u32> for BlockSize {
-    fn from(x: u32) -> BlockSize {
+impl From<usize> for BlockSize {
+    fn from(x: usize) -> BlockSize {
         BlockSize::x(x)
     }
 }
-impl From<(u32, u32)> for BlockSize {
-    fn from((x, y): (u32, u32)) -> BlockSize {
+impl From<(usize, usize)> for BlockSize {
+    fn from((x, y): (usize, usize)) -> BlockSize {
         BlockSize::xy(x, y)
     }
 }
-impl From<(u32, u32, u32)> for BlockSize {
-    fn from((x, y, z): (u32, u32, u32)) -> BlockSize {
+impl From<(usize, usize, usize)> for BlockSize {
+    fn from((x, y, z): (usize, usize, usize)) -> BlockSize {
         BlockSize::xyz(x, y, z)
     }
 }
@@ -169,52 +145,28 @@ impl From<&BlockSize> for BlockSize {
     }
 }
 #[cfg(feature = "vek")]
-impl From<vek::Vec2<u32>> for BlockSize {
-    fn from(vec: vek::Vec2<u32>) -> Self {
-        BlockSize::xy(vec.x, vec.y)
-    }
-}
-#[cfg(feature = "vek")]
-impl From<vek::Vec3<u32>> for BlockSize {
-    fn from(vec: vek::Vec3<u32>) -> Self {
-        BlockSize::xyz(vec.x, vec.y, vec.z)
-    }
-}
-#[cfg(feature = "vek")]
 impl From<vek::Vec2<usize>> for BlockSize {
     fn from(vec: vek::Vec2<usize>) -> Self {
-        BlockSize::xy(vec.x as u32, vec.y as u32)
+        BlockSize::xy(vec.x, vec.y)
     }
 }
 #[cfg(feature = "vek")]
 impl From<vek::Vec3<usize>> for BlockSize {
     fn from(vec: vek::Vec3<usize>) -> Self {
-        BlockSize::xyz(vec.x as u32, vec.y as u32, vec.z as u32)
-    }
-}
-
-#[cfg(feature = "glam")]
-impl From<glam::UVec2> for BlockSize {
-    fn from(vec: glam::UVec2) -> Self {
-        BlockSize::xy(vec.x, vec.y)
-    }
-}
-#[cfg(feature = "glam")]
-impl From<glam::UVec3> for BlockSize {
-    fn from(vec: glam::UVec3) -> Self {
         BlockSize::xyz(vec.x, vec.y, vec.z)
     }
 }
+
 #[cfg(feature = "glam")]
 impl From<glam::USizeVec2> for BlockSize {
     fn from(vec: glam::USizeVec2) -> Self {
-        BlockSize::xy(vec.x as u32, vec.y as u32)
+        BlockSize::xy(vec.x, vec.y)
     }
 }
 #[cfg(feature = "glam")]
 impl From<glam::USizeVec3> for BlockSize {
     fn from(vec: glam::USizeVec3) -> Self {
-        BlockSize::xyz(vec.x as u32, vec.y as u32, vec.z as u32)
+        BlockSize::xyz(vec.x, vec.y, vec.z)
     }
 }
 
@@ -448,7 +400,7 @@ impl Function<'_> {
         &self,
         dynamic_smem_size: usize,
         block_size_limit: BlockSize,
-    ) -> CudaResult<(u32, u32)> {
+    ) -> CudaResult<(usize, usize)> {
         let mut min_grid_size = MaybeUninit::uninit();
         let mut block_size = MaybeUninit::uninit();
 
@@ -465,8 +417,8 @@ impl Function<'_> {
             )
             .to_result()?;
             Ok((
-                min_grid_size.assume_init() as u32,
-                block_size.assume_init() as u32,
+                min_grid_size.assume_init() as usize,
+                block_size.assume_init() as usize,
             ))
         }
     }

crates/cust/src/graph.rs

@@ -85,12 +85,12 @@ impl KernelInvocation {
     pub fn to_raw(self) -> driver_sys::CUDA_KERNEL_NODE_PARAMS {
         driver_sys::CUDA_KERNEL_NODE_PARAMS {
             func: self.func,
-            gridDimX: self.grid_dim.x,
-            gridDimY: self.grid_dim.y,
-            gridDimZ: self.grid_dim.z,
-            blockDimX: self.block_dim.x,
-            blockDimY: self.block_dim.y,
-            blockDimZ: self.block_dim.z,
+            gridDimX: self.grid_dim.x as u32,
+            gridDimY: self.grid_dim.y as u32,
+            gridDimZ: self.grid_dim.z as u32,
+            blockDimX: self.block_dim.x as u32,
+            blockDimY: self.block_dim.y as u32,
+            blockDimZ: self.block_dim.z as u32,
             kernelParams: Box::into_raw(self.params),
             sharedMemBytes: self.shared_mem_bytes,
             extra: ptr::null_mut(),
@@ -109,8 +109,16 @@ impl KernelInvocation {
     pub unsafe fn from_raw(raw: driver_sys::CUDA_KERNEL_NODE_PARAMS) -> Self {
         Self {
             func: raw.func,
-            grid_dim: GridSize::xyz(raw.gridDimX, raw.gridDimY, raw.gridDimZ),
-            block_dim: BlockSize::xyz(raw.blockDimX, raw.gridDimY, raw.gridDimZ),
+            grid_dim: GridSize::xyz(
+                raw.gridDimX as usize,
+                raw.gridDimY as usize,
+                raw.gridDimZ as usize,
+            ),
+            block_dim: BlockSize::xyz(
+                raw.blockDimX as usize,
+                raw.gridDimY as usize,
+                raw.gridDimZ as usize,
+            ),
             params: Box::from_raw(raw.kernelParams),
             shared_mem_bytes: raw.sharedMemBytes,
             params_len: None,

crates/cust/src/stream.rs

@@ -274,12 +274,12 @@ impl Stream {
         unsafe {
             driver_sys::cuLaunchKernel(
                 func.to_raw(),
-                grid_size.x,
-                grid_size.y,
-                grid_size.z,
-                block_size.x,
-                block_size.y,
-                block_size.z,
+                grid_size.x as u32,
+                grid_size.y as u32,
+                grid_size.z as u32,
+                block_size.x as u32,
+                block_size.y as u32,
+                block_size.z as u32,
                 shared_mem_bytes,
                 self.inner,
                 args.as_ptr() as *mut _,

examples/gemm/src/main.rs

@@ -365,16 +365,12 @@ pub fn gemm_naive(
     // This will try to maximize how much of the GPU is used by finding the best launch configuration for the
     // current CUDA device/architecture.
     let (_, block_size) = kernel.suggested_launch_configuration(0, 0.into())?;
-    let block_size = block_size as usize;
     let (block_size_x, block_size_y) = if block_size > m * n {
-        (block_size.div_ceil(m) as u32, m as u32)
+        (block_size.div_ceil(m), m)
     } else {
-        (1, block_size as u32)
+        (1, block_size)
     };
-    let (grid_size_x, grid_size_y) = (
-        (m as u32).div_ceil(block_size_x),
-        (n as u32).div_ceil(block_size_y),
-    );
+    let (grid_size_x, grid_size_y) = (m.div_ceil(block_size_x), n.div_ceil(block_size_y));
     unsafe {
         launch!(
             kernel<<<
@@ -438,12 +434,12 @@ pub fn gemm_tiled(
     });
     let kernel = &*kernel_cell;
 
-    let (grid_size_x, grid_size_y) = (n.div_ceil(TILE_SIZE) as u32, m.div_ceil(TILE_SIZE) as u32);
+    let (grid_size_x, grid_size_y) = (n.div_ceil(TILE_SIZE), m.div_ceil(TILE_SIZE));
     unsafe {
         launch!(
             kernel<<<
                 (grid_size_x, grid_size_y),
-                (TILE_SIZE as u32, TILE_SIZE as u32),
+                (TILE_SIZE, TILE_SIZE),
                 0,
                 stream
             >>>(

examples/i128_demo/src/main.rs

@@ -119,8 +119,8 @@ fn main() -> Result<(), Box<dyn Error>> {
     let urem_gpu = DeviceBuffer::from_slice(&vec![0u128; len])?;
     let srem_gpu = DeviceBuffer::from_slice(&vec![0u128; len])?;
 
-    let block_size = 128u32;
-    let grid_size = (len as u32).div_ceil(block_size);
+    let block_size = 128usize;
+    let grid_size = len.div_ceil(block_size);
 
     unsafe {
         launch!(
@@ -241,7 +241,7 @@ fn main() -> Result<(), Box<dyn Error>> {
 
     let trap_launch = unsafe {
         launch!(
-            kernel<<<1u32, 1u32, 0, trap_stream>>>(
+            kernel<<<1, 1, 0, trap_stream>>>(
                 trap_a.as_device_ptr(),
                 trap_a.len(),
                 trap_b.as_device_ptr(),

examples/vecadd/src/main.rs

@@ -45,7 +45,7 @@ fn main() -> Result<(), Box<dyn Error>> {
     // current CUDA device/architecture.
     let (_, block_size) = vecadd.suggested_launch_configuration(0, 0.into())?;
 
-    let grid_size = (NUMBERS_LEN as u32).div_ceil(block_size);
+    let grid_size = NUMBERS_LEN.div_ceil(block_size);
 
     println!("using {grid_size} blocks and {block_size} threads per block");
 

samples/introduction/async_api/src/main.rs

@@ -37,7 +37,7 @@ fn main() -> Result<(), cust::error::CudaError> {
     let value = 26;
 
     let blocks = BlockSize::xy(512, 1);
-    let grids = GridSize::xy((N / (blocks.x as usize)).try_into().unwrap(), 1);
+    let grids = GridSize::xy(N / blocks.x, 1);
 
     let start_event = Event::new(EventFlags::DEFAULT)?;
     let stop_event = Event::new(EventFlags::DEFAULT)?;