diff --git a/README.md b/README.md index 5e2ff4d..0c7366c 100644 --- a/README.md +++ b/README.md @@ -10,10 +10,14 @@ This repo contains the exercises and the tutorials used for Unimore's HPC class ### OpenMP Exercises The exercises related to OpenMP programming model can be found in the folder `openmp`. Here the list of currectly available classes: -- `openmp\lab1`: OpenMP basics: *parallel*, *for-loop*, *sections*, and *tasking*. -- `openmp\lab2`: OpenMP Advanced: *reduction*, *tasking*, *optimizations*. +- `openmp\lab1`: OpenMP basics: *parallel*, *for-loop*, *sections*, and *tasking* +- `openmp\lab2`: OpenMP Advanced: *reduction*, *tasking*, *optimizations* - `openmp\lab3`: OpenMP 4.x+: *Accelerator Model (targeting: Nvidia GP-GPU)* ### CUDA Exercises -- `cuda\lab1`: CUDA Basics. -- `cuda\lab2`: CUDA Memory Model. +- `cuda\lab1`: CUDA Basics +- `cuda\lab2`: CUDA Memory Model +- `cuda\lab3`: CUDA Advanced Host Management + +### (Optional) +- `challenge`: Parallelize the code with everything you learned and submit the result before *21 May 2021* diff --git a/challenge/Makefile b/challenge/Makefile new file mode 100644 index 0000000..668c088 --- /dev/null +++ b/challenge/Makefile @@ -0,0 +1,55 @@ +ifndef CUDA_HOME +CUDA_HOME:=/usr/local/cuda +endif + +ifndef EXERCISE +EXERCISE=TrackColour.cpp +endif + +BUILD_DIR ?= ./build + +NVCC=$(CUDA_HOME)/bin/nvcc +CXX=g++ + +OPT:=-O2 -g +NVOPT:=-Xcompiler -fopenmp -lineinfo -arch=sm_53 --ptxas-options=-v --use_fast_math + +CXXFLAGS:=$(OPT) -I. $(EXT_CXXFLAGS) `pkg-config --cflags --libs opencv4` +LDFLAGS:=-lm -lcudart $(EXT_LDFLAGS) + +NVCFLAGS:=$(CXXFLAGS) $(NVOPT) +NVLDFLAGS:=$(LDFLAGS) -lgomp + +SRCS:= utils.c +OBJS := $(SRCS:%=$(BUILD_DIR)/%.o) $(EXERCISE:%=$(BUILD_DIR)/%.o) +EXE=$(EXERCISE:%=%.exe) + +$(EXE): $(OBJS) + $(MKDIR_P) $(dir $@) + $(NVCC) $(NVCFLAGS) $(OBJS) -o $@ $(NVLDFLAGS) + +$(BUILD_DIR)/%.cu.o: %.cu + $(MKDIR_P) $(dir $@) + $(NVCC) $(NVCFLAGS) -c $< -o $@ + +$(BUILD_DIR)/%.cpp.o: %.cpp + $(MKDIR_P) $(dir $@) + $(CXX) $(CXXFLAGS) -c $< -o $@ + +$(BUILD_DIR)/%.c.o: %.c + $(MKDIR_P) $(dir $@) + $(CXX) $(CXXFLAGS) -c $< -o $@ + +all: $(EXE) + +.PHONY: run profile clean +run: $(EXE) + ./$(EXE) + +profile: $(EXE) + sudo LD_LIBRARY_PATH=$(CUDA_HOME)/lib:/usr/ext/lib:${LD_LIBRARY_PATH} LIBRARY_PATH=/usr/ext/lib:${LIBRARY_PATH} nvprof ./$(EXE) + +clean: + -rm -fr $(BUILD_DIR) *.exe *.out *~ + +MKDIR_P ?= mkdir -p diff --git a/challenge/TrackColour.cpp b/challenge/TrackColour.cpp new file mode 100644 index 0000000..0788fa2 --- /dev/null +++ b/challenge/TrackColour.cpp @@ -0,0 +1,283 @@ +/* + * BSD 2-Clause License + * + * Copyright (c) 2020, Alessandro Capotondi + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * * Redistributions of source code must retain the above copyright notice, this + * list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright notice, + * this list of conditions and the following disclaimer in the documentation + * and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" + * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE + * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL + * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR + * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, + * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/** + * @file colorTracking.cpp + * @author Alessandro Capotondi + * @date 5 May 2020 + * @brief Color Tracking + * + * @see https://dolly.fim.unimore.it/2019/course/view.php?id=152 + */ + +#include +#include +#include +#include +#include "opencv2/imgproc/imgproc_c.h" +#include "opencv2/imgproc/imgproc.hpp" +#include +#include +#include + +using namespace cv; +using namespace std; + +#ifdef CHECK +#include "data/sample_ground_truth.h" +#endif + +extern "C" +{ +#include "utils.h" +} + +void ConversionRgb2Hsv(uint8_t *__restrict__ out, uint8_t *__restrict__ in, int width, int height, int nch) +{ + // Get Raw Image Data pointer + int pixels = width * height * nch; + + //Convert each pixel of input RGB image to HSV + for (int idx = 0; idx < pixels; idx += nch) + { + uint8_t V, S, H; + uint8_t rgb_min, rgb_max, rgb_mm; + + //Read color channels + uint8_t red = in[idx]; + uint8_t green = in[idx + 1]; + uint8_t blue = in[idx + 2]; + + rgb_max = max(max(red, green), blue); + rgb_min = min(min(red, green), blue); + rgb_mm = rgb_max - rgb_min; + + //Value computation + V = rgb_max; + if (V == 0) + H = S = 0; + else + { + //Saturation computation + S = (int)(((long)255 * (long)(rgb_mm)) / (long)V); + if (S == 0) + H = 0; + else + { + //Hue computation + if (rgb_max == red) + H = 0 + 43 * (green - blue) / rgb_mm; + else if (rgb_max == green) + H = 85 + 43 * (blue - red) / rgb_mm; + else + H = 171 + 43 * (red - green) / rgb_mm; + } + } + + //Write HSV + out[idx] = (uint8_t)H; + out[idx + 1] = (uint8_t)S; + out[idx + 2] = (uint8_t)V; + } +} + +void ImgThreashold(uint8_t *__restrict__ out, uint8_t *__restrict__ in, int32_t thLow[3], int32_t thHi[3], int width, int height, int nch) +{ + int pixels = width * height; + uint8_t lb1 = (uint8_t)thLow[0]; + uint8_t lb2 = (uint8_t)thLow[1]; + uint8_t lb3 = (uint8_t)thLow[2]; + uint8_t ub1 = (uint8_t)thHi[0]; + uint8_t ub2 = (uint8_t)thHi[1]; + uint8_t ub3 = (uint8_t)thHi[2]; + + for (int idx = 0; idx < pixels; ++idx) + { + out[idx] = ((in[idx * nch] >= lb1) && (in[idx * nch] <= ub1) && + (in[idx * nch + 1] >= lb2) && (in[idx * nch + 1] <= ub2) && + (in[idx * nch + 2] >= lb3) && (in[idx * nch + 2] <= ub3)) + ? 255 + : 0; + } +} + +void ImgCenterbyMoments(int *y, int *x, uint8_t *__restrict__ in, int width, int height, int nch) +{ + uint64_t m_00 = 0, m_01 = 0, m_10 = 0; + for (int i = 0; i < height; ++i) + { + for (int j = 0; j < width; ++j) + { + if (in[i * width + j] > 0) + { + m_00++; + m_01 += j; + m_10 += i; + } + } + } + *y = m_00 ? (double)m_01 / m_00 : 0; + *x = m_00 ? (double)m_10 / m_00 : 0; +} + +void ImgMerge(uint8_t *__restrict__ out, uint8_t *__restrict__ in1, uint8_t *__restrict__ in2, int width, int height, int nch) +{ + // Get Raw Image Data pointer + int pixels = width * height * nch; + + for (int idx = 0; idx < pixels; ++idx) + { + if(in2[idx]) + out[idx] = in2[idx]; + else + out[idx] = in1[idx]; + } +} + +int main(int argc, char *argv[]) +{ + struct timespec rt[2]; + double wt; + + //Open Video Example + VideoCapture cap("data/sample.avi"); + int width = cap.get(CAP_PROP_FRAME_WIDTH); + int height = cap.get(CAP_PROP_FRAME_HEIGHT); + int nCh = 3; + + // Upper and Lower Color Threasholds + int32_t thHi[3], thLow[3]; + Scalar _thHi, _thLow; + + // Frame Buffers + Mat frameRGB = Mat::zeros(height, width, CV_8UC3); + Mat frameHVS = Mat::zeros(height, width, CV_8UC3); + Mat frameTrack = Mat::zeros(height, width, CV_8UC3); + Mat frameMask = Mat::zeros(height, width, CV_8UC1); + + //Check Video + if (!cap.isOpened()) + { + cout << "[Error] Cannot open the video file" << endl; + return -1; + } + + if (argc > 1) + { + if (argc == 7) + { + thLow[0] = atoi(argv[1]); + thLow[1] = atoi(argv[2]); + thLow[2] = atoi(argv[3]); + thHi[0] = atoi(argv[4]); + thHi[1] = atoi(argv[5]); + thHi[2] = atoi(argv[6]); + } + else + { + cout << "[Error] Invalid arguments: usage ./cTracking [thLow(H S V) thHi(H S V)]" << endl; + return -1; + } + } + else + { + //Default Values + thLow[0] = 160; + thLow[1] = 100; + thLow[2] = 100; + thHi[0] = 180; + thHi[1] = 255; + thHi[2] = 255; + } + _thHi = Scalar(thHi[0], thHi[1], thHi[2]); + + //Print Information + printf("--------------------------------------\n"); + printf("Video Info\n"); + printf("--------------------------------------\n"); + printf("width :\t%d\n", (int)cap.get(CAP_PROP_FRAME_WIDTH)); + printf("height:\t%d\n", (int)cap.get(CAP_PROP_FRAME_HEIGHT)); + printf("fps :\t%d\n", (int)cap.get(CAP_PROP_FPS)); + printf("--------------------------------------\n"); + + int lastX = 0; + int lastY = 0; + int posX = 0; + int posY = 0; +#ifdef CHECK + int check_id = 0; +#endif + int nFrames = 0; + double time_cnt = 0.0; + while (1) + { + bool lastFrame = cap.read(frameRGB); // read a new frame from video + if (!lastFrame) + break; + + clock_gettime(CLOCK_REALTIME, rt + 0); + ConversionRgb2Hsv(frameHVS.ptr(), frameRGB.ptr(), width, height, nCh); + ImgThreashold(frameMask.ptr(), frameHVS.ptr(), thLow, thHi, width, height, nCh); + ImgCenterbyMoments(&posY, &posX, frameMask.ptr(), width, height, nCh); + + // We want to draw a line only if its a valid position + if (lastX > 0 && lastY > 0 && posX > 0 && posY > 0) + line(frameTrack, Point(lastY, lastX), Point(posY, posX), _thHi, 2, CV_8UC3, 0); + + ImgMerge(frameRGB.ptr(), frameRGB.ptr(), frameTrack.ptr(), width, height, nCh); + clock_gettime(CLOCK_REALTIME, rt + 1); + time_cnt+= (rt[1].tv_sec - rt[0].tv_sec) + 1.0e-9 * (rt[1].tv_nsec - rt[0].tv_nsec); + nFrames++; + +#ifdef DISPLAY + // Show frames + imshow("frameRGB", frameRGB); + imshow("frameMask", frameMask); + imshow("frameTrack", frameTrack); + waitKey(1); +#endif + +#ifdef CHECK + assert(ground_truth_YX[check_id++]==posY); + assert(ground_truth_YX[check_id++]==posX); +#endif + lastX = posX; + lastY = posY; + } + + printf("ColorTracking: %d frames, %9.6f s per-frame (%9.6f fps)\n", nFrames, time_cnt/nFrames, 1/(time_cnt/nFrames)); + + //Release Memory - Avoid Memory Leak! + frameRGB.release(); + frameHVS.release(); + frameTrack.release(); + frameMask.release(); + cap.release(); + + return 0; +} diff --git a/challenge/data/sample.avi b/challenge/data/sample.avi new file mode 100644 index 0000000..8e2d79b Binary files /dev/null and b/challenge/data/sample.avi differ diff --git a/challenge/data/sample_ground_truth.h b/challenge/data/sample_ground_truth.h new file mode 100644 index 0000000..65e1db7 --- /dev/null +++ b/challenge/data/sample_ground_truth.h @@ -0,0 +1,595 @@ +short ground_truth_YX[] = { + 950, 522, + 950, 522, + 950, 522, + 950, 522, + 950, 522, + 950, 522, + 950, 522, + 950, 522, + 950, 522, + 950, 522, + 950, 522, + 950, 522, + 950, 522, + 950, 522, + 950, 522, + 950, 522, + 950, 522, + 950, 522, + 950, 522, + 950, 521, + 950, 521, + 950, 521, + 951, 520, + 951, 520, + 951, 520, + 950, 520, + 949, 520, + 949, 520, + 949, 520, + 949, 520, + 949, 519, + 949, 519, + 949, 519, + 950, 518, + 949, 518, + 949, 518, + 949, 518, + 949, 518, + 948, 518, + 948, 518, + 948, 518, + 947, 518, + 946, 517, + 946, 517, + 946, 517, + 945, 516, + 945, 516, + 945, 516, + 944, 514, + 944, 512, + 944, 512, + 943, 508, + 943, 503, + 944, 497, + 945, 491, + 945, 491, + 943, 484, + 942, 478, + 942, 478, + 940, 472, + 938, 466, + 933, 460, + 933, 460, + 928, 454, + 923, 447, + 919, 440, + 919, 440, + 914, 434, + 909, 426, + 905, 419, + 905, 419, + 903, 410, + 903, 400, + 903, 400, + 901, 391, + 899, 380, + 895, 369, + 890, 358, + 890, 358, + 884, 348, + 878, 338, + 878, 338, + 871, 328, + 861, 318, + 861, 318, + 850, 308, + 839, 299, + 829, 291, + 820, 284, + 820, 284, + 811, 277, + 804, 273, + 804, 273, + 799, 271, + 796, 268, + 794, 265, + 794, 265, + 793, 265, + 790, 265, + 790, 265, + 789, 263, + 790, 264, + 791, 263, + 793, 264, + 793, 264, + 796, 263, + 802, 264, + 802, 264, + 808, 261, + 816, 257, + 816, 257, + 826, 256, + 837, 254, + 847, 251, + 857, 251, + 857, 251, + 865, 250, + 873, 250, + 873, 250, + 880, 250, + 888, 249, + 897, 248, + 897, 248, + 905, 248, + 912, 250, + 918, 249, + 918, 249, + 923, 252, + 927, 252, + 928, 253, + 928, 253, + 927, 252, + 924, 251, + 924, 251, + 920, 251, + 915, 251, + 907, 251, + 898, 250, + 898, 250, + 888, 250, + 876, 251, + 876, 251, + 860, 254, + 843, 256, + 843, 256, + 826, 259, + 809, 262, + 792, 265, + 779, 269, + 779, 269, + 767, 270, + 755, 271, + 755, 271, + 744, 273, + 732, 276, + 721, 277, + 721, 277, + 713, 278, + 704, 279, + 704, 279, + 697, 279, + 690, 279, + 685, 280, + 679, 281, + 679, 281, + 674, 283, + 666, 284, + 666, 284, + 660, 288, + 652, 294, + 652, 294, + 644, 299, + 635, 304, + 626, 309, + 617, 312, + 617, 312, + 605, 315, + 593, 319, + 593, 319, + 578, 323, + 562, 329, + 548, 335, + 548, 335, + 535, 343, + 523, 350, + 512, 359, + 512, 359, + 499, 367, + 486, 374, + 475, 381, + 475, 381, + 462, 387, + 448, 393, + 448, 393, + 433, 398, + 421, 403, + 409, 407, + 401, 411, + 401, 411, + 394, 415, + 391, 417, + 391, 417, + 387, 420, + 382, 423, + 375, 427, + 375, 427, + 371, 431, + 366, 436, + 359, 440, + 359, 440, + 353, 443, + 347, 447, + 347, 447, + 340, 450, + 336, 453, + 331, 457, + 331, 457, + 324, 460, + 318, 464, + 318, 464, + 318, 464, + 313, 469, + 309, 472, + 306, 475, + 303, 477, + 300, 481, + 299, 484, + 299, 484, + 304, 489, + 314, 495, + 323, 502, + 323, 502, + 334, 511, + 348, 519, + 366, 526, + 366, 526, + 387, 532, + 409, 536, + 433, 538, + 433, 538, + 457, 539, + 482, 538, + 482, 538, + 508, 536, + 536, 532, + 565, 527, + 565, 527, + 591, 522, + 616, 518, + 637, 516, + 637, 516, + 655, 514, + 673, 513, + 673, 513, + 689, 514, + 706, 515, + 720, 516, + 733, 517, + 733, 517, + 746, 519, + 759, 521, + 759, 521, + 771, 522, + 782, 524, + 792, 527, + 792, 527, + 802, 530, + 811, 533, + 820, 536, + 820, 536, + 830, 538, + 841, 540, + 853, 543, + 853, 543, + 867, 546, + 880, 548, + 880, 548, + 892, 550, + 905, 551, + 919, 551, + 932, 552, + 932, 552, + 945, 552, + 956, 552, + 956, 552, + 967, 550, + 978, 550, + 978, 550, + 989, 550, + 998, 550, + 1004, 551, + 1004, 551, + 1009, 551, + 1015, 552, + 1018, 553, + 1018, 553, + 1017, 552, + 1015, 550, + 1012, 547, + 1012, 547, + 1010, 542, + 1008, 537, + 1008, 537, + 1007, 532, + 1007, 526, + 1005, 520, + 1003, 512, + 1003, 512, + 1002, 504, + 999, 494, + 999, 494, + 998, 483, + 993, 472, + 993, 472, + 985, 460, + 978, 447, + 970, 435, + 962, 423, + 962, 423, + 953, 410, + 940, 397, + 940, 397, + 929, 384, + 918, 371, + 908, 358, + 908, 358, + 901, 346, + 893, 334, + 885, 321, + 885, 321, + 877, 309, + 865, 296, + 852, 284, + 852, 284, + 837, 272, + 823, 261, + 823, 261, + 813, 253, + 803, 246, + 803, 246, + 793, 241, + 787, 238, + 782, 234, + 777, 229, + 777, 229, + 774, 227, + 769, 226, + 769, 226, + 765, 227, + 762, 227, + 760, 225, + 760, 225, + 759, 227, + 759, 226, + 761, 225, + 761, 225, + 762, 225, + 762, 226, + 760, 226, + 760, 226, + 756, 224, + 748, 225, + 748, 225, + 739, 227, + 727, 226, + 713, 228, + 701, 232, + 701, 232, + 688, 235, + 674, 241, + 674, 241, + 663, 244, + 655, 246, + 655, 246, + 646, 248, + 637, 252, + 629, 253, + 620, 256, + 620, 256, + 612, 257, + 604, 259, + 604, 259, + 594, 260, + 583, 263, + 573, 264, + 573, 264, + 562, 265, + 550, 267, + 538, 267, + 538, 267, + 526, 268, + 515, 269, + 504, 269, + 504, 269, + 495, 271, + 487, 272, + 487, 272, + 477, 272, + 466, 273, + 466, 273, + 455, 275, + 444, 277, + 433, 279, + 423, 283, + 423, 283, + 414, 287, + 405, 291, + 405, 291, + 402, 296, + 402, 300, + 406, 305, + 406, 305, + 411, 311, + 420, 318, + 430, 329, + 430, 329, + 439, 340, + 449, 354, + 457, 367, + 457, 367, + 462, 380, + 466, 396, + 466, 396, + 469, 411, + 470, 428, + 472, 444, + 473, 459, + 473, 459, + 476, 473, + 479, 485, + 479, 485, + 484, 496, + 491, 508, + 499, 518, + 499, 518, + 507, 527, + 515, 536, + 520, 545, + 520, 545, + 525, 551, + 529, 555, + 529, 555, + 531, 558, + 534, 559, + 536, 559, + 536, 559, + 538, 559, + 541, 559, + 546, 558, + 546, 558, + 553, 556, + 560, 554, + 568, 553, + 568, 553, + 577, 551, + 586, 550, + 586, 550, + 598, 549, + 611, 547, + 626, 546, + 642, 545, + 642, 545, + 657, 544, + 672, 542, + 672, 542, + 689, 542, + 709, 541, + 728, 540, + 728, 540, + 745, 537, + 761, 535, + 779, 535, + 779, 535, + 798, 536, + 814, 538, + 829, 539, + 829, 539, + 844, 541, + 859, 542, + 859, 542, + 873, 542, + 888, 543, + 888, 543, + 905, 544, + 918, 544, + 931, 545, + 943, 546, + 943, 546, + 954, 545, + 960, 542, + 960, 542, + 959, 538, + 956, 535, + 953, 532, + 953, 532, + 949, 530, + 947, 528, + 946, 526, + 946, 526, + 944, 522, + 944, 517, + 943, 511, + 943, 511, + 943, 502, + 941, 493, + 941, 493, + 937, 482, + 934, 470, + 931, 458, + 926, 445, + 926, 445, + 921, 432, + 913, 420, + 913, 420, + 903, 408, + 893, 396, + 893, 396, + 884, 386, + 872, 376, + 860, 367, + 848, 358, + 848, 358, + 839, 349, + 831, 341, + 831, 341, + 823, 333, + 814, 326, + 807, 321, + 807, 321, + 801, 317, + 796, 315, + 796, 315, + 791, 313, + 789, 315, + 790, 315, + 792, 316, + 792, 316, + 795, 317, + 798, 317, + 798, 317, + 799, 318, + 803, 317, + 803, 317, + 806, 315, + 809, 315, + 814, 314, + 819, 314, + 819, 314, + 825, 313, + 830, 313, + 830, 313, + 832, 313, + 834, 312, + 834, 312, + 834, 312, + 834, 312, + 834, 312, + 834, 312, + 834, 312, + 834, 314, + 836, 314, + 838, 314, + 838, 314, + 838, 317, + 836, 320, + 836, 320, + 835, 321, + 834, 323, + 832, 325, + 832, 327, + 832, 327, + 834, 329, + 836, 331, + 836, 331, + 839, 332, + 843, 334, + 843, 334, + 846, 335, + 849, 335, + 852, 338, + 854, 339, + 854, 339, + 855, 340, + 855, 341, + 855, 341, + 855, 342, + 854, 343, + 855, 342, + 855, 342, + 854, 344, + 852, 345, + 852, 345, + 850, 346, + 849, 347}; \ No newline at end of file diff --git a/challenge/utils.c b/challenge/utils.c new file mode 100644 index 0000000..0ce0dc5 --- /dev/null +++ b/challenge/utils.c @@ -0,0 +1,138 @@ +/* + * BSD 2-Clause License + * + * Copyright (c) 2020, Alessandro Capotondi + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * * Redistributions of source code must retain the above copyright notice, this + * list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright notice, + * this list of conditions and the following disclaimer in the documentation + * and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" + * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE + * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL + * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR + * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, + * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ +/** + * @file utils.c + * @author Alessandro Capotondi + * @date 27 Mar 2020 + * @brief File containing utilities functions for HPC Unimore Class + * + * Utilities for OpenMP lab. + * + * @see http://algo.ing.unimo.it/people/andrea/Didattica/HPC/index.html + */ + +#define _POSIX_C_SOURCE 199309L +#include +#include +#include +#include +#include + +extern "C" { + +#include "utils.h" + +#define MAX_ITERATIONS 100 +static struct timespec timestampA, timestampB; +static unsigned long long statistics[MAX_ITERATIONS]; +static int iterations = 0; + +static unsigned long long __diff_ns(struct timespec start, struct timespec end) +{ + struct timespec temp; + if ((end.tv_nsec - start.tv_nsec) < 0) + { + temp.tv_sec = end.tv_sec - start.tv_sec - 1; + temp.tv_nsec = 1000000000ULL + end.tv_nsec - start.tv_nsec; + } + else + { + temp.tv_sec = end.tv_sec - start.tv_sec; + temp.tv_nsec = end.tv_nsec - start.tv_nsec; + } + + return temp.tv_nsec + temp.tv_sec * 1000000000ULL; +} + +void start_timer() +{ + asm volatile("" :: + : "memory"); + clock_gettime(CLOCK_MONOTONIC_RAW, ×tampA); + asm volatile("" :: + : "memory"); +} + +void stop_timer() +{ + unsigned long long elapsed = 0ULL; + asm volatile("" :: + : "memory"); + clock_gettime(CLOCK_MONOTONIC_RAW, ×tampB); + asm volatile("" :: + : "memory"); +} + +unsigned long long elapsed_ns() +{ + return __diff_ns(timestampA, timestampB); +} + +void start_stats() +{ + start_timer(); +} + +void collect_stats() +{ + assert(iterations < MAX_ITERATIONS); + stop_timer(); + statistics[iterations++] = elapsed_ns(); +} + +void print_stats() +{ + unsigned long long min = ULLONG_MAX; + unsigned long long max = 0LL; + double average = 0.0; + double std_deviation = 0.0; + double sum = 0.0; + + /* Compute the sum of all elements */ + for (int i = 0; i < iterations; i++) + { + if (statistics[i] > max) + max = statistics[i]; + if (statistics[i] < min) + min = statistics[i]; + sum = sum + statistics[i] / 1E6; + } + average = sum / (double)iterations; + + /* Compute variance and standard deviation */ + for (int i = 0; i < iterations; i++) + { + sum = sum + pow((statistics[i] / 1E6 - average), 2); + } + std_deviation = sqrt(sum / (double)iterations); + + printf("AvgTime\tMinTime\tMaxTime\tStdDev\n"); + printf("%.4f ms\t%.4f ms\t%.4f ms\t%.4f\n", (double)average, (double)min / 1E6, (double)max / 1E6, (double)std_deviation); +} + +} diff --git a/challenge/utils.h b/challenge/utils.h new file mode 100644 index 0000000..966281c --- /dev/null +++ b/challenge/utils.h @@ -0,0 +1,142 @@ +/* + * BSD 2-Clause License + * + * Copyright (c) 2020, Alessandro Capotondi + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * * Redistributions of source code must retain the above copyright notice, this + * list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright notice, + * this list of conditions and the following disclaimer in the documentation + * and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" + * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE + * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL + * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR + * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, + * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/** + * @file utils.h + * @author Alessandro Capotondi + * @date 27 Mar 2020 + * @brief File containing utilities functions for HPC Unimore Class + * + * The header define time functions and dummy workload used on the example tests. + * + * @see http://algo.ing.unimo.it/people/andrea/Didattica/HPC/index.html + */ +#ifndef __UTILS_H__ +#define __UTILS_H__ + +#include + +#if defined(VERBOSE) +#define DEBUG_PRINT(x, ...) printf((x), ##__VA_ARGS__) +#else +#define DEBUG_PRINT(x, ...) +#endif + +#if !defined(NTHREADS) +#define NTHREADS (4) +#endif + +extern "C" +{ + +/** + * @brief The function set the timestampA + * + * The function is used to measure elapsed time between two execution points. + * The function start_timer() sets the starting point timestamp, while the function + * stop_timer() sets the termination timestamp. The elapsed time, expressed in nanoseconds, + * between the two points can be retrieved using the function elapsed_ns(). + * + * Example usage: + * @code + * start_timer(); // Point A + * //SOME CODE HERE + * stop_timer(); // Point B + * printf("Elapsed time = %llu ns\n", elapsed_ns())); //Elapsed time between A and B + * //SOME OTHER CODE HERE + * stop_timer(); // Point C + * printf("Elapsed time = %llu ns\n", elapsed_ns())); //Elapsed time between A and C + * @endcode + * + * @return void + * @see start_timer() + * @see stop_timer() + * @see elapsed_ns() + */ + void start_timer(); + +/** + * @brief The function set the second timestamps + * + * The function is used to measure elapsed time between two execution points. + * The function start_timer() sets the starting point timestamp, while the function + * stop_timer() returns the elapsed time, expressed in nanoseconds between the last call + * of start_timer() and the current execution point. + * + * Example usage: + * @code + * start_timer(); // Point A + * //SOME CODE HERE + * stop_timer(); // Point B + * printf("Elapsed time = %llu ns\n", elapsed_ns())); //Elapsed time between A and B + * //SOME OTHER CODE HERE + * stop_timer(); // Point C + * printf("Elapsed time = %llu ns\n", elapsed_ns())); //Elapsed time between A and C + * @endcode + * + * @return void + * @see start_timer() + * @see stop_timer() + * @see elapsed_ns() + */ + void stop_timer(); + +/** + * @brief Elapsed nano seconds between start_timer() and stop_timer(). + * + * @return Elapsed nano seconds + * @see start_timer() + * @see stop_timer() + */ + unsigned long long elapsed_ns(); + +/** + * @brief The function init the starting point of stat measurement. + * + * The function is similar to start_timer(). + * + * @return void + * @see start_timer + */ + void start_stats(); + +/** + * @brief The function collects the elapsed time between the current exeuction point and the + * last call of start_stats(). + * + * @return void + */ + void collect_stats(); + +/** + * @brief The function display the collected statistics. + * @return void + */ + void print_stats(); +} +#endif /*__UTILS_H__*/ diff --git a/cuda/lab3/.solutions/gemm-v1.cu b/cuda/lab3/.solutions/gemm-v1.cu new file mode 100644 index 0000000..7f8563e --- /dev/null +++ b/cuda/lab3/.solutions/gemm-v1.cu @@ -0,0 +1,232 @@ +/* + * BSD 2-Clause License + * + * Copyright (c) 2020, Alessandro Capotondi + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * * Redistributions of source code must retain the above copyright notice, this + * list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright notice, + * this list of conditions and the following disclaimer in the documentation + * and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" + * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE + * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL + * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR + * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, + * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/** + * @file gemm.cu + * @author Alessandro Capotondi + * @date 12 May 2020 + * @brief GEMM Kernel + * + * @see https://dolly.fim.unimore.it/2019/course/view.php?id=152 + */ + +#include +#include +#include +#include + +#define gpuErrchk(ans) \ + { \ + gpuAssert((ans), __FILE__, __LINE__); \ + } +static inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) +{ + if (code != cudaSuccess) + { + fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); + if (abort) + exit(code); + } +} + +extern "C" +{ +#include "utils.h" +} + +#define TWO02 (1 << 2) +#define TWO04 (1 << 4) +#define TWO08 (1 << 8) + +#ifndef N +#define N (1 << 10) +#endif +#ifndef TILE_W +#define TILE_W 128 +#endif +#ifndef BLOCK_SIZE +#define BLOCK_SIZE 32 +#endif + +#define SM 64 +static void reorder(float *__restrict__ a, float *__restrict__ b, int n) +{ + for (int i = 0; i < SM; i++) + for (int j = 0; j < SM; j++) + b[i * SM + j] = a[i * n + j]; +} + +static void mm(float *__restrict__ a, float *__restrict__ b, float *__restrict__ c, int n) +{ + for (int i = 0; i < SM; i++) + { + for (int k = 0; k < SM; k++) + { + for (int j = 0; j < SM; j++) + { + c[i * n + j] += a[i * n + k] * b[k * SM + j]; + } + } + } +} +void gemm_host(float *a, float *b, float *c, int n) +{ + int bk = n / SM; +#pragma omp parallel for collapse(3) + for (int i = 0; i < bk; i++) + { + for (int j = 0; j < bk; j++) + { + for (int k = 0; k < bk; k++) + { + float b2[SM * SM]; + reorder(&b[SM * (k * n + j)], b2, n); + mm(&a[SM * (i * n + k)], b2, &c[SM * (i * n + j)], n); + } + } + } +} +__global__ void gemm(float *__restrict__ a, float *__restrict__ b, float *__restrict__ c, int n) +{ + __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; + __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; + + int ib = blockIdx.y; + int jb = blockIdx.x; + + int it = threadIdx.y; + int jt = threadIdx.x; + + int a_offset, b_offset, c_offset; + + float Cvalue = 0.0f; + for (int kb = 0; kb < (n / BLOCK_SIZE); ++kb) + { + a_offset = ib * n * BLOCK_SIZE + kb * BLOCK_SIZE; + b_offset = kb * n * BLOCK_SIZE + jb * BLOCK_SIZE; + As[it][jt] = a[a_offset + it * n + jt]; + Bs[it][jt] = b[b_offset + it * n + jt]; + __syncthreads(); + + for (int k = 0; k < BLOCK_SIZE; ++k) + Cvalue += As[it][k] * Bs[k][jt]; + + __syncthreads(); + } + + c_offset = ib * n * BLOCK_SIZE + jb * BLOCK_SIZE; + c[c_offset + it * n + jt] = Cvalue; +} + +int main(int argc, char *argv[]) +{ + int n = N, iret = 0; + float *a, *b, *c, *g; + struct timespec rt[2]; + double wt; // walltime + + if (argc > 1) + n = atoi(argv[1]); + + //TODO Update malloc to cudaMallocManaged + gpuErrchk(cudaMallocHost((void **)&a, sizeof(float) * n *n)); + gpuErrchk(cudaMallocHost((void **)&b, sizeof(float) * n *n)); + gpuErrchk(cudaMallocHost((void **)&c, sizeof(float) * n *n)); + if (NULL == (g = (float *)malloc(sizeof(*g) * n * n))) + { + printf("error: memory allocation for 'z'\n"); + iret = -1; + } + + if (0 != iret) + { + gpuErrchk(cudaFreeHost(a)); + gpuErrchk(cudaFreeHost(b)); + gpuErrchk(cudaFreeHost(c)); + free(g); + exit(EXIT_FAILURE); + } + + //Init Data + int _b = rand() % TWO04; + int _c = rand() % TWO08; + +#pragma omp parallel for + for (int i = 0; i < n * n; i++) + { + a[i] = _b / (float)TWO02; + b[i] = _c / (float)TWO04; + c[i] = g[i] = 0.0; + } + + clock_gettime(CLOCK_REALTIME, rt + 0); + gemm_host(a, b, g, n); + clock_gettime(CLOCK_REALTIME, rt + 1); + wt = (rt[1].tv_sec - rt[0].tv_sec) + 1.0e-9 * (rt[1].tv_nsec - rt[0].tv_nsec); + printf("GEMM (Host) : %9.3f sec %9.1f GFLOPS\n", wt, 2.0 * n * n * n / (1.0e9 * wt)); + + //TODO Remove if unecessary + float *d_a, *d_b, *d_c; + gpuErrchk(cudaMalloc((void **)&d_a, sizeof(float) * n * n)); + gpuErrchk(cudaMalloc((void **)&d_b, sizeof(float) * n * n)); + gpuErrchk(cudaMalloc((void **)&d_c, sizeof(float) * n * n)); + + clock_gettime(CLOCK_REALTIME, rt + 0); + //TODO Remove if unecessary + gpuErrchk(cudaMemcpy(d_a, a, sizeof(float) * n * n, cudaMemcpyHostToDevice)); + gpuErrchk(cudaMemcpy(d_b, b, sizeof(float) * n * n, cudaMemcpyHostToDevice)); + dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); + dim3 dimGrid((n + (BLOCK_SIZE)-1) / (BLOCK_SIZE), (n + (BLOCK_SIZE)-1) / (BLOCK_SIZE)); + gemm<<>>(d_a, d_b, d_c, n); + gpuErrchk(cudaPeekAtLastError()); + //TODO Remove if unecessary + gpuErrchk(cudaMemcpy(c, d_c, sizeof(float) * n * n, cudaMemcpyDeviceToHost)); + clock_gettime(CLOCK_REALTIME, rt + 1); + wt = (rt[1].tv_sec - rt[0].tv_sec) + 1.0e-9 * (rt[1].tv_nsec - rt[0].tv_nsec); + printf("GEMM-v1 (GPU): %9.3f sec %9.1f GFLOPS\n", wt, 2.0 * n * n * n / (1.0e9 * wt)); + + for (int i = 0; i < n * n; i++) + { + iret = *(int *)(g + i) ^ *(int *)(c + i); + assert(iret == 0); + } + + //TODO Update cudaFreeHost or cudaFree (if necessary) + gpuErrchk(cudaFreeHost(a)); + gpuErrchk(cudaFreeHost(b)); + gpuErrchk(cudaFreeHost(c)); + free(g); + //TODO Remove if unecessary + gpuErrchk(cudaFree(d_a)); + //TODO Remove if unecessary + gpuErrchk(cudaFree(d_b)); + //TODO Remove if unecessary + gpuErrchk(cudaFree(d_c)); + + return 0; +} diff --git a/cuda/lab3/.solutions/gemm-v2.cu b/cuda/lab3/.solutions/gemm-v2.cu new file mode 100644 index 0000000..b536618 --- /dev/null +++ b/cuda/lab3/.solutions/gemm-v2.cu @@ -0,0 +1,214 @@ +/* + * BSD 2-Clause License + * + * Copyright (c) 2020, Alessandro Capotondi + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * * Redistributions of source code must retain the above copyright notice, this + * list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright notice, + * this list of conditions and the following disclaimer in the documentation + * and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" + * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE + * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL + * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR + * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, + * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/** + * @file gemm.cu + * @author Alessandro Capotondi + * @date 12 May 2020 + * @brief GEMM Kernel + * + * @see https://dolly.fim.unimore.it/2019/course/view.php?id=152 + */ + +#include +#include +#include +#include + +#define gpuErrchk(ans) \ + { \ + gpuAssert((ans), __FILE__, __LINE__); \ + } +static inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) +{ + if (code != cudaSuccess) + { + fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); + if (abort) + exit(code); + } +} + +extern "C" +{ +#include "utils.h" +} + +#define TWO02 (1 << 2) +#define TWO04 (1 << 4) +#define TWO08 (1 << 8) + +#ifndef N +#define N (1 << 10) +#endif +#ifndef TILE_W +#define TILE_W 128 +#endif +#ifndef BLOCK_SIZE +#define BLOCK_SIZE 32 +#endif + +#define SM 64 +static void reorder(float *__restrict__ a, float *__restrict__ b, int n) +{ + for (int i = 0; i < SM; i++) + for (int j = 0; j < SM; j++) + b[i * SM + j] = a[i * n + j]; +} + +static void mm(float *__restrict__ a, float *__restrict__ b, float *__restrict__ c, int n) +{ + for (int i = 0; i < SM; i++) + { + for (int k = 0; k < SM; k++) + { + for (int j = 0; j < SM; j++) + { + c[i * n + j] += a[i * n + k] * b[k * SM + j]; + } + } + } +} +void gemm_host(float *a, float *b, float *c, int n) +{ + int bk = n / SM; +#pragma omp parallel for collapse(3) + for (int i = 0; i < bk; i++) + { + for (int j = 0; j < bk; j++) + { + for (int k = 0; k < bk; k++) + { + float b2[SM * SM]; + reorder(&b[SM * (k * n + j)], b2, n); + mm(&a[SM * (i * n + k)], b2, &c[SM * (i * n + j)], n); + } + } + } +} +__global__ void gemm(float *__restrict__ a, float *__restrict__ b, float *__restrict__ c, int n) +{ + __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; + __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; + + int ib = blockIdx.y; + int jb = blockIdx.x; + + int it = threadIdx.y; + int jt = threadIdx.x; + + int a_offset, b_offset, c_offset; + + float Cvalue = 0.0f; + for (int kb = 0; kb < (n / BLOCK_SIZE); ++kb) + { + a_offset = ib * n * BLOCK_SIZE + kb * BLOCK_SIZE; + b_offset = kb * n * BLOCK_SIZE + jb * BLOCK_SIZE; + As[it][jt] = a[a_offset + it * n + jt]; + Bs[it][jt] = b[b_offset + it * n + jt]; + __syncthreads(); + + for (int k = 0; k < BLOCK_SIZE; ++k) + Cvalue += As[it][k] * Bs[k][jt]; + + __syncthreads(); + } + + c_offset = ib * n * BLOCK_SIZE + jb * BLOCK_SIZE; + c[c_offset + it * n + jt] = Cvalue; +} + +int main(int argc, char *argv[]) +{ + int n = N, iret = 0; + float *a, *b, *c, *g; + struct timespec rt[2]; + double wt; // walltime + + if (argc > 1) + n = atoi(argv[1]); + + //TODO Update malloc to cudaMallocManaged + gpuErrchk(cudaMallocManaged((void **)&a, sizeof(float) * n *n)); + gpuErrchk(cudaMallocManaged((void **)&b, sizeof(float) * n *n)); + gpuErrchk(cudaMallocManaged((void **)&c, sizeof(float) * n *n)); + if (NULL == (g = (float *)malloc(sizeof(*g) * n * n))) + { + printf("error: memory allocation for 'z'\n"); + iret = -1; + } + + if (0 != iret) + { + gpuErrchk(cudaFree(a)); + gpuErrchk(cudaFree(b)); + gpuErrchk(cudaFree(c)); + free(g); + exit(EXIT_FAILURE); + } + + //Init Data + int _b = rand() % TWO04; + int _c = rand() % TWO08; + +#pragma omp parallel for + for (int i = 0; i < n * n; i++) + { + a[i] = _b / (float)TWO02; + b[i] = _c / (float)TWO04; + c[i] = g[i] = 0.0; + } + + clock_gettime(CLOCK_REALTIME, rt + 0); + gemm_host(a, b, g, n); + clock_gettime(CLOCK_REALTIME, rt + 1); + wt = (rt[1].tv_sec - rt[0].tv_sec) + 1.0e-9 * (rt[1].tv_nsec - rt[0].tv_nsec); + printf("GEMM (Host) : %9.3f sec %9.1f GFLOPS\n", wt, 2.0 * n * n * n / (1.0e9 * wt)); + + clock_gettime(CLOCK_REALTIME, rt + 0); + dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); + dim3 dimGrid((n + (BLOCK_SIZE)-1) / (BLOCK_SIZE), (n + (BLOCK_SIZE)-1) / (BLOCK_SIZE)); + gemm<<>>(a, b, c, n); + gpuErrchk(cudaPeekAtLastError()); + cudaDeviceSynchronize(); + clock_gettime(CLOCK_REALTIME, rt + 1); + wt = (rt[1].tv_sec - rt[0].tv_sec) + 1.0e-9 * (rt[1].tv_nsec - rt[0].tv_nsec); + printf("GEMM-v1 (GPU): %9.3f sec %9.1f GFLOPS\n", wt, 2.0 * n * n * n / (1.0e9 * wt)); + + for (int i = 0; i < n * n; i++) + { + iret = *(int *)(g + i) ^ *(int *)(c + i); + assert(iret == 0); + } + + gpuErrchk(cudaFree(a)); + gpuErrchk(cudaFree(b)); + gpuErrchk(cudaFree(c)); + free(g); + return 0; +} diff --git a/cuda/lab3/.solutions/saxpy-v1.cu b/cuda/lab3/.solutions/saxpy-v1.cu new file mode 100644 index 0000000..3ef821d --- /dev/null +++ b/cuda/lab3/.solutions/saxpy-v1.cu @@ -0,0 +1,175 @@ +/* + * BSD 2-Clause License + * + * Copyright (c) 2020, Alessandro Capotondi + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * * Redistributions of source code must retain the above copyright notice, this + * list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright notice, + * this list of conditions and the following disclaimer in the documentation + * and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" + * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE + * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL + * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR + * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, + * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/** + * @file saxpy.c + * @author Alessandro Capotondi + * @date 12 May 2020 + * @brief Saxpy + * + * @see https://dolly.fim.unimore.it/2019/course/view.php?id=152 + */ + +#include +#include +#include +#include +#include +#include +#include + +#define gpuErrchk(ans) \ + { \ + gpuAssert((ans), __FILE__, __LINE__); \ + } +static inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) +{ + if (code != cudaSuccess) + { + fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); + if (abort) + exit(code); + } +} + +extern "C" +{ +#include "utils.h" +} + +#define TWO02 (1 << 2) +#define TWO04 (1 << 4) +#define TWO08 (1 << 8) +#ifndef N +#define N (1 << 27) +#endif + +#ifndef BLOCK_SIZE +#define BLOCK_SIZE (512) +#endif + +/* + *SAXPY (host implementation) + * y := a * x + y + */ +void host_saxpy(float *__restrict__ y, float a, float *__restrict__ x, int n) +{ +#pragma omp parallel for simd schedule(simd \ + : static) + for (int i = 0; i < n; i++) + { + y[i] = a * x[i] + y[i]; + } +} + +__global__ void gpu_saxpy(float *__restrict__ y, float a, float *__restrict__ x, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + y[i] = a * x[i] + y[i]; +} + +int main(int argc, const char **argv) +{ + int iret = 0; + int n = N; + float *h_x, *d_x; + float *h_y, *d_y; + float *h_z; + float a = 101.0f / TWO02, + b, c; + + if (argc > 1) + n = atoi(argv[1]); + + //TODO Update malloc to cudaMallocHost or cudaMallocManaged (if necessary) + gpuErrchk(cudaMallocHost((void **)&h_x, sizeof(float) * n)); + //TODO Update malloc to cudaMallocHost or cudaMallocManaged (if necessary) + gpuErrchk(cudaMallocHost((void **)&h_y, sizeof(float) * n)); + if (NULL == (h_z = (float *)malloc(sizeof(float) * n))) + { + printf("error: memory allocation for 'z'\n"); + iret = -1; + } + if (0 != iret) + { + //TODO Update cudaFreeHost or cudaFree (if necessary) + cudaFreeHost(h_x); + //TODO Update cudaFreeHost or cudaFree (if necessary) + cudaFreeHost(h_y); + free(h_z); + exit(EXIT_FAILURE); + } + + //Init Data + b = rand() % TWO04; + c = rand() % TWO08; + for (int i = 0; i < n; i++) + { + h_x[i] = b / (float)TWO02; + h_y[i] = h_z[i] = c / (float)TWO04; + } + + //TODO Remove if unecessary + gpuErrchk(cudaMalloc((void **)&d_x, sizeof(float) * n)); + gpuErrchk(cudaMalloc((void **)&d_y, sizeof(float) * n)); + + start_timer(); + //TODO Remove if unecessary + gpuErrchk(cudaMemcpy(d_x, h_x, sizeof(float) * n, cudaMemcpyHostToDevice)); + gpuErrchk(cudaMemcpy(d_y, h_y, sizeof(float) * n, cudaMemcpyHostToDevice)); + gpu_saxpy<<<((n + BLOCK_SIZE - 1) / BLOCK_SIZE), BLOCK_SIZE>>>(d_y, a, d_x, n); + gpuErrchk(cudaPeekAtLastError()); + //TODO Remove if unecessary + gpuErrchk(cudaMemcpy(h_y, d_y, sizeof(float) * n, cudaMemcpyDeviceToHost)); + stop_timer(); + printf("saxpy (GPU): %9.3f sec %9.1f GFLOPS\n", elapsed_ns() / 1.0e9, 2 * n / ((float)elapsed_ns())); + + //Check Matematical Consistency + start_timer(); + host_saxpy(h_z, a, h_x, n); + stop_timer(); + printf("saxpy (Host): %9.3f sec %9.1f GFLOPS\n", elapsed_ns() / 1.0e9, 2 * n / ((float)elapsed_ns())); + for (int i = 0; i < n; ++i) + { + iret = *(int *)(h_y + i) ^ *(int *)(h_z + i); + assert(iret == 0); + } + + //TODO Update cudaFreeHost or cudaFree (if necessary) + gpuErrchk(cudaFreeHost((void *)h_x)); + gpuErrchk(cudaFree(d_x)); + //TODO Update cudaFreeHost or cudaFree (if necessary) + gpuErrchk(cudaFreeHost((void *)h_y)); + gpuErrchk(cudaFree(d_y)); + free(h_z); + + // CUDA exit -- needed to flush printf write buffer + cudaDeviceReset(); + return 0; +} diff --git a/cuda/lab3/.solutions/saxpy-v2.cu b/cuda/lab3/.solutions/saxpy-v2.cu new file mode 100644 index 0000000..d9bbaf7 --- /dev/null +++ b/cuda/lab3/.solutions/saxpy-v2.cu @@ -0,0 +1,167 @@ +/* + * BSD 2-Clause License + * + * Copyright (c) 2020, Alessandro Capotondi + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * * Redistributions of source code must retain the above copyright notice, this + * list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright notice, + * this list of conditions and the following disclaimer in the documentation + * and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" + * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE + * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL + * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR + * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, + * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/** + * @file saxpy.c + * @author Alessandro Capotondi + * @date 12 May 2020 + * @brief Saxpy + * + * @see https://dolly.fim.unimore.it/2019/course/view.php?id=152 + */ + +#include +#include +#include +#include +#include +#include +#include + +#define gpuErrchk(ans) \ + { \ + gpuAssert((ans), __FILE__, __LINE__); \ + } +static inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) +{ + if (code != cudaSuccess) + { + fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); + if (abort) + exit(code); + } +} + +extern "C" +{ +#include "utils.h" +} + +#define TWO02 (1 << 2) +#define TWO04 (1 << 4) +#define TWO08 (1 << 8) +#ifndef N +#define N (1 << 27) +#endif + +#ifndef BLOCK_SIZE +#define BLOCK_SIZE (512) +#endif + +/* + *SAXPY (host implementation) + * y := a * x + y + */ +void host_saxpy(float * __restrict__ y, float a, float * __restrict__ x, int n) +{ +#pragma omp parallel for simd schedule(simd: static) + for (int i = 0; i < n; i++) + { + y[i] = a * x[i] + y[i]; + } +} + +__global__ void gpu_saxpy(float * __restrict__ y, float a, float * __restrict__ x, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + y[i] = a * x[i] + y[i]; +} + +int main(int argc, const char **argv) +{ + int iret = 0; + int n = N; + float *h_x; + float *h_y; + float *h_z; + float a = 101.0f / TWO02, + b, c; + + if (argc > 1) + n = atoi(argv[1]); + + //TODO Update malloc to cudaMallocHost or cudaMallocManaged (if necessary) + gpuErrchk(cudaMallocManaged((void **)&h_x, sizeof(float) * n)); + //TODO Update malloc to cudaMallocHost or cudaMallocManaged (if necessary) + gpuErrchk(cudaMallocManaged((void **)&h_y, sizeof(float) * n)); + if (NULL == (h_z = (float *)malloc(sizeof(float) * n))) + { + printf("error: memory allocation for 'z'\n"); + iret = -1; + } + if (0 != iret) + { + //TODO Update cudaFreeHost or cudaFree (if necessary) + gpuErrchk(cudaFree(h_x)); + //TODO Update cudaFreeHost or cudaFree (if necessary) + gpuErrchk(cudaFree(h_y)); + free(h_z); + exit(EXIT_FAILURE); + } + + //Init Data + b = rand() % TWO04; + c = rand() % TWO08; + for (int i = 0; i < n; i++) + { + h_x[i] = b / (float)TWO02; + h_y[i] = h_z[i] = c / (float)TWO04; + } + + //TODO Remove if unecessary + + start_timer(); + //TODO Remove if unecessary + gpu_saxpy<<<((n + BLOCK_SIZE - 1) / BLOCK_SIZE), BLOCK_SIZE>>>(h_y, a, h_x, n); + gpuErrchk(cudaPeekAtLastError()); + //TODO Remove if unecessary + cudaDeviceSynchronize(); + stop_timer(); + printf("saxpy (GPU): %9.3f sec %9.1f GFLOPS\n", elapsed_ns() / 1.0e9, 2 * n / ((float) elapsed_ns())); + + //Check Matematical Consistency + start_timer(); + host_saxpy(h_z, a, h_x, n); + stop_timer(); + printf("saxpy (Host): %9.3f sec %9.1f GFLOPS\n", elapsed_ns() / 1.0e9, 2 * n / ((float) elapsed_ns())); + for (int i = 0; i < n; ++i) + { + iret = *(int *)(h_y + i) ^ *(int *)(h_z + i); + assert(iret == 0); + } + + //TODO Update cudaFreeHost or cudaFree (if necessary) + gpuErrchk(cudaFree(h_x)); + gpuErrchk(cudaFree(h_y)); + free(h_z); + + // CUDA exit -- needed to flush printf write buffer + cudaDeviceReset(); + return 0; +} diff --git a/cuda/lab3/.solutions/saxpy-v3.cu b/cuda/lab3/.solutions/saxpy-v3.cu new file mode 100644 index 0000000..6da8ed0 --- /dev/null +++ b/cuda/lab3/.solutions/saxpy-v3.cu @@ -0,0 +1,197 @@ +/* + * BSD 2-Clause License + * + * Copyright (c) 2020, Alessandro Capotondi + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * * Redistributions of source code must retain the above copyright notice, this + * list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright notice, + * this list of conditions and the following disclaimer in the documentation + * and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" + * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE + * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL + * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR + * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, + * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/** + * @file saxpy.c + * @author Alessandro Capotondi + * @date 12 May 2020 + * @brief Saxpy + * + * @see https://dolly.fim.unimore.it/2019/course/view.php?id=152 + */ + +#include +#include +#include +#include +#include +#include +#include + +#define gpuErrchk(ans) \ + { \ + gpuAssert((ans), __FILE__, __LINE__); \ + } +static inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) +{ + if (code != cudaSuccess) + { + fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); + if (abort) + exit(code); + } +} + +extern "C" +{ +#include "utils.h" +} + +#define TWO02 (1 << 2) +#define TWO04 (1 << 4) +#define TWO08 (1 << 8) +#ifndef N +#define N (1 << 27) +#endif + +#ifndef BLOCK_SIZE +#define BLOCK_SIZE (512) +#endif + +/* + *SAXPY (host implementation) + * y := a * x + y + */ +void host_saxpy(float * __restrict__ y, float a, float * __restrict__ x, int n) +{ +#pragma omp parallel for simd schedule(simd: static) + for (int i = 0; i < n; i++) + { + y[i] = a * x[i] + y[i]; + } +} + +__global__ void gpu_saxpy(float * __restrict__ y, float a, float * __restrict__ x, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + y[i] = a * x[i] + y[i]; +} + +int main(int argc, const char **argv) +{ + int iret = 0; + int n = N; + float *h_x, *d_x; + float *h_y, *d_y; + float *h_z; + float a = 101.0f / TWO02, + b, c; + + if (argc > 1) + n = atoi(argv[1]); + + if (NULL == (h_x = (float *)malloc(sizeof(float) * n))) + { + printf("error: memory allocation for 'x'\n"); + iret = -1; + } + if (NULL == (h_y = (float *)malloc(sizeof(float) * n))) + { + printf("error: memory allocation for 'y'\n"); + iret = -1; + } + if (NULL == (h_z = (float *)malloc(sizeof(float) * n))) + { + printf("error: memory allocation for 'z'\n"); + iret = -1; + } + if (0 != iret) + { + free(h_x); + free(h_y); + free(h_z); + exit(EXIT_FAILURE); + } + + //Init Data + b = rand() % TWO04; + c = rand() % TWO08; + for (int i = 0; i < n; i++) + { + h_x[i] = b / (float)TWO02; + h_y[i] = h_z[i] = c / (float)TWO04; + } + + //CUDA Buffer Allocation + gpuErrchk(cudaMalloc((void **)&d_x, sizeof(float) * n)); + gpuErrchk(cudaMalloc((void **)&d_y, sizeof(float) * n)); + + start_timer(); + int TILE = n / 8; + //TODO Copy the first Tile (i=0) + gpuErrchk(cudaMemcpyAsync(&d_x[0], &h_x[0], sizeof(float) * TILE, cudaMemcpyHostToDevice)); + gpuErrchk(cudaMemcpyAsync(&d_y[0], &h_y[0], sizeof(float) * TILE, cudaMemcpyHostToDevice)); + //TODO Loop over the Tiles + for (int i = 0; i < n; i += TILE) + { + //TODO Wait Tile i + cudaDeviceSynchronize(); + + //TODO Copy the out tile i-1 + if(i>0) + gpuErrchk(cudaMemcpyAsync(&h_y[i-TILE], &d_y[i-TILE], sizeof(float) * TILE, cudaMemcpyDeviceToHost)); + + //TODO Launch Kernel over tile i + gpu_saxpy<<<((TILE + BLOCK_SIZE - 1) / BLOCK_SIZE), BLOCK_SIZE>>>(&d_y[i], a, &d_x[i], TILE); + + //TODO Copy the in tile i+=TILE + if(i+TILE < n){ + gpuErrchk(cudaMemcpyAsync(&d_x[i+TILE], &h_x[i+TILE], sizeof(float) * TILE, cudaMemcpyHostToDevice)); + gpuErrchk(cudaMemcpyAsync(&d_y[i+TILE], &h_y[i+TILE], sizeof(float) * TILE, cudaMemcpyHostToDevice)); + } + } + //TODO Copy out the last tile n-TILE + gpuErrchk(cudaMemcpyAsync(&h_y[n-TILE], &d_y[n-TILE], sizeof(float) * TILE, cudaMemcpyDeviceToHost)); + //TODO Wait last tile + cudaDeviceSynchronize(); + stop_timer(); + printf("saxpy (GPU): %9.3f sec %9.1f GFLOPS\n", elapsed_ns() / 1.0e9, 2 * n / ((float) elapsed_ns())); + + //Check Matematical Consistency + start_timer(); + host_saxpy(h_z, a, h_x, n); + stop_timer(); + printf("saxpy (Host): %9.3f sec %9.1f GFLOPS\n", elapsed_ns() / 1.0e9, 2 * n / ((float) elapsed_ns())); + for (int i = 0; i < n; ++i) + { + iret = *(int *)(h_y + i) ^ *(int *)(h_z + i); + assert(iret == 0); + } + + //CUDA Buffer Allocation + free(h_x); + gpuErrchk(cudaFree(d_x)); + free(h_y); + gpuErrchk(cudaFree(d_y)); + free(h_z); + + // CUDA exit -- needed to flush printf write buffer + cudaDeviceReset(); + return 0; +} diff --git a/cuda/lab3/.solutions/saxpy-v4.cu b/cuda/lab3/.solutions/saxpy-v4.cu new file mode 100644 index 0000000..1026220 --- /dev/null +++ b/cuda/lab3/.solutions/saxpy-v4.cu @@ -0,0 +1,197 @@ +/* + * BSD 2-Clause License + * + * Copyright (c) 2020, Alessandro Capotondi + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * * Redistributions of source code must retain the above copyright notice, this + * list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright notice, + * this list of conditions and the following disclaimer in the documentation + * and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" + * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE + * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL + * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR + * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, + * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/** + * @file saxpy.c + * @author Alessandro Capotondi + * @date 12 May 2020 + * @brief Saxpy + * + * @see https://dolly.fim.unimore.it/2019/course/view.php?id=152 + */ + +#include +#include +#include +#include +#include +#include +#include + +#define gpuErrchk(ans) \ + { \ + gpuAssert((ans), __FILE__, __LINE__); \ + } +static inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) +{ + if (code != cudaSuccess) + { + fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); + if (abort) + exit(code); + } +} + +extern "C" +{ +#include "utils.h" +} + +#define TWO02 (1 << 2) +#define TWO04 (1 << 4) +#define TWO08 (1 << 8) +#ifndef N +#define N (1 << 27) +#endif + +#ifndef BLOCK_SIZE +#define BLOCK_SIZE (512) +#endif + +#ifndef N_STREAMS +#define N_STREAMS (16) +#endif + +/* + *SAXPY (host implementation) + * y := a * x + y + */ +void host_saxpy(float * __restrict__ y, float a, float * __restrict__ x, int n) +{ +#pragma omp parallel for simd schedule(simd: static) + for (int i = 0; i < n; i++) + { + y[i] = a * x[i] + y[i]; + } +} + +__global__ void gpu_saxpy(float * __restrict__ y, float a, float * __restrict__ x, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + y[i] = a * x[i] + y[i]; +} + +int main(int argc, const char **argv) +{ + int iret = 0; + int n = N; + float *h_x, *d_x; + float *h_y, *d_y; + float *h_z; + float a = 101.0f / TWO02, + b, c; + + if (argc > 1) + n = atoi(argv[1]); + + if (NULL == (h_x = (float *)malloc(sizeof(float) * n))) + { + printf("error: memory allocation for 'x'\n"); + iret = -1; + } + if (NULL == (h_y = (float *)malloc(sizeof(float) * n))) + { + printf("error: memory allocation for 'y'\n"); + iret = -1; + } + if (NULL == (h_z = (float *)malloc(sizeof(float) * n))) + { + printf("error: memory allocation for 'z'\n"); + iret = -1; + } + if (0 != iret) + { + free(h_x); + free(h_y); + free(h_z); + exit(EXIT_FAILURE); + } + + //Init Data + b = rand() % TWO04; + c = rand() % TWO08; + for (int i = 0; i < n; i++) + { + h_x[i] = b / (float)TWO02; + h_y[i] = h_z[i] = c / (float)TWO04; + } + + //CUDA Buffer Allocation + gpuErrchk(cudaMalloc((void **)&d_x, sizeof(float) * n)); + gpuErrchk(cudaMalloc((void **)&d_y, sizeof(float) * n)); + + start_timer(); + int TILE = n / N_STREAMS; + cudaStream_t stream[N_STREAMS]; + for(int i = 0; i < N_STREAMS; i++) + cudaStreamCreate(&stream[i]); + + //TODO Loop over the Tiles + for (int i = 0; i < n; i += TILE) + { + //TODO Copy in Tile i (stream i) + gpuErrchk(cudaMemcpyAsync(&d_x[i], &h_x[i], sizeof(float) * TILE, cudaMemcpyHostToDevice, stream[i/TILE])); + gpuErrchk(cudaMemcpyAsync(&d_y[i], &h_y[i], sizeof(float) * TILE, cudaMemcpyHostToDevice, stream[i/TILE])); + + //TODO Kernel Tile i (stream i) + gpu_saxpy<<<((TILE + BLOCK_SIZE - 1) / BLOCK_SIZE), BLOCK_SIZE,0,stream[i/TILE]>>>(&d_y[i], a, &d_x[i], TILE); + + //TODO Copy out Tile i (stream i) + gpuErrchk(cudaMemcpyAsync(&h_y[i], &d_y[i], sizeof(float) * TILE, cudaMemcpyDeviceToHost,stream[i/TILE])); + } + //TODO Wait all the streams... + cudaDeviceSynchronize(); + stop_timer(); + printf("saxpy (GPU): %9.3f sec %9.1f GFLOPS\n", elapsed_ns() / 1.0e9, 2 * n / ((float) elapsed_ns())); + + //Check Matematical Consistency + start_timer(); + host_saxpy(h_z, a, h_x, n); + stop_timer(); + printf("saxpy (Host): %9.3f sec %9.1f GFLOPS\n", elapsed_ns() / 1.0e9, 2 * n / ((float) elapsed_ns())); + for (int i = 0; i < n; ++i) + { + iret = *(int *)(h_y + i) ^ *(int *)(h_z + i); + assert(iret == 0); + } + + free(h_x); + gpuErrchk(cudaFree(d_x)); + free(h_y); + gpuErrchk(cudaFree(d_y)); + free(h_z); + + for (int i=0; i +#include +#include +#include +#include +#include +#include + +#define gpuErrchk(ans) \ + { \ + gpuAssert((ans), __FILE__, __LINE__); \ + } +static inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) +{ + if (code != cudaSuccess) + { + fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); + if (abort) + exit(code); + } +} + +extern "C" +{ +#include "utils.h" +} + +#define TWO02 (1 << 2) +#define TWO04 (1 << 4) +#define TWO08 (1 << 8) +#ifndef N +#define N (1 << 27) +#endif + +#ifndef BLOCK_SIZE +#define BLOCK_SIZE (128) +#endif + +#ifndef N_STREAMS +#define N_STREAMS (16) +#endif + +/* + *SAXPY (host implementation) + * y := a * x + y + */ +void host_saxpy(float *__restrict__ y, float a, float *__restrict__ x, int n) +{ +#pragma omp parallel for simd schedule(simd \ + : static) + for (int i = 0; i < n; i++) + { + y[i] = a * x[i] + y[i]; + } +} + +__global__ void gpu_saxpy(float *__restrict__ y, float a, float *__restrict__ x, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + y[i] = a * x[i] + y[i]; +} + +int main(int argc, const char **argv) +{ + int iret = 0; + int n = N; + float *h_x; + float *h_y; + float *h_z; + float a = 101.0f / TWO02, + b, c; + + if (argc > 1) + n = atoi(argv[1]); + + //CUDA Buffer Allocation + gpuErrchk(cudaMallocManaged((void **)&h_x, sizeof(float) * n)); + gpuErrchk(cudaMallocManaged((void **)&h_y, sizeof(float) * n)); + + if (NULL == (h_z = (float *)malloc(sizeof(float) * n))) + { + printf("error: memory allocation for 'z'\n"); + iret = -1; + } + if (0 != iret) + { + gpuErrchk(cudaFree(h_x)); + gpuErrchk(cudaFree(h_y)); + free(h_z); + exit(EXIT_FAILURE); + } + + //Init Data + b = rand() % TWO04; + c = rand() % TWO08; + for (int i = 0; i < n; i++) + { + h_x[i] = b / (float)TWO02; + h_y[i] = h_z[i] = c / (float)TWO04; + } + + start_timer(); + int TILE = n / N_STREAMS; + cudaStream_t stream[N_STREAMS]; + for (int i = 0; i < N_STREAMS; i++) + cudaStreamCreate(&stream[i]); + + //TODO Loop over the Tiles + for (int i = 0; i < n; i += TILE) + { + //TODO Kernel Tile i (stream i) + gpu_saxpy<<<((TILE + BLOCK_SIZE - 1) / BLOCK_SIZE), BLOCK_SIZE, 0, stream[i / TILE]>>>(&h_y[i], a, &h_x[i], TILE); + } + //TODO Wait all the streams... + cudaDeviceSynchronize(); + stop_timer(); + printf("saxpy (GPU): %9.3f sec %9.1f GFLOPS\n", elapsed_ns() / 1.0e9, 2 * n / ((float)elapsed_ns())); + + //Check Matematical Consistency + start_timer(); + host_saxpy(h_z, a, h_x, n); + stop_timer(); + printf("saxpy (Host): %9.3f sec %9.1f GFLOPS\n", elapsed_ns() / 1.0e9, 2 * n / ((float)elapsed_ns())); + for (int i = 0; i < n; ++i) + { + iret = *(int *)(h_y + i) ^ *(int *)(h_z + i); + assert(iret == 0); + } + + gpuErrchk(cudaFree(h_x)); + gpuErrchk(cudaFree(h_y)); + free(h_z); + + for (int i = 0; i < N_STREAMS; ++i) + cudaStreamDestroy(stream[i]); + + // CUDA exit -- needed to flush printf write buffer + cudaDeviceReset(); + return 0; +} diff --git a/cuda/lab3/.solutions/sobel-v1.cu b/cuda/lab3/.solutions/sobel-v1.cu new file mode 100644 index 0000000..184bdb9 --- /dev/null +++ b/cuda/lab3/.solutions/sobel-v1.cu @@ -0,0 +1,240 @@ +/* + * BSD 2-Clause License + * + * Copyright (c) 2020, Alessandro Capotondi + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * * Redistributions of source code must retain the above copyright notice, this + * list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright notice, + * this list of conditions and the following disclaimer in the documentation + * and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" + * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE + * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL + * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR + * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, + * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/** + * @file sobel.cu + * @author Alessandro Capotondi + * @date 12 May 2020 + * @brief Sobel Filtering + * + * @see https://dolly.fim.unimore.it/2019/course/view.php?id=152 + */ + +#include +#include +#include +#include +#include +#include +#include + +using namespace cv; +using namespace std; + +#ifndef BLOCK_SIZE +#define BLOCK_SIZE 16 +#endif +#ifndef NSTREAMS +#define NSTREAMS 12 +#endif + +#define gpuErrchk(ans) \ + { \ + gpuAssert((ans), __FILE__, __LINE__); \ + } +static inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) +{ + if (code != cudaSuccess) + { + fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); + if (abort) + exit(code); + } +} + +extern "C" +{ +#include "utils.h" +} + +int FILTER_HOST[3][3] = {{-1, 0, 1}, + {-2, 0, 2}, + {-1, 0, 1}}; + +void sobel_host(unsigned char *__restrict__ orig, unsigned char *__restrict__ out, int width, int height) +{ +#pragma omp parallel for collapse(2) + for (int y = 1; y < height - 1; y++) + { + for (int x = 1; x < width - 1; x++) + { + int dx = 0, dy = 0; + for (int k = -1; k <= 1; k++) + { + for (int z = -1; z <= 1; z++) + { + dx += FILTER_HOST[k + 1][z + 1] * orig[(y + k) * width + x + z]; + dy += FILTER_HOST[z + 1][k + 1] * orig[(y + k) * width + x + z]; + } + } + out[y * width + x] = sqrt((float)((dx * dx) + (dy * dy))); + } + } +} + +__constant__ int FILTER_GPU[3][3] = {{-1, 0, 1}, + {-2, 0, 2}, + {-1, 0, 1}}; + +__global__ void sobel_v1(unsigned char *__restrict__ orig, unsigned char *__restrict__ out, int width, int height) +{ + int i = threadIdx.y + blockIdx.y * blockDim.y; + int j = threadIdx.x + blockIdx.x * blockDim.x; + + if (j > 0 && i > 0 && j < width - 1 && i < height - 1) + { + int dx = 0, dy = 0; + for (int k = -1; k <= 1; k++) + { + for (int z = -1; z <= 1; z++) + { + dx += FILTER_GPU[k + 1][z + 1] * orig[(i + k) * width + j + z]; + dy += FILTER_GPU[z + 1][k + 1] * orig[(i + k) * width + j + z]; + } + } + out[i * width + j] = sqrt((float)((dx * dx) + (dy * dy))); + } +} + +int main(int argc, char *argv[]) +{ + int iret = 0; + struct timespec rt[2]; + string filename("data/sample.avi"); + + if (argc > 1) + filename = argv[1]; + + //Open Video Example + VideoCapture cap(filename); + // Check if camera opened successfully + if (!cap.isOpened()) + { + cout << "Error opening video stream or file" << endl; + return -1; + } + + int width = cap.get(CAP_PROP_FRAME_WIDTH); + int height = cap.get(CAP_PROP_FRAME_HEIGHT); + int nCh = 3; + + // Frame Buffers + Mat frameRGB = Mat::zeros(height, width, CV_8UC3); + Mat frameIn = Mat::zeros(height, width, CV_8UC1); + Mat frameOut = Mat::zeros(height, width, CV_8UC1); + + int nFrames = 0; + double time_cnt = 0.0; + while (1) + { + bool lastFrame = cap.read(frameRGB); // read a new frame from video + if (!lastFrame) + break; + + cvtColor(frameRGB, frameIn, COLOR_BGR2GRAY); + + // Compute CPU Version - Golden Model + clock_gettime(CLOCK_REALTIME, rt + 0); + sobel_host(frameIn.ptr(), frameOut.ptr(), width, height); + clock_gettime(CLOCK_REALTIME, rt + 1); + time_cnt+= (rt[1].tv_sec - rt[0].tv_sec) + 1.0e-9 * (rt[1].tv_nsec - rt[0].tv_nsec); + nFrames++; + +#ifdef DISPLAY + // Show frames + imshow("frameIn", frameIn); + imshow("frameOut", frameOut); + waitKey(1); +#endif + } + printf("Sobel (Host) : %d frames, %9.6f s per-frame (%9.6f fps)\n", nFrames, time_cnt/nFrames, 1/(time_cnt/nFrames)); + + // CUDA VERSION -------------------------------------------------- + //Open Video Example + cap = VideoCapture(filename); + // Check if camera opened successfully + if (!cap.isOpened()) + { + cout << "Error opening video stream or file" << endl; + return -1; + } + + unsigned char *d_image_in; + unsigned char *d_image_out; + gpuErrchk(cudaMalloc((void **)&d_image_in, sizeof(unsigned char) * width * height)); + gpuErrchk(cudaMalloc((void **)&d_image_out, sizeof(unsigned char) * width * height)); + gpuErrchk(cudaMemset(d_image_out, 0, sizeof(unsigned char) * width * height)); + + cudaStream_t stream[NSTREAMS]; + for(int i = 0; i < NSTREAMS; i++) + cudaStreamCreate(&stream[i]); + + nFrames = 0; + time_cnt = 0.0; + while (1) + { + bool lastFrame = cap.read(frameRGB); // read a new frame from video + if (!lastFrame) + break; + + cvtColor(frameRGB, frameIn, COLOR_BGR2GRAY); + + // Compute CPU Version - Golden Model + clock_gettime(CLOCK_REALTIME, rt + 0); + gpuErrchk(cudaMemcpyAsync(d_image_in, frameIn.ptr(), sizeof(unsigned char) * width * height, cudaMemcpyHostToDevice,stream[nFrames%NSTREAMS])); + dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); + dim3 dimGrid((width + BLOCK_SIZE - 1) / BLOCK_SIZE, (height + BLOCK_SIZE - 1) / BLOCK_SIZE); + sobel_v1<<>>(d_image_in, d_image_out, width, height); + gpuErrchk(cudaPeekAtLastError()); + gpuErrchk(cudaMemcpyAsync(frameOut.ptr(), d_image_out, sizeof(unsigned char) * width * height, cudaMemcpyDeviceToHost,stream[nFrames%NSTREAMS])); + clock_gettime(CLOCK_REALTIME, rt + 1); + time_cnt+= (rt[1].tv_sec - rt[0].tv_sec) + 1.0e-9 * (rt[1].tv_nsec - rt[0].tv_nsec); + nFrames++; + +#ifdef DISPLAY + // Show frames + imshow("frameIn", frameIn); + imshow("frameOut", frameOut); + waitKey(1); +#endif + } + cudaDeviceSynchronize(); + printf("Sobel (GPU) : %d frames, %9.6f s per-frame (%9.6f fps)\n", nFrames, time_cnt/nFrames, 1/(time_cnt/nFrames)); + + gpuErrchk(cudaFree(d_image_out)); + gpuErrchk(cudaFree(d_image_in)); + for (int i=0; i +#include +#include +#include + +#define gpuErrchk(ans) \ + { \ + gpuAssert((ans), __FILE__, __LINE__); \ + } +static inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) +{ + if (code != cudaSuccess) + { + fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); + if (abort) + exit(code); + } +} + +extern "C" +{ +#include "utils.h" +} + +#define TWO02 (1 << 2) +#define TWO04 (1 << 4) +#define TWO08 (1 << 8) + +#ifndef N +#define N (1 << 10) +#endif +#ifndef TILE_W +#define TILE_W 128 +#endif +#ifndef BLOCK_SIZE +#define BLOCK_SIZE 32 +#endif + +#define SM 64 +static void reorder(float *__restrict__ a, float *__restrict__ b, int n) +{ + for (int i = 0; i < SM; i++) + for (int j = 0; j < SM; j++) + b[i * SM + j] = a[i * n + j]; +} + +static void mm(float *__restrict__ a, float *__restrict__ b, float *__restrict__ c, int n) +{ + for (int i = 0; i < SM; i++) + { + for (int k = 0; k < SM; k++) + { + for (int j = 0; j < SM; j++) + { + c[i * n + j] += a[i * n + k] * b[k * SM + j]; + } + } + } +} +void gemm_host(float *a, float *b, float *c, int n) +{ + int bk = n / SM; +#pragma omp parallel for collapse(3) + for (int i = 0; i < bk; i++) + { + for (int j = 0; j < bk; j++) + { + for (int k = 0; k < bk; k++) + { + float b2[SM * SM]; + reorder(&b[SM * (k * n + j)], b2, n); + mm(&a[SM * (i * n + k)], b2, &c[SM * (i * n + j)], n); + } + } + } +} +__global__ void gemm(float *__restrict__ a, float *__restrict__ b, float *__restrict__ c, int n) +{ + __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; + __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; + + int ib = blockIdx.y; + int jb = blockIdx.x; + + int it = threadIdx.y; + int jt = threadIdx.x; + + int a_offset, b_offset, c_offset; + + float Cvalue = 0.0f; + for (int kb = 0; kb < (n / BLOCK_SIZE); ++kb) + { + a_offset = ib * n * BLOCK_SIZE + kb * BLOCK_SIZE; + b_offset = kb * n * BLOCK_SIZE + jb * BLOCK_SIZE; + As[it][jt] = a[a_offset + it * n + jt]; + Bs[it][jt] = b[b_offset + it * n + jt]; + __syncthreads(); + + for (int k = 0; k < BLOCK_SIZE; ++k) + Cvalue += As[it][k] * Bs[k][jt]; + + __syncthreads(); + } + + c_offset = ib * n * BLOCK_SIZE + jb * BLOCK_SIZE; + c[c_offset + it * n + jt] = Cvalue; +} + +int main(int argc, char *argv[]) +{ + int n = N, iret = 0; + float *a, *b, *c, *g; + struct timespec rt[2]; + double wt; // walltime + + if (argc > 1) + n = atoi(argv[1]); + + //TODO Update malloc to cudaMallocHost or cudaMallocManaged (if necessary) + if (NULL == (a = (float *)malloc(sizeof(*a) * n * n))) + { + printf("error: memory allocation for 'x'\n"); + iret = -1; + } + //TODO Update malloc to cudaMallocHost or cudaMallocManaged (if necessary) + if (NULL == (b = (float *)malloc(sizeof(*b) * n * n))) + { + printf("error: memory allocation for 'y'\n"); + iret = -1; + } + //TODO Update malloc to cudaMallocHost or cudaMallocManaged (if necessary) + if (NULL == (c = (float *)malloc(sizeof(*c) * n * n))) + { + printf("error: memory allocation for 'z'\n"); + iret = -1; + } + if (NULL == (g = (float *)malloc(sizeof(*g) * n * n))) + { + printf("error: memory allocation for 'z'\n"); + iret = -1; + } + + if (0 != iret) + { + //TODO Update cudaFreeHost or cudaFree (if necessary) + free(a); + //TODO Update cudaFreeHost or cudaFree (if necessary) + free(b); + //TODO Update cudaFreeHost or cudaFree (if necessary) + free(c); + free(g); + exit(EXIT_FAILURE); + } + + //Init Data + int _b = rand() % TWO04; + int _c = rand() % TWO08; + +#pragma omp parallel for + for (int i = 0; i < n * n; i++) + { + a[i] = _b / (float)TWO02; + b[i] = _c / (float)TWO04; + c[i] = g[i] = 0.0; + } + + clock_gettime(CLOCK_REALTIME, rt + 0); + gemm_host(a, b, g, n); + clock_gettime(CLOCK_REALTIME, rt + 1); + wt = (rt[1].tv_sec - rt[0].tv_sec) + 1.0e-9 * (rt[1].tv_nsec - rt[0].tv_nsec); + printf("GEMM (Host) : %9.3f sec %9.1f GFLOPS\n", wt, 2.0 * n * n * n / (1.0e9 * wt)); + + //TODO Remove if unecessary + float *d_a, *d_b, *d_c; + gpuErrchk(cudaMalloc((void **)&d_a, sizeof(float) * n * n)); + gpuErrchk(cudaMalloc((void **)&d_b, sizeof(float) * n * n)); + gpuErrchk(cudaMalloc((void **)&d_c, sizeof(float) * n * n)); + + clock_gettime(CLOCK_REALTIME, rt + 0); + //TODO Remove if unecessary + gpuErrchk(cudaMemcpy(d_a, a, sizeof(float) * n * n, cudaMemcpyHostToDevice)); + gpuErrchk(cudaMemcpy(d_b, b, sizeof(float) * n * n, cudaMemcpyHostToDevice)); + dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); + dim3 dimGrid((n + (BLOCK_SIZE)-1) / (BLOCK_SIZE), (n + (BLOCK_SIZE)-1) / (BLOCK_SIZE)); + gemm<<>>(d_a, d_b, d_c, n); + gpuErrchk(cudaPeekAtLastError()); + //TODO Remove if unecessary + gpuErrchk(cudaMemcpy(c, d_c, sizeof(float) * n * n, cudaMemcpyDeviceToHost)); + clock_gettime(CLOCK_REALTIME, rt + 1); + wt = (rt[1].tv_sec - rt[0].tv_sec) + 1.0e-9 * (rt[1].tv_nsec - rt[0].tv_nsec); + printf("GEMM-v1 (GPU): %9.3f sec %9.1f GFLOPS\n", wt, 2.0 * n * n * n / (1.0e9 * wt)); + + for (int i = 0; i < n * n; i++) + { + iret = *(int *)(g + i) ^ *(int *)(c + i); + assert(iret == 0); + } + + //TODO Update cudaFreeHost or cudaFree (if necessary) + free(a); + //TODO Update cudaFreeHost or cudaFree (if necessary) + free(b); + //TODO Update cudaFreeHost or cudaFree (if necessary) + free(c); + free(g); + //TODO Remove if unecessary + gpuErrchk(cudaFree(d_a)); + //TODO Remove if unecessary + gpuErrchk(cudaFree(d_b)); + //TODO Remove if unecessary + gpuErrchk(cudaFree(d_c)); + + return 0; +} diff --git a/cuda/lab3/saxpy.cu b/cuda/lab3/saxpy.cu new file mode 100644 index 0000000..d37b6e6 --- /dev/null +++ b/cuda/lab3/saxpy.cu @@ -0,0 +1,184 @@ +/* + * BSD 2-Clause License + * + * Copyright (c) 2020, Alessandro Capotondi + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * * Redistributions of source code must retain the above copyright notice, this + * list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright notice, + * this list of conditions and the following disclaimer in the documentation + * and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" + * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE + * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL + * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR + * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, + * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/** + * @file saxpy.c + * @author Alessandro Capotondi + * @date 12 May 2020 + * @brief Saxpy + * + * @see https://dolly.fim.unimore.it/2019/course/view.php?id=152 + */ + +#include +#include +#include +#include +#include +#include +#include + +#define gpuErrchk(ans) \ + { \ + gpuAssert((ans), __FILE__, __LINE__); \ + } +static inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) +{ + if (code != cudaSuccess) + { + fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); + if (abort) + exit(code); + } +} + +extern "C" +{ +#include "utils.h" +} + +#define TWO02 (1 << 2) +#define TWO04 (1 << 4) +#define TWO08 (1 << 8) +#ifndef N +#define N (1 << 27) +#endif + +#ifndef BLOCK_SIZE +#define BLOCK_SIZE (512) +#endif + +/* + *SAXPY (host implementation) + * y := a * x + y + */ +void host_saxpy(float * __restrict__ y, float a, float * __restrict__ x, int n) +{ +#pragma omp parallel for simd schedule(simd: static) + for (int i = 0; i < n; i++) + { + y[i] = a * x[i] + y[i]; + } +} + +__global__ void gpu_saxpy(float * __restrict__ y, float a, float * __restrict__ x, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + y[i] = a * x[i] + y[i]; +} + +int main(int argc, const char **argv) +{ + int iret = 0; + int n = N; + float *h_x, *d_x; + float *h_y, *d_y; + float *h_z; + float a = 101.0f / TWO02, + b, c; + + if (argc > 1) + n = atoi(argv[1]); + + //TODO Update malloc to cudaMallocHost or cudaMallocManaged (if necessary) + if (NULL == (h_x = (float *)malloc(sizeof(float) * n))) + { + printf("error: memory allocation for 'x'\n"); + iret = -1; + } + //TODO Update malloc to cudaMallocHost or cudaMallocManaged (if necessary) + if (NULL == (h_y = (float *)malloc(sizeof(float) * n))) + { + printf("error: memory allocation for 'y'\n"); + iret = -1; + } + if (NULL == (h_z = (float *)malloc(sizeof(float) * n))) + { + printf("error: memory allocation for 'z'\n"); + iret = -1; + } + if (0 != iret) + { + //TODO Update cudaFreeHost or cudaFree (if necessary) + free(h_x); + //TODO Update cudaFreeHost or cudaFree (if necessary) + free(h_y); + free(h_z); + exit(EXIT_FAILURE); + } + + //Init Data + b = rand() % TWO04; + c = rand() % TWO08; + for (int i = 0; i < n; i++) + { + h_x[i] = b / (float)TWO02; + h_y[i] = h_z[i] = c / (float)TWO04; + } + + //TODO Remove if unecessary + gpuErrchk(cudaMalloc((void **)&d_x, sizeof(float) * n)); + gpuErrchk(cudaMalloc((void **)&d_y, sizeof(float) * n)); + + start_timer(); + //TODO Remove if unecessary + gpuErrchk(cudaMemcpy(d_x, h_x, sizeof(float) * n, cudaMemcpyHostToDevice)); + gpuErrchk(cudaMemcpy(d_y, h_y, sizeof(float) * n, cudaMemcpyHostToDevice)); + gpu_saxpy<<<((n + BLOCK_SIZE - 1) / BLOCK_SIZE), BLOCK_SIZE>>>(d_y, a, d_x, n); + gpuErrchk(cudaPeekAtLastError()); + //TODO Remove if unecessary + gpuErrchk(cudaMemcpy(h_y, d_y, sizeof(float) * n, cudaMemcpyDeviceToHost)); + stop_timer(); + printf("saxpy (GPU): %9.3f sec %9.1f GFLOPS\n", elapsed_ns() / 1.0e9, 2 * n / ((float) elapsed_ns())); + + //Check Matematical Consistency + start_timer(); + host_saxpy(h_z, a, h_x, n); + stop_timer(); + printf("saxpy (Host): %9.3f sec %9.1f GFLOPS\n", elapsed_ns() / 1.0e9, 2 * n / ((float) elapsed_ns())); + for (int i = 0; i < n; ++i) + { + iret = *(int *)(h_y + i) ^ *(int *)(h_z + i); + assert(iret == 0); + } + + //TODO Update cudaFreeHost or cudaFree (if necessary) + free(h_x); + //TODO Remove if unecessary + gpuErrchk(cudaFree(d_x)); + //TODO Update cudaFreeHost or cudaFree (if necessary) + free(h_y); + //TODO Remove if unecessary + gpuErrchk(cudaFree(d_y)); + free(h_z); + + // CUDA exit -- needed to flush printf write buffer + cudaDeviceReset(); + return 0; +} diff --git a/cuda/lab3/sobel.cu b/cuda/lab3/sobel.cu new file mode 100644 index 0000000..a0f713d --- /dev/null +++ b/cuda/lab3/sobel.cu @@ -0,0 +1,229 @@ +/* + * BSD 2-Clause License + * + * Copyright (c) 2020, Alessandro Capotondi + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * * Redistributions of source code must retain the above copyright notice, this + * list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright notice, + * this list of conditions and the following disclaimer in the documentation + * and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" + * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE + * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL + * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR + * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, + * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/** + * @file sobel.cu + * @author Alessandro Capotondi + * @date 12 May 2020 + * @brief Sobel Filtering + * + * @see https://dolly.fim.unimore.it/2019/course/view.php?id=152 + */ + +#include +#include +#include +#include +#include +#include +#include + +using namespace cv; +using namespace std; + +#ifndef BLOCK_SIZE +#define BLOCK_SIZE 32 +#endif + +#define gpuErrchk(ans) \ + { \ + gpuAssert((ans), __FILE__, __LINE__); \ + } +static inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) +{ + if (code != cudaSuccess) + { + fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); + if (abort) + exit(code); + } +} + +extern "C" +{ +#include "utils.h" +} + +int FILTER_HOST[3][3] = {{-1, 0, 1}, + {-2, 0, 2}, + {-1, 0, 1}}; + +void sobel_host(unsigned char *__restrict__ orig, unsigned char *__restrict__ out, int width, int height) +{ +#pragma omp parallel for collapse(2) + for (int y = 1; y < height - 1; y++) + { + for (int x = 1; x < width - 1; x++) + { + int dx = 0, dy = 0; + for (int k = -1; k <= 1; k++) + { + for (int z = -1; z <= 1; z++) + { + dx += FILTER_HOST[k + 1][z + 1] * orig[(y + k) * width + x + z]; + dy += FILTER_HOST[z + 1][k + 1] * orig[(y + k) * width + x + z]; + } + } + out[y * width + x] = sqrt((float)((dx * dx) + (dy * dy))); + } + } +} + +__constant__ int FILTER_GPU[3][3] = {{-1, 0, 1}, + {-2, 0, 2}, + {-1, 0, 1}}; + +__global__ void sobel_v1(unsigned char *__restrict__ orig, unsigned char *__restrict__ out, int width, int height) +{ + int i = threadIdx.y + blockIdx.y * blockDim.y; + int j = threadIdx.x + blockIdx.x * blockDim.x; + + if (j > 0 && i > 0 && j < width - 1 && i < height - 1) + { + int dx = 0, dy = 0; + for (int k = -1; k <= 1; k++) + { + for (int z = -1; z <= 1; z++) + { + dx += FILTER_GPU[k + 1][z + 1] * orig[(i + k) * width + j + z]; + dy += FILTER_GPU[z + 1][k + 1] * orig[(i + k) * width + j + z]; + } + } + out[i * width + j] = sqrt((float)((dx * dx) + (dy * dy))); + } +} + +int main(int argc, char *argv[]) +{ + int iret = 0; + struct timespec rt[2]; + string filename("data/sample.avi"); + + if (argc > 1) + filename = argv[1]; + + //Open Video Example + VideoCapture cap(filename); + // Check if camera opened successfully + if (!cap.isOpened()) + { + cout << "Error opening video stream or file" << endl; + return -1; + } + + int width = cap.get(CAP_PROP_FRAME_WIDTH); + int height = cap.get(CAP_PROP_FRAME_HEIGHT); + int nCh = 3; + + // Frame Buffers + Mat frameRGB = Mat::zeros(height, width, CV_8UC3); + Mat frameIn = Mat::zeros(height, width, CV_8UC1); + Mat frameOut = Mat::zeros(height, width, CV_8UC1); + + int nFrames = 0; + double time_cnt = 0.0; + while (1) + { + bool lastFrame = cap.read(frameRGB); // read a new frame from video + if (!lastFrame) + break; + + cvtColor(frameRGB, frameIn, COLOR_BGR2GRAY); + + // Compute CPU Version - Golden Model + clock_gettime(CLOCK_REALTIME, rt + 0); + sobel_host(frameIn.ptr(), frameOut.ptr(), width, height); + clock_gettime(CLOCK_REALTIME, rt + 1); + time_cnt+= (rt[1].tv_sec - rt[0].tv_sec) + 1.0e-9 * (rt[1].tv_nsec - rt[0].tv_nsec); + nFrames++; + +#ifdef DISPLAY + // Show frames + imshow("frameIn", frameIn); + imshow("frameOut", frameOut); + waitKey(1); +#endif + } + printf("Sobel (Host) : %d frames, %9.6f s per-frame (%9.6f fps)\n", nFrames, time_cnt/nFrames, 1/(time_cnt/nFrames)); + + // CUDA VERSION -------------------------------------------------- + //Open Video Example + cap = VideoCapture(filename); + // Check if camera opened successfully + if (!cap.isOpened()) + { + cout << "Error opening video stream or file" << endl; + return -1; + } + + unsigned char *d_image_in; + unsigned char *d_image_out; + gpuErrchk(cudaMalloc((void **)&d_image_in, sizeof(unsigned char) * width * height)); + gpuErrchk(cudaMalloc((void **)&d_image_out, sizeof(unsigned char) * width * height)); + gpuErrchk(cudaMemset(d_image_out, 0, sizeof(unsigned char) * width * height)); + + nFrames = 0; + time_cnt = 0.0; + while (1) + { + bool lastFrame = cap.read(frameRGB); // read a new frame from video + if (!lastFrame) + break; + + cvtColor(frameRGB, frameIn, COLOR_BGR2GRAY); + + // Compute CPU Version - Golden Model + clock_gettime(CLOCK_REALTIME, rt + 0); + gpuErrchk(cudaMemcpy(d_image_in, frameIn.ptr(), sizeof(unsigned char) * width * height, cudaMemcpyHostToDevice)); + dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); + dim3 dimGrid((width + BLOCK_SIZE - 1) / BLOCK_SIZE, (height + BLOCK_SIZE - 1) / BLOCK_SIZE); + sobel_v1<<>>(d_image_in, d_image_out, width, height); + gpuErrchk(cudaPeekAtLastError()); + gpuErrchk(cudaMemcpy(frameOut.ptr(), d_image_out, sizeof(unsigned char) * width * height, cudaMemcpyDeviceToHost)); + clock_gettime(CLOCK_REALTIME, rt + 1); + time_cnt+= (rt[1].tv_sec - rt[0].tv_sec) + 1.0e-9 * (rt[1].tv_nsec - rt[0].tv_nsec); + nFrames++; + +#ifdef DISPLAY + // Show frames + imshow("frameIn", frameIn); + imshow("frameOut", frameOut); + waitKey(1); +#endif + } + printf("Sobel (GPU) : %d frames, %9.6f s per-frame (%9.6f fps)\n", nFrames, time_cnt/nFrames, 1/(time_cnt/nFrames)); + + gpuErrchk(cudaFree(d_image_out)); + gpuErrchk(cudaFree(d_image_in)); + frameOut.release(); + frameIn.release(); + frameRGB.release(); + cap.release(); + + return iret; +} diff --git a/cuda/lab3/utils.c b/cuda/lab3/utils.c new file mode 100644 index 0000000..0ce0dc5 --- /dev/null +++ b/cuda/lab3/utils.c @@ -0,0 +1,138 @@ +/* + * BSD 2-Clause License + * + * Copyright (c) 2020, Alessandro Capotondi + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * * Redistributions of source code must retain the above copyright notice, this + * list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright notice, + * this list of conditions and the following disclaimer in the documentation + * and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" + * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE + * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL + * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR + * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, + * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ +/** + * @file utils.c + * @author Alessandro Capotondi + * @date 27 Mar 2020 + * @brief File containing utilities functions for HPC Unimore Class + * + * Utilities for OpenMP lab. + * + * @see http://algo.ing.unimo.it/people/andrea/Didattica/HPC/index.html + */ + +#define _POSIX_C_SOURCE 199309L +#include +#include +#include +#include +#include + +extern "C" { + +#include "utils.h" + +#define MAX_ITERATIONS 100 +static struct timespec timestampA, timestampB; +static unsigned long long statistics[MAX_ITERATIONS]; +static int iterations = 0; + +static unsigned long long __diff_ns(struct timespec start, struct timespec end) +{ + struct timespec temp; + if ((end.tv_nsec - start.tv_nsec) < 0) + { + temp.tv_sec = end.tv_sec - start.tv_sec - 1; + temp.tv_nsec = 1000000000ULL + end.tv_nsec - start.tv_nsec; + } + else + { + temp.tv_sec = end.tv_sec - start.tv_sec; + temp.tv_nsec = end.tv_nsec - start.tv_nsec; + } + + return temp.tv_nsec + temp.tv_sec * 1000000000ULL; +} + +void start_timer() +{ + asm volatile("" :: + : "memory"); + clock_gettime(CLOCK_MONOTONIC_RAW, ×tampA); + asm volatile("" :: + : "memory"); +} + +void stop_timer() +{ + unsigned long long elapsed = 0ULL; + asm volatile("" :: + : "memory"); + clock_gettime(CLOCK_MONOTONIC_RAW, ×tampB); + asm volatile("" :: + : "memory"); +} + +unsigned long long elapsed_ns() +{ + return __diff_ns(timestampA, timestampB); +} + +void start_stats() +{ + start_timer(); +} + +void collect_stats() +{ + assert(iterations < MAX_ITERATIONS); + stop_timer(); + statistics[iterations++] = elapsed_ns(); +} + +void print_stats() +{ + unsigned long long min = ULLONG_MAX; + unsigned long long max = 0LL; + double average = 0.0; + double std_deviation = 0.0; + double sum = 0.0; + + /* Compute the sum of all elements */ + for (int i = 0; i < iterations; i++) + { + if (statistics[i] > max) + max = statistics[i]; + if (statistics[i] < min) + min = statistics[i]; + sum = sum + statistics[i] / 1E6; + } + average = sum / (double)iterations; + + /* Compute variance and standard deviation */ + for (int i = 0; i < iterations; i++) + { + sum = sum + pow((statistics[i] / 1E6 - average), 2); + } + std_deviation = sqrt(sum / (double)iterations); + + printf("AvgTime\tMinTime\tMaxTime\tStdDev\n"); + printf("%.4f ms\t%.4f ms\t%.4f ms\t%.4f\n", (double)average, (double)min / 1E6, (double)max / 1E6, (double)std_deviation); +} + +} diff --git a/cuda/lab3/utils.h b/cuda/lab3/utils.h new file mode 100644 index 0000000..966281c --- /dev/null +++ b/cuda/lab3/utils.h @@ -0,0 +1,142 @@ +/* + * BSD 2-Clause License + * + * Copyright (c) 2020, Alessandro Capotondi + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * * Redistributions of source code must retain the above copyright notice, this + * list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright notice, + * this list of conditions and the following disclaimer in the documentation + * and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" + * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE + * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL + * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR + * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, + * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/** + * @file utils.h + * @author Alessandro Capotondi + * @date 27 Mar 2020 + * @brief File containing utilities functions for HPC Unimore Class + * + * The header define time functions and dummy workload used on the example tests. + * + * @see http://algo.ing.unimo.it/people/andrea/Didattica/HPC/index.html + */ +#ifndef __UTILS_H__ +#define __UTILS_H__ + +#include + +#if defined(VERBOSE) +#define DEBUG_PRINT(x, ...) printf((x), ##__VA_ARGS__) +#else +#define DEBUG_PRINT(x, ...) +#endif + +#if !defined(NTHREADS) +#define NTHREADS (4) +#endif + +extern "C" +{ + +/** + * @brief The function set the timestampA + * + * The function is used to measure elapsed time between two execution points. + * The function start_timer() sets the starting point timestamp, while the function + * stop_timer() sets the termination timestamp. The elapsed time, expressed in nanoseconds, + * between the two points can be retrieved using the function elapsed_ns(). + * + * Example usage: + * @code + * start_timer(); // Point A + * //SOME CODE HERE + * stop_timer(); // Point B + * printf("Elapsed time = %llu ns\n", elapsed_ns())); //Elapsed time between A and B + * //SOME OTHER CODE HERE + * stop_timer(); // Point C + * printf("Elapsed time = %llu ns\n", elapsed_ns())); //Elapsed time between A and C + * @endcode + * + * @return void + * @see start_timer() + * @see stop_timer() + * @see elapsed_ns() + */ + void start_timer(); + +/** + * @brief The function set the second timestamps + * + * The function is used to measure elapsed time between two execution points. + * The function start_timer() sets the starting point timestamp, while the function + * stop_timer() returns the elapsed time, expressed in nanoseconds between the last call + * of start_timer() and the current execution point. + * + * Example usage: + * @code + * start_timer(); // Point A + * //SOME CODE HERE + * stop_timer(); // Point B + * printf("Elapsed time = %llu ns\n", elapsed_ns())); //Elapsed time between A and B + * //SOME OTHER CODE HERE + * stop_timer(); // Point C + * printf("Elapsed time = %llu ns\n", elapsed_ns())); //Elapsed time between A and C + * @endcode + * + * @return void + * @see start_timer() + * @see stop_timer() + * @see elapsed_ns() + */ + void stop_timer(); + +/** + * @brief Elapsed nano seconds between start_timer() and stop_timer(). + * + * @return Elapsed nano seconds + * @see start_timer() + * @see stop_timer() + */ + unsigned long long elapsed_ns(); + +/** + * @brief The function init the starting point of stat measurement. + * + * The function is similar to start_timer(). + * + * @return void + * @see start_timer + */ + void start_stats(); + +/** + * @brief The function collects the elapsed time between the current exeuction point and the + * last call of start_stats(). + * + * @return void + */ + void collect_stats(); + +/** + * @brief The function display the collected statistics. + * @return void + */ + void print_stats(); +} +#endif /*__UTILS_H__*/