cory-master/BabelStream
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge changes

Browse Source
This commit is contained in:
Michael Boulton 2015-07-28 11:37:20 +01:00
parent bb0dcce28b
commit b43eb9cf16
4 changed files with 485 additions and 13654 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

30
Makefile
View File

@ -1,6 +1,5 @@
LDLIBS = -l OpenCL
LIBS = -l OpenCL CXXFLAGS = -std=c++11 -O3
FLAGS = -std=c++11 -O3
PLATFORM = $(shell uname -s) PLATFORM = $(shell uname -s)
ifeq ($(PLATFORM), Darwin) ifeq ($(PLATFORM), Darwin)
@ -9,15 +8,24 @@ endif
all: gpu-stream-ocl gpu-stream-cuda all: gpu-stream-ocl gpu-stream-cuda
gpu-stream-ocl: ocl-stream.cpp gpu-stream-ocl: ocl-stream.cpp common.o Makefile
c++ $< $(FLAGS) -o $@ $(LIBS) $(CXX) $(CXXFLAGS) -Wno-deprecated-declarations common.o $< -o $@ $(LDLIBS)
gpu-stream-cuda: cuda-stream.cu common.o: common.cpp Makefile
ifeq ($(shell which nvcc > /dev/null; echo $$?), 0)
nvcc $< $(FLAGS) -o $@ ifeq ($(shell which nvcc),"")
else $(error "Cannot find nvcc, please install CUDA toolkit")
@echo "Cannot find nvcc, please install CUDA";
endif endif
gpu-stream-cuda: cuda-stream.cu common.o Makefile
ifeq ($(shell which nvcc > /dev/null; echo $$?), 0)
nvcc $(CXXFLAGS) common.o $< -o $@
else
$(error "Cannot find nvcc, please install CUDA toolkit")
endif
.PHONY: clean
clean: clean:
rm -f gpu-stream-ocl gpu-stream-cuda rm -f gpu-stream-ocl gpu-stream-cuda *.o

12906
cl.hpp
View File

File diff suppressed because it is too large Load Diff

576
cuda-stream.cu
View File

