From 3698b144047e73ebacbc5f695b439c44e5c4da5c Mon Sep 17 00:00:00 2001 From: Andrew Boessen Date: Thu, 9 Jan 2025 14:19:38 -0600 Subject: [PATCH 1/4] add tests for warp sort --- .github/workflows/makefile.yml | 7 +- makefile | 5 +- test_bitonic_sort.cu | 178 +++++++++++++++++++++++++++++++++ 3 files changed, 185 insertions(+), 5 deletions(-) create mode 100644 test_bitonic_sort.cu diff --git a/.github/workflows/makefile.yml b/.github/workflows/makefile.yml index f8ef971..51ad7ce 100644 --- a/.github/workflows/makefile.yml +++ b/.github/workflows/makefile.yml @@ -42,9 +42,8 @@ jobs: exit 1 fi - - name: Run tests (optional) + - name: Run tests run: | - # Run tests if you have them defined - # Example: ./test_program - echo "No tests defined" + # Run tests + ./test_bitonic_sort diff --git a/makefile b/makefile index 54dbc0c..d6066a7 100644 --- a/makefile +++ b/makefile @@ -20,5 +20,8 @@ main.o: main.cpp warp_bitonic_sort.cuh warp_bitonic_sort.o: warp_bitonic_sort.cu warp_bitonic_sort.cuh $(NVCC) $(NVCCFLAGS) -c $< -o $@ +test_bitonic_sort.o: test_bitonic_sort.cu warp_bitonic_sort.cuh + $(NVCC) $(NVCCFLAGS) -c $< -o $@ + clean: - rm -f *.o warp_bitonic_sort cpu_bitonic_sort + rm -f *.o warp_bitonic_sort cpu_bitonic_sort test_bitonic_sort diff --git a/test_bitonic_sort.cu b/test_bitonic_sort.cu new file mode 100644 index 0000000..ca47b73 --- /dev/null +++ b/test_bitonic_sort.cu @@ -0,0 +1,178 @@ +#include "warp_bitonic_sort.cuh" +#include +#include +#include +#include + +// Function to check if the array is sorted +bool isSorted(int *arr, int size) { + for (int i = 1; i < size; i++) { + if (arr[i] < arr[i - 1]) + return false; + } + return true; +} + +// Test fixture for CUDA Bitonic Sort +class BitonicSortTest : public ::testing::Test { +protected: + void SetUp() override { + // Initialize CUDA + cudaSetDevice(0); + } + + void TearDown() override { + // Clean up CUDA + cudaDeviceReset(); + } + + // Helper function to generate random arrays + std::vector generateRandomArray(int size) { + std::vector arr(size); + for (int i = 0; i < size; i++) { + arr[i] = rand() % 1000; // Random integers between 0 and 999 + } + return arr; + } + + // Helper function to compare two arrays + bool compareArrays(int *arr1, int *arr2, int size) { + for (int i = 0; i < size; i++) { + if (arr1[i] != arr2[i]) + return false; + } + return true; + } +}; + +// Test case for sorting a small array +TEST_F(BitonicSortTest, SmallArraySort) { + const int SIZE = 32; // Must be a multiple of 32 for warp-level sorting + std::vector h_arr = generateRandomArray(SIZE); + std::vector h_arr_sorted = h_arr; + std::sort(h_arr_sorted.begin(), h_arr_sorted.end()); + + int *d_arr; + cudaMalloc(&d_arr, SIZE * sizeof(int)); + cudaMemcpy(d_arr, h_arr.data(), SIZE * sizeof(int), cudaMemcpyHostToDevice); + + launchWarpBitonicSort(d_arr, SIZE); + + cudaMemcpy(h_arr.data(), d_arr, SIZE * sizeof(int), cudaMemcpyDeviceToHost); + + EXPECT_TRUE(isSorted(h_arr.data(), SIZE)); + EXPECT_TRUE(compareArrays(h_arr.data(), h_arr_sorted.data(), SIZE)); + + cudaFree(d_arr); +} + +// Test case for sorting a large array +TEST_F(BitonicSortTest, LargeArraySort) { + const int SIZE = 4096; // Must be a multiple of 32 for warp-level sorting + std::vector h_arr = generateRandomArray(SIZE); + std::vector h_arr_sorted = h_arr; + std::sort(h_arr_sorted.begin(), h_arr_sorted.end()); + + int *d_arr; + cudaMalloc(&d_arr, SIZE * sizeof(int)); + cudaMemcpy(d_arr, h_arr.data(), SIZE * sizeof(int), cudaMemcpyHostToDevice); + + launchWarpBitonicSort(d_arr, SIZE); + + cudaMemcpy(h_arr.data(), d_arr, SIZE * sizeof(int), cudaMemcpyDeviceToHost); + + EXPECT_TRUE(isSorted(h_arr.data(), SIZE)); + EXPECT_TRUE(compareArrays(h_arr.data(), h_arr_sorted.data(), SIZE)); + + cudaFree(d_arr); +} + +// Test case for sorting an already sorted array +TEST_F(BitonicSortTest, AlreadySortedArray) { + const int SIZE = 256; // Must be a multiple of 32 for warp-level sorting + std::vector h_arr(SIZE); + for (int i = 0; i < SIZE; i++) { + h_arr[i] = i; + } + + int *d_arr; + cudaMalloc(&d_arr, SIZE * sizeof(int)); + cudaMemcpy(d_arr, h_arr.data(), SIZE * sizeof(int), cudaMemcpyHostToDevice); + + launchWarpBitonicSort(d_arr, SIZE); + + cudaMemcpy(h_arr.data(), d_arr, SIZE * sizeof(int), cudaMemcpyDeviceToHost); + + EXPECT_TRUE(isSorted(h_arr.data(), SIZE)); + + cudaFree(d_arr); +} + +// Test case for sorting a reverse-sorted array +TEST_F(BitonicSortTest, ReverseSortedArray) { + const int SIZE = 512; // Must be a multiple of 32 for warp-level sorting + std::vector h_arr(SIZE); + for (int i = 0; i < SIZE; i++) { + h_arr[i] = SIZE - i; + } + + int *d_arr; + cudaMalloc(&d_arr, SIZE * sizeof(int)); + cudaMemcpy(d_arr, h_arr.data(), SIZE * sizeof(int), cudaMemcpyHostToDevice); + + launchWarpBitonicSort(d_arr, SIZE); + + cudaMemcpy(h_arr.data(), d_arr, SIZE * sizeof(int), cudaMemcpyDeviceToHost); + + EXPECT_TRUE(isSorted(h_arr.data(), SIZE)); + + cudaFree(d_arr); +} + +// Test case for sorting an array with duplicate elements +TEST_F(BitonicSortTest, ArrayWithDuplicates) { + const int SIZE = 1024; // Must be a multiple of 32 for warp-level sorting + std::vector h_arr = generateRandomArray(SIZE); + for (int i = 0; i < SIZE / 2; i++) { + h_arr[i] = h_arr[i + SIZE / 2]; // Introduce duplicates + } + + std::vector h_arr_sorted = h_arr; + std::sort(h_arr_sorted.begin(), h_arr_sorted.end()); + + int *d_arr; + cudaMalloc(&d_arr, SIZE * sizeof(int)); + cudaMemcpy(d_arr, h_arr.data(), SIZE * sizeof(int), cudaMemcpyHostToDevice); + + launchWarpBitonicSort(d_arr, SIZE); + + cudaMemcpy(h_arr.data(), d_arr, SIZE * sizeof(int), cudaMemcpyDeviceToHost); + + EXPECT_TRUE(isSorted(h_arr.data(), SIZE)); + EXPECT_TRUE(compareArrays(h_arr.data(), h_arr_sorted.data(), SIZE)); + + cudaFree(d_arr); +} + +// Test case for sorting an array with a single element +TEST_F(BitonicSortTest, SingleElementArray) { + const int SIZE = 1; + std::vector h_arr = {42}; + + int *d_arr; + cudaMalloc(&d_arr, SIZE * sizeof(int)); + cudaMemcpy(d_arr, h_arr.data(), SIZE * sizeof(int), cudaMemcpyHostToDevice); + + launchWarpBitonicSort(d_arr, SIZE); + + cudaMemcpy(h_arr.data(), d_arr, SIZE * sizeof(int), cudaMemcpyDeviceToHost); + + EXPECT_TRUE(isSorted(h_arr.data(), SIZE)); + + cudaFree(d_arr); +} + +int main(int argc, char **argv) { + ::testing::InitGoogleTest(&argc, argv); + return RUN_ALL_TESTS(); +} From afd26512b5d4b07f86eea3a7c572a1baf35b612c Mon Sep 17 00:00:00 2001 From: Andrew Boessen Date: Thu, 9 Jan 2025 14:23:59 -0600 Subject: [PATCH 2/4] add tests to all in makefile --- makefile | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/makefile b/makefile index d6066a7..2304894 100644 --- a/makefile +++ b/makefile @@ -5,8 +5,9 @@ NVCCFLAGS = -O2 CUDA_PATH = /opt/cuda INCLUDES = -I$(CUDA_PATH)/include LDFLAGS = -L$(CUDA_PATH)/lib64 -lcudart +GTEST_LDFLAGS = -lgtest -lgtest_main -lpthread -all: warp_bitonic_sort cpu_bitonic_sort +all: warp_bitonic_sort cpu_bitonic_sort test_bitonic_sort warp_bitonic_sort: main.o warp_bitonic_sort.o $(CXX) $^ -o $@ $(LDFLAGS) @@ -14,6 +15,9 @@ warp_bitonic_sort: main.o warp_bitonic_sort.o cpu_bitonic_sort: cpu_bitonic_sort.cpp $(CXX) $^ -o $@ +test_bitonic_sort: test_bitonic_sort.o warp_bitonic_sort.o + $(NVCC) $^ -o $@ $(LDFLAGS) $(GTEST_LDFLAGS) + main.o: main.cpp warp_bitonic_sort.cuh $(CXX) $(CXXFLAGS) $(INCLUDES) -c $< -o $@ From 5a4b79a130b2af3583c9d25bc41048ae44c3b923 Mon Sep 17 00:00:00 2001 From: Andrew Boessen Date: Thu, 9 Jan 2025 14:39:23 -0600 Subject: [PATCH 3/4] install test deps --- .github/workflows/makefile.yml | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/.github/workflows/makefile.yml b/.github/workflows/makefile.yml index 51ad7ce..11e2f43 100644 --- a/.github/workflows/makefile.yml +++ b/.github/workflows/makefile.yml @@ -27,6 +27,7 @@ jobs: # Install other necessary dependencies like build-essential sudo apt-get install -y build-essential + sudo apt-get install libgtest-dev # Verify that CUDA is installed nvcc --version @@ -45,5 +46,6 @@ jobs: - name: Run tests run: | # Run tests - ./test_bitonic_sort + #./test_bitonic_sort + echo "No tests to run" From 26b61db51c75955946f898dd0ebd7ae4bb78f7c1 Mon Sep 17 00:00:00 2001 From: Andrew Boessen Date: Thu, 9 Jan 2025 16:06:27 -0600 Subject: [PATCH 4/4] remove tests from all --- makefile | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/makefile b/makefile index 2304894..9a474c7 100644 --- a/makefile +++ b/makefile @@ -7,7 +7,7 @@ INCLUDES = -I$(CUDA_PATH)/include LDFLAGS = -L$(CUDA_PATH)/lib64 -lcudart GTEST_LDFLAGS = -lgtest -lgtest_main -lpthread -all: warp_bitonic_sort cpu_bitonic_sort test_bitonic_sort +all: warp_bitonic_sort cpu_bitonic_sort warp_bitonic_sort: main.o warp_bitonic_sort.o $(CXX) $^ -o $@ $(LDFLAGS)