rustyverse/rustytorch
2
0
Fork 0
Code Issues Pull Requests Actions 46 Packages Projects Releases Wiki Activity

rtx-cfd legacy GPU modules: cudarc 0.18 API — unsafe launches, the stream accessor, memcpy_htod argument order; the renamed k-ε constants and ::standard(); get_module clones; CfdConfig test literals take ..Default
CI / Build (macos-latest) (push) Waiting to run
Details
CI / Test (macos-latest) (push) Blocked by required conditions
Details
CI / Test (ubuntu-latest) (push) Blocked by required conditions
Details
CI / Python Bindings (maturin) (macos-latest) (push) Blocked by required conditions
Details
CI / Python Bindings (maturin) (ubuntu-latest) (push) Blocked by required conditions
Details
CI / WASM Build + Size Check (push) Blocked by required conditions
Details
CI / Distributed Training Tests (push) Blocked by required conditions
Details
CI / CI Success (push) Blocked by required conditions
Details
CI / Build (ubuntu-latest) (push) Failing after 5s
Details
CI / Clippy Check (push) Failing after 5s
Details
Documentation / Build User Guide (push) Successful in 6s
Details
Documentation / Build API Documentation (push) Failing after 17s
Details
CI / Build CPU-Only (Explicit) (push) Failing after 32s
Details
CI / Format Check (push) Failing after 42s
Details
Performance Benchmarks / Run Benchmarks (push) Successful in 1m59s
Details

