     ASSERT_EQ(err, GDF_SUCCESS);
 }
 
-
-TEST(gdf_foo_sample_TEST, case1) {
-    standard_context_t context;
-    cudaEvent_t start, stop;
-    cudaEventCreate(&start);
-    cudaEventCreate(&stop);
-    float milliseconds=0,sortms=0,hashms = 0;
-    for (int countSize=1e3; countSize <=1e8; countSize*=10){ 
-        int countA=countSize;
-        int countB=countSize;
-        for(int maxkey=1e4; maxkey<=1e8; maxkey*=10){
-            mem_t<int> dataA = fill_random(0, maxkey, countA, false, context);
-            mem_t<int> dataB = fill_random(0, maxkey, countB, false, context);
-            cudaEventRecord(start); 
-            mergesort(dataA.data(), countA, less_t<int>(), context);
-            mergesort(dataB.data(), countB, less_t<int>(), context);
-            cudaEventRecord(stop);
-            cudaEventSynchronize(stop);
-            cudaEventElapsedTime(&sortms, start, stop);
-            // printf("Sorting, %10d, %10d, %8.5f, %10.0f\n", 
-            printf("Sorting,%d,%d,%f,%f\n", 
-                countA, maxkey, sortms, float(1000*countA)/(float)sortms);
-            mem_t<int> common;
-            cudaEventRecord(start);
-            //common = inner_join(dataA.data(), countA, dataB.data(), countB, less_t<int>() , context);
-            cudaEventRecord(stop);
-            cudaEventSynchronize(stop);
-            cudaEventElapsedTime(&milliseconds, start, stop);
-            // printf("time: %f  - common elements %d\n", milliseconds,common.size());
-            mem_t<int> dataA1 = fill_random(0, maxkey, countA, false, context);
-            mem_t<int> dataA2 = fill_random(0, maxkey, countA, false, context);
-            mem_t<int> dataA3 = fill_random(0, maxkey, countA, false, context);
-            mem_t<int> dataB1 = fill_random(0, maxkey, countB, false, context);
-            mem_t<int> dataB2 = fill_random(0, maxkey, countB, false, context);
-            mem_t<int> dataB3 = fill_random(0, maxkey, countB, false, context);
-            gdf_column gdl0 = create_gdf_column(dataA1);
-            gdf_column gdl1 = create_gdf_column(dataA2);
-            gdf_column gdl2 = create_gdf_column(dataA3);
-            gdf_column gdr0 = create_gdf_column(dataB1);
-            gdf_column gdr1 = create_gdf_column(dataB2);
-            gdf_column gdr2 = create_gdf_column(dataB3);
-            gdf_column* gl[3] = {&gdl0, &gdl1, &gdl2};
-            gdf_column* gr[3] = {&gdr0, &gdr1, &gdr2};
-            gdf_join_result_type *out;
-            // gdf_error err = gdf_multi_left_join_generic(1, gl, gr, &out);
-            cudaEventRecord(start); 
-            gdf_error err = gdf_multi_left_join_generic(1, gl, gr, &out);
-            cudaEventRecord(stop);
-            cudaEventSynchronize(stop);
-            cudaEventElapsedTime(&hashms, start, stop);
-            // printf("Hashing, %10d, %10d, %8.5f, %10.0f\n", 
-            printf("Hashing,%d,%d,%f,%f\n", 
-                countA,  maxkey, hashms, float(1000*countA)/(float)hashms);
-                          
-            // gdf_error err = gdf_inner_join_i32(&gdl0, &gdr0 , &out);
-        }
-    }
-}
-
+TEST(join_TEST, gdf_inner_join) {
+    std::vector<int> l{0, 0, 1, 2, 3};
+    std::vector<int> r{0, 1, 2, 2, 3};
+    thrust::device_vector<int> dl, dr, l_pos, r_pos;
+
+    auto err = call_gdf_single_column_test(l, r, dl, dr, l_pos, r_pos, gdf_inner_join_generic);
+
+    thrust::device_vector<int> l_idx(l_pos.size());
+    thrust::device_vector<int> r_idx(r_pos.size());
+
+    thrust::gather(l_pos.begin(), l_pos.end(), dl.begin(), l_idx.begin());
+    thrust::gather(r_pos.begin(), r_pos.end(), dr.begin(), r_idx.begin());
+
+    EXPECT_THAT(host_vec(l_idx), ElementsAreArray(host_vec(r_idx)));
+    EXPECT_THAT(host_vec(l_pos), ElementsAre(0, 1, 2, 3, 3, 4));
+    EXPECT_THAT(host_vec(r_pos), ElementsAre(0, 0, 1, 2, 3, 4));
+
+    ASSERT_EQ(err, GDF_SUCCESS);
+}
+
+TEST(join_TEST, gdf_left_join) {
+    std::vector<int> l{0, 0, 4, 5, 5};
+    std::vector<int> r{0, 0, 2, 3, 5};
+    thrust::device_vector<int> dl, dr, l_pos, r_pos;
+
+    auto err = call_gdf_single_column_test(l, r, dl, dr, l_pos, r_pos, gdf_left_join_generic);
+
+    thrust::device_vector<int> l_idx(l_pos.size());
+    thrust::device_vector<int> r_idx(r_pos.size());
+
+    thrust::gather(l_pos.begin(), l_pos.end(), dl.begin(), l_idx.begin());
+    thrust::gather(r_pos.begin(), r_pos.end(), dr.begin(), r_idx.begin());
+
+    EXPECT_THAT(host_vec(l_idx), ElementsAre(0, 0, 0, 0, 4, 5, 5));
+    EXPECT_THAT(host_vec(l_pos), ElementsAre(0, 0, 1, 1, 2, 3, 4));
+    EXPECT_THAT(host_vec(r_pos), ElementsAre(0, 1, 0, 1, -1, 4, 4));
+
+    ASSERT_EQ(err, GDF_SUCCESS);
+}

    std::vector<int> data(dev_vec.size());
    thrust::copy(dev_vec.begin(), dev_vec.end(), data.begin());
    return data;
}

gdf_error

#include <cstdlib>
#include <iostream>
#include <vector>
#include <functional>

#include <thrust/device_vector.h>
#include <thrust/sort.h>
