HW3: RC1 - hpc slurm creation and submittion scripts added

homework_3/Makefile

@@ -22,15 +22,15 @@
 PROJECT         := PDS_homework_3
 
 # Excecutable's name
-TARGET          := bitonic
+TARGET          := bitonicCUDA
 
 # Source directories list(space seperated). Makefile-relative path, UNDER current directory.
-SRC_DIR_LIST    := src test test/gtest
+SRC_DIR_LIST    := src #test test/gtest
 
 # Include directories list(space seperated). Makefile-relative path.
-INC_DIR_LIST    := src \
-                   test \
-                   test/gtest/ \
+INC_DIR_LIST    := src
+#                   test \
+#                   test/gtest/ \
 
 
 # Exclude files list(space seperated). Filenames only.
@@ -45,26 +45,26 @@ OUTPUT_DIR      := out
 
 # ========== Compiler settings ==========
 # Compiler flags for debug and release
-DEB_CFLAGS      := -DDEBUG -g3 -Wall -Wextra -std=c11 -fopenmp
-REL_CFLAGS      := -Wall -Wextra -O3 -std=c11 -fopenmp
-DEB_CXXFLAGS    := -DDEBUG -g3 -Wall -Wextra -std=c++17 -fopenmp
-REL_CXXFLAGS    := -Wall -Wextra -O3 -std=c++17 -fopenmp
+DEB_CFLAGS      := -DDEBUG -std=c11 -Xcompiler "-Wall -Wextra -g -DDEBUG"
+REL_CFLAGS      := -O3 -std=c11 -Xcompiler "-Wall -Wextra"
+DEB_CXXFLAGS    := -DDEBUG -std=c++17 -Xcompiler "-Wall -Wextra -g -DDEBUG" 
+REL_CXXFLAGS    := -O3 -std=c++17 -Xcompiler "-Wall -Wextra"
 
 # Pre-defines
 # PRE_DEFS := MYCAB=1729 SUPER_MODE
-PRE_DEFS        :=
+PRE_DEFS        := TARGET=$(TARGET)
 
 # ============== Linker settings ==============
 # Linker flags (example: -pthread -lm)
-LDFLAGS         := -pthread
+LDFLAGS         :=
 
 # Map output file
-MAP_FILE        := output.map
-MAP_FLAG        := -Xlinker -Map=$(BUILD_DIR)/$(MAP_FILE)
+MAP_FILE        := # output.map
+MAP_FLAG        := # -Xlinker -Map=$(BUILD_DIR)/$(MAP_FILE)
 
 # ============== Docker settings ==============
 # We need:
-#  - Bind the entire project directory(the dir that icludes all the code) as volume.
+#  - Bind the entire project directory(the dir that includes all the code) as volume.
 #  - In docker instance, change to working directory(where the makefile is).
 DOCKER_VOL_DIR  := $(shell pwd)
 DOCKER_WRK_DIR  :=
@@ -83,16 +83,14 @@ DOCKER          :=
 CSIZE           := size
 CFLAGS          := $(DEB_CFLAGS)
 CXXFLAGS        := $(DEB_CXXFLAGS)
-CXX             := g++ #mpic++
-CC              := gcc #mpicc
+CXX             := g++
+CC              := gcc
+LINKER          := g++
 
 #
 # =========== Main body and Patterns ===========
 #
 