@ -4,24 +4,11 @@
#include <vector> #include <vector>
#include <chrono> #include <chrono>
#include <cfloat> #include <cfloat>
#include <iomanip>
#include <cmath> #include <cmath>
#include <cuda.h> #include <cuda.h>
#include "common.h"
#define DATATYPE double
unsigned int ARRAY_SIZE = 50000000;
unsigned int NTIMES = 10;
size_t DATATYPE_SIZE = sizeof(double);
bool useFloat = false;
#define MIN(a,b) ((a) < (b)) ? (a) : (b)
#define MAX(a,b) ((a) > (b)) ? (a) : (b)
#define VERSION_STRING "0.0"
void parseArguments(int argc, char *argv[]);
std::string getDeviceName(int device); std::string getDeviceName(int device);
struct invaliddevice : public std::exception struct invaliddevice : public std::exception
@ -54,75 +41,6 @@ void check_cuda_error(void)
} }
} }
void check_solution(void* a, void* b, void* c)
{
// Generate correct solution
double golda = 1.0;
double goldb = 2.0;
double goldc = 0.0;
float goldaf = 1.0;
float goldbf = 2.0;
float goldcf = 0.0;
const double scalar = 3.0;
const float scalarf = 3.0;
for (unsigned int i = 0; i < NTIMES; i++)
{
// Double
goldc = golda;
goldb = scalar * goldc;
goldc = golda + goldb;
golda = goldb + scalar * goldc;
// Float
goldcf = goldaf;
goldbf = scalarf * goldcf;
goldcf = goldaf + goldbf;
goldaf = goldbf + scalarf * goldcf;
}
// Calculate average error
double erra = 0.0;
double errb = 0.0;
double errc = 0.0;
for (unsigned int i = 0; i < ARRAY_SIZE; i++)
{
if (useFloat)
{
erra += fabsf(((float*)a)[i] - goldaf);
errb += fabsf(((float*)b)[i] - goldbf);
errc += fabsf(((float*)c)[i] - goldcf);
}
else
{
erra += fabs(((double*)a)[i] - (double)golda);
errb += fabs(((double*)b)[i] - (double)goldb);
errc += fabs(((double*)c)[i] - (double)goldc);
}
}
erra /= (double)ARRAY_SIZE;
errb /= (double)ARRAY_SIZE;
errc /= (double)ARRAY_SIZE;
double epsi;
if (useFloat) epsi = 1.0E-6;
else epsi = 1.0E-13;
if (erra > epsi)
std::cout
<< "Validation failed on a[]. Average error " << erra
<< std::endl;
if (errb > epsi)
std::cout
<< "Validation failed on b[]. Average error " << errb
<< std::endl;
if (errc > epsi)
std::cout
<< "Validation failed on c[]. Average error " << errc
<< std::endl;
}
template <typename T> template <typename T>
__global__ void copy(const T * a, T * c) __global__ void copy(const T * a, T * c)
{ {
@ -153,8 +71,6 @@ __global__ void triad(T * a, const T * b, const T * c)
a[i] = b[i] + scalar * c[i]; a[i] = b[i] + scalar * c[i];
} }
int deviceIndex = 0;
int main(int argc, char *argv[]) int main(int argc, char *argv[])
{ {
@ -164,218 +80,230 @@ int main(int argc, char *argv[])
<< "Version: " << VERSION_STRING << std::endl << "Version: " << VERSION_STRING << std::endl
<< "Implementation: CUDA" << std::endl; << "Implementation: CUDA" << std::endl;
try parseArguments(argc, argv);
if (NTIMES < 2) throw badntimes();
std::cout << "Precision: ";
if (useFloat) std::cout << "float";
else std::cout << "double";
std::cout << std::endl << std::endl;
if (ARRAY_SIZE % 1024 != 0)
{ {
parseArguments(argc, argv); unsigned int OLD_ARRAY_SIZE = ARRAY_SIZE;
ARRAY_SIZE -= ARRAY_SIZE % 1024;
std::cout
<< "Warning: array size must divide 1024" << std::endl
<< "Resizing array from " << OLD_ARRAY_SIZE
<< " to " << ARRAY_SIZE << std::endl;
}
if (NTIMES < 2) throw badntimes(); // Get precision (used to reset later)
std::streamsize ss = std::cout.precision();
std::cout << "Precision: "; size_t DATATYPE_SIZE;
if (useFloat) std::cout << "float";
else std::cout << "double";
std::cout << std::endl << std::endl;
if (ARRAY_SIZE % 1024 != 0) if (useFloat)
{
DATATYPE_SIZE = sizeof(float);
}
else
{
DATATYPE_SIZE = sizeof(double);
}
// Display number of bytes in array
std::cout << std::setprecision(1) << std::fixed
<< "Array size: " << ARRAY_SIZE*DATATYPE_SIZE/1024.0/1024.0 << " MB"
<< " (=" << ARRAY_SIZE*DATATYPE_SIZE/1024.0/1024.0/1024.0 << " GB)"
<< std::endl;
std::cout << "Total size: " << 3*ARRAY_SIZE*DATATYPE_SIZE/1024.0/1024.0 << " MB"
<< " (=" << 3*ARRAY_SIZE*DATATYPE_SIZE/1024.0/1024.0/1024.0 << " GB)"
<< std::endl;
// Reset precision
std::cout.precision(ss);
// Check device index is in range
int count;
cudaGetDeviceCount(&count);
check_cuda_error();
if (deviceIndex >= count) throw invaliddevice();
cudaSetDevice(deviceIndex);
check_cuda_error();
// Print out device name
std::cout << "Using CUDA device " << getDeviceName(deviceIndex) << std::endl;
// Create host vectors
void * h_a = malloc(ARRAY_SIZE*DATATYPE_SIZE);
void * h_b = malloc(ARRAY_SIZE*DATATYPE_SIZE);
void * h_c = malloc(ARRAY_SIZE*DATATYPE_SIZE);
// Initilise arrays
for (unsigned int i = 0; i < ARRAY_SIZE; i++)
{
if (useFloat)
{ {
unsigned int OLD_ARRAY_SIZE = ARRAY_SIZE; ((float*)h_a)[i] = 1.0;
ARRAY_SIZE -= ARRAY_SIZE % 1024; ((float*)h_b)[i] = 2.0;
std::cout ((float*)h_c)[i] = 0.0;
<< "Warning: array size must divide 1024" << std::endl
<< "Resizing array from " << OLD_ARRAY_SIZE
<< " to " << ARRAY_SIZE << std::endl;
} }
else
// Get precision (used to reset later)
std::streamsize ss = std::cout.precision();
// Display number of bytes in array
std::cout << std::setprecision(1) << std::fixed
<< "Array size: " << ARRAY_SIZE*DATATYPE_SIZE/1024.0/1024.0 << " MB"
<< " (=" << ARRAY_SIZE*DATATYPE_SIZE/1024.0/1024.0/1024.0 << " GB)"
<< std::endl;
std::cout << "Total size: " << 3*ARRAY_SIZE*DATATYPE_SIZE/1024.0/1024.0 << " MB"
<< " (=" << 3*ARRAY_SIZE*DATATYPE_SIZE/1024.0/1024.0/1024.0 << " GB)"
<< std::endl;
// Reset precision
std::cout.precision(ss);
// Check device index is in range
int count;
cudaGetDeviceCount(&count);
check_cuda_error();
if (deviceIndex >= count) throw invaliddevice();
cudaSetDevice(deviceIndex);
check_cuda_error();
// Print out device name
std::cout << "Using CUDA device " << getDeviceName(deviceIndex) << std::endl;
// Create host vectors
void * h_a = malloc(ARRAY_SIZE*DATATYPE_SIZE);
void * h_b = malloc(ARRAY_SIZE*DATATYPE_SIZE);
void * h_c = malloc(ARRAY_SIZE*DATATYPE_SIZE);
// Initilise arrays
for (unsigned int i = 0; i < ARRAY_SIZE; i++)
{ {
if (useFloat) ((double*)h_a)[i] = 1.0;
{ ((double*)h_b)[i] = 2.0;
((float*)h_a)[i] = 1.0; ((double*)h_c)[i] = 0.0;
((float*)h_b)[i] = 2.0;
((float*)h_c)[i] = 0.0;
}
else
{
((double*)h_a)[i] = 1.0;
((double*)h_b)[i] = 2.0;
((double*)h_c)[i] = 0.0;
}
} }
}
// Create device buffers // Create device buffers
void * d_a, * d_b, *d_c; void * d_a, * d_b, *d_c;
cudaMalloc(&d_a, ARRAY_SIZE*DATATYPE_SIZE); cudaMalloc(&d_a, ARRAY_SIZE*DATATYPE_SIZE);
check_cuda_error(); check_cuda_error();
cudaMalloc(&d_b, ARRAY_SIZE*DATATYPE_SIZE); cudaMalloc(&d_b, ARRAY_SIZE*DATATYPE_SIZE);
check_cuda_error(); check_cuda_error();
cudaMalloc(&d_c, ARRAY_SIZE*DATATYPE_SIZE); cudaMalloc(&d_c, ARRAY_SIZE*DATATYPE_SIZE);
check_cuda_error(); check_cuda_error();
// Copy host memory to device // Copy host memory to device
cudaMemcpy(d_a, h_a, ARRAY_SIZE*DATATYPE_SIZE, cudaMemcpyHostToDevice); cudaMemcpy(d_a, h_a, ARRAY_SIZE*DATATYPE_SIZE, cudaMemcpyHostToDevice);
check_cuda_error(); check_cuda_error();
cudaMemcpy(d_b, h_b, ARRAY_SIZE*DATATYPE_SIZE, cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, ARRAY_SIZE*DATATYPE_SIZE, cudaMemcpyHostToDevice);
check_cuda_error(); check_cuda_error();
cudaMemcpy(d_c, h_c, ARRAY_SIZE*DATATYPE_SIZE, cudaMemcpyHostToDevice); cudaMemcpy(d_c, h_c, ARRAY_SIZE*DATATYPE_SIZE, cudaMemcpyHostToDevice);
check_cuda_error(); check_cuda_error();
// Make sure the copies are finished // Make sure the copies are finished
cudaDeviceSynchronize();
check_cuda_error();
// List of times
std::vector< std::vector<double> > timings;
// Declare timers
std::chrono::high_resolution_clock::time_point t1, t2;
// Main loop
for (unsigned int k = 0; k < NTIMES; k++)
{
std::vector<double> times;
t1 = std::chrono::high_resolution_clock::now();
if (useFloat)
copy<<<ARRAY_SIZE/1024, 1024>>>((float*)d_a, (float*)d_c);
else
copy<<<ARRAY_SIZE/1024, 1024>>>((double*)d_a, (double*)d_c);
check_cuda_error();
cudaDeviceSynchronize(); cudaDeviceSynchronize();
check_cuda_error(); check_cuda_error();
t2 = std::chrono::high_resolution_clock::now();
// List of times times.push_back(std::chrono::duration_cast<std::chrono::duration<double> >(t2 - t1).count());
std::vector< std::vector<double> > timings;
// Declare timers
std::chrono::high_resolution_clock::time_point t1, t2;
// Main loop
for (unsigned int k = 0; k < NTIMES; k++)
{
std::vector<double> times;
t1 = std::chrono::high_resolution_clock::now();
if (useFloat)
copy<<<ARRAY_SIZE/1024, 1024>>>((float*)d_a, (float*)d_c);
else
copy<<<ARRAY_SIZE/1024, 1024>>>((double*)d_a, (double*)d_c);
check_cuda_error();
cudaDeviceSynchronize();
check_cuda_error();
t2 = std::chrono::high_resolution_clock::now();
times.push_back(std::chrono::duration_cast<std::chrono::duration<double> >(t2 - t1).count());
t1 = std::chrono::high_resolution_clock::now(); t1 = std::chrono::high_resolution_clock::now();
if (useFloat) if (useFloat)
mul<<<ARRAY_SIZE/1024, 1024>>>((float*)d_b, (float*)d_c); mul<<<ARRAY_SIZE/1024, 1024>>>((float*)d_b, (float*)d_c);
else else
mul<<<ARRAY_SIZE/1024, 1024>>>((double*)d_b, (double*)d_c); mul<<<ARRAY_SIZE/1024, 1024>>>((double*)d_b, (double*)d_c);
check_cuda_error();
cudaDeviceSynchronize();
check_cuda_error();
t2 = std::chrono::high_resolution_clock::now();
times.push_back(std::chrono::duration_cast<std::chrono::duration<double> >(t2 - t1).count());
t1 = std::chrono::high_resolution_clock::now();
if (useFloat)
add<<<ARRAY_SIZE/1024, 1024>>>((float*)d_a, (float*)d_b, (float*)d_c);
else
add<<<ARRAY_SIZE/1024, 1024>>>((double*)d_a, (double*)d_b, (double*)d_c);
check_cuda_error();
cudaDeviceSynchronize();
check_cuda_error();
t2 = std::chrono::high_resolution_clock::now();
times.push_back(std::chrono::duration_cast<std::chrono::duration<double> >(t2 - t1).count());
t1 = std::chrono::high_resolution_clock::now();
if (useFloat)
triad<<<ARRAY_SIZE/1024, 1024>>>((float*)d_a, (float*)d_b, (float*)d_c);
else
triad<<<ARRAY_SIZE/1024, 1024>>>((double*)d_a, (double*)d_b, (double*)d_c);
check_cuda_error();
cudaDeviceSynchronize();
check_cuda_error();
t2 = std::chrono::high_resolution_clock::now();
times.push_back(std::chrono::duration_cast<std::chrono::duration<double> >(t2 - t1).count());
timings.push_back(times);
}
// Check solutions
cudaMemcpy(h_a, d_a, ARRAY_SIZE*DATATYPE_SIZE, cudaMemcpyDeviceToHost);
check_cuda_error(); check_cuda_error();
cudaMemcpy(h_b, d_b, ARRAY_SIZE*DATATYPE_SIZE, cudaMemcpyDeviceToHost); cudaDeviceSynchronize();
check_cuda_error(); check_cuda_error();
cudaMemcpy(h_c, d_c, ARRAY_SIZE*DATATYPE_SIZE, cudaMemcpyDeviceToHost); t2 = std::chrono::high_resolution_clock::now();
times.push_back(std::chrono::duration_cast<std::chrono::duration<double> >(t2 - t1).count());
t1 = std::chrono::high_resolution_clock::now();
if (useFloat)
add<<<ARRAY_SIZE/1024, 1024>>>((float*)d_a, (float*)d_b, (float*)d_c);
else
add<<<ARRAY_SIZE/1024, 1024>>>((double*)d_a, (double*)d_b, (double*)d_c);
check_cuda_error(); check_cuda_error();
check_solution(h_a, h_b, h_c); cudaDeviceSynchronize();
check_cuda_error();
t2 = std::chrono::high_resolution_clock::now();
times.push_back(std::chrono::duration_cast<std::chrono::duration<double> >(t2 - t1).count());
// Crunch results
size_t sizes[4] = {
2 * DATATYPE_SIZE * ARRAY_SIZE,
2 * DATATYPE_SIZE * ARRAY_SIZE,
3 * DATATYPE_SIZE * ARRAY_SIZE,
3 * DATATYPE_SIZE * ARRAY_SIZE
};
double min[4] = {DBL_MAX, DBL_MAX, DBL_MAX, DBL_MAX};
double max[4] = {0.0, 0.0, 0.0, 0.0};
double avg[4] = {0.0, 0.0, 0.0, 0.0};
// Ignore first result
for (unsigned int i = 1; i < NTIMES; i++)
{
for (int j = 0; j < 4; j++)
{
avg[j] += timings[i][j];
min[j] = MIN(min[j], timings[i][j]);
max[j] = MAX(max[j], timings[i][j]);
}
}
for (int j = 0; j < 4; j++)
avg[j] /= (double)(NTIMES-1);
// Display results t1 = std::chrono::high_resolution_clock::now();
std::string labels[] = {"Copy", "Mul", "Add", "Triad"}; if (useFloat)
std::cout triad<<<ARRAY_SIZE/1024, 1024>>>((float*)d_a, (float*)d_b, (float*)d_c);
<< std::left << std::setw(12) << "Function" else
<< std::left << std::setw(12) << "MBytes/sec" triad<<<ARRAY_SIZE/1024, 1024>>>((double*)d_a, (double*)d_b, (double*)d_c);
<< std::left << std::setw(12) << "Min (sec)" check_cuda_error();
<< std::left << std::setw(12) << "Max" cudaDeviceSynchronize();
<< std::left << std::setw(12) << "Average" check_cuda_error();
<< std::endl; t2 = std::chrono::high_resolution_clock::now();
for (int j = 0; j < 4; j++) times.push_back(std::chrono::duration_cast<std::chrono::duration<double> >(t2 - t1).count());
{
std::cout timings.push_back(times);
<< std::left << std::setw(12) << labels[j]
<< std::left << std::setw(12) << std::setprecision(3) << 1.0E-06 * sizes[j]/min[j]
<< std::left << std::setw(12) << std::setprecision(5) << min[j]
<< std::left << std::setw(12) << std::setprecision(5) << max[j]
<< std::left << std::setw(12) << std::setprecision(5) << avg[j]
<< std::endl;
}
} }
catch (std::exception& e)
// Check solutions
cudaMemcpy(h_a, d_a, ARRAY_SIZE*DATATYPE_SIZE, cudaMemcpyDeviceToHost);
check_cuda_error();
cudaMemcpy(h_b, d_b, ARRAY_SIZE*DATATYPE_SIZE, cudaMemcpyDeviceToHost);
check_cuda_error();
cudaMemcpy(h_c, d_c, ARRAY_SIZE*DATATYPE_SIZE, cudaMemcpyDeviceToHost);
check_cuda_error();
if (useFloat)
{ {
std::cerr check_solution<float>(h_a, h_b, h_c);
<< "Error: " }
<< e.what() else
{
check_solution<double>(h_a, h_b, h_c);
}
// Crunch results
size_t sizes[4] = {
2 * DATATYPE_SIZE * ARRAY_SIZE,
2 * DATATYPE_SIZE * ARRAY_SIZE,
3 * DATATYPE_SIZE * ARRAY_SIZE,
3 * DATATYPE_SIZE * ARRAY_SIZE
};
double min[4] = {DBL_MAX, DBL_MAX, DBL_MAX, DBL_MAX};
double max[4] = {0.0, 0.0, 0.0, 0.0};
double avg[4] = {0.0, 0.0, 0.0, 0.0};
// Ignore first result
for (unsigned int i = 1; i < NTIMES; i++)
{
for (int j = 0; j < 4; j++)
{
avg[j] += timings[i][j];
min[j] = std::min(min[j], timings[i][j]);
max[j] = std::max(max[j], timings[i][j]);
}
}
for (int j = 0; j < 4; j++)
avg[j] /= (double)(NTIMES-1);
// Display results
std::string labels[] = {"Copy", "Mul", "Add", "Triad"};
std::cout
<< std::left << std::setw(12) << "Function"
<< std::left << std::setw(12) << "MBytes/sec"
<< std::left << std::setw(12) << "Min (sec)"
<< std::left << std::setw(12) << "Max"
<< std::left << std::setw(12) << "Average"
<< std::endl;
for (int j = 0; j < 4; j++)
{
std::cout
<< std::left << std::setw(12) << labels[j]
<< std::left << std::setw(12) << std::setprecision(3) << 1.0E-06 * sizes[j]/min[j]
<< std::left << std::setw(12) << std::setprecision(5) << min[j]
<< std::left << std::setw(12) << std::setprecision(5) << max[j]
<< std::left << std::setw(12) << std::setprecision(5) << avg[j]
<< std::endl; << std::endl;
} }
} }
std::string getDeviceName(int device) std::string getDeviceName(int device)
@ -386,98 +314,28 @@ std::string getDeviceName(int device)
return std::string(prop.name); return std::string(prop.name);
} }
void listDevices(void)
int parseUInt(const char *str, unsigned int *output)
{ {
char *next; // Get number of devices
*output = strtoul(str, &next, 10); int count;
return !strlen(next); cudaGetDeviceCount(&count);
} check_cuda_error();
int parseInt(const char *str, int *output) // Print device names
{ if (count == 0)
char *next;
*output = strtol(str, &next, 10);
return !strlen(next);
}
void parseArguments(int argc, char *argv[])
{
for (int i = 1; i < argc; i++)
{ {
if (!strcmp(argv[i], "--list")) std::cout << "No devices found." << std::endl;
}
else
{
std::cout << std::endl;
std::cout << "Devices:" << std::endl;
for (int i = 0; i < count; i++)
{ {
// Get number of devices std::cout << i << ": " << getDeviceName(i) << std::endl;
int count;
cudaGetDeviceCount(&count);
check_cuda_error(); check_cuda_error();
// Print device names
if (count == 0)
{
std::cout << "No devices found." << std::endl;
}
else
{
std::cout << std::endl;
std::cout << "Devices:" << std::endl;
for (int i = 0; i < count; i++)
{
std::cout << i << ": " << getDeviceName(i) << std::endl;
check_cuda_error();
}
std::cout << std::endl;
}
exit(0);
}
else if (!strcmp(argv[i], "--device"))
{
if (++i >= argc || !parseInt(argv[i], &deviceIndex))
{
std::cout << "Invalid device index" << std::endl;
exit(1);
}
}
else if (!strcmp(argv[i], "--arraysize") || !strcmp(argv[i], "-s"))
{
if (++i >= argc || !parseUInt(argv[i], &ARRAY_SIZE))
{
std::cout << "Invalid array size" << std::endl;
exit(1);
}
}
else if (!strcmp(argv[i], "--numtimes") || !strcmp(argv[i], "-n"))
{
if (++i >= argc || !parseUInt(argv[i], &NTIMES))
{
std::cout << "Invalid number of times" << std::endl;
exit(1);
}
}
else if (!strcmp(argv[i], "--float"))
{
useFloat = true;
DATATYPE_SIZE = sizeof(float);
}
else if (!strcmp(argv[i], "--help") || !strcmp(argv[i], "-h"))
{
std::cout << std::endl;
std::cout << "Usage: ./gpu-stream-cuda [OPTIONS]" << std::endl << std::endl;
std::cout << "Options:" << std::endl;
std::cout << " -h --help Print the message" << std::endl;
std::cout << " --list List available devices" << std::endl;
std::cout << " --device INDEX Select device at INDEX" << std::endl;
std::cout << " -s --arraysize SIZE Use SIZE elements in the array" << std::endl;
std::cout << " -n --numtimes NUM Run the test NUM times (NUM >= 2)" << std::endl;
std::cout << " --float Use floats (rather than doubles)" << std::endl;
std::cout << std::endl;
exit(0);
}
else
{
std::cout << "Unrecognized argument '" << argv[i] << "' (try '--help')"
<< std::endl;
exit(1);
} }
std::cout << std::endl;
} }
} }

