wu-arch/kernels
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

adding opencl convolution benchmark

Browse Source
This commit is contained in:
Blaise Tine
2023-11-14 22:31:30 -08:00
parent 4e7a536918
commit 61e3442ef8
16 changed files with 490 additions and 170 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
tests/opencl/matmul/Makefile
View File

@@ -2,6 +2,6 @@ PROJECT = matmul
SRCS = main.cc
OPTS ?= -n16
OPTS ?= -n32
include ../common.mk

59
tests/opencl/matmul/kernel.cl
View File

@@ -7,43 +7,41 @@ __kernel void matmul(__global float *A,
{
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 localRow = get_local_id(1);
int localCol = get_local_id(0);
int localSize = get_local_size(0); // assuming square local size
float sum = 0.0f;
// Load initial blocks of A and B into local memory
int k = 0;
localA[localRow * localSize + localCol] = A[globalRow * N + k + localCol];
localB[localRow * localSize + localCol] = B[(k + localRow) * N + globalCol];
// Loop over all blocks of both matrices
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];
// Iterate over blocks
for (k = 0; k < N; k += 16) {
// Ensure the initial block is loaded
// Load block of matrix B to local memory, adjusting for column-major access
localB[localRow * localSize + localCol] = B[(k + localRow) * N + globalCol];
// Synchronize to make sure the tiles are loaded
barrier(CLK_LOCAL_MEM_FENCE);
// Compute multiplication for this block
for (int j = 0; j < 16; j++) {
// Multiply the two matrix blocks and accumulate result
for (int j = 0; j < localSize; j++) {
sum += localA[localRow * localSize + j] * localB[j * localSize + localCol];
}
// Load the next block of matrix A into local memory
if (k + 16 < N) {
localA[localRow * localSize + localCol] = A[globalRow * N + k + 16 + localCol];
localB[localRow * localSize + localCol] = B[(k + 16 + localRow) * N + globalCol];
}
}
C[globalRow * N + globalCol] = sum;
}
/*__kernel void matmul(__global float *A, __global float *B, __global float *C, const unsigned int N)
/*__kernel void matmul(__global float *A,
__global float *B,
__global float *C,
const unsigned int N)
{
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 localRow = get_local_id(1);
int localCol = get_local_id(0);
// Static local memory declaration
__local float localA[16][16];
@@ -51,26 +49,21 @@ __kernel void matmul(__global float *A,
float sum = 0.0f;
// Load initial blocks of A and B into local memory
int k = 0;
localA[localRow][localCol] = A[globalRow * N + k + localCol];
localB[localRow][localCol] = B[(k + localRow) * N + globalCol];
// Iterate over blocks
for (k = 0; k < N; k += 16) {
// Ensure the initial block is loaded
for (int k = 0; k < N; k += 16) {
// Load a block of matrix A into local memory
localA[localRow][localCol] = A[globalRow * N + k + localCol];
// Load a block of matrix B into local memory
localB[localRow][localCol] = B[(k + localRow) * N + globalCol];
// Ensure the entire block is loaded
barrier(CLK_LOCAL_MEM_FENCE);
// Compute multiplication for this block
for (int j = 0; j < 16; j++) {
sum += localA[localRow][j] * localB[j][localCol];
}
// Load the next block of matrix A into local memory
if (k + 16 < N) {
localA[localRow][localCol] = A[globalRow * N + k + 16 + localCol];
localB[localRow][localCol] = B[(k + 16 + localRow) * N + globalCol];
}
}
C[globalRow * N + globalCol] = sum;

51
tests/opencl/matmul/main.cc
View File

@@ -10,6 +10,8 @@
#define LOCAL_SIZE 16
#define FLOAT_ULP 6
#define KERNEL_NAME "matmul"
#define CL_CHECK(_expr) \
@@ -56,15 +58,16 @@ static int read_kernel_file(const char* filename, uint8_t** data, size_t* size)
return 0;
}
static bool compare_equal(float a, float b, int ulp = 21) {
union fi_t { int i; float f; };
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;
return std::abs(fa.i - fb.i) <= ulp;
auto d = std::abs(fa.i - fb.i);
return d <= FLOAT_ULP;
}
static void matrix_multiply_cpu(float *A, float *B, float *C, int N) {
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;
@@ -98,7 +101,7 @@ static void cleanup() {
if (kernel_bin) free(kernel_bin);
}
int size = 64;
int size = 32;
static void show_usage() {
printf("Usage: [-n size] [-h: help]\n");
@@ -106,7 +109,7 @@ static void show_usage() {
static void parse_args(int argc, char **argv) {
int c;
while ((c = getopt(argc, argv, "fn:h?")) != -1) {
while ((c = getopt(argc, argv, "n:h?")) != -1) {
switch (c) {
case 'n':
size = atoi(optarg);
@@ -127,6 +130,8 @@ int main (int argc, char **argv) {
// parse command arguments
parse_args(argc, argv);
uint32_t num_points = 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);
@@ -148,7 +153,7 @@ int main (int argc, char **argv) {
printf("Using device: %s\n", device_string);
printf("Allocate device buffers\n");
size_t nbytes = size * size * sizeof(float);
size_t nbytes = num_points * 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));
@@ -176,32 +181,26 @@ int main (int argc, char **argv) {
// Create kernel
kernel = CL_CHECK2(clCreateKernel(program, KERNEL_NAME, &_err));
size_t local_size[2] = {LOCAL_SIZE, LOCAL_SIZE};
size_t global_size[2] = {size, 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, 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));
//CL_CHECK(clSetKernelArg(kernel, 4, local_size[0]*local_size[1]*sizeof(float), NULL));
//CL_CHECK(clSetKernelArg(kernel, 5, local_size[0]*local_size[1]*sizeof(float), NULL));
CL_CHECK(clSetKernelArg(kernel, 4, local_size[0]*local_size[1]*sizeof(float), NULL));
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(size * size);
std::vector<float> h_b(size * size);
std::vector<float> h_c(size * size);
std::vector<float> h_a(num_points);
std::vector<float> h_b(num_points);
std::vector<float> h_c(num_points);
// Initialize values for array members.
for (int i = 0; i < (size * size); ++i) {
#ifdef USE_FLOAT
h_a[i] = (float)rand() / (float)RAND_MAX;
h_b[i] = (float)rand() / (float)RAND_MAX;
#else
h_a[i] = rand();
h_b[i] = rand();
#endif
h_c[i] = 0xdeadbeef;
// 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;
}
// Creating command queue
@@ -223,10 +222,10 @@ 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(size * size);
matrix_multiply_cpu(h_a.data(), h_b.data(), ref_vec.data(), size);
std::vector<float> ref_vec(num_points);
matmul_cpu(ref_vec.data(), h_a.data(), h_b.data(), size);
int errors = 0;
for (int i = 0; i < (size * size); i++) {
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]);