Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Aleaverfay/kinforest from stencil #299

Open
wants to merge 55 commits into
base: master
Choose a base branch
from
Open
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
55 commits
Select commit Hold shift + click to select a range
495be6c
Sketch the creation of KinForests from precomputed restype stencils
aleaverfay May 30, 2024
515958e
Save progress
aleaverfay Jun 3, 2024
e2b4aa7
Merge branch 'master' into aleaverfay/kinforest_from_stencil
aleaverfay Aug 8, 2024
1f777a9
Add a working ground-truth kin-forest that will be the target we try …
aleaverfay Aug 12, 2024
4d6c7ae
Fixed parent definition for residue 2 to go along w/ N-conn input
aleaverfay Aug 12, 2024
4355f01
Move some code out of the unit test file into tmol/kinematics
aleaverfay Aug 15, 2024
167a555
Fix unit tests following code shuffle
aleaverfay Aug 15, 2024
54bded2
Move scan_types to its own, CUDA-independent declaration
aleaverfay Sep 10, 2024
43ec5e4
Add C++ implementation of fix-jump-nodes
aleaverfay Sep 10, 2024
5cb1e28
Okay, small section of working C++ code for a few tensor dereferencin…
aleaverfay Sep 16, 2024
25c2fee
Add parent-atom lookup logic
aleaverfay Sep 16, 2024
3d24e8f
Construct parent-atom tensor + child-span lists.
aleaverfay Sep 17, 2024
4c16934
Construct id, frame_x, frame_y, and frame_z tensors w/ C++
aleaverfay Sep 18, 2024
066a861
Add draft of nearly-complete algorithm for kinforest nodes/scans/gens…
aleaverfay Sep 26, 2024
575251e
Save incremental progress; code now compiles but crashes
aleaverfay Sep 26, 2024
9e708b1
Add correct generational path decomposition algorithm for fold forest
aleaverfay Sep 26, 2024
614352d
Add a kinforest-edge delay calculation test w/ multiple poses in a Po…
aleaverfay Sep 26, 2024
5c10581
Add topological sort of FF edges
aleaverfay Sep 27, 2024
0006bab
Uncommenting some code and reformatting
aleaverfay Sep 27, 2024
845c481
Add code to record input- and 1st-output connection for each block in…
aleaverfay Sep 27, 2024
1c729a0
Add very nearly (but not!) working alg for nodes tensor creaion
aleaverfay Oct 2, 2024
3e4423c
Fix bug where wrong variable is being used in offset calculation
aleaverfay Oct 3, 2024
973104a
Write the correct indices for the parent atoms of jumps
aleaverfay Oct 3, 2024
16878e3
Rename the paths within blocks "scan path segments" and paths within …
aleaverfay Oct 8, 2024
bed3d45
Okay, compute scans and gens tensors.
aleaverfay Oct 9, 2024
4098b23
Fix logic about scan-path-continuation
aleaverfay Oct 10, 2024
29f8e67
Remove debugging statements
aleaverfay Oct 10, 2024
2a29c7d
Add code for constructing reverse paths up the FoldForest for derivat…
aleaverfay Oct 11, 2024
fda174c
Begin code rearrangement
aleaverfay Oct 11, 2024
38279c9
Moving kinforest data construction out of unit tests
aleaverfay Oct 11, 2024
8c64b38
Add smoke test for PoseStackKinematicsOp with forward + backward trav…
aleaverfay Nov 4, 2024
28b2015
Fix bug in handling "inter-block" status for kinematic leaf residues
aleaverfay Nov 5, 2024
4306488
Fix bug in the idea of certain scan-path segments being "inter residue"
aleaverfay Nov 6, 2024
dbd1e45
Remove separate tensor for "is_sps_inter_res" as it is unnecessary
aleaverfay Nov 6, 2024
7d66203
Fix CUDA compilation of kinematics code
aleaverfay Nov 6, 2024
a390043
Merge branch 'aleaverfay/kinforest_from_stencil' of https://github.co…
aleaverfay Nov 6, 2024
c221962
Fix bug in launching kernel with no work to do
aleaverfay Nov 6, 2024
2bfadad
Saving debugging progress
aleaverfay Nov 7, 2024
6bc17a9
Fix bug in identifying the first block in a backwards scan path
aleaverfay Nov 8, 2024
3da6fd2
Comment out debugging code
aleaverfay Nov 12, 2024
c9bc906
Remove remaining debug statements in kinematics; Kinematics script mo…
aleaverfay Nov 12, 2024
9f93861
Ugh. Re-enable debugging code
aleaverfay Nov 12, 2024
34efa48
Fix bug in determining which SPSs are roots of BW scan paths
aleaverfay Nov 12, 2024
8386ae6
Fix bug in pointing jump_ff_edge_rooted_at_scan_path_seg_bw TView at …
aleaverfay Nov 14, 2024
31b58ee
Flesh out edge-delay calculation unit tests.
aleaverfay Nov 14, 2024
dd3ae89
Add another topology for fold-tree unit tests: multiple jumps leaving…
aleaverfay Nov 14, 2024
21bc108
Turn duplicated test harness code into fixtures
aleaverfay Nov 14, 2024
fb0c0da
Fix device for several unit tests
aleaverfay Nov 14, 2024
d456b61
Remove debugging statements
aleaverfay Nov 14, 2024
6587711
Merge branch 'aleaverfay/kinforest_from_stencil' of github.com:uw-ipd…
aleaverfay Nov 14, 2024
1d8f004
Add some more unit tests
aleaverfay Nov 15, 2024
006ce5d
Add "minimize all phi_c and all jump DOFs" version of kinematics mini…
aleaverfay Dec 16, 2024
b979290
First attempt at adding torsion-location data to bt-scan-path annotat…
aleaverfay Dec 17, 2024
f03b368
First pass implementation of movemap --> "minimizer map" conversion
aleaverfay Dec 20, 2024
0ad241d
Add tests for jump-index correctness to check_fold_forest
aleaverfay Dec 23, 2024
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion pytest.ini
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,7 @@
ext = .cpp,.hh,.h,.cu,.py,.yaml

