Annotate all examples, function to generate just JSON, small improvments

For every example there is now an annotated version (which is also in the makefile) and the original version of that code. (It is mostly sufficient to change about 12 lines of code).

The cuda_mav class has been improved, by being able to directly produce JSON files and strings as output.
The currentSize and originalSize have been added to the output, so you know if it was not possible to store all data which was actually accessed. Instead of a tuple, a dedicated struct is now used for memory regions. 

Templates have been moved to their own dedicated directory.

To reduce git load times a correct gitignore was already added to the html folder. This will be solved later by the merge into master anyways, but eases local development.

Author: NetroScript

examples/basic/CMakeLists.txt

@@ -13,16 +13,6 @@ set_target_properties(${_TARGET_NORMAL} PROPERTIES
 # add include path to the cuda_mav.cuh header
 target_link_libraries(${_TARGET_NORMAL} PRIVATE CUDA_MAV_HEADER)
 
-set(_TARGET_ANNOTATED annotated)
-
-add_executable(${_TARGET_ANNOTATED})
-target_sources(${_TARGET_ANNOTATED}
-   PRIVATE
-   annotated.cu)
-set_target_properties(${_TARGET_ANNOTATED} PROPERTIES
-  CUDA_CXX_STANDARD 17
-)
-
 set(_TARGET_RANDOM_2D_MATRIX random2dMatrix)
 
 add_executable(${_TARGET_RANDOM_2D_MATRIX})

examples/basic/normal.cu

@@ -49,7 +49,8 @@ int main(){
     checkCudaError(cudaGetLastError());
     cudaDeviceSynchronize();
 
-    memAccessStorage->generateOutput("../../../html/template.json", "../../../out/basic.json", CudaMemAccessStorage<int>::parseDataForJSPage);
+    memAccessStorage->generateTemplatedOutput("../../../templates/basic_template.html", "../../../out/basic_html.html",
+                                              CudaMemAccessStorage<int>::parseDataForStaticHTML);
 
     checkCudaError(cudaMemcpy(h_output.data(), d_output, sizeof(int)*prob_size, cudaMemcpyDeviceToHost));
 

examples/basic/normal_original.cu

@@ -0,0 +1,55 @@
+#include <vector>
+#include <numeric>
+#include <iostream>
+
+inline void checkCudaError(cudaError_t err) {
+    if (err != cudaSuccess) {
+        std::cerr << "\rCuda Error " << err << ": " << cudaGetErrorString(err) << std::endl;
+        std::cerr << "Aborting..." << std::endl;
+        exit(1);
+    }
+}
+
+
+__global__ void kernel(int prob_size, int * const input, int * output){
+    int id = threadIdx.x + blockIdx.x * blockDim.x;
+    if(id < prob_size){
+        output[id] = input[id];
+    }
+}
+
+int main(){
+    constexpr int prob_size = 100;
+
+    std::vector<int> h_input(prob_size);
+    std::iota(h_input.begin(), h_input.end(), 0);
+    int * d_input = nullptr;
+    checkCudaError(cudaMalloc((void**) &d_input, sizeof(int)*prob_size));
+
+    std::vector<int> h_output(prob_size, 0);
+    int * d_output = nullptr;
+    checkCudaError(cudaMalloc((void**) &d_output, sizeof(int)*prob_size));
+
+    checkCudaError(cudaMemcpy(d_input, h_input.data(), sizeof(int)* prob_size, cudaMemcpyHostToDevice));
+
+    constexpr int threads = 32;
+    constexpr int blocks = (prob_size/threads)+1;
+
+    kernel<<<blocks, threads>>>(prob_size, d_input, d_output);
+    checkCudaError(cudaGetLastError());
+
+    checkCudaError(cudaMemcpy(h_output.data(), d_output, sizeof(int)*prob_size, cudaMemcpyDeviceToHost));
+
+    for(auto i = 0; i < h_input.size(); ++i){
+        if(h_input[i] != h_output[i]){
+            std::cerr << "Element at position " << i << "is not equal (input - output): " << h_input[i] << " != " << h_output[i] << std::endl;
+            std::exit(1);
+        }
+    }
+
+    checkCudaError(cudaFree(d_input));
+    checkCudaError(cudaFree(d_output));
+
+    std::cout << "kernel finished successful" << std::endl;
+    return 0;
+}

examples/basic/random2dMatrix.cu

@@ -5,6 +5,7 @@
 #include <vector>
 #include <random>
 #include <algorithm>
+#include "../../include/cuda_mav.cuh"
 
 // The wrapper macro is required, that __LINE__ is correct pointing to the line, where the check fails
 #define checkCudaError(ans)                          \
@@ -22,17 +23,17 @@ inline void checkCudaErrorFunc(cudaError_t err, const char *file, int line)
    }
 }
 
