Finishing touches

Primitives/Primitives.cpp

@@ -59,14 +59,14 @@ void add_tests(std::vector<TestCase*>* tests) {
 
 int main()
 {
-	capi();
+	//capi();
 
 	OpenCLHandler handler;
 	std::vector<TestCase*> tests;
 
 	Square s;
-	Histogram h1(false, 32, 4096);
+	Histogram h1(true, 32, 4096);
-	Histogram h2(true, 32, 4096);
+	Histogram h2(false, 32, 4096);
 	ReduceAdd r1(handler.get_max_size());
 	Compact c(1024);
 	ExclusiveScan e(handler.get_max_size());
@@ -75,9 +75,9 @@ int main()
 	tests.push_back(&h1);
 	tests.push_back(&h2);
 	tests.push_back(&r1);
-	tests.push_back(&c);
 	tests.push_back(&e);
-
+	tests.push_back(&c);
+	
 
 	for (size_t i = 0; i < tests.size(); i++) {
 		std::cout << tests[i]->description() << std::endl;

kernels/programs.cl

@@ -7,7 +7,7 @@ __kernel void square(__global float* inputData,
 }
 
 // TODO
-//
+// id  := get_global_id(0)
 // histogram[data[id]] := histogram[data[id]] + 1
 //
 // SYNCHRONIZATION!
@@ -27,7 +27,7 @@ void histogram_global(__global int* data, __global int* histogram)
 //  lhistogram[LID] := 0
 // BARRIER
 //
-// Add data to local histogram
+// lhistogram[data[id]] := lhistogram[data[id]] + 1
 //
 // BARRIER
 // 
@@ -47,7 +47,7 @@ void histogram_local(__global int* data, __global int* histogram, __local int* l
 
 	atomic_add(&lhistogram[data[id]], 1.0f);
 
-	barrier(CLK_GLOBAL_MEM_FENCE); // LOCAL???
+	barrier(CLK_LOCAL_MEM_FENCE); // LOCAL???
 
 	if (lid < histogramSize) {
 		atomic_add(&histogram[lid], lhistogram[lid]);
@@ -69,7 +69,7 @@ void reduce_global(__global float* data)
 {
 	int id = get_global_id(0);
 
-	// one right shift is = divide number by two
+	// one right shift = divide number by two
 	for (size_t s = get_global_size(0) / 2; s > 0; s >>= 1) {
 		if (id < s) {
 			data[id] = data[id] +  data[id + s];
@@ -175,7 +175,10 @@ void compact_exscan(__global int* pred, __global int* prefSum)
 // VALUE := data[ID]
 // BARRIER
 // IF pred[ID] == 1 THEN
-//  data[prefSum[ID]] = VALUE
+//   IF ID == 0 THEN
+//     data[0] = VALUE
+//   ELSE
+//     data[prefSum[ID]] = VALUE
 __kernel
 void compact_compact(__global int* data, __global int* pred, __global int* prefSum)
 {