From 94a24a26e9053ba1913a4f3ef459ad9fbf913279 Mon Sep 17 00:00:00 2001
From: trinitas <gydani2@gmail.com>
Date: Wed, 2 Mar 2022 21:18:28 +0100
Subject: [PATCH] Finishing touches
---
Primitives/Primitives.cpp | 10 +++++-----
kernels/programs.cl | 13 ++++++++-----
2 files changed, 13 insertions(+), 10 deletions(-)
diff --git a/Primitives/Primitives.cpp b/Primitives/Primitives.cpp
index 80c38d1..d9b0bcf 100644
--- a/Primitives/Primitives.cpp
+++ b/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 h2(true, 32, 4096);
+ Histogram h1(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;
diff --git a/kernels/programs.cl b/kernels/programs.cl
index 6fd24f2..6f41201 100644
--- a/kernels/programs.cl
+++ b/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)
{