-__global__ void decrement(unsigned int const size, unsigned int *data, unsigned int *control)
+__global__ void decrement(unsigned int const size, CudaMemAccessLogger<unsigned int> *data, CudaMemAccessLogger<unsigned int> *control)
 {
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;
 
    for (int i = index; i < size; i += stride)
    {
-      while (data[i] > 0)
+      while ((*data)[i] > 0)
       {
-         data[i] = data[i] - 1;
-         control[i] = control[i] + 1;
+         (*data)[i] = (*data)[i] - 1;
+         (*control)[i] = (*control)[i] + 1;
       }
    }
 }
@@ -82,10 +83,10 @@ int main(int argc, char **argv)
        { return distribution(generator); });
 
    // enable me, to create a hot spot area
-   // the hot spot area should looks interessting in the memory access visualization
-   if (false)
+   // the hot spot area should looks interesting in the memory access visualization
+   if (true)
    {
-      hot_spot(h_data, dim, 8, 10, 10, 3);
+      hot_spot(h_data, dim, 8, 10, 10, 20);
    }
 
    // enable me, to print the matrix
@@ -120,15 +121,27 @@ int main(int argc, char **argv)
    // copy h_controll to initialize all values with 0 on the GPU
    checkCudaError(cudaMemcpy(d_control, h_control.data(), buffer_size_byte, cudaMemcpyHostToDevice));
 
+   // Define amount of accesses you want to log and create a memory object which stores them
+   auto* memAccessStorage = new CudaMemAccessStorage<unsigned int>(dim * dim * 50);
+
+    // Wrap the data classes with the custom logging class
+   auto* data = new CudaMemAccessLogger<unsigned int>(d_data, dim*dim, "Decremented Data", memAccessStorage);
+   auto* control = new CudaMemAccessLogger<unsigned int>(d_control, dim*dim, "Control Data", memAccessStorage);
+
    // change me and look, how the visulization looks like
    int const blockSize = 32;
    int const numBlocks = ((dim * dim) + blockSize - 1) / blockSize;
 
-   decrement<<<numBlocks, blockSize>>>(dim * dim, d_data, d_control);
+   decrement<<<numBlocks, blockSize>>>(dim * dim, data, control);
+
    checkCudaError(cudaGetLastError());
 
    checkCudaError(cudaMemcpy(h_control.data(), d_control, buffer_size_byte, cudaMemcpyDeviceToHost));
 
+   // Generate the output json file for visualization
+   // It needs to be below the memory copy or alternatively after a synchronize to make sure the kernel is finished
+   memAccessStorage->generateJSONOutput("../../../out/random2DMatrix.json");
+
    bool success = true;
 
    for (auto y = 0; y < dim; ++y)
@@ -167,5 +180,10 @@ int main(int argc, char **argv)
    checkCudaError(cudaFree(d_data));
    checkCudaError(cudaFree(d_control));
 
+   // Free the memory of the custom logging classes
+   delete memAccessStorage;
+   delete data;
+   delete control;
+
    return 0;
 }

examples/basic/random2dMatrix_original.cu

