summaryrefslogtreecommitdiff
path: root/GEMV
diff options
context:
space:
mode:
Diffstat (limited to 'GEMV')
-rw-r--r--GEMV/Makefile44
-rw-r--r--GEMV/baselines/cpu/Makefile7
-rw-r--r--GEMV/baselines/cpu/README9
-rw-r--r--GEMV/baselines/cpu/gemv_openmp.c78
-rw-r--r--GEMV/baselines/cpu/gemv_utils.h29
-rw-r--r--GEMV/baselines/gpu/Makefile5
-rw-r--r--GEMV/baselines/gpu/README9
-rw-r--r--GEMV/baselines/gpu/gemv.cu152
-rw-r--r--GEMV/dpu/task.c171
-rw-r--r--GEMV/host/app.c269
-rwxr-xr-xGEMV/support/common.h41
-rw-r--r--GEMV/support/params.h56
-rwxr-xr-xGEMV/support/timer.h62
13 files changed, 932 insertions, 0 deletions
diff --git a/GEMV/Makefile b/GEMV/Makefile
new file mode 100644
index 0000000..fbba9b8
--- /dev/null
+++ b/GEMV/Makefile
@@ -0,0 +1,44 @@
+DPU_DIR := dpu
+HOST_DIR := host
+BUILDDIR ?= bin
+NR_TASKLETS ?= 16
+BL ?= 10
+NR_DPUS ?= 1
+
+define conf_filename
+ ${BUILDDIR}/.NR_DPUS_$(1)_NR_TASKLETS_$(2)_BL_$(3).conf
+endef
+CONF := $(call conf_filename,${NR_DPUS},${NR_TASKLETS},${BL})
+
+HOST_TARGET := ${BUILDDIR}/gemv_host
+DPU_TARGET := ${BUILDDIR}/gemv_dpu
+
+COMMON_INCLUDES := support
+HOST_SOURCES := $(wildcard ${HOST_DIR}/*.c)
+DPU_SOURCES := $(wildcard ${DPU_DIR}/*.c)
+
+.PHONY: all clean test
+
+__dirs := $(shell mkdir -p ${BUILDDIR})
+
+COMMON_FLAGS := -Wall -Wextra -g -I${COMMON_INCLUDES}
+HOST_FLAGS := ${COMMON_FLAGS} -std=c11 -O3 `dpu-pkg-config --cflags --libs dpu` -DNR_TASKLETS=${NR_TASKLETS} -DNR_DPUS=${NR_DPUS} -DBL=${BL}
+DPU_FLAGS := ${COMMON_FLAGS} -O2 -DNR_TASKLETS=${NR_TASKLETS} -DBL=${BL}
+
+all: ${HOST_TARGET} ${DPU_TARGET}
+
+${CONF}:
+ $(RM) $(call conf_filename,*,*)
+ touch ${CONF}
+
+${HOST_TARGET}: ${HOST_SOURCES} ${COMMON_INCLUDES} ${CONF}
+ $(CC) -o $@ ${HOST_SOURCES} ${HOST_FLAGS}
+
+${DPU_TARGET}: ${DPU_SOURCES} ${COMMON_INCLUDES} ${CONF}
+ dpu-upmem-dpurte-clang ${DPU_FLAGS} -o $@ ${DPU_SOURCES}
+
+clean:
+ $(RM) -r $(BUILDDIR)
+
+test: all
+ ./${HOST_TARGET} -m 1024 -n 1024
diff --git a/GEMV/baselines/cpu/Makefile b/GEMV/baselines/cpu/Makefile
new file mode 100644
index 0000000..c779651
--- /dev/null
+++ b/GEMV/baselines/cpu/Makefile
@@ -0,0 +1,7 @@
+all:
+ gcc -o gemv -fopenmp gemv_openmp.c
+
+clean:
+ rm gemv
+
+
diff --git a/GEMV/baselines/cpu/README b/GEMV/baselines/cpu/README
new file mode 100644
index 0000000..92906c3
--- /dev/null
+++ b/GEMV/baselines/cpu/README
@@ -0,0 +1,9 @@
+Matrix-Vector Multiplication (GEMV)
+
+Compilation instructions:
+
+ make
+
+Execution instructions
+
+ ./gemv
diff --git a/GEMV/baselines/cpu/gemv_openmp.c b/GEMV/baselines/cpu/gemv_openmp.c
new file mode 100644
index 0000000..307e03b
--- /dev/null
+++ b/GEMV/baselines/cpu/gemv_openmp.c
@@ -0,0 +1,78 @@
+#include <stdlib.h>
+#include <stdio.h>
+#include "../../support/timer.h"
+#include "gemv_utils.h"
+
+int main(int argc, char *argv[])
+{
+ const size_t rows = 20480;
+ const size_t cols = 8192;
+
+ double **A, *b, *x;
+
+ b = (double*) malloc(sizeof(double)*rows);
+ x = (double*) malloc(sizeof(double)*cols);
+
+ allocate_dense(rows, cols, &A);
+
+ make_hilbert_mat(rows,cols, &A);
+
+#pragma omp parallel
+ {
+#pragma omp for
+ for (size_t i = 0; i < cols; i++) {
+ x[i] = (double) i+1 ;
+ }
+
+#pragma omp for
+ for (size_t i = 0; i < rows; i++) {
+ b[i] = (double) 0.0;
+ }
+ }
+
+ Timer timer;
+ start(&timer, 0, 0);
+
+
+ gemv(A, x, rows, cols, &b);
+
+ stop(&timer, 0);
+
+
+ printf("Kernel ");
+ print(&timer, 0, 1);
+ printf("\n");
+
+#if 0
+ print_vec(x, rows);
+ print_mat(A, rows, cols);
+ print_vec(b, rows);
+#endif
+
+ printf("sum(x) = %f, sum(Ax) = %f\n", sum_vec(x,cols), sum_vec(b,rows));
+ return 0;
+}
+
+void gemv(double** A, double* x, size_t rows, size_t cols, double** b) {
+#pragma omp parallel for
+ for (size_t i = 0; i < rows; i ++ )
+ for (size_t j = 0; j < cols; j ++ ) {
+ (*b)[i] = (*b)[i] + A[i][j]*x[j];
+ }
+}
+
+void make_hilbert_mat(size_t rows, size_t cols, double*** A) {
+#pragma omp parallel for
+ for (size_t i = 0; i < rows; i++) {
+ for (size_t j = 0; j < cols; j++) {
+ (*A)[i][j] = 1.0/( (double) i + (double) j + 1.0);
+ }
+ }
+}
+
+double sum_vec(double* vec, size_t rows) {
+ double sum = 0.0;
+#pragma omp parallel for reduction(+:sum)
+ for (int i = 0; i < rows; i++) sum = sum + vec[i];
+ return sum;
+}
diff --git a/GEMV/baselines/cpu/gemv_utils.h b/GEMV/baselines/cpu/gemv_utils.h
new file mode 100644
index 0000000..605f148
--- /dev/null
+++ b/GEMV/baselines/cpu/gemv_utils.h
@@ -0,0 +1,29 @@
+void allocate_dense(size_t rows,size_t cols, double*** dense) {
+
+ *dense = malloc(sizeof(double)*rows);
+ **dense = malloc(sizeof(double)*rows*cols);
+
+ for (size_t i=0; i < rows; i++ ) {
+ (*dense)[i] = (*dense)[0] + i*cols;
+ }
+
+}
+
+void print_mat(double** A, size_t rows, size_t cols) {
+ for (size_t i = 0; i < rows; i++) {
+ for (size_t j = 0; j < cols; j++) {
+ printf("%f ", A[i][j]);
+ }
+ printf("\n");
+ }
+}
+
+void print_vec(double* b, size_t rows) {
+ for (size_t i = 0; i < rows; i++) {
+ printf("%f\n", b[i]);
+ }
+}
+
+void gemv(double** A, double* x, size_t rows, size_t cols, double** b);
+void make_hilbert_mat(size_t rows, size_t cols, double*** A);
+double sum_vec(double* vec, size_t rows);
diff --git a/GEMV/baselines/gpu/Makefile b/GEMV/baselines/gpu/Makefile
new file mode 100644
index 0000000..bd1be55
--- /dev/null
+++ b/GEMV/baselines/gpu/Makefile
@@ -0,0 +1,5 @@
+all:
+ /usr/local/cuda/bin/nvcc gemv.cu -I/usr/local/cuda/include -lm -o gemv
+
+clean:
+ rm gemv
diff --git a/GEMV/baselines/gpu/README b/GEMV/baselines/gpu/README
new file mode 100644
index 0000000..92906c3
--- /dev/null
+++ b/GEMV/baselines/gpu/README
@@ -0,0 +1,9 @@
+Matrix-Vector Multiplication (GEMV)
+
+Compilation instructions:
+
+ make
+
+Execution instructions
+
+ ./gemv
diff --git a/GEMV/baselines/gpu/gemv.cu b/GEMV/baselines/gpu/gemv.cu
new file mode 100644
index 0000000..5c0e240
--- /dev/null
+++ b/GEMV/baselines/gpu/gemv.cu
@@ -0,0 +1,152 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include <sys/time.h>
+#include <cuda.h>
+
+#define THREAD 128
+
+#define T int
+
+__global__ void gemv(int m, int n, T *adim, T *b, T *d_ans);
+
+void cgemv(int m, int n, T *adim, T *b, T *d_ans);
+
+double gettime()
+{
+struct timeval tv;
+gettimeofday(&tv, NULL);
+return tv.tv_sec + (double)tv.tv_usec*1.0e-6;
+}
+
+int main(int argc, char **argv)
+{
+/* for CPU */
+int i, j;
+int *bdim, *c, *ans, *h_ans;
+//double start, stop;
+//double cpu_time, gpu_time;
+int n = 8192;
+int m = 20480;
+
+bdim = (T*)malloc(sizeof(T) *m*n);
+c = (T*)malloc(sizeof(T) *n);
+ans = (T*)malloc(sizeof(T) *m);
+h_ans = (T*)malloc(sizeof(T) *m);
+
+/* for GPU */
+T *d_bdim, *d_c, *d_ans;
+cudaMalloc((void **)&d_bdim, sizeof(T)*m*n);
+cudaMalloc((void **)&d_c, sizeof(T)*n);
+cudaMalloc((void **)&d_ans, sizeof(T)*m);
+
+for(i = 0; i < n; i++)
+{
+c[i] = 1;
+for(j = 0; j < m; j++)
+bdim[i*m+j] = 1;
+}
+
+//start = gettime();
+cgemv(m, n, bdim, c, ans);
+//stop = gettime();
+//cpu_time=stop - start;
+
+// Event creation
+cudaEvent_t start, stop;
+cudaEventCreate(&start);
+cudaEventCreate(&stop);
+float time1 = 0;
+
+
+cudaMemcpy(d_bdim, bdim, sizeof(T)*m*n, cudaMemcpyHostToDevice);
+cudaMemcpy(d_c, c, sizeof(T)*n, cudaMemcpyHostToDevice);
+
+// Start timer
+cudaEventRecord( start, 0 );
+//start = gettime();
+gemv<<<m, THREAD>>>(m, n, d_bdim, d_c, d_ans);
+//stop = gettime();
+// End timer
+cudaEventRecord( stop, 0 );
+cudaEventSynchronize( stop );
+cudaEventElapsedTime( &time1, start, stop );
+
+//gpu_time=stop - start;
+
+cudaMemcpy(h_ans, d_ans, sizeof(T)*m, cudaMemcpyDeviceToHost);
+
+//printf("cpu_time : %.6f[sec]\n",cpu_time);
+//printf("gpu_time : %.6f[sec]\n",gpu_time);
+//printf("%f x\n", cpu_time / gpu_time);
+
+
+for(i = 0; i < m; i++)
+printf("%d -- %d\n", ans[i], h_ans[i]);
+
+printf("Execution time = %f ms\n", time1);
+
+
+free(bdim);
+free(c);
+free(ans);
+free(h_ans);
+cudaFree(d_bdim);
+cudaFree(d_c);
+cudaFree(d_ans);
+
+return 0;
+}
+
+__global__ void gemv(int m, int n, T* adim, T* b, T* d_ans)
+{
+int i;
+int div = n/THREAD;
+__shared__ T tmp[THREAD];
+
+tmp[threadIdx.x] = 0.0;
+
+for(i = 0; i < div; i++)
+{
+tmp[threadIdx.x] += adim[blockIdx.x*n+i*THREAD+threadIdx.x] * b[i * THREAD + threadIdx.x];
+}
+if(threadIdx.x < m%THREAD)
+tmp[threadIdx.x] += adim[blockIdx.x*n+THREAD*div+threadIdx.x] * b[THREAD * div + threadIdx.x];
+
+__syncthreads();
+
+for(i = THREAD / 2; i > 31; i = i / 2)
+{
+if(threadIdx.x < i)
+tmp[threadIdx.x] += tmp[threadIdx.x + i];
+__syncthreads();
+}
+
+if(threadIdx.x < 16)
+{
+tmp[threadIdx.x] += tmp[threadIdx.x + 16];
+__syncthreads();
+tmp[threadIdx.x] += tmp[threadIdx.x + 8];
+__syncthreads();
+tmp[threadIdx.x] += tmp[threadIdx.x + 4];
+__syncthreads();
+tmp[threadIdx.x] += tmp[threadIdx.x + 2];
+__syncthreads();
+tmp[threadIdx.x] += tmp[threadIdx.x + 1];
+__syncthreads();
+}
+
+
+if(threadIdx.x == 0)
+d_ans[blockIdx.x] = tmp[0];
+
+}
+
+void cgemv(int m, int n, T *adim, T *b, T *d_ans)
+{
+int i, j;
+
+for(i = 0; i < m; i++)
+for(j = 0; j < n; j++)
+d_ans[i] += adim[i*n+j] * b[j];
+
+}
diff --git a/GEMV/dpu/task.c b/GEMV/dpu/task.c
new file mode 100644
index 0000000..de3e554
--- /dev/null
+++ b/GEMV/dpu/task.c
@@ -0,0 +1,171 @@
+/*
+ * Matrix vector multiplication with multiple tasklet
+ *
+ */
+#include <stdint.h>
+#include <stdio.h>
+#include <defs.h>
+#include <mram.h>
+#include <alloc.h>
+#include <barrier.h>
+#include <seqread.h>
+
+#include "../support/common.h"
+
+__host dpu_arguments_t DPU_INPUT_ARGUMENTS;
+
+// GEMV
+static void gemv(T *bufferC, T *bufferA, T *bufferB, int pos) {
+ for (unsigned int i = 0; i < BLOCK_SIZE / sizeof(T); i++) {
+ bufferC[pos] += bufferA[i] * bufferB[i];
+ }
+ return;
+}
+
+// Barrier
+BARRIER_INIT(my_barrier, NR_TASKLETS);
+
+// main
+int main() {
+ unsigned int tasklet_id = me();
+#if PRINT
+ printf("tasklet_id = %u\n", tasklet_id);
+#endif
+ if (tasklet_id == 0){ // Initialize once the cycle counter
+ mem_reset(); // Reset the heap
+ }
+ // Barrier
+ barrier_wait(&my_barrier);
+
+ int32_t n_size = DPU_INPUT_ARGUMENTS.n_size;
+ int32_t n_size_pad = DPU_INPUT_ARGUMENTS.n_size_pad;
+ uint32_t nr_rows = DPU_INPUT_ARGUMENTS.nr_rows;
+ uint32_t max_rows = DPU_INPUT_ARGUMENTS.max_rows;
+
+
+ unsigned int nrows = nr_rows;
+ unsigned int rows_per_tasklet;
+ unsigned int start_row;
+ unsigned int chunks = nrows / (NR_TASKLETS + NR_TASKLETS);
+ unsigned int dbl_chunks = chunks + chunks;
+ rows_per_tasklet = dbl_chunks;
+ unsigned int rest_rows = nrows % (NR_TASKLETS + NR_TASKLETS);
+
+ if ((tasklet_id + tasklet_id) < rest_rows)
+ rows_per_tasklet += 2;
+ if (rest_rows > 0) {
+ if ((tasklet_id + tasklet_id) >= rest_rows) {
+ unsigned int hlf_rest_rows = rest_rows >> 1;
+ if ((rest_rows & 1) == 1)
+ start_row = (hlf_rest_rows + 1) * (dbl_chunks + 2) + (tasklet_id - 1 - hlf_rest_rows) * dbl_chunks;
+ else
+ start_row = (hlf_rest_rows) * (dbl_chunks + 2) + (tasklet_id - hlf_rest_rows) * dbl_chunks;
+ } else
+ start_row = tasklet_id * (dbl_chunks + 2);
+ } else {
+ start_row = tasklet_id * (dbl_chunks);
+ }
+
+ // Address of the current row in MRAM
+ uint32_t mram_base_addr_A = (uint32_t) (DPU_MRAM_HEAP_POINTER + start_row * n_size * sizeof(T));
+ uint32_t mram_base_addr_B = (uint32_t) (DPU_MRAM_HEAP_POINTER + max_rows * n_size_pad * sizeof(T));
+ uint32_t mram_base_addr_C = (uint32_t) (DPU_MRAM_HEAP_POINTER + max_rows * n_size_pad * sizeof(T) + n_size_pad * sizeof(T) + start_row * sizeof(T));
+ uint32_t mram_temp_addr_A = mram_base_addr_A;
+ uint32_t mram_temp_addr_B = mram_base_addr_B;
+
+ // Inititalize a local cache to store the MRAM block
+ T *cache_A = (T *) mem_alloc(BLOCK_SIZE + 8);
+ T *cache_A_aux = (T *) mem_alloc(8);
+ T *cache_B = (T *) mem_alloc(BLOCK_SIZE);
+ T *cache_C = (T *) mem_alloc(8);
+
+ int offset = 0;
+
+ // Iterate over nr_rows
+ for (unsigned int i = start_row; i < start_row + rows_per_tasklet; i += 2) {
+
+ mram_temp_addr_A = (uint32_t) (DPU_MRAM_HEAP_POINTER + i * n_size * sizeof(T));
+ mram_temp_addr_B = mram_base_addr_B;
+
+ cache_C[0] = 0;
+ cache_C[1] = 0;
+ for(unsigned int pos = 0; pos < 2 && i + pos < nr_rows; pos++){
+ int n = 0, j;
+ for (n = 0; n < (int32_t) (n_size - (BLOCK_SIZE/sizeof(T))); n += (BLOCK_SIZE / sizeof(T)))
+ {
+
+ mram_read((__mram_ptr void const*) (mram_temp_addr_A), cache_A, BLOCK_SIZE);
+ mram_read((__mram_ptr void const*) (mram_temp_addr_B), cache_B, BLOCK_SIZE);
+
+ if(offset)
+ {
+
+ for(unsigned int off = 0; off < (BLOCK_SIZE / sizeof(T)) - 1; off++)
+ {
+ cache_A[off] = cache_A[off + 1];
+ }
+
+ mram_read((__mram_ptr void const*) (mram_temp_addr_A + BLOCK_SIZE), cache_A_aux, 8);
+
+ cache_A[BLOCK_SIZE / sizeof(T) - 1] = cache_A_aux[0];
+ }
+
+ // Compute GEMV
+ gemv(cache_C, cache_A, cache_B, pos);
+
+ // Update memory addresses
+ mram_temp_addr_A += BLOCK_SIZE;
+ mram_temp_addr_B += BLOCK_SIZE;
+ }
+
+ mram_read((__mram_ptr void const*) (mram_temp_addr_A), cache_A, BLOCK_SIZE);
+
+
+ if(offset)
+ {
+ for(unsigned int off = 0; off < (BLOCK_SIZE / sizeof(T)) -1; off++)
+ {
+
+ cache_A[off] = cache_A[off + 1];
+ }
+
+ mram_read((__mram_ptr void const*) (mram_temp_addr_A + BLOCK_SIZE ), cache_A_aux, 8);
+
+ cache_A[BLOCK_SIZE / sizeof(T) - 1] = cache_A_aux[0];
+ }
+
+
+ mram_read((__mram_ptr void const*) (mram_temp_addr_B), cache_B, BLOCK_SIZE);
+
+ for (j = 0; j < (int) (n_size - n); j++) {
+ // Compute GEMV
+ if(j >= (int)(BLOCK_SIZE / sizeof(T))){
+ printf("error\n");
+ break;
+ }
+ cache_C[pos] += cache_A[j] * cache_B[j];
+ }
+
+
+ mram_temp_addr_A += (BLOCK_SIZE - ((BLOCK_SIZE / sizeof(T)) - (n_size - n)) * sizeof(T));
+ mram_temp_addr_B = mram_base_addr_B;
+
+ if(mram_temp_addr_A % 8 != 0)
+ {
+ offset = 1;
+ }
+ else
+ {
+ offset = 0;
+ }
+ }
+ // Write cache to current MRAM block
+ mram_write(cache_C, (__mram_ptr void *) (mram_base_addr_C), 8);
+
+ // Update memory address
+ mram_base_addr_C += 2 * sizeof(T);
+
+ }
+
+ return 0;
+}
diff --git a/GEMV/host/app.c b/GEMV/host/app.c
new file mode 100644
index 0000000..de2e847
--- /dev/null
+++ b/GEMV/host/app.c
@@ -0,0 +1,269 @@
+/**
+ * app.c
+ * GEMV Host Application Source File
+ *
+ */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdbool.h>
+#include <string.h>
+#include <dpu.h>
+#include <dpu_log.h>
+#include <unistd.h>
+#include <getopt.h>
+#include <assert.h>
+
+#if ENERGY
+#include <dpu_probe.h>
+#endif
+
+#include "../support/common.h"
+#include "../support/timer.h"
+#include "../support/params.h"
+
+// Define the DPU Binary path as DPU_BINARY here
+#ifndef DPU_BINARY
+#define DPU_BINARY "./bin/gemv_dpu"
+#endif
+
+static T* A;
+static T* B;
+static T* C;
+static T* C_dpu;
+
+// Create input arrays
+static void init_data(T* A, T* B, unsigned int m_size, unsigned int n_size) {
+ srand(0);
+
+ for (unsigned int i = 0; i < m_size * n_size; i++)
+ {
+ A[i] = (unsigned int) (rand()%50);
+ }
+
+ for (unsigned int i = 0; i < n_size; i++)
+ {
+ B[i] = (unsigned int) (rand()%50);
+ }
+}
+
+// Compute output in the host
+static void gemv_host(T* C, T* A, T* B, unsigned int m_size, unsigned int n_size) {
+ for (unsigned int i = 0; i < m_size; i++)
+ {
+ C[i] = 0;
+ }
+
+ for (unsigned int m = 0; m < m_size; m++) {
+ for (unsigned int n = 0; n < n_size; n++)
+ {
+ C[m] += A[m * n_size + n] * B[n];
+ }
+ }
+}
+
+// Main of the Host Application
+int main(int argc, char **argv) {
+
+ struct Params p = input_params(argc, argv);
+
+ struct dpu_set_t dpu_set, dpu;
+ uint32_t nr_of_dpus;
+
+ // Allocate DPUs and load binary
+ DPU_ASSERT(dpu_alloc(NR_DPUS, NULL, &dpu_set));
+ DPU_ASSERT(dpu_load(dpu_set, DPU_BINARY, NULL));
+ DPU_ASSERT(dpu_get_nr_dpus(dpu_set, &nr_of_dpus));
+
+#if ENERGY
+ struct dpu_probe_t probe;
+ DPU_ASSERT(dpu_probe_init("energy_probe", &probe));
+#endif
+
+ unsigned int i;
+ unsigned int m_size = p.m_size;
+ unsigned int n_size = p.n_size;
+
+ // Initialize help data
+ dpu_info = (struct dpu_info_t *) malloc(nr_of_dpus * sizeof(struct dpu_info_t));
+ dpu_arguments_t *input_args = (dpu_arguments_t *) malloc(nr_of_dpus * sizeof(dpu_arguments_t));
+ uint32_t max_rows_per_dpu = 0;
+ uint32_t n_size_pad = n_size;
+ if(n_size % 2 == 1)
+ {
+ n_size_pad++;
+ }
+
+ i = 0;
+ DPU_FOREACH(dpu_set, dpu, i) {
+ uint32_t rows_per_dpu;
+ uint32_t prev_rows_dpu = 0;
+ uint32_t chunks = m_size / nr_of_dpus;
+ rows_per_dpu = chunks;
+ uint32_t rest_rows = m_size % nr_of_dpus;
+ if (i < rest_rows)
+ rows_per_dpu++;
+ if (rest_rows > 0) {
+ if (i >= rest_rows)
+ prev_rows_dpu = rest_rows * (chunks + 1) + (i - rest_rows) * chunks;
+ else
+ prev_rows_dpu = i * (chunks + 1);
+ } else {
+ prev_rows_dpu = i * chunks;
+ }
+
+ // Keep max rows for parallel transfers
+ uint32_t rows_per_dpu_pad = rows_per_dpu;
+ if (rows_per_dpu_pad % 2 == 1) // 4-byte elements
+ rows_per_dpu_pad++;
+ if (rows_per_dpu_pad > max_rows_per_dpu)
+ max_rows_per_dpu = rows_per_dpu_pad;
+
+ dpu_info[i].rows_per_dpu = rows_per_dpu;
+ dpu_info[i].rows_per_dpu_pad = rows_per_dpu_pad;
+ dpu_info[i].prev_rows_dpu = prev_rows_dpu;
+
+ // Copy input arguments to DPU
+ input_args[i].n_size = n_size;
+ input_args[i].n_size_pad = n_size_pad;
+ input_args[i].nr_rows = rows_per_dpu;
+ }
+
+ A = malloc(max_rows_per_dpu * nr_of_dpus * n_size_pad * sizeof(T));
+ B = malloc(n_size_pad * sizeof(T));
+ C = malloc(max_rows_per_dpu * nr_of_dpus * sizeof(T));
+
+ // Initialize data with arbitrary data
+ init_data(A, B, m_size, n_size);
+
+ // Timer
+ Timer timer;
+
+ // Compute output on CPU (performance comparison and verification purposes)
+ start(&timer, 0, 0);
+ gemv_host(C, A, B, m_size, n_size);
+ stop(&timer, 0);
+ for (unsigned int rep = 0; rep < p.n_warmup + p.n_reps; rep++) {
+
+
+
+ if (rep >= p.n_warmup)
+ start(&timer, 1, rep - p.n_warmup);
+ // Input arguments
+ i = 0;
+ DPU_FOREACH(dpu_set, dpu, i) {
+ // Copy input arguments to DPU
+ input_args[i].max_rows = max_rows_per_dpu;
+
+ DPU_ASSERT(dpu_prepare_xfer(dpu, input_args + i));
+ }
+
+ DPU_ASSERT(dpu_push_xfer(dpu_set, DPU_XFER_TO_DPU, "DPU_INPUT_ARGUMENTS", 0, sizeof(dpu_arguments_t), DPU_XFER_DEFAULT));
+
+ // Copy input array and vector
+ i = 0;
+ DPU_FOREACH(dpu_set, dpu, i) {
+ DPU_ASSERT(dpu_prepare_xfer(dpu, A + dpu_info[i].prev_rows_dpu * n_size));
+ }
+ DPU_ASSERT(dpu_push_xfer(dpu_set, DPU_XFER_TO_DPU, DPU_MRAM_HEAP_POINTER_NAME, 0, max_rows_per_dpu * n_size_pad * sizeof(T), DPU_XFER_DEFAULT));
+ DPU_FOREACH(dpu_set, dpu, i) {
+ DPU_ASSERT(dpu_prepare_xfer(dpu, B));
+ }
+ DPU_ASSERT(dpu_push_xfer(dpu_set, DPU_XFER_TO_DPU, DPU_MRAM_HEAP_POINTER_NAME, max_rows_per_dpu * n_size_pad * sizeof(T) , n_size_pad * sizeof(T), DPU_XFER_DEFAULT));
+
+ if (rep >= p.n_warmup)
+ stop(&timer, 1);
+
+ // Run kernel on DPUs
+ if (rep >= p.n_warmup)
+ {
+ start(&timer, 2, rep - p.n_warmup);
+#if ENERGY
+ DPU_ASSERT(dpu_probe_start(&probe));
+#endif
+ }
+
+ DPU_ASSERT(dpu_launch(dpu_set, DPU_SYNCHRONOUS));
+
+ if (rep >= p.n_warmup)
+ {
+ stop(&timer, 2);
+#if ENERGY
+ DPU_ASSERT(dpu_probe_stop(&probe));
+#endif
+ }
+#if PRINT
+ // Display DPU Logs
+ DPU_FOREACH(dpu_set, dpu) {
+ DPU_ASSERT(dpulog_read_for_dpu(dpu.dpu, stdout));
+ }
+#endif
+
+ // Retrieve results
+ C_dpu = malloc(max_rows_per_dpu * nr_of_dpus * sizeof(T));
+ if (rep >= p.n_warmup)
+ start(&timer, 3, rep - p.n_warmup);
+ i = 0;
+ DPU_FOREACH(dpu_set, dpu, i) {
+ DPU_ASSERT(dpu_prepare_xfer(dpu, C_dpu + i * max_rows_per_dpu));
+ }
+ DPU_ASSERT(dpu_push_xfer(dpu_set, DPU_XFER_FROM_DPU, DPU_MRAM_HEAP_POINTER_NAME, max_rows_per_dpu * n_size_pad * sizeof(T) + n_size_pad * sizeof(T), max_rows_per_dpu * sizeof(T), DPU_XFER_DEFAULT));
+ if(rep >= p.n_warmup)
+ stop(&timer, 3);
+ }
+#if ENERGY
+ double acc_energy, avg_energy, acc_time, avg_time;
+ DPU_ASSERT(dpu_probe_get(&probe, DPU_ENERGY, DPU_ACCUMULATE, &acc_energy));
+ DPU_ASSERT(dpu_probe_get(&probe, DPU_ENERGY, DPU_AVERAGE, &avg_energy));
+ DPU_ASSERT(dpu_probe_get(&probe, DPU_TIME, DPU_ACCUMULATE, &acc_time));
+ DPU_ASSERT(dpu_probe_get(&probe, DPU_TIME, DPU_AVERAGE, &avg_time));
+#endif
+
+ // Print timing results
+ printf("CPU Version Time (ms): ");
+ print(&timer, 0, 1);
+ printf("CPU-DPU Time (ms): ");
+ print(&timer, 1, p.n_reps);
+ printf("DPU Kernel Time (ms): ");
+ print(&timer, 2, p.n_reps);
+ printf("DPU-CPU Time (ms): ");
+ print(&timer, 3, p.n_reps);
+
+#if ENERGY
+ printf("Energy (J): %f J\t", avg_energy);
+#endif
+
+ // Check output
+ bool status = true;
+ unsigned int n,j;
+ i = 0;
+ for (n = 0; n < nr_of_dpus; n++) {
+ for (j = 0; j < dpu_info[n].rows_per_dpu; j++) {
+ if(C[i] != C_dpu[n * max_rows_per_dpu + j]) {
+ status = false;
+#if PRINT
+ // printf("%d: %d -- %d\n", i, C[i], C_dpu[n * max_rows_per_dpu + j]);
+#endif
+ }
+ i++;
+ }
+ }
+ if (status) {
+ printf("[" ANSI_COLOR_GREEN "OK" ANSI_COLOR_RESET "] Outputs are equal\n");
+ } else {
+ printf("[" ANSI_COLOR_RED "ERROR" ANSI_COLOR_RESET "] Outputs differ!\n");
+ }
+
+ // Deallocation
+ free(A);
+ free(B);
+ free(C);
+ free(C_dpu);
+ DPU_ASSERT(dpu_free(dpu_set));
+
+#if ENERGY
+ DPU_ASSERT(dpu_probe_deinit(&probe));
+#endif
+
+ return status ? 0 : -1;
+}
diff --git a/GEMV/support/common.h b/GEMV/support/common.h
new file mode 100755
index 0000000..0deebcb
--- /dev/null
+++ b/GEMV/support/common.h
@@ -0,0 +1,41 @@
+#ifndef _COMMON_H_
+#define _COMMON_H_
+
+// Structures used by both the host and the dpu to communicate information
+typedef struct {
+ uint32_t n_size;
+ uint32_t n_size_pad;
+ uint32_t nr_rows;
+ uint32_t max_rows;
+} dpu_arguments_t;
+
+// Specific information for each DPU
+struct dpu_info_t {
+ uint32_t rows_per_dpu;
+ uint32_t rows_per_dpu_pad;
+ uint32_t prev_rows_dpu;
+};
+struct dpu_info_t *dpu_info;
+
+// Transfer size between MRAM and WRAM
+#ifdef BL
+#define BLOCK_SIZE_LOG2 BL
+#define BLOCK_SIZE (1 << BLOCK_SIZE_LOG2)
+#else
+#define BLOCK_SIZE_LOG2 8
+#define BLOCK_SIZE (1 << BLOCK_SIZE_LOG2)
+#define BL BLOCK_SIZE_LOG2
+#endif
+
+// Data type
+#define T uint32_t
+
+#ifndef ENERGY
+#define ENERGY 0
+#endif
+#define PRINT 0
+
+#define ANSI_COLOR_RED "\x1b[31m"
+#define ANSI_COLOR_GREEN "\x1b[32m"
+#define ANSI_COLOR_RESET "\x1b[0m"
+#endif
diff --git a/GEMV/support/params.h b/GEMV/support/params.h
new file mode 100644
index 0000000..526c71c
--- /dev/null
+++ b/GEMV/support/params.h
@@ -0,0 +1,56 @@
+#ifndef _PARAMS_H_
+#define _PARAMS_H_
+
+#include "common.h"
+
+typedef struct Params {
+ unsigned int m_size;
+ unsigned int n_size;
+ unsigned int n_warmup;
+ unsigned int n_reps;
+}Params;
+
+static void usage() {
+ fprintf(stderr,
+ "\nUsage: ./program [options]"
+ "\n"
+ "\nGeneral options:"
+ "\n -h help"
+ "\n -w <W> # of untimed warmup iterations (default=1)"
+ "\n -e <E> # of timed repetition iterations (default=3)"
+ "\n"
+ "\nBenchmark-specific options:"
+ "\n -m <I> m_size (default=8192 elements)"
+ "\n -n <I> n_size (default=8192 elements)"
+ "\n");
+}
+
+struct Params input_params(int argc, char **argv) {
+ struct Params p;
+ p.m_size = 8192;
+ p.n_size = 8192;
+ p.n_warmup = 1;
+ p.n_reps = 3;
+
+ int opt;
+ while((opt = getopt(argc, argv, "hm:n:w:e:")) >= 0) {
+ switch(opt) {
+ case 'h':
+ usage();
+ exit(0);
+ break;
+ case 'm': p.m_size = atoi(optarg); break;
+ case 'n': p.n_size = atoi(optarg); break;
+ case 'w': p.n_warmup = atoi(optarg); break;
+ case 'e': p.n_reps = atoi(optarg); break;
+ default:
+ fprintf(stderr, "\nUnrecognized option!\n");
+ usage();
+ exit(0);
+ }
+ }
+ assert(NR_DPUS > 0 && "Invalid # of dpus!");
+
+ return p;
+}
+#endif
diff --git a/GEMV/support/timer.h b/GEMV/support/timer.h
new file mode 100755
index 0000000..2ea4119
--- /dev/null
+++ b/GEMV/support/timer.h
@@ -0,0 +1,62 @@
+/*
+ * Copyright (c) 2016 University of Cordoba and University of Illinois
+ * All rights reserved.
+ *
+ * Developed by: IMPACT Research Group
+ * University of Cordoba and University of Illinois
+ * http://impact.crhc.illinois.edu/
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * with the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * > Redistributions of source code must retain the above copyright notice,
+ * this list of conditions and the following disclaimers.
+ * > Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimers in the
+ * documentation and/or other materials provided with the distribution.
+ * > Neither the names of IMPACT Research Group, University of Cordoba,
+ * University of Illinois nor the names of its contributors may be used
+ * to endorse or promote products derived from this Software without
+ * specific prior written permission.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS WITH
+ * THE SOFTWARE.
+ *
+ */
+
+#include <sys/time.h>
+
+typedef struct Timer{
+
+ struct timeval startTime[4];
+ struct timeval stopTime[4];
+ double time[4];
+
+}Timer;
+
+void start(Timer *timer, int i, int rep) {
+ if(rep == 0) {
+ timer->time[i] = 0.0;
+ }
+ gettimeofday(&timer->startTime[i], NULL);
+}
+
+void stop(Timer *timer, int i) {
+ gettimeofday(&timer->stopTime[i], NULL);
+ timer->time[i] += (timer->stopTime[i].tv_sec - timer->startTime[i].tv_sec) * 1000000.0 +
+ (timer->stopTime[i].tv_usec - timer->startTime[i].tv_usec);
+ //printf("Time (ms): %f\t",((timer->stopTime[i].tv_sec - timer->startTime[i].tv_sec) * 1000000.0 +
+ // (timer->stopTime[i].tv_usec - timer->startTime[i].tv_usec)) / 1000);
+
+}
+
+void print(Timer *timer, int i, int REP) { printf("%f\t", timer->time[i] / (1000 * REP)); }
generated by cgit v1.2.3 (git 2.39.1) at 2024-11-25 21:20:28 +0000