-#ifeq ($(OS), Windows_NT)
-#	TARGET := $(TARGET).exe
-#endif
 INC     := $(foreach dir,$(INC_DIR_LIST),-I$(dir))
 DEF     := $(foreach def,$(PRE_DEFS),-D$(def))
 EXC     := $(foreach fil,$(EXC_FILE_LIST),                              \
@@ -110,44 +108,22 @@ OBJ     := $(foreach file,$(SRC:%.cpp=%.o),$(OBJ_DIR)/$(file))
 DEP     := $(foreach file,$(SRC:%.cpp=%.d),$(DEP_DIR)/$(file))
 
 
-# Make Dependencies pattern.
-# This little trick enables recompilation only when dependencies change
-# and it does so for changes both in source AND header files ;)
-# 
-# It is based on Tom Tromey's method.
-# 
-# Invoke cpp to create makefile rules with dependencies for each source file
-$(DEP_DIR)/%.d: %.c
-	@mkdir -p $(@D)
-	@$(DOCKER) $(CC) -E $(CFLAGS) $(INC) $(DEF) -MM -MT $(OBJ_DIR)/$(<:.c=.o) -MF $@ $<
-
 # c file objects depent on .c AND dependency files, which have an empty recipe 
-$(OBJ_DIR)/%.o: %.c $(DEP_DIR)/%.d
+$(OBJ_DIR)/%.o: %.c
 	@mkdir -p $(@D)
-	@$(DOCKER) $(CC) -c $(CFLAGS) $(INC) $(DEF) -o $@ $<
+	$(DOCKER) $(CC) -c $(CFLAGS) $(INC) $(DEF) -o $@ $<
 
-$(DEP_DIR)/%.d: %.cpp
+# cpp file objects depend on .cpp AND dependency files, which have an empty recipe
+$(OBJ_DIR)/%.o: %.cpp
 	@mkdir -p $(@D)
-	@$(DOCKER) $(CXX) -E $(CXXFLAGS) $(INC) $(DEF) -MM -MT $(OBJ_DIR)/$(<:.cpp=.o) -MF $@ $<
-
-# cpp file objects depent on .cpp AND dependency files, which have an empty recipe 
-$(OBJ_DIR)/%.o: %.cpp $(DEP_DIR)/%.d
-	@mkdir -p $(@D)
-	@$(DOCKER) $(CXX) -c $(CXXFLAGS) $(INC) $(DEF) -o $@ $<
-
-# empty recipe for dependency files. This prevents make errors
-$(DEP):
-
-# now include all dependencies
-# After all they are makefile dependency rules ;)
-include $(wildcard $(DEP))
+	$(DOCKER) $(CXX) -c $(CXXFLAGS) $(INC) $(DEF) -o $@ $<
 
 # main target rule
 $(BUILD_DIR)/$(TARGET): $(OBJ)
 	@mkdir -p $(@D)
 	@echo Linking to target: $(TARGET)
-	@echo $(DOCKER) $(CXX) '$$(OBJ)' $(LDFLAGS) $(MAP_FLAG) -o $(@D)/$(TARGET)
-	@$(DOCKER) $(CXX) $(OBJ) $(LDFLAGS) $(MAP_FLAG) -o $(@D)/$(TARGET)
+	@echo $(DOCKER) $(LINKER) '$$(OBJ)' $(LDFLAGS) $(MAP_FLAG) -o $(@D)/$(TARGET)
+	@$(DOCKER) $(LINKER) $(OBJ) $(LDFLAGS) $(MAP_FLAG) -o $(@D)/$(TARGET)
 	@echo
 	@echo Print size information
 	@$(CSIZE) $(@D)/$(TARGET)
@@ -178,24 +154,79 @@ release: $(BUILD_DIR)/$(TARGET)
 # ================ Build rules =================
 #
 
+bitonic_v0deb: CC := nvcc -G -g -x cu
+bitonic_v0deb: CXX := nvcc -G -g -x cu
+bitonic_v0deb: LINKER := nvcc
+bitonic_v0deb: CFLAGS := $(DEB_CFLAGS) -DCODE_VERSION=V0
+bitonic_v0deb: CXXFLAGS := $(DEB_CXXFLAGS) -DCODE_VERSION=V0
+bitonic_v0deb: OUTPUT_DIR := $(OUTPUT_DIR)/v0
+bitonic_v0deb: $(BUILD_DIR)/$(TARGET)
+	@mkdir -p $(OUTPUT_DIR)
+	cp $(BUILD_DIR)/$(TARGET) $(OUTPUT_DIR)/$(TARGET)
 
-bitonic_v0: CC := nvcc
-bitonic_v0: CXX := nvcc
+
+bitonic_v1deb: CC := nvcc -G -g -x cu
+bitonic_v1deb: CXX := nvcc -G -g -x cu
+bitonic_v1deb: LINKER := nvcc
+bitonic_v1deb: CFLAGS := $(DEB_CFLAGS) -DCODE_VERSION=V1
+bitonic_v1deb: CXXFLAGS := $(DEB_CXXFLAGS) -DCODE_VERSION=V1
+bitonic_v1deb: OUTPUT_DIR := $(OUTPUT_DIR)/v1
+bitonic_v1deb: $(BUILD_DIR)/$(TARGET)
+	@mkdir -p $(OUTPUT_DIR)
+	cp $(BUILD_DIR)/$(TARGET) $(OUTPUT_DIR)/$(TARGET)
+
+
+bitonic_v2deb: CC := nvcc -G -g -x cu
+bitonic_v2deb: CXX := nvcc -G -g -x cu
+bitonic_v2deb: LINKER := nvcc
+bitonic_v2deb: CFLAGS := $(DEB_CFLAGS) -DCODE_VERSION=V2
+bitonic_v2deb: CXXFLAGS := $(DEB_CXXFLAGS) -DCODE_VERSION=V2
+bitonic_v2deb: OUTPUT_DIR := $(OUTPUT_DIR)/v2
+bitonic_v2deb: $(BUILD_DIR)/$(TARGET)
+	@mkdir -p $(OUTPUT_DIR)
+	cp $(BUILD_DIR)/$(TARGET) $(OUTPUT_DIR)/$(TARGET)
+
+
+
+bitonic_v0: CC := nvcc -x cu
+bitonic_v0: CXX := nvcc -x cu
+bitonic_v0: LINKER := nvcc
 bitonic_v0: CFLAGS := $(REL_CFLAGS) -DCODE_VERSION=V0
 bitonic_v0: CXXFLAGS := $(REL_CXXFLAGS) -DCODE_VERSION=V0
-bitonic_v0: TARGET := bitonic_v0
+bitonic_v0: OUTPUT_DIR := $(OUTPUT_DIR)/v0
 bitonic_v0: $(BUILD_DIR)/$(TARGET)
 	@mkdir -p $(OUTPUT_DIR)
 	cp $(BUILD_DIR)/$(TARGET) $(OUTPUT_DIR)/$(TARGET)
 
 
+bitonic_v1: CC := nvcc -x cu
+bitonic_v1: CXX := nvcc -x cu
+bitonic_v1: LINKER := nvcc
+bitonic_v1: CFLAGS := $(REL_CFLAGS) -DCODE_VERSION=V1
+bitonic_v1: CXXFLAGS := $(REL_CXXFLAGS) -DCODE_VERSION=V1
+bitonic_v1: OUTPUT_DIR := $(OUTPUT_DIR)/v1
+bitonic_v1: $(BUILD_DIR)/$(TARGET)
+	@mkdir -p $(OUTPUT_DIR)
+	cp $(BUILD_DIR)/$(TARGET) $(OUTPUT_DIR)/$(TARGET)
+
+bitonic_v2: CC := nvcc -x cu
+bitonic_v2: CXX := nvcc -x cu
+bitonic_v2: LINKER := nvcc
+bitonic_v2: CFLAGS := $(REL_CFLAGS) -DCODE_VERSION=V2
+bitonic_v2: CXXFLAGS := $(REL_CXXFLAGS) -DCODE_VERSION=V2
+bitonic_v2: OUTPUT_DIR := $(OUTPUT_DIR)/v2
+bitonic_v2: $(BUILD_DIR)/$(TARGET)
+	@mkdir -p $(OUTPUT_DIR)
+	cp $(BUILD_DIR)/$(TARGET) $(OUTPUT_DIR)/$(TARGET)
+
+
 hpc-build:
 	make clean
-	make distbubbletonic
+	make bitonic_v0
 	make clean
-	make distbitonic
+	make bitonic_v1
 	make clean
-	make tests
+	make bitonic_v2
 
 
 all: debug bitonic_v0

homework_3/exersize.md

@@ -0,0 +1,33 @@
+Parallel & Distributed Computer Systems HW3
+
+January, 2025
+
+Write a program that sorts $N$ integers in ascending order, using CUDA.
+
+The program must perform the following tasks:
+
+- The user specifies a positive integers $q$.
+
+- Start a process with an array of $N = 2^q$ random integers is each processes.
+
+- Sort all $N$ elements int ascending order.
+
+- Check the correctness of the final result.
+
+Your implementation should be based on the following steps:
+
+V0. A kernel where each thread only compares and exchanges. This "eliminates" the 1:n innermost loop. Easy to write, but too many function calls and global synchronizations.
+
+V1. Include the k inner loop in the kernel function. How do we handle the synchronization? Fewer calls, fewer global synchronizations. Faster than V0!
+
+V2. Modify the kernel of V1 to work with local memory instead of global.
+
+You must deliver:
+- A report (about $3-4$ pages) that describes your parallel algorithm and implementation.
+
+- Your comments on the speed of your parallel program compared to the serial sort, after trying you program on aristotelis for $q = [20:27]$.
+
+- The source code of your program uploaded online.
+
+Ethics: If you use code found on the web or by an LLM, you should mention your source and the changes you made. You may work in pairs; both partners must submit a single report with both names.
+Deadline: 2 February, $2025$.

homework_3/hpc/makeSlurmScripts.sh

@@ -0,0 +1,26 @@
+#!/usr/bin/env bash
+
+# Parameters
+versions=("v0" "v1" "v2")
+q_values=(20 21 22 23 24 25 26 27 28 29 30)
+
+# Make scripts
+for version in "${versions[@]}"; do
+  for q in "${q_values[@]}"; do
+    filename="Bitnc${version^^}Q${q}.sh"  # Convert v0 -> V0 etc...
+    cat > "$filename" <<EOL
+#! /usr/bin/env bash
+
+#SBATCH --job-name=Bitnc${version^^}Q${q}
+#SBATCH --nodes=1
+#SBATCH --gres=gpu:1
+#SBATCH --time=10:00
+
+module load gcc/9.2.0 cuda/11.1.0
+
+./out/${version}/bitonicCUDA -v --validation --perf 7 -b 512 -q ${q}
+
+EOL
+    echo "Create: $filename"
+  done
+done

homework_3/hpc/submitJobs.sh

@@ -0,0 +1,26 @@
+#!/usr/bin/env bash
+
+# Submission parameters
+QOS="small"
+PARTITION="ampere"
+SCRIPT_DIR="hpc"  # Directory containing the job scripts
+
+# Range of values for the -q parameter
+VERSIONS=("V0" "V1" "V2")
+Q_START=20
+Q_END=30
+
+# Submitting the jobs
+for version in "${VERSIONS[@]}"; do
+  for ((q = Q_START; q <= Q_END; q++)); do
+    script_name="Bitnc${version}Q${q}.sh"
+    script_path="${SCRIPT_DIR}/${script_name}"
+
+    if [[ -f "$script_path" ]]; then
+      sbatch --qos="$QOS" -p "$PARTITION" "$script_path"
+      echo "Submitted: $script_path"
+    else
+      echo "Warning: File not found - $script_path"
+    fi
+  done
+done

homework_3/src/bitonicsort.hpp

@@ -0,0 +1,456 @@
+/*!
+ * \file
+ * \brief   Bitonic sort CUDA implementation header
+ *
+ * \author
+ *    Christos Choutouridis AEM:8997
+ *    <cchoutou@ece.auth.gr>
+ */
+
+#ifndef BITONICSORTCUDA_H_
+#define BITONICSORTCUDA_H_
+
+#include <cuda_runtime.h>
+#include <vector>
+#include <cmath>
+#include <cstdint>
+#include <utility>
+
+#include "utils.hpp"
+
+/*
+ * Exported timers
+ */
+extern Timing Timer_total, Timer_memory, Timer_sorting;
+
+using threadId_t = size_t;
+
+
+/*
+ * ============================== Sort utilities ==============================
+ */
+
+/*!
+ * Returns the ascending or descending configuration (up/down phase) of the thread id
+ * depending on the current depth
+ *
+ * @param tid   [threadId_t] The current thread
+ * @param stage [size_t]     The current stage of the sorting network (same for each step)
+ * @return      [bool]       True if we need ascending configuration, false otherwise
+ */
+__device__ inline bool ascending(threadId_t tid, size_t stage) noexcept {
+    return !(tid & (1 << stage));
+}
+
+/*!
+ * Returns the thread's partner for data exchange during the sorting network iterations
+ * of Bitonic
+ *
+ * @param tid   [threadId_t] The current node
+ * @param step  [size_t]     The step of the sorting network
+ * @return      [threadId_t] The node id of the partner for data exchange
+ */
+__device__ inline threadId_t partner(threadId_t tid, size_t step) noexcept {
+    return (tid ^ (1 << step));
+}
+
+
+/*!
+ * Predicate to check if a node keeps the small numbers during the bitonic sort network exchange.
+ *
+ * @param tid       [threadId_t] The node for which we check
+ * @param partner   [threadId_t] The partner of the data exchange
+ * @param stage     [size_t]     The current stage of the sorting network (same for each step)
+ * @return          [bool]       True if the node should keep the small values, false otherwise
+ */
+
+__device__ inline bool keepSmall(threadId_t tid, threadId_t partner, size_t stage) {
+    return ascending(tid, stage) == (tid < partner);
+}
+
+
+
+/*
+ * ============================== Sort algorithms ==============================
+ */
+
+/*!
+ * Each thread can handle 2 points in the array. For each of these 2 points it may
+ * - compare and exchange if needed
+ * - copy data to local and back if needed
+ */
+static constexpr size_t SizeToThreadsRatio = 2;
+
+/*!
+ * Calculates the blocks needed for the entire sorting process
+ *
+ * @note
+ *  This "redundant" little trick makes sure blocks are allocated for arraySizes that are not exact
+ *  multipliers of config.blockSize.
+ *  Even if we don't need it, we keep it in case we experiment with weird sizes in the future!
+ *
+ * @param arraySize [ArraySize_t]   The size of the entire array (in points)
+ * @return          [size_t]        The number of blocks
+ */
+inline size_t NBlocks(ArraySize_t arraySize) {
+	return (((arraySize + config.blockSize - 1) / config.blockSize) / SizeToThreadsRatio);
+}
+
+
+/*!
+ * Exchange utility
+ *
+ * @tparam  ValueT  The underlying data type of the array items
+ *
+ * @param data      [ValueT*]    Pointer to data array
+ * @param tid       [threadId_t] Current thread's index to data
+ * @param pid       [threadId_t] Parents's index to data
+ * @param keepSmall [bool]       Flag to indicate if current threads is keeping the small
+ */
+template <typename ValueT>
+__device__ void exchange(ValueT* data, threadId_t tid, threadId_t partner, bool keepSmall) {
+    if (( keepSmall && (data[tid] > data[partner])) ||
+        (!keepSmall && (data[tid] < data[partner])) ) {
+        ValueT temp = data[tid];
+        data[tid] = data[partner];
+        data[partner] = temp;
+    }
+}
+
+#if CODE_VERSION == V0
+
+/*!
+ * This is the body of each thread. This function compare and exchange data
+ *
+ * @tparam ValueT   The underlying data type of the array items
+ * @param data      [ValueT*]   Pointer to data array
+ * @param n         [size_t]    The total size of the array
+ * @param step      [size_t]    The current step of the current stage of bitonic sort
+ * @param stage     [size_t]    The current stage of bitonic sort
+ */
+template <typename ValueT>
+__global__ void bitonicStep(ValueT* data, size_t n, size_t step, size_t stage) {
+    threadId_t tid = threadIdx.x + blockIdx.x * blockDim.x; // Keep contiguous addressing to the first half of the array
+    threadId_t pid = partner(tid, step);
+    if (tid > pid) {
+        // Shift to the other half of the array for global data
+        tid += n / SizeToThreadsRatio;
+        pid += n / SizeToThreadsRatio;
+    }
+    if ((tid < n) && (pid < n)) {       // Boundary check
+        bool keep = keepSmall(tid, pid, stage);
+        exchange(data, tid, pid, keep);
+    }
+}
+
+
+/*!
+ * A CUDA version of the Bitonic sort algorithm.
+ *
+ * @tparam DataT    A container type to hold data array. Should have .data() and .size() methods
+ * @param data      [DataT&]    Reference to the container to sort
+ */
+template <typename DataT>
+void bitonicSort(DataT& data) {
+    using value_t = typename DataT::value_type;
+
+    value_t* dev_data;
+    auto size = data.size();
+
+    Timer_memory.start();
+    if (cudaMalloc(&dev_data, size * sizeof(value_t)) != cudaSuccess)
+        throw std::runtime_error("[CUDA] - Can not allocate memory\n");
+    if (cudaMemcpy(dev_data, data.data(), size * sizeof(value_t), cudaMemcpyHostToDevice) != cudaSuccess)
+        throw std::runtime_error("[CUDA] - Can not copy memory to device\n");
+    Timer_memory.stop();
+
+    size_t Nth = config.blockSize;
+    size_t Nbl = NBlocks(size);
+
+    size_t Stages = static_cast<size_t>(log2(size));
+    Timer_sorting.start();
+    for (size_t stage = 1; stage <= Stages; ++stage) {
+        for (size_t step = stage; step > 0; ) {
+            --step;
+            bitonicStep<<<Nbl, Nth>>>(dev_data, size, step, stage);
+            cudaDeviceSynchronize();
+        }
+    }
+    Timer_sorting.stop();
+
+    Timer_memory.start();
+    if (cudaMemcpy(data.data(), dev_data, size * sizeof(value_t), cudaMemcpyDeviceToHost) != cudaSuccess)
+        throw std::runtime_error("[CUDA] - Can not copy memory from device\n");
+    cudaFree(dev_data);
+    Timer_memory.stop();
+}
+
+#elif CODE_VERSION == V1
+
+/*!
+ * This is the body of each thread. This function compare and exchange data
+ *
+ * @tparam ValueT   The underlying data type of the array items
+ * @param data      [ValueT*]   Pointer to data array
+ * @param n         [size_t]    The total size of the array
+ * @param step      [size_t]    The current step of the current stage of bitonic sort
+ * @param stage     [size_t]    The current stage of bitonic sort
+ */
+template <typename ValueT>
+__device__ void interBlockStep_(ValueT* data, size_t n, size_t step, size_t stage) {
+    /*
+     * Here we skip blocks every time (one for SizeToThreadsRatio = 2)
+     * And we use the neighbor block address indices for the other half of the threads
+     */
+    threadId_t tid = threadIdx.x + SizeToThreadsRatio * blockIdx.x * blockDim.x;
+    threadId_t pid = partner(tid, step);
+    if (tid > pid) {
+        // Shift to the other half of the array for global data
+        tid += blockDim.x;
+        pid += blockDim.x;
+    }
+    if ((tid < n) && (pid < n)) {       // Boundary check
+        bool keep = keepSmall(tid, pid, stage);
+        exchange(data, tid, pid, keep);
+    }
+}
+
+/*!
+ * This is the version of the body that is called outside of the loop unrolling
+ *
+ * @tparam ValueT   The underlying data type of the array items
+ * @param data      [ValueT*]   Pointer to data array
+ * @param n         [size_t]    The total size of the array
+ * @param step      [size_t]    The current step of the current stage of bitonic sort
+ * @param stage     [size_t]    The current stage of bitonic sort
+ */
+template <typename ValueT>
+__global__ void interBlockStep(ValueT* data, size_t n, size_t step, size_t stage) {
+    interBlockStep_(data, n, step, stage);
+}
+
+/*!
+ * This is unrolled part of the bitonic double loop.
+ *
+ * @tparam ValueT   The underlying data type of the array items
+ * @param data      [ValueT*]   Pointer to data array
+ * @param n         [size_t]    The total size of the array
+ * @param step      [size_t]    The current step of the current stage of bitonic sort
+ * @param stage     [size_t]    The current stage of bitonic sort
+ */
+template <typename ValueT>
+__global__ void inBlockStep(ValueT* data, size_t n, size_t innerSteps, size_t stage) {
+    for (size_t step = innerSteps + 1; step > 0; ) {
+        --step;
+        interBlockStep_(data, n, step, stage);
+        __syncthreads();
+    }
+}
+
+/*!
+ * A CUDA version of the Bitonic sort algorithm.
+ *
+ * @tparam DataT    A container type to hold data array. Should have .data() and .size() methods
+ * @param data      [DataT&]    Reference to the container to sort
+ */
+template <typename DataT>
+void bitonicSort(DataT& data) {
+    using value_t = typename DataT::value_type;
+
+    value_t* dev_data;
+    auto size = data.size();
+
+    Timer_memory.start();
+    if (cudaMalloc(&dev_data, size * sizeof(value_t)) != cudaSuccess)
+        throw std::runtime_error("[CUDA] - Can not allocate memory\n");
+    if (cudaMemcpy(dev_data, data.data(), size * sizeof(value_t), cudaMemcpyHostToDevice) != cudaSuccess)
+        throw std::runtime_error("[CUDA] - Can not copy memory to device\n");
+    Timer_memory.stop();
+
+    size_t Nth = config.blockSize;
+    size_t Nbl = NBlocks(size);
+
+    auto Stages          = static_cast<size_t>(log2(size));
+    auto InnerBlockSteps = static_cast<size_t>(log2(Nth));    //
+    Timer_sorting.start();
+    for (size_t stage = 1; stage <= Stages; ++stage) {
+        size_t step = stage - 1;
+        for ( ; step > InnerBlockSteps; --step) {
+            interBlockStep<<<Nbl, Nth>>>(dev_data, size, step, stage);
+            cudaDeviceSynchronize();
+        }
+        inBlockStep<<<Nbl, Nth>>>(dev_data, size, step, stage);
+        cudaDeviceSynchronize();
+    }
+    Timer_sorting.stop();
+
+    Timer_memory.start();
+    if (cudaMemcpy(data.data(), dev_data, size * sizeof(value_t), cudaMemcpyDeviceToHost) != cudaSuccess)
+        throw std::runtime_error("[CUDA] - Can not copy memory from device\n");
+    cudaFree(dev_data);
+    Timer_memory.stop();
+}
+
+
+
+#elif CODE_VERSION == V2
+
+/*!
+ * @return The memory that each block local threads can affect.
+ *
+ * @note
+ *  Each block thread collection can exchange twice the size of data points.
+ */
+inline size_t effectiveBlockSize() { return SizeToThreadsRatio * config.blockSize; }
+
+
+
+/*!
+ * Converts the global address of the data to the local shared memory array which is used
+ * as cached memory to the unrolled part of the bitonic sort loop.
+ *
+ * @note
+ *  Each block's thread collection can exchange twice the size of data points.
+ *  These points get copied (cached) in the shared memory location. We use contiguous blocks
+ *  both in global data memory and the shared memory buffer.
+ *
+ * @param gIndex    The global array index
+ * @param blockDim  The block size (threads per block)
+ * @return          The equivalent local address of the shared memory
+ */
+__device__ inline size_t toLocal(size_t gIndex, size_t blockDim) {
+    return gIndex % (SizeToThreadsRatio * blockDim);
+}
+
+/*!
+ * This is the version of the body that is called outside of the loop unrolling
+ *
+ * @tparam ValueT   The underlying data type of the array items
+ * @param data      [ValueT*]   Pointer to data array
+ * @param n         [size_t]    The total size of the array
+ * @param step      [size_t]    The current step of the current stage of bitonic sort
+ * @param stage     [size_t]    The current stage of bitonic sort
+ */
+template <typename ValueT>
+__global__ void interBlockStep(ValueT* data, size_t n, size_t step, size_t stage) {
+    threadId_t tid = threadIdx.x + blockIdx.x * blockDim.x;     // Keep contiguous addressing to the first half of the array
+    threadId_t pid = partner(tid, step);
+    if (tid > pid) {
+        // Shift to the other half of the array for global data
+        tid += n / SizeToThreadsRatio;
+        pid += n / SizeToThreadsRatio;
+    }
+    if ((tid < n) && (pid < n)) {       // Boundary check
+        bool keep = keepSmall(tid, pid, stage);
+        exchange(data, tid, pid, keep);
+    }
+}
+
+/*!
+ * This is unrolled part of the bitonic double loop.
+ *
+ * First each thread caches its corresponding data point from the current and the following data block.
+ * After that we execute the loop unrolling on the local data and then we write back to global memory.
+ *
+ * @tparam ValueT   The underlying data type of the array items
+ * @param data      [ValueT*]   Pointer to data array
+ * @param n         [size_t]    The total size of the array
+ * @param step      [size_t]    The current step of the current stage of bitonic sort
+ * @param stage     [size_t]    The current stage of bitonic sort
+ */
+template <typename ValueT>
+__global__ void inBlockStep(ValueT* data, size_t n, size_t innerSteps, size_t stage) {
+    extern __shared__ ValueT shared_data[];
+
+    /*
+     * Global and local(shared) memory indices (calculated once)
+     * Here we skip blocks every time (one for SizeToThreadsRatio = 2)
+     * And we cache the neighbor block address indexes in local (shared) memory
+     */
+    threadId_t gIdx0 = threadIdx.x + SizeToThreadsRatio * blockIdx.x * blockDim.x;
+    threadId_t lIdx0 = toLocal(gIdx0, blockDim.x);
+
+    if (gIdx0 + blockDim.x >= n)   // Boundary check
+        return;
+
+    // Fetch to local memory the entire effective block size (2 positions for each thread)
+    shared_data[lIdx0]              = data[gIdx0];
+    shared_data[lIdx0 + blockDim.x] = data[gIdx0 + blockDim.x];
+    __syncthreads();
+
+    for (size_t step = innerSteps + 1; step > 0; ) {
+        --step;
+
+        // Init thread global and local indices
+        threadId_t gIdx = gIdx0;
+        threadId_t lIdx = lIdx0;
+        // Find partner and keep-small configuration based on the global data positions
+        threadId_t pIdx  = partner(gIdx, step);
+        if (gIdx > pIdx) {
+            // Shift inside effective block
+            gIdx += blockDim.x; // global
+            pIdx += blockDim.x;
+            lIdx += blockDim.x;	// local
+        }
+        bool keep = keepSmall(gIdx, pIdx, stage);
+
+        // Exchange data on local(shared) copy
+        threadId_t lpIdx = toLocal(pIdx, blockDim.x);
+        exchange(shared_data, lIdx, lpIdx, keep);
+        __syncthreads();
+    }
+
+    // Write back to global memory
+    data[gIdx0]              = shared_data[lIdx0];
+    data[gIdx0 + blockDim.x] = shared_data[lIdx0 + blockDim.x];
+    __syncthreads();
+}
+
+/*!
+ * A CUDA version of the Bitonic sort algorithm.
+ *
+ * @tparam DataT    A container type to hold data array. Should have .data() and .size() methods
+ * @param data      [DataT&]    Reference to the container to sort
+ */
+template <typename DataT>
+void bitonicSort(DataT& data) {
+    using value_t = typename DataT::value_type;
+
+    value_t* dev_data;
+    auto size = data.size();
+
+    Timer_memory.start();
+    if (cudaMalloc(&dev_data, size * sizeof(value_t)) != cudaSuccess)
+        throw std::runtime_error("[CUDA] - Can not allocate memory\n");
+    if (cudaMemcpy(dev_data, data.data(), size * sizeof(value_t), cudaMemcpyHostToDevice) != cudaSuccess)
+        throw std::runtime_error("[CUDA] - Can not copy memory to device\n");
+    Timer_memory.stop();
+
+    size_t Nth = config.blockSize;
+    size_t Nbl = NBlocks(size);
+    size_t kernelMemSize = effectiveBlockSize() * sizeof(value_t);
+
+    auto Stages          = static_cast<size_t>(log2(size));
+    auto InnerBlockSteps = static_cast<size_t>(log2(Nth));
+    Timer_sorting.start();
+    for (size_t stage = 1; stage <= Stages; ++stage) {
+        size_t step = stage - 1;
+        for ( ; step > InnerBlockSteps; --step) {
+            interBlockStep<<<Nbl, Nth>>>(dev_data, size, step, stage);
+            cudaDeviceSynchronize();
+        }
+        inBlockStep<<<Nbl, Nth, kernelMemSize>>>(dev_data, size, step, stage);
+        cudaDeviceSynchronize();
+    }
+    Timer_sorting.stop();
+
+    Timer_memory.start();
+    if (cudaMemcpy(data.data(), dev_data, size * sizeof(value_t), cudaMemcpyDeviceToHost) != cudaSuccess)
+        throw std::runtime_error("[CUDA] - Can not copy memory from device\n");
+    cudaFree(dev_data);
+    Timer_memory.stop();
+}
+
+#endif
+
+#endif //BITONICSORTCUDA_H_

homework_3/src/config.h

@@ -1,6 +1,6 @@
 /*!
  * \file
- * \brief   Build configuration file.
+ * \brief   Build and runtime configuration file.
  *
  * \author
  *    Christos Choutouridis AEM:8997
@@ -11,31 +11,35 @@
 #define CONFIG_H_
 
 #include <cstdint>
+#include <cuda_runtime.h>
 
 /*
  * Versioning:
- * - RC1:
+ * - RC1: First version to test on HPC
  */
-static constexpr char version[] = "0.0";
+static constexpr char version[] = "0.1";
 
 /*
  * Defines for different version of the exercise
  */
-#define  V0        (0)
-#define  V1        (1)
-#define  V2        (2)
+#define  V0        0
+#define  V1        1
+#define  V2        2
 
 // Fail-safe version selection
 #if !defined CODE_VERSION
-#define CODE_VERSION   V0
+#define CODE_VERSION   V2
 #endif
 
 // Default Data size (in case -q <N> is not present)
 static constexpr size_t DEFAULT_DATA_SIZE   = 1 << 16;
 
+// Placeholder default (actual default comes from device properties read at initialization)
+static constexpr size_t THREADS_PER_BLOCK   = 1024;
+
 
 /*!
- * Value type selection
+ * Value and Buffer type selection
  *
  * We support the following compiler types or the <cstdint> that translate to them:
  *  char       -  unsigned char
@@ -46,7 +50,13 @@ static constexpr size_t DEFAULT_DATA_SIZE   = 1 << 16;
  *  float
  *  double
  */
-using   distValue_t = uint32_t;
+using Value_t = uint32_t;
+using Data_t  = std::vector<Value_t>;
+
+/*!
+ * In theory we can support large arrays ;)
+ */
+using ArraySize_t = uint64_t;
 
 /*!
  * Session option for each invocation of the executable.
@@ -55,13 +65,11 @@ using   distValue_t = uint32_t;
  *  The values of the members are set from the command line.
  */
 struct config_t {
-    size_t  arraySize{DEFAULT_DATA_SIZE};   //!< The array size of the local data to sort.
-    bool    exchangeOpt{false};             //!< Flag to request the exchange optimization
-    size_t  pipeline{1UL};                  //!< Pipeline stages (1 to disable)
+	ArraySize_t	arraySize{DEFAULT_DATA_SIZE};   //!< The array size of the local data to sort.
+	size_t      blockSize{THREADS_PER_BLOCK};	//!< The block size (threads per block) for the session.
     bool        validation{false};              //!< Request a full validation at the end, performed by process rank 0.
-    bool    ndebug{false};                  //!< Skips debug trap on DEBUG builds.
-    size_t  perf{1};                        //!< Enable performance timing measurements and prints and repeat
-                                            //!< the sorting <perf> times.
+    size_t      perf{1};                        //!< Enable performance timing measurements and prints. Repeat
+                                                //!< the sorting <perf> times to do so.
     bool        verbose{false};                 //!< Flag to enable verbose output to stdout.
 };
 
@@ -69,6 +77,6 @@ struct config_t {
  * Exported data types
  */
 extern config_t 		config;
-
+extern cudaDeviceProp   device;
 
 #endif /* CONFIG_H_ */

homework_3/src/distsort.cpp

@@ -1,51 +0,0 @@
-/*!
- * \file
- * \brief   Distributed sort implementation
- *
- * \author
- *    Christos Choutouridis AEM:8997
- *    <cchoutou@ece.auth.gr>
- */
-#include "utils.hpp"
-#include "distsort.hpp"
-
-
-/*!
- * Returns the ascending or descending configuration of the node's sequence based on
- * the current node (MPI process) and the depth of the sorting network
- *
- * @param node      [mpi_id_t] The current node (MPI process)
- * @param depth     [size_t]   The total depth of the sorting network (same for each step for a given network)
- * @return          [bool]     True if we need ascending configuration, false otherwise
- */
-bool ascending(mpi_id_t node, size_t depth) noexcept {
-    return !(node & (1 << depth));
-}
-
-/*!
- * Returns the node's partner for data exchange during the sorting network iterations
- * of Bitonic
- *
- * @param node      [mpi_id_t] The current node
- * @param step      [size_t]   The step of the sorting network
- * @return          [mpi_id_t] The node id of the partner for data exchange
- */
-mpi_id_t partner(mpi_id_t node, size_t step) noexcept {
-    return (node ^ (1 << step));
-}
-
-
-/*!
- * Predicate to check if a node keeps the small numbers during the bitonic sort network exchange.
- *
- * @param node      [mpi_id_t] The node for which we check
- * @param partner   [mpi_id_t] The partner of the data exchange
- * @param depth     [size_t]   The total depth of the sorting network (same for each step for a given network)
- * @return          [bool]     True if the node should keep the small values, false otherwise
- */
-
-bool keepSmall(mpi_id_t node, mpi_id_t partner, size_t depth) {
-    if (node == partner)
-        throw std::runtime_error("(keepSmall) Node and Partner can not be the same\n");
-    return ascending(node, depth) == (node < partner);
-}

homework_3/src/distsort.hpp

@@ -1,223 +0,0 @@
-/*!
- * \file
- * \brief   Distributed sort implementation header
- *
- * \author
- *    Christos Choutouridis AEM:8997
- *    <cchoutou@ece.auth.gr>
- */
-
-#ifndef DISTBITONIC_H_
-#define DISTBITONIC_H_
-
-#include <vector>
-#include <algorithm>
-#include <parallel/algorithm>
-#include <cmath>
-#include <cstdint>
-#if !defined DEBUG
-#define NDEBUG
-#endif
-#include <cassert>
-
-#include "utils.hpp"
-
-/*
- * Exported timers
- */
-extern Timing Timer_total;
-extern Timing Timer_fullSort;
-extern Timing Timer_exchange;
-extern Timing Timer_minmax;
-extern Timing Timer_elbowSort;
-
-
-
-/*
- * ============================== Sort utilities ==============================
- */
-
-
-
-/*!
- * Returns the ascending or descending configuration of the node's sequence based on
- * the current node (MPI process) and the depth of the sorting network
- *
- * @param node      [mpi_id_t] The current node (MPI process)
- * @param depth     [size_t]   The total depth of the sorting network (same for each step for a given network)
- * @return          [bool]     True if we need ascending configuration, false otherwise
- */
-bool ascending(mpi_id_t node, size_t depth);
-
-/*!
- * Returns the node's partner for data exchange during the sorting network iterations
- * of Bitonic
- *
- * @param node      [mpi_id_t] The current node
- * @param step      [size_t]   The step of the sorting network
- * @return          [mpi_id_t] The node id of the partner for data exchange
- */
-mpi_id_t partner(mpi_id_t node, size_t step);
-
-
-/*!
- * Predicate to check if a node keeps the small numbers during the bitonic sort network exchange.
- *
- * @param node      [mpi_id_t] The node for which we check
- * @param partner   [mpi_id_t] The partner of the data exchange
- * @param depth     [size_t]   The total depth of the sorting network (same for each step for a given network)
- * @return          [bool]     True if the node should keep the small values, false otherwise
- */
-bool keepSmall(mpi_id_t node, mpi_id_t partner, size_t depth);
-
-
-
-/*
- * ============================== Data utilities ==============================
- */
-
-/*!
- * Sort a range using the build-in O(Nlog(N)) algorithm
- *
- * @tparam RangeT   A range type with random access iterator
- *
- * @param data      [RangeT] The data to be sorted
- * @param ascending [bool]   Flag to indicate the sorting order
- */
-template<typename RangeT>
-void fullSort(RangeT& data, bool ascending) noexcept {
-    // Use introsort from stdlib++ here, unless ... __gnu_parallel
-    if (ascending) {
-        __gnu_parallel::sort(data.begin(), data.end(), std::less<>());
-    }
-    else {
-        __gnu_parallel::sort(data.begin(), data.end(), std::greater<>());
-    }
-}
-
-/*!
- * Core functionality of sort for shadowed buffer types using
- * the "elbow sort" algorithm.
- *
- * @note:
- *  This algorithm can not work "in place".
- *  We use the active buffer as source and the shadow as target.
- *  At the end we switch which buffer is active and which is the shadow.
- * @note
- *  This is the core functionality. Use the elbowSort() function instead
- *
- * @tparam ShadowedDataT    A Shadowed buffer type with random access iterator.
- * @tparam CompT            A Comparison type for binary operation comparisons
- *
- * @param data          [ShadowedDataT] The data to sort
- * @param ascending     [bool]          Flag to indicate the sorting order
- * @param comp          [CompT]         The binary operator object
- */
-template<typename ShadowedDataT, typename CompT>
-void elbowSortCore(ShadowedDataT& data, bool ascending, CompT comp) noexcept {
-    auto& active = data.getActive(); // Get the source vector (the data to sort)
-    auto& shadow = data.getShadow(); // Get the target vector (the sorted data)
-
-    size_t N = data.size();         // The total size is the same or both vectors
-    size_t left = std::distance(
-            active.begin(),
-            (ascending) ?
-                std::min_element(active.begin(), active.end()) :
-                std::max_element(active.begin(), active.end())
-    );                              // start 'left' from elbow of the bitonic
-    size_t right = (left == N-1) ? 0 : left + 1;
-
-    // Walk in opposite directions from elbow and insert-sort to target vector
-    for (size_t i = 0 ; i<N ; ++i) {
-        if (comp(active[left], active[right])) {
-            shadow[i] = active[left];
-            left = (left == 0) ? N-1 : left -1; // cycle decrease
-        }
-        else {
-            shadow[i] = active[right];
-            right = (right + 1) % N;            // cycle increase
-        }
-    }
-    data.switch_active();           // Switch active-shadow buffers
-}
-
-/*!
- * Sort a shadowed buffer using the "elbow sort" algorithm.
- *
- * @tparam ShadowedDataT    A Shadowed buffer type with random access iterator.
- *
- * @param data          [ShadowedDataT] The data to sort
- * @param ascending     [bool]          Flag to indicate the sorting order
- */
-template<typename ShadowedDataT>
-void elbowSort(ShadowedDataT& data, bool ascending) noexcept {
-    if (ascending)
-        elbowSortCore(data, ascending, std::less<>());
-    else
-        elbowSortCore(data, ascending, std::greater<>());
-}
-
-
-/*!
- * Takes two sequences and selects either the larger or the smaller items
- * in one-to-one comparison between them. If the initial sequences are bitonic, then
- * the result is a bitonic sequence too!
- *
- * @tparam ValueT   The underlying type of the sequences
- *
- * @param local     [ValueT*]       Pointer to the local sequence
- * @param remote    [const ValueT*] Pointer to the remote sequence (copied locally by MPI)
- * @param count     [size_t]        The number of items to process
- * @param keepSmall [bool]          Flag to indicate if we keep the small items in local sequence
- */
-template<typename ValueT>
-void keepMinOrMax(ValueT* local, const ValueT* remote, size_t count, bool keepSmall) noexcept {
-    std::transform(
-            local, local + count,
-            remote,
-            local,
-            [&keepSmall](const ValueT& a, const ValueT& b){
-                return (keepSmall) ? std::min(a, b) : std::max(a, b);
-            });
-}
-
-/*
- * ============================== Sort algorithms ==============================
- */
-
-
-/*!
- * A distributed version of the Bitonic sort algorithm.
- *
- * @note
- *  Each MPI process should run an instance of this function.
- *
- * @tparam ShadowedDataT    A Shadowed buffer type with random access iterator.
- *
- * @param data          [ShadowedDataT] The local to MPI process data to sort
- * @param Processes     [mpi_id_t]      The total number of MPI processes
- * @param rank          [mpi_id_t]      The current process id
- */
-template<typename ShadowedDataT>
-void distBitonic(ShadowedDataT& data) {
-    // Initially sort to create a half part of a bitonic sequence
-    timeCall(Timer_fullSort, fullSort, data, ascending(rank, 0));
-
-    // Run through sort network using elbow-sort ( O(LogN * LogN) iterations )
-    auto p = static_cast<uint32_t>(std::log2(Processes));
-    for (size_t depth = 1; depth <= p; ++depth) {
-        for (size_t step = depth; step > 0;) {
-            --step;
-            // Find out exchange configuration
-            auto part = partner(rank, step);
-            auto ks = keepSmall(rank, part, depth);
-            // Exchange with partner, keep nim-or-max
-            exchange(data, part, ks, tag);
-
-        }
-        // sort - O(N)
-        timeCall(Timer_elbowSort, elbowSort, data, ascending(rank, depth));
-    }
-}
-
-#endif //DISTBITONIC_H_

homework_3/src/main.cpp

@@ -11,37 +11,33 @@
 #include <iostream>
 #include <algorithm>
 #include <random>
+#include <cuda_runtime.h>
 
 #include "utils.hpp"
 #include "config.h"
-#include "distsort.hpp"
+#include "bitonicsort.hpp"
 
 
 // Global session data
+Data_t          Data;
 config_t        config;
-distBuffer_t    Data;
 Log             logger;
-
+cudaDeviceProp 	device;
 
 // Mersenne seeded from hw if possible. range: [type_min, type_max]
 std::random_device  rd;
 std::mt19937        gen(rd());
 
 //! Performance timers for each one of the "costly" functions
-Timing Timer_total;
-Timing Timer_fullSort;
-Timing Timer_exchange;
-Timing Timer_minmax;
-Timing Timer_elbowSort;
+Timing Timer_total, Timer_memory, Timer_sorting;
+
 
 //! Init timing objects for extra rounds
 void measurements_init() {
     if (config.perf > 1) {
         Timer_total.init(config.perf);
-        Timer_fullSort.init(config.perf);
-        Timer_exchange.init(config.perf);
-        Timer_minmax.init(config.perf);
-        Timer_elbowSort.init(config.perf);
+        Timer_memory.init(config.perf);
+        Timer_sorting.init(config.perf);
     }
 }
 
@@ -49,10 +45,8 @@ void measurements_init() {
 void measurements_next() {
     if (config.perf > 1) {
         Timer_total.next();
-        Timer_fullSort.next();
-        Timer_exchange.next();
-        Timer_minmax.next();
-        Timer_elbowSort.next();
+        Timer_memory.next();
+        Timer_sorting.next();
     }
 }
 
@@ -69,7 +63,15 @@ bool get_options(int argc, char* argv[]){
 
         if (arg == "-q" || arg == "--array-size") {
             if (i+1 < argc) {
-                config.arraySize = 1 << atoi(argv[++i]);
+                config.arraySize = (ArraySize_t)1 << atoi(argv[++i]);
+            }
+            else {
+                status = false;
+            }
+        }
+        else if (arg == "-b" || arg == "--block-size") {
+            if (i+1 < argc) {
+                config.blockSize = atoi(argv[++i]);
             }
             else {
                 status = false;
@@ -86,32 +88,34 @@ bool get_options(int argc, char* argv[]){
                 status = false;
             }
         }
-        else if (arg == "--ndebug") {
-            config.ndebug = true;
-        }
         else if (arg == "-v" || arg == "--verbose") {
             config.verbose = true;
         }
         else if (arg == "--version") {
-            std::cout << "bitonic - A GPU accelerated sort utility\n";
+            std::cout << STR(TARGET) << " - A GPU accelerated bitonic sort utility (V" << STR(CODE_VERSION)<< ") \n";
             std::cout << "version: " << version << "\n\n";
             exit(0);
         }
         else if (arg == "-h" || arg == "--help") {
-            std::cout << "distbitonic - A distributed sort utility\n\n";
-            std::cout << "  distbitonic -q <N> [--validation] [--perf <N>] [--ndebug] [-v]\n";
-            std::cout << "  distbitonic -h\n";
+            std::cout << STR(TARGET) << " - A GPU accelerated bitonic sort utility (V" << STR(CODE_VERSION)<< ") \n\n";
+            std::cout << "  " << STR(TARGET) << " -q <N> -b <N> [--validation] [--perf <N>] [-v]\n";
+            std::cout << "  " << STR(TARGET) << " -h\n";
+            std::cout << "  " << STR(TARGET) << " --version\n";
             std::cout << '\n';
             std::cout << "Options:\n\n";
             std::cout << "   -q | --array-size <N>\n";
-            std::cout << "      Selects the array size according to size = 2^N\n\n";
+            std::cout << "      Selects the array size according to size = 2^N\n";
+            std::cout << "      [Size must be larger than 2 * blockSize]\n";
+            std::cout << "      [Default is 2^16]\n\n";
+            std::cout << "   -b | --block-size <N>\n";
+            std::cout << "      Selects the number of CUDA threads per block\n";
+            std::cout << "      [Size has to be multiple of device's warp size (usually 32)\n";
+            std::cout << "      [Default is the maximum device supported number. For ex: (GTX 1650) block-size=1024]\n\n";
             std::cout << "   --validation\n";
-            std::cout << "      Request a full validation at the end, performed by process rank 0\n\n";
+            std::cout << "      Request a full validation at the end\n\n";
             std::cout << "   --perf <N> \n";
             std::cout << "      Enable performance timing measurements and prints, and repeat\n";
             std::cout << "      the sorting <N> times.\n\n";
-            std::cout << "   --ndebug\n";
-            std::cout << "      Skip debug breakpoint when on debug build.\n\n";
             std::cout << "   -v | --verbose\n";
             std::cout << "      Request a more verbose output to stdout.\n\n";
             std::cout << "   -h | --help\n";
@@ -119,8 +123,12 @@ bool get_options(int argc, char* argv[]){
             std::cout << "   --version\n";
             std::cout << "      Prints version and exit.\n\n";
             std::cout << "Examples:\n\n";
-            std::cout << "   bitonic -q 24\n";
-            std::cout << "      Runs bitonic with GPU acceleration with 2^24 array points\n\n";
+            std::cout << "  " << STR(TARGET) << " -q 24\n";
+            std::cout << "      Runs bitonic sort on an 2^24 points array, using GPU acceleration\n\n";
+            std::cout << "  " << STR(TARGET) << " --validation --perf 5 -b 512 -q 26\n";
+            std::cout << "      Runs bitonic sort on an 2^26 points array 5 times, using GPU acceleration with\n";
+            std::cout << "      512 threads per block, performs a validation check at the end and prints the time\n";
+            std::cout << "      of the median.\n\n";
 
             exit(0);
         }
@@ -130,26 +138,31 @@ bool get_options(int argc, char* argv[]){
         }
     }
 
+    // Check configuration requirements
+    if (config.blockSize % device.warpSize)
+    	throw std::runtime_error("[Config] - Number of threads per block is not an exact multiple of warp size\n");
+    if (config.arraySize < 2*config.blockSize)
+    	throw std::runtime_error("[Config] - Unsupported array size (smaller than "
+                                 + std::to_string(SizeToThreadsRatio*config.blockSize) + ")\n");
+    if (device.totalGlobalMem < config.arraySize * sizeof(Value_t))
+    	throw std::runtime_error("[CUDA] - Unsupported array size: "
+                                 + std::to_string(config.arraySize * sizeof(Value_t))
+                                 + " (larger than GPU's: " + std::to_string(device.totalGlobalMem) + ")\n");
+
     return status;
 }
 
 /*!
  * A simple validator for the entire distributed process
  *
- * @tparam ShadowedDataT    A Shadowed buffer type with random access iterator.
+ * @tparam DataT    A buffer type with random access iterator.
  *
- * @param data          [ShadowedDataT] The local to MPI process
- * @param Processes     [mpi_id_t]      The total number of MPI processes
- * @param rank          [mpi_id_t]      The current process id
- *
- * @return              [bool]          True if all are sorted and in total ascending order
+ * @param data      [DataT] The data
+ * @return          [bool]  True if sorted in ascending order
  */
-template<typename ShadowedDataT>
-bool validator(ShadowedDataT& data) {
-    using value_t = typename ShadowedDataT::value_type;
-    bool ret = true;    // Have faith!
-
-    return ret;
+template<typename DataT>
+bool validator(DataT& data) {
+    return std::is_sorted(data.begin(), data.end());
 }
 
 /*!
@@ -159,6 +172,13 @@ bool validator(ShadowedDataT& data) {
  * @param argv  [char***]   POINTER to main's argv argument
  */
 void init(int* argc, char*** argv) {
+
+    // Get device configuration
+    if (cudaGetDeviceProperties(&device, 0) != cudaSuccess)
+        throw std::runtime_error("[CUDA] - Can not read GPU");
+
+    config.blockSize = static_cast<size_t>(device.maxThreadsPerBlock);
+
     // try to read command line
     if (!get_options(*argc, *argv))
         exit(1);
@@ -177,18 +197,25 @@ int main(int argc, char* argv[]) try {
     // Init everything
     init(&argc, &argv);
 
+    logger << "Array size:     " << config.arraySize << " (Q=" << static_cast<size_t>(log2(config.arraySize))<< ")" << logger.endl;
+    logger << "Repeated sorts: " << config.perf << logger.endl;
+    logger << "GPU:            " << device.name << logger.endl;
+    logger << "Block size:     " << config.blockSize << logger.endl;
+
     for (size_t it = 0 ; it < config.perf ; ++it) {
         // Initialize local data
-        logger << "Initialize local array of " << config.arraySize << " elements" << logger.endl;
-        std::uniform_int_distribution<distValue_t > dis(
-                std::numeric_limits<distValue_t>::min(),
-                std::numeric_limits<distValue_t>::max()
+        logger << "Initialize array ... ";
+        std::uniform_int_distribution<Value_t > dis(
+                std::numeric_limits<Value_t>::min(),
+                std::numeric_limits<Value_t>::max()
         );
         std::generate(Data.begin(), Data.end(), [&]() { return dis(gen); });
+        logger << " Done." << logger.endl;
+
         // Run distributed sort
-        logger << "Starting distributed sorting ... ";
+        logger << "Start sorting ...    ";
         Timer_total.start();
-        distBitonic(Data);
+        bitonicSort(Data);
         Timer_total.stop();
         measurements_next();
         logger << " Done." << logger.endl;
@@ -196,17 +223,15 @@ int main(int argc, char* argv[]) try {
 
     // Print-outs and validation
     if (config.perf > 1) {
-        Timing::print_duration(Timer_total.median(),    "Total     ", 0);
-        Timing::print_duration(Timer_fullSort.median(), "Full-Sort ", 0);
-        Timing::print_duration(Timer_exchange.median(), "Exchange  ", 0);
-        Timing::print_duration(Timer_minmax.median(),   "Min-Max   ", 0);
-        Timing::print_duration(Timer_elbowSort.median(),"Elbow-Sort", 0);
+        Timing::print_duration(Timer_total.median(),  "Total   ");
+        Timing::print_duration(Timer_memory.median(), "Mem-xch ");
+        Timing::print_duration(Timer_sorting.median(),"Sorting ");
     }
     if (config.validation) {
         // If requested, we have the chance to fail!
         std::cout << "[Validation] Results validation ...";
         bool val = validator(Data);
-        std::cout << ((val) ? "\x1B[32m [PASSED] \x1B[0m\n" : " \x1B[32m [FAILED] \x1B[0m\n");
+        std::cout << ((val) ? "\x1B[32m [PASSED] \x1B[0m\n" : " \x1B[31m [FAILED] \x1B[0m\n");
     }
     return 0;
 }

homework_3/src/utils.hpp

@@ -17,124 +17,11 @@
 
 #include "config.h"
 
-
 /*!
- * @brief A std::vector wrapper with 2 vectors, an active and a shadow.
- *
- * This type exposes the standard vector functionality of the active vector.
- * The shadow can be used when we need to use the vector as mutable
- * data in algorithms that can not support "in-place" editing (like elbow-sort for example)
- *
- * @tparam Value_t  the underlying data type of the vectors
+ * Stringify preprocessor util
  */
-template <typename Value_t>
-struct ShadowedVec_t {
-    // STL requirements
-    using value_type     = Value_t;
-    using iterator       = typename std::vector<Value_t>::iterator;
-    using const_iterator = typename std::vector<Value_t>::const_iterator;
-    using size_type      = typename std::vector<Value_t>::size_type;
-
-    // Default constructor
-    ShadowedVec_t() = default;
-
-    // Constructor from an std::vector
-    explicit ShadowedVec_t(const std::vector<Value_t>& vec)
-            : North(vec), South(), active(north) {
-        South.resize(North.size());
-    }
-
-    explicit ShadowedVec_t(std::vector<Value_t>&& vec)
-            : North(std::move(vec)), South(), active(north) {
-        South.resize(North.size());
-    }
-
-    // Copy assignment operator
-    ShadowedVec_t& operator=(const ShadowedVec_t& other) {
-        if (this != &other) { // Avoid self-assignment
-            North = other.North;
-            South = other.South;
-            active = other.active;
-        }
-        return *this;
-    }
-
-    // Move assignment operator
-    ShadowedVec_t& operator=(ShadowedVec_t&& other) noexcept {
-        if (this != &other) { // Avoid self-assignment
-            North = std::move(other.North);
-            South = std::move(other.South);
-            active = other.active;
-
-            // There is no need to zero out other since it is valid but in a non-defined state
-        }
-        return *this;
-    }
-
-    // Type accessors
-    std::vector<Value_t>& getActive() { return (active == north) ? North : South; }
-    std::vector<Value_t>& getShadow() { return (active == north) ? South : North; }
-    const std::vector<Value_t>& getActive() const { return (active == north) ? North : South; }
-    const std::vector<Value_t>& getShadow() const { return (active == north) ? South : North; }
-
-    // Swap vectors
-    void switch_active() { active = (active == north) ? south : north; }
-
-    // Dispatch vector functionality to active vector
-    Value_t& operator[](size_type index) { return getActive()[index]; }
-    const Value_t& operator[](size_type index) const { return getActive()[index]; }
-
-    Value_t& at(size_type index) { return getActive().at(index); }
-    const Value_t& at(size_type index) const { return getActive().at(index); }
-
-    void push_back(const Value_t& value) { getActive().push_back(value); }
-    void push_back(Value_t&& value)      { getActive().push_back(std::move(value)); }
-    void pop_back()                      { getActive().pop_back(); }
-    Value_t& front() { return getActive().front(); }
-    Value_t& back()  { return getActive().back(); }
-    const Value_t& front() const { return getActive().front(); }
-    const Value_t& back()  const { return getActive().back(); }
-
-    iterator begin() { return getActive().begin(); }
-    const_iterator begin() const { return getActive().begin(); }
-    iterator end() { return getActive().end(); }
-    const_iterator end() const { return getActive().end(); }
-
-    size_type size() const { return getActive().size(); }
-    void resize(size_t new_size) {
-        North.resize(new_size);
-        South.resize(new_size);
-    }
-
-    void reserve(size_t new_capacity) {
-        North.reserve(new_capacity);
-        South.reserve(new_capacity);
-    }
-    [[nodiscard]] size_t capacity() const { return getActive().capacity(); }
-    [[nodiscard]] bool empty() const { return getActive().empty(); }
-
-    void clear() { getActive().clear(); }
-    void swap(std::vector<Value_t>& other) { getActive().swap(other); }
-
-    // Comparisons
-    bool operator== (const ShadowedVec_t& other) { return getActive() == other.getActive(); }
-    bool operator!= (const ShadowedVec_t& other) { return getActive() != other.getActive(); }
-    bool operator== (const std::vector<value_type>& other) { return getActive() == other; }
-    bool operator!= (const std::vector<value_type>& other) { return getActive() != other; }
-
-private:
-    std::vector<Value_t> North{};       //!< Actual buffer to be used either as active or shadow
-    std::vector<Value_t> South{};       //!< Actual buffer to be used either as active or shadow
-    enum {
-        north, south
-    } active{north};                    //!< Flag to select between North and South buffer
-};
-
-/*
- * Exported data types
- */
-using distBuffer_t = ShadowedVec_t<distValue_t>;
-extern distBuffer_t Data;
+#define STR(s)		STR_(s)
+#define STR_(s)		#s
 
 /*!
  * A Logger for entire program.

homework_3/test/tests.cpp

@@ -25,8 +25,7 @@ protected:
 
 
 /*
- * MPI: SysTest (acceptance)
- * Each process executes distBubbletonic for uin8_t [16]
+ *
  */
 TEST_F(TCUDAbitonic, test1) {