diff --git a/gpu-slo/main b/gpu-slo/main
deleted file mode 100755
index 38eeea5..0000000
Binary files a/gpu-slo/main and /dev/null differ
diff --git a/gpu-slo/main.cu b/gpu-slo/main.cu
index c18c059..57c26ae 100644
--- a/gpu-slo/main.cu
+++ b/gpu-slo/main.cu
@@ -1,6 +1,6 @@
 #include <thrust/device_vector.h>
 #include <thrust/host_vector.h>
-#include <thrust/for_each.h>
+#include <thrust/logical.h>
 
 #include <vector>
 #include <iostream>
@@ -14,8 +14,10 @@ struct Row {
 
     int from, to;
 
+    bool learning;
+
     __host__ __device__
-    Row() : rel(0), l(0), r(0), from(0), to(0) {}
+    Row() : rel(0), l(0), r(0), from(0), to(0), learning(true) {}
 
     __device__
     Row(int rel, int cos, int size) {
@@ -25,11 +27,13 @@ struct Row {
         from = to = cos;
         
         this->rel = rel;
+
+        learning = true;
     }
 };
 
 std::ostream &operator<<(std::ostream &o, const Row &r) {
-    return o << "Row[" << r.rel << "]{" << r.l << ":" << r.from << "-" << r.to << ":" << r.r << "}";
+    return o << "Row[" << r.rel << "]{" << r.l << ":" << r.from << "-" << r.to << ":" << r.r << "}(" << r.learning << ")";
 }
 
 struct Rel {
@@ -52,7 +56,10 @@ struct Solver {
     
     __device__
     void operator()(Row &r) {
-        if (r.r - r.l <= 1) return;
+        if (r.r - r.l <= 1) {
+            r.learning = false;
+            return;
+        }
         
         while (r.r - r.l > 1) {
             int gen = rels[r.rel].gens[r.l & 1];
@@ -74,7 +81,12 @@ struct Solver {
             int gen = rels[r.rel].gens[r.l & 1];
             cosets[r.from * ngens + gen] = r.to;
             cosets[r.to * ngens + gen] = r.from;
+
+            r.learning = true;
+            return;
         }
+
+        r.learning = false;
     }
 };
 
@@ -113,29 +125,49 @@ struct RowIncomplete {
     }
 };
 
+struct Relearn {
+    __device__
+    void operator()(Row &r) {
+        r.learning = true;
+    }
+};
+
+struct Learning {
+    __device__
+    bool operator()(Row r) {
+        return r.learning;
+    }
+};
+
 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(
+bool 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++;
+    while (cosets[*hint] >= 0) {
+        *hint = *hint + 1;
+        if (*hint >= cosets.size()) 
+            return true;
+    }
     int from = *hint / ngens;
     int gen = *hint % ngens;
     
+    add_row(ngens, cosets);
+    
     cosets[*hint] = *coset;
     cosets[*coset * ngens + gen] = from;
+
+    return false;
 }
 
 void gen_rows(
@@ -171,26 +203,60 @@ thrust::device_vector<int> solve(
     int coset = 0;
     int hint = 0;
 
-    // the main loop should go here.
+    while (true) {
+        thrust::for_each(
+                thrust::device, 
+                rows.begin(), 
+                rows.end(),
+                Relearn());
 
-    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);
-    }
+        while (true) {
+            thrust::for_each(
+                    thrust::device,
+                    rows.begin(), rows.end(),
+                    solve);
+            bool r = thrust::any_of(
+                    thrust::device,
+                    rows.begin(), rows.end(),
+                    Learning());
+            if (!r) break;
+        }
 
-    auto cut = thrust::partition(
-            thrust::device, 
-            rows.begin(), rows.end(), 
-            RowIncomplete());
-    rows.erase(cut, rows.end());
+        bool done = add_coset(
+                ngens,
+                &coset, &hint,
+                cosets);
+
+        if (done) break;
+
+        gen_rows(coset, rels, rows);
+
+        auto cut = thrust::partition(
+                thrust::device, 
+                rows.begin(), rows.end(), 
+                RowIncomplete());
+
+        rows.erase(cut, rows.end());
+    }
 
     return cosets;
 }
 
 
 int main(int argc, char* argv[]) {
+    // int ngens = 4;
+    // std::vector<Rel> rels = {
+    //     {0, 1, 3},
+    //     {1, 2, 3},
+    //     {2, 3, 3},
+
+    //     {0, 2, 2},
+    //     {1, 2, 2},
+    //     {1, 3, 2},
+    // };
+    // std::vector<int> subs = {};
+
     int ngens = 4;
     std::vector<Rel> rels = {
         {0, 1, 4},
@@ -198,14 +264,19 @@ int main(int argc, char* argv[]) {
         {2, 3, 3},
 
         {0, 2, 2},
-        {1, 2, 2},
         {1, 3, 2},
+        {0, 3, 2},
     };
-    std::vector<int> subs = {1, 3};
+    std::vector<int> subs = {};
 
     thrust::host_vector<int> cosets = solve(ngens, subs, rels);
 
-    std::cout << cosets << std::endl;
+    std::cout << cosets.size() / ngens << " cosets" << std::endl;
+    for (int c = 0; c < cosets.size(); c += ngens) {
+        for (int g = c; g < c + ngens; g++ ) {
+            std::cout << cosets[g] << " ";
+        } std::cout << std::endl;
+    }
 
     return 0;
 }