@@ -0,0 +1,171 @@
+// The applicaiton creates a 2D matrix and initialize each element randomly with a value between 0 and 10.
+// The kernel is simply decrementing each element until 0 in a very ineffective way.
+
+#include <iostream>
+#include <vector>
+#include <random>
+#include <algorithm>
+
+// The wrapper macro is required, that __LINE__ is correct pointing to the line, where the check fails
+#define checkCudaError(ans)                          \
+   {                                                 \
+      checkCudaErrorFunc((ans), __FILE__, __LINE__); \
+   }
+
+inline void checkCudaErrorFunc(cudaError_t err, const char *file, int line)
+{
+   if (err != cudaSuccess)
+   {
+      std::cout << "\r" << file << ":" << line << " -> Cuda Error " << err << ": " << cudaGetErrorString(err) << std::endl;
+      std::cout << "Aborting..." << std::endl;
+      exit(0);
+   }
+}
+
+__global__ void decrement(unsigned int const size, unsigned int *data, unsigned int *control)
+{
+   int index = threadIdx.x + blockIdx.x * blockDim.x;
+   int stride = blockDim.x * gridDim.x;
+
+   for (int i = index; i < size; i += stride)
+   {
+      while (data[i] > 0)
+      {
+         data[i] = data[i] - 1;
+         control[i] = control[i] + 1;
+      }
+   }
+}
+
+/// @brief Increment all values in a specific area by the value of increment. The maximum value of an entry is clamp to 10.
+/// @param data Data to increment.
+/// @param dim Dimensions of the 2D matrix.
+/// @param y_start Y start coordinate of the area to increment.
+/// @param x_start X start coordinate of the area to increment.
+/// @param size Size of the Y and X direction of the area to increment. 
+/// @param increment Value to increment.
+void hot_spot(std::vector<unsigned int> &data, unsigned int const dim, unsigned int const y_start, unsigned int const x_start, unsigned int const size, unsigned int const increment)
+{
+   for (auto y = y_start; y < y_start + size; ++y)
+   {
+      for (auto x = x_start; x < x_start + size; ++x)
+      {
+         if (data[y * dim + x] + increment > 10)
+         {
+            data[y * dim + x] = 10;
+         }
+         else
+         {
+            data[y * dim + x] += increment;
+         }
+      }
+   }
+}
+
+int main(int argc, char **argv)
+{
+   unsigned int dim = 100;
+
+   std::vector<unsigned int> h_data(dim * dim);
+   // create a 2D matrix where all elements are 0
+   std::vector<unsigned int> h_control(dim * dim, 0);
+
+   // initialize data matrix with random numbers betweem 0 and 10
+   std::uniform_int_distribution<unsigned int> distribution(
+       0,
+       10);
+   std::default_random_engine generator;
+   std::generate(
+       h_data.begin(),
+       h_data.end(),
+       [&distribution, &generator]()
+       { return distribution(generator); });
+
+   // enable me, to create a hot spot area
+   // the hot spot area should looks interessting in the memory access visualization
+   if (false)
+   {
+      hot_spot(h_data, dim, 8, 10, 10, 3);
+   }
+
+   // enable me, to print the matrix
+   if (false)
+   {
+      for (auto y = 0; y < dim; ++y)
+      {
+         for (auto x = 0; x < dim; ++x)
+         {
+            if (h_data[y * dim + x] < 10)
+            {
+               std::cout << " " << h_data[y * dim + x] << " ";
+            }
+            else
+            {
+               std::cout << h_data[y * dim + x] << " ";
+            }
+         }
+         std::cout << std::endl;
+      }
+   }
+
+   unsigned int *d_data = nullptr;
+   unsigned int *d_control = nullptr;
+
+   size_t const buffer_size_byte = sizeof(unsigned int) * dim * dim;
+
+   checkCudaError(cudaMalloc((void **)&d_data, buffer_size_byte));
+   checkCudaError(cudaMalloc((void **)&d_control, buffer_size_byte));
+
+   checkCudaError(cudaMemcpy(d_data, h_data.data(), buffer_size_byte, cudaMemcpyHostToDevice));
+   // copy h_controll to initialize all values with 0 on the GPU
+   checkCudaError(cudaMemcpy(d_control, h_control.data(), buffer_size_byte, cudaMemcpyHostToDevice));
+
+   // change me and look, how the visulization looks like
+   int const blockSize = 32;
+   int const numBlocks = ((dim * dim) + blockSize - 1) / blockSize;
+
+   decrement<<<numBlocks, blockSize>>>(dim * dim, d_data, d_control);
+   checkCudaError(cudaGetLastError());
+
+   checkCudaError(cudaMemcpy(h_control.data(), d_control, buffer_size_byte, cudaMemcpyDeviceToHost));
+
+   bool success = true;
+
+   for (auto y = 0; y < dim; ++y)
+   {
+      for (auto x = 0; x < dim; ++x)
+      {
+         if (h_control[y * dim + x] != h_data[y * dim + x])
+         {
+            std::cout << "h_control[" << y << ", " << x << "] != h_data[" << y << ", " << x << "]" << std::endl;
+            std::cout << h_control[y * dim + x] << " != " << h_data[y * dim + x] << std::endl;
+            success = false;
+         }
+      }
+   }
+
+   checkCudaError(cudaMemcpy(h_data.data(), d_data, buffer_size_byte, cudaMemcpyDeviceToHost));
+
+   for (auto y = 0; y < dim; ++y)
+   {
+      for (auto x = 0; x < dim; ++x)
+      {
+         if (h_data[y * dim + x] != 0)
+         {
+            std::cout << "h_data[" << y << ", " << x << "] != 0" << std::endl;
+            std::cout << "value is: " << h_data[y * dim + x] << std::endl;
+            success = false;
+         }
+      }
+   }
+
+   if (success)
+   {
+      std::cout << "The kernel worked correctly" << std::endl;
+   }
+
+   checkCudaError(cudaFree(d_data));
+   checkCudaError(cudaFree(d_control));
+
+   return 0;
+}