github-mirrors/vortex
1
0
Fork 0
mirror of https://github.com/vortexgpgpu/vortex.git synced 2025-04-23 21:39:10 -04:00
Code Issues Projects Releases Packages Wiki Activity Actions

opencl tests update

Browse source
This commit is contained in:
Blaise Tine 2024-05-06 01:23:25 -07:00
parent 311799b423
commit 0cabd24f08
17 changed files with 375 additions and 546 deletions
Download patch file Download diff file Expand all files Collapse all files

114
tests/opencl/sgemm/main.cc
View file

@ -4,7 +4,7 @@
#include <CL/opencl.h>
#include <string.h>
#include <time.h>
#include <unistd.h>
#include <unistd.h>
#include <chrono>
#include <vector>
#include "common.h"
@ -44,10 +44,10 @@ public:
static const char* type_str() {
return "integer";
}
static int generate() {
return rand();
static int generate() {
return rand();
}
static bool compare(int a, int b, int index, int errors) {
static bool compare(int a, int b, int index, int errors) {
if (a != b) {
if (errors < 100) {
printf("*** error: [%d] expected=%d, actual=%d\n", index, a, b);
@ -55,7 +55,7 @@ public:
return false;
}
return true;
}
}
};
template <>
@ -64,10 +64,10 @@ public:
static const char* type_str() {
return "float";
}
static int generate() {
static int generate() {
return static_cast<float>(rand()) / RAND_MAX;
}
static bool compare(float a, float b, int index, int errors) {
static bool compare(float a, float b, int index, int errors) {
union fi_t { float f; int32_t i; };
fi_t fa, fb;
fa.f = a;
@ -80,10 +80,10 @@ public:
return false;
}
return true;
}
}
};
/*static void sgemm_cpu(TYPE *C, const TYPE* A, const TYPE *B, int M, int N, int K) {
static void sgemm_cpu(TYPE *C, const TYPE* A, const TYPE *B, int M, int N, int K) {
for (int m = 0; m < M; ++m) {
for (int n = 0; n < N; ++n) {
TYPE acc = 0;
@ -93,7 +93,7 @@ public:
C[n * M + m] = acc;
}
}
}*/
}
static int read_kernel_file(const char* filename, uint8_t** data, size_t* size) {
if (nullptr == filename || nullptr == data || 0 == size)
@ -110,9 +110,9 @@ static int read_kernel_file(const char* filename, uint8_t** data, size_t* size)
*data = (uint8_t*)malloc(fsize);
*size = fread(*data, 1, fsize, fp);
fclose(fp);
return 0;
}
@ -123,10 +123,7 @@ cl_program program = NULL;
cl_kernel kernel = NULL;
cl_mem a_memobj = NULL;
cl_mem b_memobj = NULL;
cl_mem c_memobj = NULL;
TYPE *h_a = NULL;
TYPE *h_b = NULL;
TYPE *h_c = NULL;
cl_mem c_memobj = NULL;
uint8_t *kernel_bin = NULL;
static void cleanup() {
@ -135,14 +132,11 @@ static void cleanup() {
if (program) clReleaseProgram(program);
if (a_memobj) clReleaseMemObject(a_memobj);
if (b_memobj) clReleaseMemObject(b_memobj);
if (c_memobj) clReleaseMemObject(c_memobj);
if (c_memobj) clReleaseMemObject(c_memobj);
if (context) clReleaseContext(context);
if (device_id) clReleaseDevice(device_id);
if (kernel_bin) free(kernel_bin);
if (h_a) free(h_a);
if (h_b) free(h_b);
if (h_c) free(h_c);
}
int size = 32;
@ -181,13 +175,13 @@ int main (int argc, char **argv) {
// parse command arguments
parse_args(argc, argv);
uint32_t num_points = size * size;
uint32_t size_sq = size * size;
cl_platform_id platform_id;
size_t kernel_size;
srand(50);
// Getting platform and device information
CL_CHECK(clGetPlatformIDs(1, &platform_id, NULL));
CL_CHECK(clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, NULL));
@ -196,7 +190,7 @@ int main (int argc, char **argv) {
context = CL_CHECK2(clCreateContext(NULL, 1, &device_id, NULL, NULL, &_err));
// Allocate device buffers
size_t nbytes = num_points * sizeof(TYPE);
size_t nbytes = size_sq * sizeof(TYPE);
a_memobj = CL_CHECK2(clCreateBuffer(context, CL_MEM_READ_ONLY, nbytes, NULL, &_err));
b_memobj = CL_CHECK2(clCreateBuffer(context, CL_MEM_READ_ONLY, nbytes, NULL, &_err));
c_memobj = CL_CHECK2(clCreateBuffer(context, CL_MEM_WRITE_ONLY, nbytes, NULL, &_err));
@ -206,7 +200,7 @@ int main (int argc, char **argv) {
if (0 != read_kernel_file("kernel.cl", &kernel_bin, &kernel_size))
return -1;
program = CL_CHECK2(clCreateProgramWithSource(
context, 1, (const char**)&kernel_bin, &kernel_size, &_err));
context, 1, (const char**)&kernel_bin, &kernel_size, &_err));
#else
if (0 != read_kernel_file("kernel.pocl", &kernel_bin, &kernel_size))
return -1;
@ -216,24 +210,23 @@ int main (int argc, char **argv) {
// Build program
CL_CHECK(clBuildProgram(program, 1, &device_id, NULL, NULL, NULL));
// Create kernel
kernel = CL_CHECK2(clCreateKernel(program, KERNEL_NAME, &_err));
// Set kernel arguments
int width = size;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&a_memobj));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&b_memobj));
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&a_memobj));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&b_memobj));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&c_memobj));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(width), (void*)&width));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(size), (void*)&size));
// Allocate memories for input arrays and output arrays.
h_a = (TYPE*)malloc(nbytes);
h_b = (TYPE*)malloc(nbytes);
h_c = (TYPE*)malloc(nbytes);
// Generate input values
for (uint32_t i = 0; i < num_points; ++i) {
// Allocate memories for input arrays and output arrays.
std::vector<TYPE> h_a(size_sq);
std::vector<TYPE> h_b(size_sq);
std::vector<TYPE> h_c(size_sq);
// Generate input values
for (uint32_t i = 0; i < size_sq; ++i) {
h_a[i] = Comparator<TYPE>::generate();
h_b[i] = Comparator<TYPE>::generate();
}
@ -242,39 +235,14 @@ int main (int argc, char **argv) {
size_t global_work_size[2] = {size, size};
size_t local_work_size[2] = {1, 1};
std::vector<float> ref_vec(num_points);
// reference generation
size_t num_groups_y = global_work_size[1] / local_work_size[1];
size_t num_groups_x = global_work_size[0] / local_work_size[0];
for (size_t workgroup_id_y = 0; workgroup_id_y < num_groups_y; ++workgroup_id_y) {
for (size_t workgroup_id_x = 0; workgroup_id_x < num_groups_x; ++workgroup_id_x) {
for (size_t local_id_y = 0; local_id_y < local_work_size[1]; ++local_id_y) {
for (size_t local_id_x = 0; local_id_x < local_work_size[0]; ++local_id_x) {
// Calculate global ID for the work-item
int global_id_x = global_offset[0] + local_work_size[0] * workgroup_id_x + local_id_x;
int global_id_y = global_offset[1] + local_work_size[1] * workgroup_id_y + local_id_y;
// kernel operation
int r = global_id_x;
int c = global_id_y;
TYPE acc = 0;
for (int k = 0; k < width; k++) {
acc += h_a[k * width + r] * h_b[c * width + k];
}
ref_vec[c * width + r] = acc;
}
}
}
}
// Creating command queue
commandQueue = CL_CHECK2(clCreateCommandQueue(context, device_id, 0, &_err));
commandQueue = CL_CHECK2(clCreateCommandQueue(context, device_id, 0, &_err));
printf("Upload source buffers\n");
CL_CHECK(clEnqueueWriteBuffer(commandQueue, a_memobj, CL_TRUE, 0, nbytes, h_a, 0, NULL, NULL));
CL_CHECK(clEnqueueWriteBuffer(commandQueue, b_memobj, CL_TRUE, 0, nbytes, h_b, 0, NULL, NULL));
CL_CHECK(clEnqueueWriteBuffer(commandQueue, a_memobj, CL_TRUE, 0, nbytes, h_a.data(), 0, NULL, NULL));
CL_CHECK(clEnqueueWriteBuffer(commandQueue, b_memobj, CL_TRUE, 0, nbytes, h_b.data(), 0, NULL, NULL));
printf("Execute the kernel\n");
printf("Execute the kernel\n");
auto time_start = std::chrono::high_resolution_clock::now();
CL_CHECK(clEnqueueNDRangeKernel(commandQueue, kernel, 2, global_offset, global_work_size, local_work_size, 0, NULL, NULL));
CL_CHECK(clFinish(commandQueue));
@ -283,22 +251,24 @@ int main (int argc, char **argv) {
printf("Elapsed time: %lg ms\n", elapsed);
printf("Download destination buffer\n");
CL_CHECK(clEnqueueReadBuffer(commandQueue, c_memobj, CL_TRUE, 0, nbytes, h_c, 0, NULL, NULL));
CL_CHECK(clEnqueueReadBuffer(commandQueue, c_memobj, CL_TRUE, 0, nbytes, h_c.data(), 0, NULL, NULL));
printf("Verify result\n");
std::vector<TYPE> h_ref(size_sq);
sgemm_cpu(h_ref.data(), h_a.data(), h_b.data(), size, size, size);
int errors = 0;
for (uint32_t i = 0; i < num_points; ++i) {
if (!Comparator<TYPE>::compare(h_c[i], ref_vec[i], i, errors)) {
for (uint32_t i = 0; i < size_sq; ++i) {
if (!Comparator<TYPE>::compare(h_c[i], h_ref[i], i, errors)) {
++errors;
}
}
if (errors != 0) {
printf("FAILED! - %d errors\n", errors);
printf("FAILED! - %d errors\n", errors);
} else {
printf("PASSED!\n");
}
// Clean up
cleanup();
// Clean up
cleanup();
return errors;
}

6
tests/opencl/sgemm2/Makefile
View file

@ -7,14 +7,12 @@ SRC_DIR := $(VORTEX_HOME)/tests/opencl/$(PROJECT)
SRCS := $(SRC_DIR)/main.cc
CXXFLAGS += -I$(SRC_DIR)
common.h: $(SRC_DIR)/common.h
common.h: $(SRC_DIR)/common.h
cp $< $@
setup: common.h
USE_SETUP := yes
OPTS ?= -n32
OPTS ?= -n16
include ../common.mk

23
tests/opencl/sgemm2/common.h
View file

@ -1,15 +1,10 @@
// Copyright 2019-2023
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#ifndef COMMON_H
#define COMMON_H
#define LOCAL_SIZE 16
#define LOCAL_SIZE 4
#ifndef TYPE
#define TYPE float
#endif
#endif // COMMON_H

24
tests/opencl/sgemm2/kernel.cl
View file

@ -1,8 +1,8 @@
#include "common.h"
__kernel void sgemm2(__global float *A,
__global float *B,
__global float *C,
__kernel void sgemm2(__global TYPE *A,
__global TYPE *B,
__global TYPE *C,
const unsigned int N)
{
int globalRow = get_global_id(1);
@ -11,20 +11,16 @@ __kernel void sgemm2(__global float *A,
int localCol = get_local_id(0);
// Static local memory declaration
__local float localA[LOCAL_SIZE][LOCAL_SIZE];
__local float localB[LOCAL_SIZE][LOCAL_SIZE];
__local TYPE localA[LOCAL_SIZE][LOCAL_SIZE];
__local TYPE localB[LOCAL_SIZE][LOCAL_SIZE];
float sum = 0.0f;
//printf("l=(%d, %d), g=(%d, %d)\n", localCol, localRow, globalCol, globalRow);
TYPE sum = 0;
// Iterate over blocks
for (int k = 0; k < N; k += LOCAL_SIZE) {
float a = A[globalRow * N + k + localCol];
float b = B[(k + localRow) * N + globalCol];
localA[localRow][localCol] = a;
localB[localRow][localCol] = b;
// Load block of matrix A & B to local memory
localA[localRow][localCol] = A[globalRow * N + (k + localCol)];
localB[localRow][localCol] = B[(k + localRow) * N + globalCol];
// Ensure the entire block is loaded
barrier(CLK_LOCAL_MEM_FENCE);
@ -36,8 +32,6 @@ __kernel void sgemm2(__global float *A,
// Ensure computation is done before loading next block
barrier(CLK_LOCAL_MEM_FENCE);
//printf("k=%d, a=%f, b=%f, sum=%f\n", k, a, b, sum);
}
C[globalRow * N + globalCol] = sum;

137
tests/opencl/sgemm2/main.cc
View file

@ -4,7 +4,7 @@
#include <CL/opencl.h>
#include <string.h>
#include <time.h>
#include <unistd.h>
#include <unistd.h>
#include <chrono>
#include <vector>
#include "common.h"
@ -44,38 +44,77 @@ static int read_kernel_file(const char* filename, uint8_t** data, size_t* size)
fprintf(stderr, "Failed to load kernel.");
return -1;
}
fseek(fp , 0 , SEEK_END);
long fsize = ftell(fp);
rewind(fp);
*data = (uint8_t*)malloc(fsize);
*size = fread(*data, 1, fsize, fp);
fclose(fp);
return 0;
}
static bool compare_equal(float a, float b) {
union fi_t { float f; int32_t i; };
fi_t fa, fb;
fa.f = a;
fb.f = b;
auto d = std::abs(fa.i - fb.i);
return d <= FLOAT_ULP;
}
template <typename Type>
class Comparator {};
static void matmul_cpu(float *C, float *A, float *B, int N) {
for (int i = 0; i < N; i++) {
for (int j = 0; j < N; j++) {
float sum = 0.0f;
for (int k = 0; k < N; k++) {
sum += A[i * N + k] * B[k * N + j];
}
C[i * N + j] = sum;
}
template <>
class Comparator<int> {
public:
static const char* type_str() {
return "integer";
}
static int generate() {
return rand();
}
static bool compare(int a, int b, int index, int errors) {
if (a != b) {
if (errors < 100) {
printf("*** error: [%d] expected=%d, actual=%d\n", index, a, b);
}
return false;
}
return true;
}
};
template <>
class Comparator<float> {
public:
static const char* type_str() {
return "float";
}
static int generate() {
return static_cast<float>(rand()) / RAND_MAX;
}
static bool compare(float a, float b, int index, int errors) {
union fi_t { float f; int32_t i; };
fi_t fa, fb;
fa.f = a;
fb.f = b;
auto d = std::abs(fa.i - fb.i);
if (d > FLOAT_ULP) {
if (errors < 100) {
printf("*** error: [%d] expected=%f, actual=%f\n", index, a, b);
}
return false;
}
return true;
}
};
static void sgemm_cpu(TYPE *C, const TYPE* A, const TYPE *B, int M, int N, int K) {
for (int m = 0; m < M; ++m) {
for (int n = 0; n < N; ++n) {
TYPE acc = 0;
for (int k = 0; k < K; ++k) {
acc += A[k * M + m] * B[n * K + k];
}
C[n * M + m] = acc;
}
}
}
cl_device_id device_id = NULL;
@ -96,11 +135,11 @@ static void cleanup() {
if (b_memobj) clReleaseMemObject(b_memobj);
if (c_memobj) clReleaseMemObject(c_memobj);
if (context) clReleaseContext(context);
if (device_id) clReleaseDevice(device_id);
if (device_id) clReleaseDevice(device_id);
if (kernel_bin) free(kernel_bin);
}
int size = 32;
int size = 16;
static void show_usage() {
printf("Usage: [-n size] [-h: help]\n");
@ -129,17 +168,17 @@ int main (int argc, char **argv) {
// parse command arguments
parse_args(argc, argv);
uint32_t num_points = size * size;
uint32_t size_sq = size * size;
printf("Matrix size=%d\n", size);
if ((size / LOCAL_SIZE) * LOCAL_SIZE != size) {
printf("Error: matrix size must be a multiple of %d\n", LOCAL_SIZE);
return -1;
}
cl_platform_id platform_id;
size_t kernel_size;
// Getting platform and device information
CL_CHECK(clGetPlatformIDs(1, &platform_id, NULL));
CL_CHECK(clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, NULL));
@ -152,7 +191,7 @@ int main (int argc, char **argv) {
printf("Using device: %s\n", device_string);
printf("Allocate device buffers\n");
size_t nbytes = num_points * sizeof(float);
size_t nbytes = size_sq * sizeof(TYPE);
a_memobj = CL_CHECK2(clCreateBuffer(context, CL_MEM_READ_ONLY, nbytes, NULL, &_err));
b_memobj = CL_CHECK2(clCreateBuffer(context, CL_MEM_READ_ONLY, nbytes, NULL, &_err));
c_memobj = CL_CHECK2(clCreateBuffer(context, CL_MEM_WRITE_ONLY, nbytes, NULL, &_err));
@ -162,13 +201,13 @@ int main (int argc, char **argv) {
if (0 != read_kernel_file("kernel.cl", &kernel_bin, &kernel_size))
return -1;
program = CL_CHECK2(clCreateProgramWithSource(
context, 1, (const char**)&kernel_bin, &kernel_size, &_err));
context, 1, (const char**)&kernel_bin, &kernel_size, &_err));
#else
if (0 != read_kernel_file("kernel.pocl", &kernel_bin, &kernel_size))
return -1;
program = CL_CHECK2(clCreateProgramWithBinary(
context, 1, &device_id, &kernel_size, (const uint8_t**)&kernel_bin, NULL, &_err));
#endif
#endif
if (program == NULL) {
cleanup();
return -1;
@ -176,32 +215,32 @@ int main (int argc, char **argv) {
// Build program
CL_CHECK(clBuildProgram(program, 1, &device_id, NULL, NULL, NULL));
// Create kernel
kernel = CL_CHECK2(clCreateKernel(program, KERNEL_NAME, &_err));
size_t global_size[2] = {size, size};
size_t local_size[2] = {LOCAL_SIZE, LOCAL_SIZE};
size_t local_size[2] = {LOCAL_SIZE, LOCAL_SIZE};
// Set kernel arguments
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&a_memobj));
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&a_memobj));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&b_memobj));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&c_memobj));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(uint32_t), &size));
// Allocate memories for input arrays and output arrays.
std::vector<float> h_a(num_points);
std::vector<float> h_b(num_points);
std::vector<float> h_c(num_points);
// Allocate memories for input arrays and output arrays.
std::vector<TYPE> h_a(size_sq);
std::vector<TYPE> h_b(size_sq);
std::vector<TYPE> h_c(size_sq);
// Generate input values
for (uint32_t i = 0; i < num_points; ++i) {
h_a[i] = static_cast<float>(rand()) / RAND_MAX;
h_b[i] = static_cast<float>(rand()) / RAND_MAX;
for (uint32_t i = 0; i < size_sq; ++i) {
h_a[i] = Comparator<TYPE>::generate();
h_b[i] = Comparator<TYPE>::generate();
}
// Creating command queue
commandQueue = CL_CHECK2(clCreateCommandQueue(context, device_id, 0, &_err));
commandQueue = CL_CHECK2(clCreateCommandQueue(context, device_id, 0, &_err));
printf("Upload source buffers\n");
CL_CHECK(clEnqueueWriteBuffer(commandQueue, a_memobj, CL_TRUE, 0, nbytes, h_a.data(), 0, NULL, NULL));
@ -219,24 +258,22 @@ int main (int argc, char **argv) {
CL_CHECK(clEnqueueReadBuffer(commandQueue, c_memobj, CL_TRUE, 0, nbytes, h_c.data(), 0, NULL, NULL));
printf("Verify result\n");
std::vector<float> ref_vec(num_points);
matmul_cpu(ref_vec.data(), h_a.data(), h_b.data(), size);
std::vector<TYPE> h_ref(size_sq);
sgemm_cpu(h_ref.data(), h_a.data(), h_b.data(), size, size, size);
int errors = 0;
for (uint32_t i = 0; i < num_points; ++i) {
if (!compare_equal(h_c[i], ref_vec[i])) {
if (errors < 100)
printf("*** error: [%d] expected=%f, actual=%f\n", i, ref_vec[i], h_c[i]);
for (uint32_t i = 0; i < size_sq; ++i) {
if (!Comparator<TYPE>::compare(h_c[i], h_ref[i], i, errors)) {
++errors;
}
}
if (errors != 0) {
printf("FAILED! - %d errors\n", errors);
printf("FAILED! - %d errors\n", errors);
} else {
printf("PASSED!\n");
}
// Clean up
cleanup();
// Clean up
cleanup();
return errors;
}

10
tests/opencl/sgemm3/Makefile
View file

@ -7,6 +7,12 @@ SRC_DIR := $(VORTEX_HOME)/tests/opencl/$(PROJECT)
SRCS := $(SRC_DIR)/main.cc
OPTS ?= -n32
common.h: $(SRC_DIR)/common.h
cp $< $@
include ../common.mk
setup: common.h
USE_SETUP := yes
OPTS ?= -n16
include ../common.mk

8
tests/opencl/sgemm3/common.h Normal file
View file

@ -0,0 +1,8 @@
#ifndef COMMON_H
#define COMMON_H
#ifndef TYPE
#define TYPE float
#endif
#endif // COMMON_H

26
tests/opencl/sgemm3/kernel.cl
View file

@ -1,24 +1,24 @@
__kernel void sgemm3(__global float *A,
__global float *B,
__global float *C,
const unsigned int N,
__local float *localA,
__local float *localB)
#include "common.h"
__kernel void sgemm3(__global TYPE *A,
__global TYPE *B,
__global TYPE *C,
const unsigned int N,
__local TYPE *localA,
__local TYPE *localB)
{
int globalRow = get_global_id(1);
int globalCol = get_global_id(0);
int localRow = get_local_id(1);
int localCol = get_local_id(0);
int localSize = get_local_size(0); // assuming square local size
int localSize = get_local_size(0);
float sum = 0.0f;
TYPE sum = 0;
// Loop over all blocks of both matrices
// Iterate over blocks
for (int k = 0; k < N; k += localSize) {
// Load block of matrix A to local memory
localA[localRow * localSize + localCol] = A[globalRow * N + k + localCol];
// Load block of matrix B to local memory, adjusting for column-major access
// Load block of matrix A & B to local memory
localA[localRow * localSize + localCol] = A[globalRow * N + (k + localCol)];
localB[localRow * localSize + localCol] = B[(k + localRow) * N + globalCol];
// Synchronize to make sure the tiles are loaded

138
tests/opencl/sgemm3/main.cc
View file

@ -4,11 +4,12 @@
#include <CL/opencl.h>
#include <string.h>
#include <time.h>
#include <unistd.h>
#include <unistd.h>
#include <chrono>
#include <vector>
#include "common.h"
#define LOCAL_SIZE 16
#define LOCAL_SIZE 4
#define FLOAT_ULP 6
@ -45,38 +46,77 @@ static int read_kernel_file(const char* filename, uint8_t** data, size_t* size)
fprintf(stderr, "Failed to load kernel.");
return -1;
}
fseek(fp , 0 , SEEK_END);
long fsize = ftell(fp);
rewind(fp);
*data = (uint8_t*)malloc(fsize);
*size = fread(*data, 1, fsize, fp);
fclose(fp);
return 0;
}
static bool compare_equal(float a, float b) {
union fi_t { float f; int32_t i; };
fi_t fa, fb;
fa.f = a;
fb.f = b;
auto d = std::abs(fa.i - fb.i);
return d <= FLOAT_ULP;
}
template <typename Type>
class Comparator {};
static void matmul_cpu(float *C, float *A, float *B, int N) {
for (int i = 0; i < N; i++) {
for (int j = 0; j < N; j++) {
float sum = 0.0f;
for (int k = 0; k < N; k++) {
sum += A[i * N + k] * B[k * N + j];
}
C[i * N + j] = sum;
}
template <>
class Comparator<int> {
public:
static const char* type_str() {
return "integer";
}
static int generate() {
return rand();
}
static bool compare(int a, int b, int index, int errors) {
if (a != b) {
if (errors < 100) {
printf("*** error: [%d] expected=%d, actual=%d\n", index, a, b);
}
return false;
}
return true;
}
};
template <>
class Comparator<float> {
public:
static const char* type_str() {
return "float";
}
static int generate() {
return static_cast<float>(rand()) / RAND_MAX;
}
static bool compare(float a, float b, int index, int errors) {
union fi_t { float f; int32_t i; };
fi_t fa, fb;
fa.f = a;
fb.f = b;
auto d = std::abs(fa.i - fb.i);
if (d > FLOAT_ULP) {
if (errors < 100) {
printf("*** error: [%d] expected=%f, actual=%f\n", index, a, b);
}
return false;
}
return true;
}
};
static void sgemm_cpu(TYPE *C, const TYPE* A, const TYPE *B, int M, int N, int K) {
for (int m = 0; m < M; ++m) {
for (int n = 0; n < N; ++n) {
TYPE acc = 0;
for (int k = 0; k < K; ++k) {
acc += A[k * M + m] * B[n * K + k];
}
C[n * M + m] = acc;
}
}
}
cl_device_id device_id = NULL;
@ -97,11 +137,11 @@ static void cleanup() {
if (b_memobj) clReleaseMemObject(b_memobj);
if (c_memobj) clReleaseMemObject(c_memobj);
if (context) clReleaseContext(context);
if (device_id) clReleaseDevice(device_id);
if (device_id) clReleaseDevice(device_id);
if (kernel_bin) free(kernel_bin);
}
int size = 32;
int size = 16;
static void show_usage() {
printf("Usage: [-n size] [-h: help]\n");
@ -130,17 +170,17 @@ int main (int argc, char **argv) {
// parse command arguments
parse_args(argc, argv);
uint32_t num_points = size * size;
uint32_t size_sq = size * size;
printf("Matrix size=%d\n", size);
if ((size / LOCAL_SIZE) * LOCAL_SIZE != size) {
printf("Error: matrix size must be a multiple of %d\n", LOCAL_SIZE);
return -1;
}
cl_platform_id platform_id;
size_t kernel_size;
// Getting platform and device information
CL_CHECK(clGetPlatformIDs(1, &platform_id, NULL));
CL_CHECK(clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, NULL));
@ -153,7 +193,7 @@ int main (int argc, char **argv) {
printf("Using device: %s\n", device_string);
printf("Allocate device buffers\n");
size_t nbytes = num_points * sizeof(float);
size_t nbytes = size_sq * sizeof(float);
a_memobj = CL_CHECK2(clCreateBuffer(context, CL_MEM_READ_ONLY, nbytes, NULL, &_err));
b_memobj = CL_CHECK2(clCreateBuffer(context, CL_MEM_READ_ONLY, nbytes, NULL, &_err));
c_memobj = CL_CHECK2(clCreateBuffer(context, CL_MEM_WRITE_ONLY, nbytes, NULL, &_err));
@ -163,13 +203,13 @@ int main (int argc, char **argv) {
if (0 != read_kernel_file("kernel.cl", &kernel_bin, &kernel_size))
return -1;
program = CL_CHECK2(clCreateProgramWithSource(
context, 1, (const char**)&kernel_bin, &kernel_size, &_err));
context, 1, (const char**)&kernel_bin, &kernel_size, &_err));
#else
if (0 != read_kernel_file("kernel.pocl", &kernel_bin, &kernel_size))
return -1;
program = CL_CHECK2(clCreateProgramWithBinary(
context, 1, &device_id, &kernel_size, (const uint8_t**)&kernel_bin, NULL, &_err));
#endif
#endif
if (program == NULL) {
cleanup();
return -1;
@ -177,15 +217,15 @@ int main (int argc, char **argv) {
// Build program
CL_CHECK(clBuildProgram(program, 1, &device_id, NULL, NULL, NULL));
// Create kernel
kernel = CL_CHECK2(clCreateKernel(program, KERNEL_NAME, &_err));
size_t global_size[2] = {size, size};
size_t local_size[2] = {LOCAL_SIZE, LOCAL_SIZE};
size_t local_size[2] = {LOCAL_SIZE, LOCAL_SIZE};
// Set kernel arguments
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&a_memobj));
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&a_memobj));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&b_memobj));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&c_memobj));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(uint32_t), &size));
@ -193,18 +233,18 @@ int main (int argc, char **argv) {
CL_CHECK(clSetKernelArg(kernel, 5, local_size[0]*local_size[1]*sizeof(float), NULL));
// Allocate memories for input arrays and output arrays.
std::vector<float> h_a(num_points);
std::vector<float> h_b(num_points);
std::vector<float> h_c(num_points);
std::vector<float> h_a(size_sq);
std::vector<float> h_b(size_sq);
std::vector<float> h_c(size_sq);
// Generate input values
for (uint32_t i = 0; i < num_points; ++i) {
h_a[i] = static_cast<float>(rand()) / RAND_MAX;
h_b[i] = static_cast<float>(rand()) / RAND_MAX;
for (uint32_t i = 0; i < size_sq; ++i) {
h_a[i] = Comparator<TYPE>::generate();
h_b[i] = Comparator<TYPE>::generate();
}
// Creating command queue
commandQueue = CL_CHECK2(clCreateCommandQueue(context, device_id, 0, &_err));
commandQueue = CL_CHECK2(clCreateCommandQueue(context, device_id, 0, &_err));
printf("Upload source buffers\n");
CL_CHECK(clEnqueueWriteBuffer(commandQueue, a_memobj, CL_TRUE, 0, nbytes, h_a.data(), 0, NULL, NULL));
@ -222,24 +262,22 @@ int main (int argc, char **argv) {
CL_CHECK(clEnqueueReadBuffer(commandQueue, c_memobj, CL_TRUE, 0, nbytes, h_c.data(), 0, NULL, NULL));
printf("Verify result\n");
std::vector<float> ref_vec(num_points);
matmul_cpu(ref_vec.data(), h_a.data(), h_b.data(), size);
std::vector<TYPE> h_ref(size_sq);
sgemm_cpu(h_ref.data(), h_a.data(), h_b.data(), size, size, size);
int errors = 0;
for (uint32_t i = 0; i < num_points; ++i) {
if (!compare_equal(h_c[i], ref_vec[i])) {
if (errors < 100)
printf("*** error: [%d] expected=%f, actual=%f\n", i, ref_vec[i], h_c[i]);
for (uint32_t i = 0; i < size_sq; ++i) {
if (!Comparator<TYPE>::compare(h_c[i], h_ref[i], i, errors)) {
++errors;
}
}
if (errors != 0) {
printf("FAILED! - %d errors\n", errors);
printf("FAILED! - %d errors\n", errors);
} else {
printf("PASSED!\n");
}
// Clean up
cleanup();
// Clean up
cleanup();
return errors;
}

6
tests/opencl/vecadd/Makefile
View file

@ -7,6 +7,12 @@ SRC_DIR := $(VORTEX_HOME)/tests/opencl/$(PROJECT)
SRCS := $(SRC_DIR)/main.cc
common.h: $(SRC_DIR)/common.h
cp $< $@
setup: common.h
USE_SETUP := yes
OPTS ?= -n64
include ../common.mk

8
tests/opencl/vecadd/common.h Normal file
View file

@ -0,0 +1,8 @@
#ifndef COMMON_H
#define COMMON_H
#ifndef TYPE
#define TYPE float
#endif
#endif // COMMON_H

9
tests/opencl/vecadd/kernel.cl
View file

@ -1,6 +1,9 @@
__kernel void vecadd (__global const float *A,
__global const float *B,
__global float *C)
#include "common.h"
__kernel void vecadd (__global const TYPE *A,
__global const TYPE *B,
__global TYPE *C)
{
int gid = get_global_id(0);
C[gid] = A[gid] + B[gid];

135
tests/opencl/vecadd/main.cc
View file

@ -3,12 +3,16 @@
#include <assert.h>
#include <math.h>
#include <CL/opencl.h>
#include <unistd.h>
#include <unistd.h>
#include <string.h>
#include <chrono>
#include <vector>
#include "common.h"
#define KERNEL_NAME "vecadd"
#define FLOAT_ULP 6
#define CL_CHECK(_expr) \
do { \
cl_int _err = _expr; \
@ -46,18 +50,64 @@ static int read_kernel_file(const char* filename, uint8_t** data, size_t* size)
*data = (uint8_t*)malloc(fsize);
*size = fread(*data, 1, fsize, fp);
fclose(fp);
return 0;
}
static bool almost_equal(float a, float b, int ulp = 4) {
union fi_t { int i; float f; };
fi_t fa, fb;
fa.f = a;
fb.f = b;
return std::abs(fa.i - fb.i) <= ulp;
template <typename Type>
class Comparator {};
template <>
class Comparator<int> {
public:
static const char* type_str() {
return "integer";
}
static int generate() {
return rand();
}
static bool compare(int a, int b, int index, int errors) {
if (a != b) {
if (errors < 100) {
printf("*** error: [%d] expected=%d, actual=%d\n", index, a, b);
}
return false;
}
return true;
}
};
template <>
class Comparator<float> {
public:
static const char* type_str() {
return "float";
}
static int generate() {
return static_cast<float>(rand()) / RAND_MAX;
}
static bool compare(float a, float b, int index, int errors) {
union fi_t { float f; int32_t i; };
fi_t fa, fb;
fa.f = a;
fb.f = b;
auto d = std::abs(fa.i - fb.i);
if (d > FLOAT_ULP) {
if (errors < 100) {
printf("*** error: [%d] expected=%f, actual=%f\n", index, a, b);
}
return false;
}
return true;
}
};
static void vecadd_cpu(TYPE *C, const TYPE* A, const TYPE *B, int N) {
for (int i = 0; i < N; ++i) {
C[i] = A[i] + B[i];
}
}
cl_device_id device_id = NULL;
@ -67,10 +117,7 @@ cl_program program = NULL;
cl_kernel kernel = NULL;
cl_mem a_memobj = NULL;
cl_mem b_memobj = NULL;
cl_mem c_memobj = NULL;
float *h_a = NULL;
float *h_b = NULL;
float *h_c = NULL;
cl_mem c_memobj = NULL;
uint8_t *kernel_bin = NULL;
static void cleanup() {
@ -79,14 +126,11 @@ static void cleanup() {
if (program) clReleaseProgram(program);
if (a_memobj) clReleaseMemObject(a_memobj);
if (b_memobj) clReleaseMemObject(b_memobj);
if (c_memobj) clReleaseMemObject(c_memobj);
if (c_memobj) clReleaseMemObject(c_memobj);
if (context) clReleaseContext(context);
if (device_id) clReleaseDevice(device_id);
if (kernel_bin) free(kernel_bin);
if (h_a) free(h_a);
if (h_b) free(h_b);
if (h_c) free(h_c);
}
int size = 64;
@ -119,10 +163,10 @@ static void parse_args(int argc, char **argv) {
int main (int argc, char **argv) {
// parse command arguments
parse_args(argc, argv);
cl_platform_id platform_id;
size_t kernel_size;
// Getting platform and device information
CL_CHECK(clGetPlatformIDs(1, &platform_id, NULL));
CL_CHECK(clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, NULL));
@ -131,7 +175,7 @@ int main (int argc, char **argv) {
context = CL_CHECK2(clCreateContext(NULL, 1, &device_id, NULL, NULL, &_err));
printf("Allocate device buffers\n");
size_t nbytes = size * sizeof(float);
size_t nbytes = size * sizeof(TYPE);
a_memobj = CL_CHECK2(clCreateBuffer(context, CL_MEM_READ_ONLY, nbytes, NULL, &_err));
b_memobj = CL_CHECK2(clCreateBuffer(context, CL_MEM_READ_ONLY, nbytes, NULL, &_err));
c_memobj = CL_CHECK2(clCreateBuffer(context, CL_MEM_WRITE_ONLY, nbytes, NULL, &_err));
@ -141,7 +185,7 @@ int main (int argc, char **argv) {
if (0 != read_kernel_file("kernel.cl", &kernel_bin, &kernel_size))
return -1;
program = CL_CHECK2(clCreateProgramWithSource(
context, 1, (const char**)&kernel_bin, &kernel_size, &_err));
context, 1, (const char**)&kernel_bin, &kernel_size, &_err));
#else
if (0 != read_kernel_file("kernel.pocl", &kernel_bin, &kernel_size))
return -1;
@ -151,35 +195,35 @@ int main (int argc, char **argv) {
// Build program
CL_CHECK(clBuildProgram(program, 1, &device_id, NULL, NULL, NULL));
// Create kernel
kernel = CL_CHECK2(clCreateKernel(program, KERNEL_NAME, &_err));
// Set kernel arguments
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&a_memobj));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&b_memobj));
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&a_memobj));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&b_memobj));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&c_memobj));
// Allocate memories for input arrays and output arrays.
h_a = (float*)malloc(nbytes);
h_b = (float*)malloc(nbytes);
h_c = (float*)malloc(nbytes);
// Allocate memories for input arrays and output arrays.
std::vector<TYPE> h_a(size);
std::vector<TYPE> h_b(size);
std::vector<TYPE> h_c(size);
// Generate input values
for (int i = 0; i < size; ++i) {
h_a[i] = sinf(i)*sinf(i);
h_b[i] = cosf(i)*cosf(i);
for (uint32_t i = 0; i < size; ++i) {
h_a[i] = Comparator<TYPE>::generate();
h_b[i] = Comparator<TYPE>::generate();
}
// Creating command queue
commandQueue = CL_CHECK2(clCreateCommandQueue(context, device_id, 0, &_err));
commandQueue = CL_CHECK2(clCreateCommandQueue(context, device_id, 0, &_err));
printf("Upload source buffers\n");
CL_CHECK(clEnqueueWriteBuffer(commandQueue, a_memobj, CL_TRUE, 0, nbytes, h_a, 0, NULL, NULL));
CL_CHECK(clEnqueueWriteBuffer(commandQueue, b_memobj, CL_TRUE, 0, nbytes, h_b, 0, NULL, NULL));
CL_CHECK(clEnqueueWriteBuffer(commandQueue, a_memobj, CL_TRUE, 0, nbytes, h_a.data(), 0, NULL, NULL));
CL_CHECK(clEnqueueWriteBuffer(commandQueue, b_memobj, CL_TRUE, 0, nbytes, h_b.data(), 0, NULL, NULL));
printf("Execute the kernel\n");
size_t global_work_size[1] = {size};
size_t global_work_size[1] = {size};
size_t local_work_size[1] = {1};
auto time_start = std::chrono::high_resolution_clock::now();
CL_CHECK(clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL));
@ -189,26 +233,25 @@ int main (int argc, char **argv) {
printf("Elapsed time: %lg ms\n", elapsed);
printf("Download destination buffer\n");
CL_CHECK(clEnqueueReadBuffer(commandQueue, c_memobj, CL_TRUE, 0, nbytes, h_c, 0, NULL, NULL));
CL_CHECK(clEnqueueReadBuffer(commandQueue, c_memobj, CL_TRUE, 0, nbytes, h_c.data(), 0, NULL, NULL));
printf("Verify result\n");
std::vector<TYPE> h_ref(size);
vecadd_cpu(h_ref.data(), h_a.data(), h_b.data(), size);
int errors = 0;
for (int i = 0; i < size; ++i) {
float ref = h_a[i] + h_b[i];
if (!almost_equal(h_c[i], ref)) {
if (errors < 100)
printf("*** error: [%d] expected=%f, actual=%f, a=%f, b=%f\n", i, ref, h_c[i], h_a[i], h_b[i]);
for (uint32_t i = 0; i < size; ++i) {
if (!Comparator<TYPE>::compare(h_c[i], h_ref[i], i, errors)) {
++errors;
}
}
if (0 == errors) {
printf("PASSED!\n");
} else {
printf("FAILED! - %d errors\n", errors);
printf("FAILED! - %d errors\n", errors);
}
// Clean up
cleanup();
// Clean up
cleanup();
return errors;
}

14
tests/regression/lmem/Makefile
View file

@ -1,14 +0,0 @@
ROOT_DIR := $(realpath ../../..)
include $(ROOT_DIR)/config.mk
PROJECT := lmem
SRC_DIR := $(VORTEX_HOME)/tests/regression/$(PROJECT)
SRCS := $(SRC_DIR)/main.cpp
VX_SRCS := $(SRC_DIR)/kernel.cpp
OPTS ?= -n64
include ../common.mk

16
tests/regression/lmem/common.h
View file

@ -1,16 +0,0 @@
#ifndef _COMMON_H_
#define _COMMON_H_
#ifndef TYPE
#define TYPE float
#endif
typedef struct {
uint32_t num_tasks;
uint32_t task_size;
uint64_t src0_addr;
uint64_t src1_addr;
uint64_t dst_addr;
} kernel_arg_t;
#endif

23
tests/regression/lmem/kernel.cpp
View file

@ -1,23 +0,0 @@
#include <stdint.h>
#include <vx_intrinsics.h>
#include <vx_spawn.h>
#include "common.h"
void kernel_body(int task_id, kernel_arg_t* __UNIFORM__ arg) {
auto src0_ptr = reinterpret_cast<TYPE*>(arg->src0_addr);
auto src1_ptr = reinterpret_cast<TYPE*>(arg->src1_addr);
auto dst_ptr = reinterpret_cast<TYPE*>(arg->dst_addr);
uint32_t count = arg->task_size;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
dst_ptr[offset+i] = src0_ptr[offset+i] + src1_ptr[offset+i];
}
}
int main() {
kernel_arg_t* arg = (kernel_arg_t*)csr_read(VX_CSR_MSCRATCH);
vx_spawn_tasks(arg->num_tasks, (vx_spawn_tasks_cb)kernel_body, arg);
return 0;
}

224
tests/regression/lmem/main.cpp
View file

@ -1,224 +0,0 @@
#include <iostream>
#include <unistd.h>
#include <string.h>
#include <vector>
#include <vortex.h>
#include "common.h"
#define FLOAT_ULP 6
#define RT_CHECK(_expr) \
do { \
int _ret = _expr; \
if (0 == _ret) \
break; \
printf("Error: '%s' returned %d!\n", #_expr, (int)_ret); \
cleanup(); \
exit(-1); \
} while (false)
///////////////////////////////////////////////////////////////////////////////
template <typename Type>
class Comparator {};
template <>
class Comparator<int> {
public:
static const char* type_str() {
return "integer";
}
static int generate() {
return rand();
}
static bool compare(int a, int b, int index, int errors) {
if (a != b) {
if (errors < 100) {
printf("*** error: [%d] expected=%d, actual=%d\n", index, a, b);
}
return false;
}
return true;
}
};
template <>
class Comparator<float> {
private:
union Float_t { float f; int i; };
public:
static const char* type_str() {
return "float";
}
static int generate() {
return static_cast<float>(rand()) / RAND_MAX;
}
static bool compare(float a, float b, int index, int errors) {
union fi_t { float f; int32_t i; };
fi_t fa, fb;
fa.f = a;
fb.f = b;
auto d = std::abs(fa.i - fb.i);
if (d > FLOAT_ULP) {
if (errors < 100) {
printf("*** error: [%d] expected=%f, actual=%f\n", index, a, b);
}
return false;
}
return true;
}
};
const char* kernel_file = "kernel.vxbin";
uint32_t count = 16;
vx_device_h device = nullptr;
vx_buffer_h src0_buffer = nullptr;
vx_buffer_h src1_buffer = nullptr;
vx_buffer_h dst_buffer = nullptr;
vx_buffer_h krnl_buffer = nullptr;
vx_buffer_h args_buffer = nullptr;
kernel_arg_t kernel_arg = {};
static void show_usage() {
std::cout << "Vortex Test." << std::endl;
std::cout << "Usage: [-k: kernel] [-n words] [-h: help]" << std::endl;
}
static void parse_args(int argc, char **argv) {
int c;
while ((c = getopt(argc, argv, "n:k:h?")) != -1) {
switch (c) {
case 'n':
count = atoi(optarg);
break;
case 'k':
kernel_file = optarg;
break;
case 'h':
case '?': {
show_usage();
exit(0);
} break;
default:
show_usage();
exit(-1);
}
}
}
void cleanup() {
if (device) {
vx_mem_free(src0_buffer);
vx_mem_free(src1_buffer);
vx_mem_free(dst_buffer);
vx_mem_free(krnl_buffer);
vx_mem_free(args_buffer);
vx_dev_close(device);
}
}
int main(int argc, char *argv[]) {
// parse command arguments
parse_args(argc, argv);
std::srand(50);
// open device connection
std::cout << "open device connection" << std::endl;
RT_CHECK(vx_dev_open(&device));
uint64_t num_cores, num_warps, num_threads;
RT_CHECK(vx_dev_caps(device, VX_CAPS_NUM_CORES, &num_cores));
RT_CHECK(vx_dev_caps(device, VX_CAPS_NUM_WARPS, &num_warps));
RT_CHECK(vx_dev_caps(device, VX_CAPS_NUM_THREADS, &num_threads));
uint32_t total_threads = num_cores * num_warps * num_threads;
uint32_t num_points = count * total_threads;
uint32_t buf_size = num_points * sizeof(TYPE);
std::cout << "data type: " << Comparator<TYPE>::type_str() << std::endl;
std::cout << "number of points: " << num_points << std::endl;
std::cout << "buffer size: " << buf_size << " bytes" << std::endl;
kernel_arg.num_tasks = total_threads;
kernel_arg.task_size = count;
// allocate device memory
std::cout << "allocate device memory" << std::endl;
RT_CHECK(vx_mem_alloc(device, buf_size, VX_MEM_READ, &src0_buffer));
RT_CHECK(vx_mem_address(src0_buffer, &kernel_arg.src0_addr));
RT_CHECK(vx_mem_alloc(device, buf_size, VX_MEM_READ, &src1_buffer));
RT_CHECK(vx_mem_address(src1_buffer, &kernel_arg.src1_addr));
RT_CHECK(vx_mem_alloc(device, buf_size, VX_MEM_WRITE, &dst_buffer));
RT_CHECK(vx_mem_address(dst_buffer, &kernel_arg.dst_addr));
std::cout << "dev_src0=0x" << std::hex << kernel_arg.src0_addr << std::endl;
std::cout << "dev_src1=0x" << std::hex << kernel_arg.src1_addr << std::endl;
std::cout << "dev_dst=0x" << std::hex << kernel_arg.dst_addr << std::endl;
// allocate host buffers
std::cout << "allocate host buffers" << std::endl;
std::vector<TYPE> h_src0(num_points);
std::vector<TYPE> h_src1(num_points);
std::vector<TYPE> h_dst(num_points);
// generate source data
for (uint32_t i = 0; i < num_points; ++i) {
h_src0[i] = Comparator<TYPE>::generate();
h_src1[i] = Comparator<TYPE>::generate();
}
// upload source buffer0
std::cout << "upload source buffer0" << std::endl;
RT_CHECK(vx_copy_to_dev(src0_buffer, h_src0.data(), 0, buf_size));
// upload source buffer1
std::cout << "upload source buffer1" << std::endl;
RT_CHECK(vx_copy_to_dev(src1_buffer, h_src1.data(), 0, buf_size));
// upload program
std::cout << "upload program" << std::endl;
RT_CHECK(vx_upload_kernel_file(device, kernel_file, &krnl_buffer));
// upload kernel argument
std::cout << "upload kernel argument" << std::endl;
RT_CHECK(vx_upload_bytes(device, &kernel_arg, sizeof(kernel_arg_t), &args_buffer));
// start device
std::cout << "start device" << std::endl;
RT_CHECK(vx_start(device, krnl_buffer, args_buffer));
// wait for completion
std::cout << "wait for completion" << std::endl;
RT_CHECK(vx_ready_wait(device, VX_MAX_TIMEOUT));
// download destination buffer
std::cout << "download destination buffer" << std::endl;
RT_CHECK(vx_copy_from_dev(h_dst.data(), dst_buffer, 0, buf_size));
// verify result
std::cout << "verify result" << std::endl;
int errors = 0;
for (uint32_t i = 0; i < num_points; ++i) {
auto ref = h_src0[i] + h_src1[i];
auto cur = h_dst[i];
if (!Comparator<TYPE>::compare(cur, ref, i, errors)) {
++errors;
}
}
// cleanup
std::cout << "cleanup" << std::endl;
cleanup();
if (errors != 0) {
std::cout << "Found " << std::dec << errors << " errors!" << std::endl;
std::cout << "FAILED!" << std::endl;
return errors;
}
std::cout << "PASSED!" << std::endl;
return 0;
}