Project 2: Rony Edde

stream_compaction/CMakeLists.txt

@@ -9,6 +9,8 @@ set(SOURCE_FILES
     "efficient.cu"
     "thrust.h"
     "thrust.cu"
+    "radix.h"
+    "radix.cu"
     )
 
 cuda_add_library(stream_compaction

stream_compaction/common.cu

@@ -24,6 +24,12 @@ namespace Common {
  */
 __global__ void kernMapToBoolean(int n, int *bools, const int *idata) {
     // TODO
+    int index = (blockIdx.x * blockDim.x) + threadIdx.x;
+    if (index >= n)
+        return;
+
+    bools[index] = idata[index] != 0;
+
 }
 
 /**
@@ -33,6 +39,12 @@ __global__ void kernMapToBoolean(int n, int *bools, const int *idata) {
 __global__ void kernScatter(int n, int *odata,
         const int *idata, const int *bools, const int *indices) {
     // TODO
+    int index = (blockIdx.x * blockDim.x) + threadIdx.x;
+    if (index >= n)
+        return;
+
+    if (bools[index] == 1)
+        odata[indices[index]] = idata[index];
 }
 
 }

stream_compaction/cpu.cu

@@ -9,7 +9,12 @@ namespace CPU {
  */
 void scan(int n, int *odata, const int *idata) {
     // TODO
-    printf("TODO\n");
+    //printf("TODO\n");
+    odata[0] = 0;
+    for (int i = 1; i < n; i++)
+    {
+        odata[i] = odata[i-1] + idata[i-1];
+    }
 }
 
 /**
@@ -19,7 +24,18 @@ void scan(int n, int *odata, const int *idata) {
  */
 int compactWithoutScan(int n, int *odata, const int *idata) {
     // TODO
-    return -1;
+    int count = 0;
+    odata[0] = 0;
+
+    for (int i = 0; i < n; i++)
+    {
+        if (idata[i] != 0)
+        {
+            odata[count] = idata[i];
+            count++;
+        }
+    }
+    return count;
 }
 
 /**
@@ -29,7 +45,52 @@ int compactWithoutScan(int n, int *odata, const int *idata) {
  */
 int compactWithScan(int n, int *odata, const int *idata) {
     // TODO
-    return -1;
+    //return -1;
+
+    int* mapdata = new int[n];
+    int* scandata = new int[n];
+
+    memset(mapdata, 0, n*sizeof(int));
+    memset(scandata, 0, n*sizeof(int));
+
+
+    for (int i = 0; i < n; i++)
+    {
+        if (idata[i] != 0)
+            mapdata[i] = 1;
+    }
+
+    scan(n, scandata, mapdata);
+
+    int count = 0;
+    for (int i = 0; i < n; i++)
+    {
+        //odata[count] = mapdata[i] * idata[scandata[i]];
+        //count += mapdata[i];
+        if (mapdata[i] != 0)
+        {
+            odata[scandata[i]] = idata[i];
+        }
+    }
+    count = scandata[n - 1] + mapdata[n-1];
+
+    /*
+    printf("\n%-10s", "idata: "); for (int i = 0; i < 20; i++)
+       printf("%d ", idata[i]);
+    printf("\n%-10s", "mapdata: "); for (int i = 0; i < 20; i++)
+        printf("%d ", mapdata[i]);
+    printf("\n%-10s", "scandata: "); for (int i = 0; i < 20; i++)
+        printf("%d ", scandata[i]);
+    printf("\n%-10s", "odata: "); for (int i = 0; i < 20; i++)
+        printf("%d ", odata[i]);
+    printf("\n");
+    */
+
+    delete[] mapdata;
+    delete[] scandata;
+
+    return count;
+
 }
 
 }