mirror of
https://github.com/Steffo99/unimore-hpc-assignments.git
synced 2024-11-25 09:34:23 +00:00
HPC CUDA Lab 3
This commit is contained in:
Alessandro Capotondi
parent
809d881f2f
commit
9570af5b66
26 changed files with 3823 additions and 4 deletions
12
README.md
12
README.md
|
|
@ -10,10 +10,14 @@ This repo contains the exercises and the tutorials used for Unimore's HPC class
|
||||||
|
|
||||||
### OpenMP Exercises
|
### OpenMP Exercises
|
||||||
The exercises related to OpenMP programming model can be found in the folder `openmp`. Here the list of currectly available classes:
|
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\lab1`: OpenMP basics: *parallel*, *for-loop*, *sections*, and *tasking*
|
||||||
- `openmp\lab2`: OpenMP Advanced: *reduction*, *tasking*, *optimizations*.
|
- `openmp\lab2`: OpenMP Advanced: *reduction*, *tasking*, *optimizations*
|
||||||
- `openmp\lab3`: OpenMP 4.x+: *Accelerator Model (targeting: Nvidia GP-GPU)*
|
- `openmp\lab3`: OpenMP 4.x+: *Accelerator Model (targeting: Nvidia GP-GPU)*
|
||||||
|
|
||||||
### CUDA Exercises
|
### CUDA Exercises
|
||||||
- `cuda\lab1`: CUDA Basics.
|
- `cuda\lab1`: CUDA Basics
|
||||||
- `cuda\lab2`: CUDA Memory Model.
|
- `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*
|
||||||
|
|
|
||||||
55
challenge/Makefile
Normal file
55
challenge/Makefile
Normal file
|
|
@ -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
|
||||||
283
challenge/TrackColour.cpp
Normal file
283
challenge/TrackColour.cpp
Normal file
|
|
@ -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 <stdio.h>
|
||||||
|
#include <time.h>
|
||||||
|
#include <sys/time.h>
|
||||||
|
#include <opencv2/opencv.hpp>
|
||||||
|
#include "opencv2/imgproc/imgproc_c.h"
|
||||||
|
#include "opencv2/imgproc/imgproc.hpp"
|
||||||
|
#include <opencv2/imgcodecs/imgcodecs.hpp>
|
||||||
|
#include <opencv2/objdetect/objdetect.hpp>
|
||||||
|
#include <opencv2/highgui/highgui.hpp>
|
||||||
|
|
||||||
|
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;
|
||||||
|
}
|
||||||
BIN
challenge/data/sample.avi
Normal file
BIN
challenge/data/sample.avi
Normal file
Binary file not shown.
595
challenge/data/sample_ground_truth.h
Normal file
595
challenge/data/sample_ground_truth.h
Normal file
|
|
@ -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};
|
||||||
138
challenge/utils.c
Normal file
138
challenge/utils.c
Normal file
|
|
@ -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 <time.h>
|
||||||
|
#include <limits.h>
|
||||||
|
#include <math.h>
|
||||||
|
#include <stdio.h>
|
||||||
|
#include <assert.h>
|
||||||
|
|
||||||
|
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);
|
||||||
|
}
|
||||||
|
|
||||||
|
}
|
||||||
142
challenge/utils.h
Normal file
142
challenge/utils.h
Normal file
|
|
@ -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 <stdarg.h>
|
||||||
|
|
||||||
|
#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__*/
|
||||||
232
cuda/lab3/.solutions/gemm-v1.cu
Normal file
232
cuda/lab3/.solutions/gemm-v1.cu
Normal file
|
|
@ -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 <assert.h>
|
||||||
|
#include <stdio.h>
|
||||||
|
#include <stdlib.h>
|
||||||
|
#include <time.h>
|
||||||
|
|
||||||
|
#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<<<dimGrid, dimBlock>>>(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;
|
||||||
|
}
|
||||||
214
cuda/lab3/.solutions/gemm-v2.cu
Normal file
214
cuda/lab3/.solutions/gemm-v2.cu
Normal file
|
|
@ -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 <assert.h>
|
||||||
|
#include <stdio.h>
|
||||||
|
#include <stdlib.h>
|
||||||
|
#include <time.h>
|
||||||
|
|
||||||
|
#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<<<dimGrid, dimBlock>>>(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;
|
||||||
|
}
|
||||||
175
cuda/lab3/.solutions/saxpy-v1.cu
Normal file
175
cuda/lab3/.solutions/saxpy-v1.cu
Normal file
|
|
@ -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 <assert.h>
|
||||||
|
#include <time.h>
|
||||||
|
#include <stdlib.h>
|
||||||
|
#include <stdio.h>
|
||||||
|
#include <string.h>
|
||||||
|
#include <math.h>
|
||||||
|
#include <cuda_runtime.h>
|
||||||
|
|
||||||
|
#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;
|
||||||
|
}
|
||||||
167
cuda/lab3/.solutions/saxpy-v2.cu
Normal file
167
cuda/lab3/.solutions/saxpy-v2.cu
Normal file
|
|
@ -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 <assert.h>
|
||||||
|
#include <time.h>
|
||||||
|
#include <stdlib.h>
|
||||||
|
#include <stdio.h>
|
||||||
|
#include <string.h>
|
||||||
|
#include <math.h>
|
||||||
|
#include <cuda_runtime.h>
|
||||||
|
|
||||||
|
#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;
|
||||||
|
}
|
||||||
197
cuda/lab3/.solutions/saxpy-v3.cu
Normal file
197
cuda/lab3/.solutions/saxpy-v3.cu
Normal file
|
|
@ -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 <assert.h>
|
||||||
|
#include <time.h>
|
||||||
|
#include <stdlib.h>
|
||||||
|
#include <stdio.h>
|
||||||
|
#include <string.h>
|
||||||
|
#include <math.h>
|
||||||
|
#include <cuda_runtime.h>
|
||||||
|
|
||||||
|
#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;
|
||||||
|
}
|
||||||
197
cuda/lab3/.solutions/saxpy-v4.cu
Normal file
197
cuda/lab3/.solutions/saxpy-v4.cu
Normal file
|
|
@ -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 <assert.h>
|
||||||
|
#include <time.h>
|
||||||
|
#include <stdlib.h>
|
||||||
|
#include <stdio.h>
|
||||||
|
#include <string.h>
|
||||||
|
#include <math.h>
|
||||||
|
#include <cuda_runtime.h>
|
||||||
|
|
||||||
|
#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<N_STREAMS; ++i)
|
||||||
|
cudaStreamDestroy(stream[i]);
|
||||||
|
|
||||||
|
|
||||||
|
// CUDA exit -- needed to flush printf write buffer
|
||||||
|
cudaDeviceReset();
|
||||||
|
return 0;
|
||||||
|
}
|
||||||
178
cuda/lab3/.solutions/saxpy-v5.cu
Normal file
178
cuda/lab3/.solutions/saxpy-v5.cu
Normal file
|
|
@ -0,0 +1,178 @@
|
||||||
|
/*
|
||||||
|
* 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 <assert.h>
|
||||||
|
#include <time.h>
|
||||||
|
#include <stdlib.h>
|
||||||
|
#include <stdio.h>
|
||||||
|
#include <string.h>
|
||||||
|
#include <math.h>
|
||||||
|
#include <cuda_runtime.h>
|
||||||
|
|
||||||
|
#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;
|
||||||
|
}
|
||||||
240
cuda/lab3/.solutions/sobel-v1.cu
Normal file
240
cuda/lab3/.solutions/sobel-v1.cu
Normal file
|
|
@ -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 <stdio.h>
|
||||||
|
#include <time.h>
|
||||||
|
#include <sys/time.h>
|
||||||
|
#include <opencv2/opencv.hpp>
|
||||||
|
#include <opencv2/imgcodecs/imgcodecs.hpp>
|
||||||
|
#include <opencv2/objdetect/objdetect.hpp>
|
||||||
|
#include <opencv2/highgui/highgui.hpp>
|
||||||
|
|
||||||
|
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<<<dimGrid, dimBlock,0,stream[nFrames%NSTREAMS]>>>(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<NSTREAMS; ++i)
|
||||||
|
gpuErrchk(cudaStreamDestroy(stream[i]));
|
||||||
|
|
||||||
|
frameOut.release();
|
||||||
|
frameIn.release();
|
||||||
|
frameRGB.release();
|
||||||
|
cap.release();
|
||||||
|
|
||||||
|
return iret;
|
||||||
|
}
|
||||||
58
cuda/lab3/Makefile
Executable file
58
cuda/lab3/Makefile
Executable file
|
|
@ -0,0 +1,58 @@
|
||||||
|
ifndef CUDA_HOME
|
||||||
|
CUDA_HOME:=/usr/local/cuda
|
||||||
|
endif
|
||||||
|
|
||||||
|
ifndef EXERCISE
|
||||||
|
EXERCISE=exercise1.cu
|
||||||
|
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 `pkg-config --cflags --libs opencv4`
|
||||||
|
|
||||||
|
CXXFLAGS:=$(OPT) -I. $(EXT_CXXFLAGS)
|
||||||
|
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:.cu=.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)
|
||||||
|
|
||||||
|
metrics: $(EXE)
|
||||||
|
sudo LD_LIBRARY_PATH=$(CUDA_HOME)/lib:/usr/ext/lib:${LD_LIBRARY_PATH} LIBRARY_PATH=/usr/ext/lib:${LIBRARY_PATH} nvprof --print-gpu-trace --metrics "eligible_warps_per_cycle,achieved_occupancy,sm_efficiency,ipc" ./$(EXE)
|
||||||
|
|
||||||
|
clean:
|
||||||
|
-rm -fr $(BUILD_DIR) *.exe *.out *~
|
||||||
|
|
||||||
|
MKDIR_P ?= mkdir -p
|
||||||
BIN
cuda/lab3/data/buzz.jpg
Normal file
BIN
cuda/lab3/data/buzz.jpg
Normal file
Binary file not shown.
|
After Width: | Height: | Size: 2.4 MiB |
BIN
cuda/lab3/data/daisy.jpg
Normal file
BIN
cuda/lab3/data/daisy.jpg
Normal file
Binary file not shown.
|
After Width: | Height: | Size: 66 KiB |
BIN
cuda/lab3/data/earth_rise.jpg
Normal file
BIN
cuda/lab3/data/earth_rise.jpg
Normal file
Binary file not shown.
|
After Width: | Height: | Size: 2.8 MiB |
BIN
cuda/lab3/data/fiore.jpg
Normal file
BIN
cuda/lab3/data/fiore.jpg
Normal file
Binary file not shown.
|
After Width: | Height: | Size: 261 KiB |
BIN
cuda/lab3/data/sample.avi
Normal file
BIN
cuda/lab3/data/sample.avi
Normal file
Binary file not shown.
251
cuda/lab3/gemm.cu
Normal file
251
cuda/lab3/gemm.cu
Normal file
|
|
@ -0,0 +1,251 @@
|
||||||
|
/*
|
||||||
|
* 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 <assert.h>
|
||||||
|
#include <stdio.h>
|
||||||
|
#include <stdlib.h>
|
||||||
|
#include <time.h>
|
||||||
|
|
||||||
|
#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<<<dimGrid, dimBlock>>>(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;
|
||||||
|
}
|
||||||
184
cuda/lab3/saxpy.cu
Normal file
184
cuda/lab3/saxpy.cu
Normal file
|
|
@ -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 <assert.h>
|
||||||
|
#include <time.h>
|
||||||
|
#include <stdlib.h>
|
||||||
|
#include <stdio.h>
|
||||||
|
#include <string.h>
|
||||||
|
#include <math.h>
|
||||||
|
#include <cuda_runtime.h>
|
||||||
|
|
||||||
|
#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;
|
||||||
|
}
|
||||||
229
cuda/lab3/sobel.cu
Normal file
229
cuda/lab3/sobel.cu
Normal file
|
|
@ -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 <stdio.h>
|
||||||
|
#include <time.h>
|
||||||
|
#include <sys/time.h>
|
||||||
|
#include <opencv2/opencv.hpp>
|
||||||
|
#include <opencv2/imgcodecs/imgcodecs.hpp>
|
||||||
|
#include <opencv2/objdetect/objdetect.hpp>
|
||||||
|
#include <opencv2/highgui/highgui.hpp>
|
||||||
|
|
||||||
|
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<<<dimGrid, dimBlock>>>(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;
|
||||||
|
}
|
||||||
138
cuda/lab3/utils.c
Normal file
138
cuda/lab3/utils.c
Normal file
|
|
@ -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 <time.h>
|
||||||
|
#include <limits.h>
|
||||||
|
#include <math.h>
|
||||||
|
#include <stdio.h>
|
||||||
|
#include <assert.h>
|
||||||
|
|
||||||
|
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);
|
||||||
|
}
|
||||||
|
|
||||||
|
}
|
||||||
142
cuda/lab3/utils.h
Normal file
142
cuda/lab3/utils.h
Normal file
|
|
@ -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 <stdarg.h>
|
||||||
|
|
||||||
|
#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__*/
|
||||||
Loading…
Reference in a new issue