Custom cuda and metal kernels for system validation#274
Conversation
Luthaf
force-pushed
the
kernels
branch
3 times, most recently
from
c2584f3 to
607c96a
Compare
| let cell_array: ArrayView2<f64> = cell.try_into()?; | ||
| for i in 0..3 { | ||
| if !pbc_array[i] && !cell_array.row(i).iter().all(|&x| x == 0.0) { | ||
| return Err(Error::InvalidParameter(format!( | ||
| "invalid cell: for non-periodic dimensions, the corresponding \ | ||
| cell vector must be zero, but cell[{}] contains non-zero values", | ||
| i | ||
| ))); | ||
| } | ||
| } |
There was a problem hiding this comment.
This is essentially the same as above, but with a different type for ArrayView2. An helper function would avoid code duplication (and make it easier to add support to more exotic types down the line if needed).
There was a problem hiding this comment.
fixed with a macro (rust generics don't support unconstrained "template parameters", they always require "concepts" i.e. traits for all parameters)
| int i = threadIdx.x; | ||
| if (i < 3) { | ||
| if (!pbc[i]) { | ||
| if (cell[3 * i] != T(0) || cell[3 * i + 1] != T(0) || cell[3 * i + 2] != T(0)) { |
There was a problem hiding this comment.
Given this indexing, worth checking that DLPack tensors are contiguous in memory before calling the kernel? If I understand correctly, they could potentially be strided in general.
There was a problem hiding this comment.
This was quite painful and added a bunch of code, but should be fine now
| eprintln!( | ||
| "is_equal_i32 for device {:?} is not implemented", | ||
| tensor.device() | ||
| ); | ||
| Ok(()) |
There was a problem hiding this comment.
I don't know since these are only checks that don't change the model behavior, I would rather skip the checks than error here, to make it easier to add new devices in metatomic without having to add code everywhere
| eprintln!( | ||
| "Cell/PBC validation for device {:?} is not implemented", | ||
| pbc.device() | ||
| ); | ||
| Ok(()) |
There was a problem hiding this comment.
As above, shouldn't this be an error?
Luthaf
force-pushed
the
kernels
branch
3 times, most recently
from
7d61507 to
d25d254
Compare
2 participants
Add custom cuda and metal kernels for system validation. The intention is to add more kernels for unit conversion & any other work that needs to be done inside
mta_execute_model.No tests other than making sure the code compile for now, they will be added once we have Python bindings.
Contributor (creator of pull-request) checklist
Tests updated (for new features and bugfixes)?Issue referenced (for PRs that solve an issue)?Reviewer checklist
CHANGELOG updated with public API or any other important changes?