[pytest]
addopts = --benchmark-disable --benchmark-columns=ops,mean,iqr
# addopts = --benchmark-disable --benchmark-columns=ops,mean,iqr
filterwarnings =
ignore:(?s).*is not compatible with the compiler Pytorch.*:
ignore:(?s).*Benchmark fixture was not used.*:
Expand Down
1 change: 1 addition & 0 deletions tmol/chemical/patched_chemdb.py
Original file line number Diff line number Diff line change
Expand Up @@ -510,6 +510,7 @@ def do_patch(res, variant, resgraph, patchgraph, marked):
icoors=res.icoors,
properties=res.properties,
chi_samples=res.chi_samples,
default_jump_connection_atom=res.default_jump_connection_atom,
)

# 1. remove atoms
Expand Down
14 changes: 14 additions & 0 deletions tmol/chemical/restypes.py
Original file line number Diff line number Diff line change
Expand Up @@ -155,6 +155,10 @@ def _setup_bond_indices(self):
bond_array.flags.writeable = False
return bond_array

@property
def n_conn(self):
return len(self.connections)

# The index of the atom for a given inter-residue connection point
connection_to_idx: Mapping[str, AtomIndex] = attr.ib()

Expand Down Expand Up @@ -259,6 +263,10 @@ def _setup_ordered_torsions(self):
)
return ordered_torsions

@property
def n_torsions(self):
return self.ordered_torsions.shape[0]

path_distance: numpy.ndarray = attr.ib()

@path_distance.default
Expand Down Expand Up @@ -448,6 +456,12 @@ def _setup_icoors_geom(self):
def compute_ideal_coords(self):
return build_coords_from_icoors(self.icoors_ancestors, self.icoors_geom)

default_jump_connection_atom_index: int = attr.ib()

@default_jump_connection_atom_index.default
def get_default_jump_connection_atom_index(self):
return self.atom_to_idx[self.default_jump_connection_atom]


@attr.s(auto_attribs=True)
class ResidueTypeSet:
Expand Down
4 changes: 4 additions & 0 deletions tmol/database/chemical/__init__.py
Original file line number Diff line number Diff line change
Expand Up @@ -136,6 +136,10 @@ class RawResidueType:
icoors: Tuple[Icoor, ...]
properties: ChemicalProperties
chi_samples: Tuple[ChiSamples, ...]
default_jump_connection_atom: str

def atom_name(self, index):
return self.atoms[index].name