Browse Source
This commit is contained in:
Omar Sobh
2026-09-16 07:16:39 -05:00
parent de435eb91c
commit 5e402918bb
6 changed files with 432 additions and 378 deletions
Download Patch File Download Diff File Expand all files Collapse all files
crates/specialized/rtx-cfd/src/solvers/lbm/d3q19_gpu.rs
+93 -78
View File
@@ -395,17 +395,19 @@ __global__ void d3q19_bounce_back_boundaries(
shared_mem_bytes: 0,
};
self.kernel_manager
.stream()
.launch_builder(&func)
.arg(&mut buffers.f.clone())
.arg(&buffers.f_eq)
.arg(&(omega as f32))
.arg(&(buffers.nx as i32))
.arg(&(buffers.ny as i32))
.arg(&(buffers.nz as i32))
.launch(config)
.map_err(|e| CfdError::gpu_error(&format!("Collision kernel launch failed: {}", e)))?;
unsafe {
self.kernel_manager
.stream()
.launch_builder(&func)
.arg(&mut buffers.f.clone())
.arg(&buffers.f_eq)
.arg(&(omega as f32))
.arg(&(buffers.nx as i32))
.arg(&(buffers.ny as i32))
.arg(&(buffers.nz as i32))
.launch(config)
.map_err(|e| CfdError::gpu_error(&format!("Collision kernel launch failed: {}", e)))?;
}
self.kernel_manager.synchronize()?;
Ok(())
@@ -436,16 +438,18 @@ __global__ void d3q19_bounce_back_boundaries(
shared_mem_bytes: 0,
};
self.kernel_manager
.stream()
.launch_builder(&func)
.arg(&mut buffers.f_temp.clone())
.arg(&buffers.f)
.arg(&(buffers.nx as i32))
.arg(&(buffers.ny as i32))
.arg(&(buffers.nz as i32))
.launch(config)
.map_err(|e| CfdError::gpu_error(&format!("Streaming kernel launch failed: {}", e)))?;
unsafe {
self.kernel_manager
.stream()
.launch_builder(&func)
.arg(&mut buffers.f_temp.clone())
.arg(&buffers.f)
.arg(&(buffers.nx as i32))
.arg(&(buffers.ny as i32))
.arg(&(buffers.nz as i32))
.launch(config)
.map_err(|e| CfdError::gpu_error(&format!("Streaming kernel launch failed: {}", e)))?;
}
// Swap buffers: f = f_temp
// In real implementation, would swap buffer pointers
@@ -483,24 +487,26 @@ __global__ void d3q19_bounce_back_boundaries(
shared_mem_bytes: 0,
};
self.kernel_manager
.stream()
.launch_builder(&func)
.arg(&mut buffers.density.clone())
.arg(&mut buffers.velocity_x.clone())
.arg(&mut buffers.velocity_y.clone())
.arg(&mut buffers.velocity_z.clone())
.arg(&buffers.f)
.arg(&(buffers.nx as i32))
.arg(&(buffers.ny as i32))
.arg(&(buffers.nz as i32))
.launch(config)
.map_err(|e| {
CfdError::gpu_error(&format!(
"Macroscopic variables kernel launch failed: {}",
e
))
})?;
unsafe {
self.kernel_manager
.stream()
.launch_builder(&func)
.arg(&mut buffers.density.clone())
.arg(&mut buffers.velocity_x.clone())
.arg(&mut buffers.velocity_y.clone())
.arg(&mut buffers.velocity_z.clone())
.arg(&buffers.f)
.arg(&(buffers.nx as i32))
.arg(&(buffers.ny as i32))
.arg(&(buffers.nz as i32))
.launch(config)
.map_err(|e| {
CfdError::gpu_error(&format!(
"Macroscopic variables kernel launch failed: {}",
e
))
})?;
}
self.kernel_manager.synchronize()?;
Ok(())
@@ -531,21 +537,23 @@ __global__ void d3q19_bounce_back_boundaries(
shared_mem_bytes: 0,
};
self.kernel_manager
.stream()
.launch_builder(&func)
.arg(&mut buffers.f_eq.clone())
.arg(&buffers.density)
.arg(&buffers.velocity_x)
.arg(&buffers.velocity_y)
.arg(&buffers.velocity_z)
.arg(&(buffers.nx as i32))
.arg(&(buffers.ny as i32))
.arg(&(buffers.nz as i32))
.launch(config)
.map_err(|e| {
CfdError::gpu_error(&format!("Equilibrium kernel launch failed: {}", e))
})?;
unsafe {
self.kernel_manager
.stream()
.launch_builder(&func)
.arg(&mut buffers.f_eq.clone())
.arg(&buffers.density)
.arg(&buffers.velocity_x)
.arg(&buffers.velocity_y)
.arg(&buffers.velocity_z)
.arg(&(buffers.nx as i32))
.arg(&(buffers.ny as i32))
.arg(&(buffers.nz as i32))
.launch(config)
.map_err(|e| {
CfdError::gpu_error(&format!("Equilibrium kernel launch failed: {}", e))
})?;
}
self.kernel_manager.synchronize()?;
Ok(())
@@ -575,15 +583,17 @@ __global__ void d3q19_bounce_back_boundaries(
shared_mem_bytes: 0,
};
self.kernel_manager
.stream()
.launch_builder(&func)
.arg(&mut buffers.f.clone())
.arg(&(buffers.nx as i32))
.arg(&(buffers.ny as i32))
.arg(&(buffers.nz as i32))
.launch(config)
.map_err(|e| CfdError::gpu_error(&format!("Boundary kernel launch failed: {}", e)))?;
unsafe {
self.kernel_manager
.stream()
.launch_builder(&func)
.arg(&mut buffers.f.clone())
.arg(&(buffers.nx as i32))
.arg(&(buffers.ny as i32))
.arg(&(buffers.nz as i32))
.launch(config)
.map_err(|e| CfdError::gpu_error(&format!("Boundary kernel launch failed: {}", e)))?;
}
self.kernel_manager.synchronize()?;
Ok(())
@@ -628,21 +638,23 @@ __global__ void d3q19_bounce_back_boundaries(
shared_mem_bytes: 0,
};
self.kernel_manager
.stream()
.launch_builder(&func)
.arg(&mut buffers.f.clone())
.arg(&(density as f32))
.arg(&(velocity.x as f32))
.arg(&(velocity.y as f32))
.arg(&(velocity.z as f32))
.arg(&(buffers.nx as i32))
.arg(&(buffers.ny as i32))
.arg(&(buffers.nz as i32))
.launch(config)
.map_err(|e| {
CfdError::gpu_error(&format!("Initialization kernel launch failed: {}", e))
})?;
unsafe {
self.kernel_manager
.stream()
.launch_builder(&func)
.arg(&mut buffers.f.clone())
.arg(&(density as f32))
.arg(&(velocity.x as f32))
.arg(&(velocity.y as f32))
.arg(&(velocity.z as f32))
.arg(&(buffers.nx as i32))
.arg(&(buffers.ny as i32))
.arg(&(buffers.nz as i32))
.launch(config)
.map_err(|e| {
CfdError::gpu_error(&format!("Initialization kernel launch failed: {}", e))
})?;
}
self.kernel_manager.synchronize()?;
Ok(())
@@ -771,6 +783,7 @@ mod tests {
viscosity: 0.01,
density: 1.0,
device_id: 0,
..CfdConfig::default()
};
let params = D3Q19Parameters::default();
@@ -799,6 +812,7 @@ mod tests {
viscosity: 0.01,
density: 1.0,
device_id: 0,
..CfdConfig::default()
};
let params = D3Q19Parameters::default();
@@ -839,6 +853,7 @@ mod tests {
viscosity: 0.01,
density: 1.0,
device_id: 0,
..CfdConfig::default()
};
let params = D3Q19Parameters::default();