627
ocl-stream.cpp
View File

@ -4,24 +4,12 @@
#include <vector> #include <vector>
#include <chrono> #include <chrono>
#include <cfloat> #include <cfloat>
#include <iomanip>
#include <cmath> #include <cmath>
#define __CL_ENABLE_EXCEPTIONS #define __CL_ENABLE_EXCEPTIONS
#include "cl.hpp" #include "CL/cl.hpp"
#include "common.h"
unsigned int ARRAY_SIZE = 50000000;
unsigned int NTIMES = 10;
size_t DATATYPE_SIZE = sizeof(double);
bool useFloat = false;
#define MIN(a,b) ((a) < (b)) ? (a) : (b)
#define MAX(a,b) ((a) > (b)) ? (a) : (b)
#define VERSION_STRING "0.0"
void parseArguments(int argc, char *argv[]);
std::string getDeviceName(const cl::Device& device); std::string getDeviceName(const cl::Device& device);
unsigned getDeviceList(std::vector<cl::Device>& devices); unsigned getDeviceList(std::vector<cl::Device>& devices);
@ -50,76 +38,6 @@ struct badntimes : public std::exception
}; };
void check_solution(void* a, void* b, void* c)
{
// Generate correct solution
double golda = 1.0;
double goldb = 2.0;
double goldc = 0.0;
float goldaf = 1.0;
float goldbf = 2.0;
float goldcf = 0.0;
const double scalar = 3.0;
const float scalarf = 3.0;
for (unsigned int i = 0; i < NTIMES; i++)
{
// Double
goldc = golda;
goldb = scalar * goldc;
goldc = golda + goldb;
golda = goldb + scalar * goldc;
// Float
goldcf = goldaf;
goldbf = scalarf * goldcf;
goldcf = goldaf + goldbf;
goldaf = goldbf + scalarf * goldcf;
}
// Calculate average error
double erra = 0.0;
double errb = 0.0;
double errc = 0.0;
for (unsigned int i = 0; i < ARRAY_SIZE; i++)
{
if (useFloat)
{
erra += fabsf(((float*)a)[i] - goldaf);
errb += fabsf(((float*)b)[i] - goldbf);
errc += fabsf(((float*)c)[i] - goldcf);
}
else
{
erra += fabs(((double*)a)[i] - (double)golda);
errb += fabs(((double*)b)[i] - (double)goldb);
errc += fabs(((double*)c)[i] - (double)goldc);
}
}
erra /= (double)ARRAY_SIZE;
errb /= (double)ARRAY_SIZE;
errc /= (double)ARRAY_SIZE;
double epsi;
if (useFloat) epsi = 1.0E-6;
else epsi = 1.0E-13;
if (erra > epsi)
std::cout
<< "Validation failed on a[]. Average error " << erra
<< std::endl;
if (errb > epsi)
std::cout
<< "Validation failed on b[]. Average error " << errb
<< std::endl;
if (errc > epsi)
std::cout
<< "Validation failed on c[]. Average error " << errc
<< std::endl;
}
cl_uint deviceIndex = 0;
int main(int argc, char *argv[]) int main(int argc, char *argv[])
{ {
@ -129,236 +47,252 @@ int main(int argc, char *argv[])
<< "Version: " << VERSION_STRING << std::endl << "Version: " << VERSION_STRING << std::endl
<< "Implementation: OpenCL" << std::endl; << "Implementation: OpenCL" << std::endl;
try parseArguments(argc, argv);
if (NTIMES < 2) throw badntimes();
std::cout << "Precision: ";
if (useFloat) std::cout << "float";
else std::cout << "double";
std::cout << std::endl << std::endl;
if (ARRAY_SIZE % 1024 != 0)
{ {
parseArguments(argc, argv); unsigned int OLD_ARRAY_SIZE = ARRAY_SIZE;
ARRAY_SIZE -= ARRAY_SIZE % 1024;
std::cout
<< "Warning: array size must divide 1024" << std::endl
<< "Resizing array from " << OLD_ARRAY_SIZE
<< " to " << ARRAY_SIZE << std::endl;
}
if (NTIMES < 2) throw badntimes(); // Get precision (used to reset later)
std::streamsize ss = std::cout.precision();
std::cout << "Precision: "; size_t DATATYPE_SIZE;
if (useFloat) std::cout << "float";
else std::cout << "double";
std::cout << std::endl << std::endl;
// Get precision (used to reset later) if (useFloat)
std::streamsize ss = std::cout.precision(); {
DATATYPE_SIZE = sizeof(float);
}
else
{
DATATYPE_SIZE = sizeof(double);
}
// Display number of bytes in array // Display number of bytes in array
std::cout << std::setprecision(1) << std::fixed std::cout << std::setprecision(1) << std::fixed
<< "Array size: " << ARRAY_SIZE*DATATYPE_SIZE/1024.0/1024.0 << " MB" << "Array size: " << ARRAY_SIZE*DATATYPE_SIZE/1024.0/1024.0 << " MB"
<< " (=" << ARRAY_SIZE*DATATYPE_SIZE/1024.0/1024.0/1024.0 << " GB)" << " (=" << ARRAY_SIZE*DATATYPE_SIZE/1024.0/1024.0/1024.0 << " GB)"
<< std::endl; << std::endl;
std::cout << "Total size: " << 3*ARRAY_SIZE*DATATYPE_SIZE/1024.0/1024.0 << " MB" std::cout << "Total size: " << 3*ARRAY_SIZE*DATATYPE_SIZE/1024.0/1024.0 << " MB"
<< " (=" << 3*ARRAY_SIZE*DATATYPE_SIZE/1024.0/1024.0/1024.0 << " GB)" << " (=" << 3*ARRAY_SIZE*DATATYPE_SIZE/1024.0/1024.0/1024.0 << " GB)"
<< std::endl; << std::endl;
// Reset precision
std::cout.precision(ss);
// Open the Kernel source // Reset precision
std::cout.precision(ss);
// Open the Kernel source
std::string kernels;
{
std::ifstream in("ocl-stream-kernels.cl"); std::ifstream in("ocl-stream-kernels.cl");
if (!in.is_open()) throw badfile(); if (!in.is_open()) throw badfile();
std::string kernels(std::istreambuf_iterator<char>(in), (std::istreambuf_iterator<char>())); kernels = std::string (std::istreambuf_iterator<char>(in), (std::istreambuf_iterator<char>()));
}
// Setup OpenCL
// Setup OpenCL
// Get list of devices
std::vector<cl::Device> devices; // Get list of devices
getDeviceList(devices); std::vector<cl::Device> devices;
getDeviceList(devices);
// Check device index is in range
if (deviceIndex >= devices.size()) throw invaliddevice(); // Check device index is in range
if (deviceIndex >= devices.size()) throw invaliddevice();
cl::Device device = devices[deviceIndex];
cl::Device device = devices[deviceIndex];
cl::Context context(device);
cl::CommandQueue queue(context); cl::Context context(device);
cl::Program program(context, kernels); cl::CommandQueue queue(context);
cl::Program program(context, kernels);
// Print out device name
std::string name = getDeviceName(device); // Print out device name
std::cout << "Using OpenCL device " << name << std::endl; std::string name = getDeviceName(device);
std::cout << "Using OpenCL device " << name << std::endl;
try
{ try
std::string options = ""; {
if (useFloat) std::string options = "";
options = "-DFLOAT"; if (useFloat)
program.build(options.c_str()); options = "-DFLOAT";
} program.build(options.c_str());
catch (cl::Error& e)
{
std::vector<cl::Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();
std::string buildlog = program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(devices[0]);
std::cerr
<< "Build error:"
<< buildlog
<< std::endl;
throw e;
}
cl::make_kernel<cl::Buffer, cl::Buffer> copy(program, "copy");
cl::make_kernel<cl::Buffer, cl::Buffer> mul(program, "mul");
cl::make_kernel<cl::Buffer, cl::Buffer, cl::Buffer> add(program, "add");
cl::make_kernel<cl::Buffer, cl::Buffer, cl::Buffer> triad(program, "triad");
// Create host vectors
void *h_a = malloc(ARRAY_SIZE * DATATYPE_SIZE);
void *h_b = malloc(ARRAY_SIZE * DATATYPE_SIZE);
void *h_c = malloc(ARRAY_SIZE * DATATYPE_SIZE);
// Initilise arrays
for (unsigned int i = 0; i < ARRAY_SIZE; i++)
{
if (useFloat)
{
((float*)h_a)[i] = 1.0;
((float*)h_b)[i] = 2.0;
((float*)h_c)[i] = 0.0;
}
else
{
((double*)h_a)[i] = 1.0;
((double*)h_b)[i] = 2.0;
((double*)h_c)[i] = 0.0;
}
}
// Create device buffers
cl::Buffer d_a(context, CL_MEM_READ_WRITE, DATATYPE_SIZE * ARRAY_SIZE);
cl::Buffer d_b(context, CL_MEM_READ_WRITE, DATATYPE_SIZE * ARRAY_SIZE);
cl::Buffer d_c(context, CL_MEM_READ_WRITE, DATATYPE_SIZE * ARRAY_SIZE);
// Copy host memory to device
queue.enqueueWriteBuffer(d_a, CL_FALSE, 0, ARRAY_SIZE*DATATYPE_SIZE, h_a);
queue.enqueueWriteBuffer(d_b, CL_FALSE, 0, ARRAY_SIZE*DATATYPE_SIZE, h_b);
queue.enqueueWriteBuffer(d_c, CL_FALSE, 0, ARRAY_SIZE*DATATYPE_SIZE, h_c);
// Make sure the copies are finished
queue.finish();
// List of times
std::vector< std::vector<double> > timings;
// Declare timers
std::chrono::high_resolution_clock::time_point t1, t2;
// Main loop
for (unsigned int k = 0; k < NTIMES; k++)
{
std::vector<double> times;
t1 = std::chrono::high_resolution_clock::now();
copy(
cl::EnqueueArgs(
queue,
cl::NDRange(ARRAY_SIZE)),
d_a, d_c);
queue.finish();
t2 = std::chrono::high_resolution_clock::now();
times.push_back(std::chrono::duration_cast<std::chrono::duration<double> >(t2 - t1).count());
t1 = std::chrono::high_resolution_clock::now();
mul(
cl::EnqueueArgs(
queue,
cl::NDRange(ARRAY_SIZE)),
d_b, d_c);
queue.finish();
t2 = std::chrono::high_resolution_clock::now();
times.push_back(std::chrono::duration_cast<std::chrono::duration<double> >(t2 - t1).count());
t1 = std::chrono::high_resolution_clock::now();
add(
cl::EnqueueArgs(
queue,
cl::NDRange(ARRAY_SIZE)),
d_a, d_b, d_c);
queue.finish();
t2 = std::chrono::high_resolution_clock::now();
times.push_back(std::chrono::duration_cast<std::chrono::duration<double> >(t2 - t1).count());
t1 = std::chrono::high_resolution_clock::now();
triad(
cl::EnqueueArgs(
queue,
cl::NDRange(ARRAY_SIZE)),
d_a, d_b, d_c);
queue.finish();
t2 = std::chrono::high_resolution_clock::now();
times.push_back(std::chrono::duration_cast<std::chrono::duration<double> >(t2 - t1).count());
timings.push_back(times);
}
// Check solutions
queue.enqueueReadBuffer(d_a, CL_FALSE, 0, ARRAY_SIZE*DATATYPE_SIZE, h_a);
queue.enqueueReadBuffer(d_b, CL_FALSE, 0, ARRAY_SIZE*DATATYPE_SIZE, h_b);
queue.enqueueReadBuffer(d_c, CL_FALSE, 0, ARRAY_SIZE*DATATYPE_SIZE, h_c);
queue.finish();
check_solution(h_a, h_b, h_c);
// Crunch results
size_t sizes[4] = {
2 * DATATYPE_SIZE * ARRAY_SIZE,
2 * DATATYPE_SIZE * ARRAY_SIZE,
3 * DATATYPE_SIZE * ARRAY_SIZE,
3 * DATATYPE_SIZE * ARRAY_SIZE
};
double min[4] = {DBL_MAX, DBL_MAX, DBL_MAX, DBL_MAX};
double max[4] = {0.0, 0.0, 0.0, 0.0};
double avg[4] = {0.0, 0.0, 0.0, 0.0};
// Ignore first result
for (unsigned int i = 1; i < NTIMES; i++)
{
for (int j = 0; j < 4; j++)
{
avg[j] += timings[i][j];
min[j] = MIN(min[j], timings[i][j]);
max[j] = MAX(max[j], timings[i][j]);
}
}
for (int j = 0; j < 4; j++)
avg[j] /= (double)(NTIMES-1);
// Display results
std::string labels[] = {"Copy", "Mul", "Add", "Triad"};
std::cout
<< std::left << std::setw(12) << "Function"
<< std::left << std::setw(12) << "MBytes/sec"
<< std::left << std::setw(12) << "Min (sec)"
<< std::left << std::setw(12) << "Max"
<< std::left << std::setw(12) << "Average"
<< std::endl;
for (int j = 0; j < 4; j++)
{
std::cout
<< std::left << std::setw(12) << labels[j]
<< std::left << std::setw(12) << std::setprecision(3) << 1.0E-06 * sizes[j]/min[j]
<< std::left << std::setw(12) << std::setprecision(5) << min[j]
<< std::left << std::setw(12) << std::setprecision(5) << max[j]
<< std::left << std::setw(12) << std::setprecision(5) << avg[j]
<< std::endl;
}
} }
// Catch OpenCL Errors and display information
catch (cl::Error& e) catch (cl::Error& e)
{ {
std::vector<cl::Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();
std::string buildlog = program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(devices[0]);
std::cerr std::cerr
<< "Error: " << "Build error:"
<< e.what() << buildlog
<< "(" << e.err() << ")"
<< std::endl; << std::endl;
throw e;
} }
catch (std::exception& e)
cl::make_kernel<cl::Buffer, cl::Buffer> copy(program, "copy");
cl::make_kernel<cl::Buffer, cl::Buffer> mul(program, "mul");
cl::make_kernel<cl::Buffer, cl::Buffer, cl::Buffer> add(program, "add");
cl::make_kernel<cl::Buffer, cl::Buffer, cl::Buffer> triad(program, "triad");
// Create host vectors
void * h_a = malloc(ARRAY_SIZE*DATATYPE_SIZE);
void * h_b = malloc(ARRAY_SIZE*DATATYPE_SIZE);
void * h_c = malloc(ARRAY_SIZE*DATATYPE_SIZE);
// Initilise arrays
for (unsigned int i = 0; i < ARRAY_SIZE; i++)
{ {
std::cerr if (useFloat)
<< "Error: " {
<< e.what() ((float*)h_a)[i] = 1.0;
((float*)h_b)[i] = 2.0;
((float*)h_c)[i] = 0.0;
}
else
{
((double*)h_a)[i] = 1.0;
((double*)h_b)[i] = 2.0;
((double*)h_c)[i] = 0.0;
}
}
// Create device buffers
cl::Buffer d_a(context, CL_MEM_READ_WRITE, DATATYPE_SIZE * ARRAY_SIZE);
cl::Buffer d_b(context, CL_MEM_READ_WRITE, DATATYPE_SIZE * ARRAY_SIZE);
cl::Buffer d_c(context, CL_MEM_READ_WRITE, DATATYPE_SIZE * ARRAY_SIZE);
// Copy host memory to device
queue.enqueueWriteBuffer(d_a, CL_FALSE, 0, ARRAY_SIZE*DATATYPE_SIZE, h_a);
queue.enqueueWriteBuffer(d_b, CL_FALSE, 0, ARRAY_SIZE*DATATYPE_SIZE, h_b);
queue.enqueueWriteBuffer(d_c, CL_FALSE, 0, ARRAY_SIZE*DATATYPE_SIZE, h_c);
// Make sure the copies are finished
queue.finish();
// List of times
std::vector< std::vector<double> > timings;
// Declare timers
std::chrono::high_resolution_clock::time_point t1, t2;
// Main loop
for (unsigned int k = 0; k < NTIMES; k++)
{
std::vector<double> times;
t1 = std::chrono::high_resolution_clock::now();
copy(
cl::EnqueueArgs(
queue,
cl::NDRange(ARRAY_SIZE)),
d_a, d_c);
queue.finish();
t2 = std::chrono::high_resolution_clock::now();
times.push_back(std::chrono::duration_cast<std::chrono::duration<double> >(t2 - t1).count());
t1 = std::chrono::high_resolution_clock::now();
mul(
cl::EnqueueArgs(
queue,
cl::NDRange(ARRAY_SIZE)),
d_b, d_c);
queue.finish();
t2 = std::chrono::high_resolution_clock::now();
times.push_back(std::chrono::duration_cast<std::chrono::duration<double> >(t2 - t1).count());
t1 = std::chrono::high_resolution_clock::now();
add(
cl::EnqueueArgs(
queue,
cl::NDRange(ARRAY_SIZE)),
d_a, d_b, d_c);
queue.finish();
t2 = std::chrono::high_resolution_clock::now();
times.push_back(std::chrono::duration_cast<std::chrono::duration<double> >(t2 - t1).count());
t1 = std::chrono::high_resolution_clock::now();
triad(
cl::EnqueueArgs(
queue,
cl::NDRange(ARRAY_SIZE)),
d_a, d_b, d_c);
queue.finish();
t2 = std::chrono::high_resolution_clock::now();
times.push_back(std::chrono::duration_cast<std::chrono::duration<double> >(t2 - t1).count());
timings.push_back(times);
}
// Check solutions
queue.enqueueReadBuffer(d_a, CL_FALSE, 0, ARRAY_SIZE*DATATYPE_SIZE, h_a);
queue.enqueueReadBuffer(d_b, CL_FALSE, 0, ARRAY_SIZE*DATATYPE_SIZE, h_b);
queue.enqueueReadBuffer(d_c, CL_FALSE, 0, ARRAY_SIZE*DATATYPE_SIZE, h_c);
queue.finish();
if (useFloat)
{
check_solution<float>(h_a, h_b, h_c);
}
else
{
check_solution<double>(h_a, h_b, h_c);
}
// Crunch results
size_t sizes[4] = {
2 * DATATYPE_SIZE * ARRAY_SIZE,
2 * DATATYPE_SIZE * ARRAY_SIZE,
3 * DATATYPE_SIZE * ARRAY_SIZE,
3 * DATATYPE_SIZE * ARRAY_SIZE
};
double min[4] = {DBL_MAX, DBL_MAX, DBL_MAX, DBL_MAX};
double max[4] = {0.0, 0.0, 0.0, 0.0};
double avg[4] = {0.0, 0.0, 0.0, 0.0};
// Ignore first result
for (unsigned int i = 1; i < NTIMES; i++)
{
for (int j = 0; j < 4; j++)
{
avg[j] += timings[i][j];
min[j] = std::min(min[j], timings[i][j]);
max[j] = std::max(max[j], timings[i][j]);
}
}
for (int j = 0; j < 4; j++)
avg[j] /= (double)(NTIMES-1);
// Display results
std::string labels[] = {"Copy", "Mul", "Add", "Triad"};
std::cout
<< std::left << std::setw(12) << "Function"
<< std::left << std::setw(12) << "MBytes/sec"
<< std::left << std::setw(12) << "Min (sec)"
<< std::left << std::setw(12) << "Max"
<< std::left << std::setw(12) << "Average"
<< std::endl;
for (int j = 0; j < 4; j++)
{
std::cout
<< std::left << std::setw(12) << labels[j]
<< std::left << std::setw(12) << std::setprecision(3) << 1.0E-06 * sizes[j]/min[j]
<< std::left << std::setw(12) << std::setprecision(5) << min[j]
<< std::left << std::setw(12) << std::setprecision(5) << max[j]
<< std::left << std::setw(12) << std::setprecision(5) << avg[j]
<< std::endl; << std::endl;
} }
} }
@ -396,89 +330,26 @@ std::string getDeviceName(const cl::Device& device)
return name; return name;
} }
void listDevices(void)
int parseUInt(const char *str, cl_uint *output)
{ {
char *next; // Get list of devices
*output = strtoul(str, &next, 10); std::vector<cl::Device> devices;
return !strlen(next); getDeviceList(devices);
}
void parseArguments(int argc, char *argv[]) // Print device names
{ if (devices.size() == 0)
for (int i = 1; i < argc; i++)
{ {
if (!strcmp(argv[i], "--list")) std::cout << "No devices found." << std::endl;
}
else
{
std::cout << std::endl;
std::cout << "Devices:" << std::endl;
for (unsigned i = 0; i < devices.size(); i++)
{ {
// Get list of devices std::cout << i << ": " << getDeviceName(devices[i]) << std::endl;
std::vector<cl::Device> devices;
getDeviceList(devices);
// Print device names
if (devices.size() == 0)
{
std::cout << "No devices found." << std::endl;
}
else
{
std::cout << std::endl;
std::cout << "Devices:" << std::endl;
for (unsigned i = 0; i < devices.size(); i++)
{
std::cout << i << ": " << getDeviceName(devices[i]) << std::endl;
}
std::cout << std::endl;
}
exit(0);
}
else if (!strcmp(argv[i], "--device"))
{
if (++i >= argc || !parseUInt(argv[i], &deviceIndex))
{
std::cout << "Invalid device index" << std::endl;
exit(1);
}
}
else if (!strcmp(argv[i], "--arraysize") || !strcmp(argv[i], "-s"))
{
if (++i >= argc || !parseUInt(argv[i], &ARRAY_SIZE))
{
std::cout << "Invalid array size" << std::endl;
exit(1);
}
}
else if (!strcmp(argv[i], "--numtimes") || !strcmp(argv[i], "-n"))
{
if (++i >= argc || !parseUInt(argv[i], &NTIMES))
{
std::cout << "Invalid number of times" << std::endl;
exit(1);
}
}
else if (!strcmp(argv[i], "--float"))
{
useFloat = true;
DATATYPE_SIZE = sizeof(float);
}
else if (!strcmp(argv[i], "--help") || !strcmp(argv[i], "-h"))
{
std::cout << std::endl;
std::cout << "Usage: ./gpu-stream-ocl [OPTIONS]" << std::endl << std::endl;
std::cout << "Options:" << std::endl;
std::cout << " -h --help Print the message" << std::endl;
std::cout << " --list List available devices" << std::endl;
std::cout << " --device INDEX Select device at INDEX" << std::endl;
std::cout << " -s --arraysize SIZE Use SIZE elements in the array" << std::endl;
std::cout << " -n --numtimes NUM Run the test NUM times (NUM >= 2)" << std::endl;
std::cout << " --float Use floats (rather than doubles)" << std::endl;
std::cout << std::endl;
exit(0);
}
else
{
std::cout << "Unrecognized argument '" << argv[i] << "' (try '--help')"
<< std::endl;
exit(1);
} }
std::cout << std::endl;
} }
} }