@attr.s(auto_attribs=True, frozen=True, slots=True)
Expand Down
23 changes: 23 additions & 0 deletions tmol/database/default/chemical/chemical.yaml
Original file line number Diff line number Diff line change
Expand Up @@ -143,6 +143,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: ARG
base_name: ARG
name3: ARG
Expand Down Expand Up @@ -277,6 +278,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: ASN
base_name: ASN
name3: ASN
Expand Down Expand Up @@ -373,6 +375,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: ASP
base_name: ASP
name3: ASP
Expand Down Expand Up @@ -461,6 +464,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: CYS
base_name: CYS
name3: CYS
Expand Down Expand Up @@ -549,6 +553,7 @@ residues:
- chi_dihedral: chi2
samples: [60, 180, 300]
expansions: []
default_jump_connection_atom: CA
- name: CYD
base_name: CYD
name3: CYS
Expand Down Expand Up @@ -631,6 +636,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: GLN
base_name: GLN
name3: GLN
Expand Down Expand Up @@ -739,6 +745,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: GLU
base_name: GLU
name3: GLU
Expand Down Expand Up @@ -839,6 +846,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: GLY
base_name: GLY
name3: GLY
Expand Down Expand Up @@ -910,6 +918,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: HIS
base_name: HIS
name3: HIS
Expand Down Expand Up @@ -1016,6 +1025,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: HIS_D
base_name: HIS_D
name3: HIS
Expand Down Expand Up @@ -1121,6 +1131,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: ILE
base_name: ILE
name3: ILE
Expand Down Expand Up @@ -1236,6 +1247,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: LEU
base_name: LEU
name3: LEU
Expand Down Expand Up @@ -1351,6 +1363,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: LYS
base_name: LYS
name3: LYS
Expand Down Expand Up @@ -1480,6 +1493,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: MET
base_name: MET
name3: MET
Expand Down Expand Up @@ -1589,6 +1603,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: PHE
base_name: PHE
name3: PHE
Expand Down Expand Up @@ -1702,6 +1717,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: PRO
base_name: PRO
name3: PRO
Expand Down Expand Up @@ -1802,6 +1818,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: SER
base_name: SER
name3: SER
Expand Down Expand Up @@ -1890,6 +1907,7 @@ residues:
- chi_dihedral: chi2
samples: [0, 20, 40, 60, 80, 100, 120, 140, 160, 180, 200, 220, 240, 260, 280, 300, 320, 340]
expansions: []
default_jump_connection_atom: CA
- name: THR
base_name: THR
name3: THR
Expand Down Expand Up @@ -1988,6 +2006,7 @@ residues:
- chi_dihedral: chi2
samples: [0, 20, 40, 60, 80, 100, 120, 140, 160, 180, 200, 220, 240, 260, 280, 300, 320, 340]
expansions: []
default_jump_connection_atom: CA
- name: TRP
base_name: TRP
name3: TRP
Expand Down Expand Up @@ -2114,6 +2133,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: TYR
base_name: TYR
name3: TYR
Expand Down Expand Up @@ -2235,6 +2255,7 @@ residues:
- chi_dihedral: chi3
samples: [0, 180]
expansions: [20]
default_jump_connection_atom: CA
- name: VAL
base_name: VAL
name3: VAL
Expand Down Expand Up @@ -2339,6 +2360,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: CA
- name: HOH
base_name: HOH
name3: HOH
Expand Down Expand Up @@ -2374,6 +2396,7 @@ residues:
pH: 7
virtual: []
chi_samples: []
default_jump_connection_atom: O

variants:
- name: CarboxyTerminus
Expand Down
46 changes: 21 additions & 25 deletions tmol/extern/moderngpu/cta_scan.hxx
Original file line number Diff line number Diff line change
Expand Up @@ -2,14 +2,10 @@
#pragma once
#include "loadstore.hxx"
#include "intrinsics.hxx"
#include "scan_types.hxx"

BEGIN_MGPU_NAMESPACE

enum scan_type_t {
scan_type_exc,
scan_type_inc
};

