diff --git a/gpu-slo/main.cu b/gpu-slo/main.cu
index c6fdb48..77ff72e 100644
--- a/gpu-slo/main.cu
+++ b/gpu-slo/main.cu
@@ -7,23 +7,6 @@
 
 #include "util.h"
 
-// struct Cosets {
-//     int width;
-//     thrust::device_vector<int> data{};
-//     
-//     __host__
-//     Cosets(int ngens) : width(ngens) {
-//     }
-//     
-//     void add_row() {
-//         data.resize(data.size() + width, -1);
-//     }
-//     
-//     thrust::host_vector<int> get_data() {
-//         return data;
-//     }
-// };
-
 struct Row {
     int rel;
 
@@ -123,20 +106,83 @@ struct RowGen {
     }
 };
 
+void add_row(
+        int ngens,
+        thrust::device_vector<int> &cosets) {
+    cosets.resize(cosets.size() + ngens, -1);
+}
+
+// todo: this part is _real_ slow.
+void add_coset(
+        int ngens,
+        int *coset,
+        int *hint,
+        thrust::device_vector<int> &cosets) {
+    *coset = cosets.size() / ngens;
+
+    add_row(ngens, cosets);
+    
+    // todo: this part especially.
+    while (cosets[*hint] >= 0)  *hint++;
+    int from = *hint / ngens;
+    int gen = *hint % ngens;
+    
+    cosets[*hint] = *coset;
+    cosets[*coset * ngens + gen] = from;
+}
+
+void gen_rows(
+        int coset,
+        thrust::device_vector<Rel> &rels,
+        thrust::device_vector<Row> &rows) {
+    rows.resize(rows.size() + rels.size());
+
+    thrust::counting_iterator<int> counter(0);
+    thrust::transform(
+            thrust::device,
+            counter, counter + rels.size(),
+            rows.end() - rels.size(),
+            RowGen(coset, rels));
+}
+
 thrust::device_vector<int> solve(
         int ngens,
         thrust::device_vector<int> subs,
         thrust::device_vector<Rel> rels) {
     
     thrust::device_vector<int> cosets;
-    cosets.resize(cosets.size() + ngens, -1);
-    int lastCoset = 0;
-
-    thrust::for_each(subs.begin(), subs.end(), 
-        CosetInitializer(cosets));
-
     thrust::device_vector<Row> rows;
 
+    add_row(ngens, cosets);
+    thrust::for_each(
+            thrust::device,
+            subs.begin(), subs.end(), 
+            CosetInitializer(cosets));
+
+    gen_rows(0, rels, rows);
+
+    int coset = 0;
+    int hint = 0;
+
+    // the main loop should go here.
+
+    std::cout << thrust::host_vector<Row>(rows) << std::endl;
+    
+    for (int i = 0; i < 4; i++) {
+        Solver solve(ngens, cosets, rels);
+        thrust::for_each(rows.begin(), rows.end(), solve);
+    }
+
+    std::cout << thrust::host_vector<Row>(rows) << std::endl;
+
+    /*
+
+    add_coset(ngens, &coset, &hint, cosets);
+
+    std::cout << coset << " " << hint << std::endl;
+
+    std::cout << rows << std::endl;
+
     thrust::counting_iterator<int> counter(0);
 
     thrust::device_vector<Row> new_rows(rels.size());
@@ -151,6 +197,7 @@ thrust::device_vector<int> solve(
     std::cout << thrust::host_vector<Row>(rows) << std::endl;
     thrust::for_each(rows.begin(), rows.end(), solv);
     std::cout << thrust::host_vector<Row>(rows) << std::endl;
+    */
 
     return cosets;
 }