template<typename type_t, int vt = 0, bool is_array = (vt > 0)>
struct scan_result_t {
type_t scan;
Expand All @@ -32,7 +28,7 @@ struct cta_scan_t {
struct { type_t threads[nt], warps[num_warps]; };
};

#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 300
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 300

//////////////////////////////////////////////////////////////////////////////
// Optimized CTA scan code that uses warp shfl intrinsics.
Expand All @@ -41,7 +37,7 @@ struct cta_scan_t {

template<typename op_t = plus_t<type_t> >
MGPU_DEVICE scan_result_t<type_t>
scan(int tid, type_t x, storage_t& storage, int count = nt, op_t op = op_t(),
scan(int tid, type_t x, storage_t& storage, int count = nt, op_t op = op_t(),
type_t init = type_t(), scan_type_t type = scan_type_exc) const {

int warp = tid / warp_size;
Expand All @@ -61,7 +57,7 @@ struct cta_scan_t {
__syncthreads();

// Scan the warp reductions.
if(tid < num_warps) {
if(tid < num_warps) {
type_t cta_scan = storage.warps[tid];
iterate<s_log2(num_warps)>([&](int pass) {
cta_scan = shfl_up_op(cta_scan, 1<< pass, op, num_warps);
Expand All @@ -78,10 +74,10 @@ struct cta_scan_t {
if(warp > 0) scan = op(scan, storage.warps[warp - 1]);

type_t reduction = storage.warps[div_up(count, warp_size) - 1];
scan_result_t<type_t> result {
tid < count ? scan : reduction,
reduction

scan_result_t<type_t> result {
tid < count ? scan : reduction,
reduction
};
__syncthreads();

Expand All @@ -91,11 +87,11 @@ struct cta_scan_t {
#else

//////////////////////////////////////////////////////////////////////////////
// Standard CTA scan code that does not use shfl intrinsics.
// Standard CTA scan code that does not use shfl intrinsics.

template<typename op_t = plus_t<type_t> >
MGPU_DEVICE scan_result_t<type_t>
scan(int tid, type_t x, storage_t& storage, int count = nt, op_t op = op_t(),
MGPU_DEVICE scan_result_t<type_t>
scan(int tid, type_t x, storage_t& storage, int count = nt, op_t op = op_t(),
type_t init = type_t(), scan_type_t type = scan_type_exc) const {

int first = 0;
Expand All @@ -113,7 +109,7 @@ struct cta_scan_t {

scan_result_t<type_t> result;
result.reduction = storage.data[first + count - 1];
result.scan = (tid < count) ?
result.scan = (tid < count) ?
(scan_type_inc == type ? x :
(tid ? storage.data[first + tid - 1] : init)) :
result.reduction;
Expand All @@ -122,16 +118,16 @@ struct cta_scan_t {
return result;
}

#endif
#endif

//////////////////////////////////////////////////////////////////////////////
// CTA vectorized scan. Accepts multiple values per thread and adds in
// CTA vectorized scan. Accepts multiple values per thread and adds in
// optional global carry-in.

template<int vt, typename op_t = plus_t<type_t> >
MGPU_DEVICE scan_result_t<type_t, vt>
scan(int tid, array_t<type_t, vt> x, storage_t& storage,
type_t carry_in = type_t(), bool use_carry_in = false,
scan(int tid, array_t<type_t, vt> x, storage_t& storage,
type_t carry_in = type_t(), bool use_carry_in = false,
int count = nt, op_t op = op_t(), type_t init = type_t(),
scan_type_t type = scan_type_exc) const {

Expand All @@ -143,14 +139,14 @@ struct cta_scan_t {
} else {
iterate<vt>([&](int i) {
int index = vt * tid + i;
x[i] = i ?
x[i] = i ?
((index < count) ? op(x[i], x[i - 1]) : x[i - 1]) :
(x[i] = (index < count) ? x[i] : init);
});
}

// Scan the thread-local reductions for a carry-in for each thread.
scan_result_t<type_t> result = scan(tid, x[vt - 1], storage,
scan_result_t<type_t> result = scan(tid, x[vt - 1], storage,
div_up(count, vt), op, init, scan_type_exc);

// Perform the scan downsweep and add both the global carry-in and the
Expand Down Expand Up @@ -185,7 +181,7 @@ struct cta_scan_t<nt, bool> {
int warps[num_warps];
};

MGPU_DEVICE scan_result_t<int> scan(int tid, bool x,
MGPU_DEVICE scan_result_t<int> scan(int tid, bool x,
storage_t& storage) const {

// Store the bit totals for each warp.
Expand All @@ -207,7 +203,7 @@ struct cta_scan_t<nt, bool> {
}
__syncthreads();
#else

if(0 == tid) {
// Inclusive scan of partial reductions..
int scan = 0;
Expand All @@ -217,7 +213,7 @@ struct cta_scan_t<nt, bool> {
}
__syncthreads();

#endif
#endif

int scan = ((warp > 0) ? storage.warps[warp - 1] : 0) +
popc(bfe(bits, 0, lane));
Expand Down
Loading