From 116f001e9842e9a371dac2853d7deb387fff6092 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Thu, 6 Aug 2026 06:12:56 -0700 Subject: [PATCH 1/7] [sundials-omp] add the example --- src/sundials-omp/CMakeLists.txt | 10 ++ src/sundials-omp/LICENSE | 24 +++++ src/sundials-omp/Makefile | 68 ++++++++++++++ src/sundials-omp/Makefile.aomp | 72 +++++++++++++++ src/sundials-omp/Makefile.nvc | 69 ++++++++++++++ src/sundials-omp/kernels.cpp | 158 ++++++++++++++++++++++++++++++++ 6 files changed, 401 insertions(+) create mode 100644 src/sundials-omp/CMakeLists.txt create mode 100644 src/sundials-omp/LICENSE create mode 100644 src/sundials-omp/Makefile create mode 100644 src/sundials-omp/Makefile.aomp create mode 100644 src/sundials-omp/Makefile.nvc create mode 100644 src/sundials-omp/kernels.cpp diff --git a/src/sundials-omp/CMakeLists.txt b/src/sundials-omp/CMakeLists.txt new file mode 100644 index 0000000000..2db61d95d1 --- /dev/null +++ b/src/sundials-omp/CMakeLists.txt @@ -0,0 +1,10 @@ +# sundials-omp/CMakeLists.txt + +set(INC_DIR "${CMAKE_CURRENT_LIST_DIR}/../sundials-cuda") + +add_hecbench_benchmark( + NAME sundials + MODEL omp + SOURCES kernels.cpp ${INC_DIR}/utils.cpp ${INC_DIR}/main.cpp + CATEGORIES simulation algorithms math +) diff --git a/src/sundials-omp/LICENSE b/src/sundials-omp/LICENSE new file mode 100644 index 0000000000..d835b2db6d --- /dev/null +++ b/src/sundials-omp/LICENSE @@ -0,0 +1,24 @@ +Copyright (c) 2020 - 2021 HPCAS Lab (https://passlab.github.io) +from University of North Carolina at Charlotte. All rights reserved. Funding for this research and +development was provided by the National Science Foundation +under award number CISE SHF-1551182 and CISE SHF-2015254. + +Redistribution and use in source and binary forms, with or without modification, +are permitted provided that the following conditions are met: + +1. Redistributions of source code must retain the above copyright notice, + this list of conditions and the following disclaimer. +2. Redistributions in binary form must reproduce the above copyright notice, + this list of conditions and the following disclaimer in the documentation + and/or other materials provided with the distribution. +3. Neither the name of the copyright holder nor the names of its contributors may + be used to endorse or promote products derived from this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR +IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND +FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR +CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, +WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. diff --git a/src/sundials-omp/Makefile b/src/sundials-omp/Makefile new file mode 100644 index 0000000000..7f3d19cc2b --- /dev/null +++ b/src/sundials-omp/Makefile @@ -0,0 +1,68 @@ +#=============================================================================== +# User Options +#=============================================================================== + +# Compiler can be set below, or via environment variable +CC = icpx +OPTIMIZE = yes +DEBUG = no +DEVICE = gpu +LAUNCHER = + +#=============================================================================== +# Program name & source code list +#=============================================================================== + +program = main + +source = main.cpp kernels.cpp utils.cpp + +obj = $(source:.cpp=.o) + +#=============================================================================== +# Sets Flags +#=============================================================================== + +# Standard Flags +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall -I../sundials-cuda + +# Linker Flags +LDFLAGS = + +# Debug Flags +ifeq ($(DEBUG),yes) + CFLAGS += -g + LDFLAGS += -g +endif + +# Optimization Flags +ifeq ($(OPTIMIZE),yes) + CFLAGS += -O3 +endif + +ifeq ($(DEVICE),gpu) + CFLAGS +=-fiopenmp -fopenmp-targets=spir64 -D__STRICT_ANSI__ +else + CFLAGS +=-qopenmp +endif +#=============================================================================== +# Targets to Build +#=============================================================================== + +$(program): $(obj) Makefile + $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) + +main.o: ../sundials-cuda/main.cpp ../sundials-cuda/sundials.h Makefile + $(CC) $(CFLAGS) -c $< -o $@ + +utils.o: ../sundials-cuda/utils.cpp ../sundials-cuda/sundials.h Makefile + $(CC) $(CFLAGS) -c $< -o $@ + +%.o: %.cpp ../sundials-cuda/sundials.h Makefile + $(CC) $(CFLAGS) -c $< -o $@ + +clean: + rm -rf $(program) $(obj) + +run: $(program) + $(LAUNCHER) ./$(program) 32 20000 6 1000 3 diff --git a/src/sundials-omp/Makefile.aomp b/src/sundials-omp/Makefile.aomp new file mode 100644 index 0000000000..451b47e83c --- /dev/null +++ b/src/sundials-omp/Makefile.aomp @@ -0,0 +1,72 @@ +#=============================================================================== +# User Options +#=============================================================================== + +# Compiler can be set below, or via environment variable +CC = clang++ +OPTIMIZE = yes +DEBUG = no +DEVICE = gpu +ARCH = gfx942 +LAUNCHER = + +#=============================================================================== +# Program name & source code list +#=============================================================================== + +program = main + +source = main.cpp kernels.cpp utils.cpp + +obj = $(source:.cpp=.o) + +#=============================================================================== +# Sets Flags +#=============================================================================== + +# Standard Flags +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall -I../sundials-cuda + +# Linker Flags +LDFLAGS = + +# Debug Flags +ifeq ($(DEBUG),yes) + CFLAGS += -g + LDFLAGS += -g +endif + +# Optimization Flags +ifeq ($(OPTIMIZE),yes) + CFLAGS += -O3 +endif + +ifeq ($(DEVICE),gpu) + CFLAGS += -target x86_64-pc-linux-gnu \ + -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa \ + -Xopenmp-target=amdgcn-amd-amdhsa \ + -march=$(ARCH) +else + CFLAGS +=-fopenmp +endif +#=============================================================================== +# Targets to Build +#=============================================================================== + +$(program): $(obj) Makefile.aomp + $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) + +main.o: ../sundials-cuda/main.cpp ../sundials-cuda/sundials.h Makefile.aomp + $(CC) $(CFLAGS) -c $< -o $@ + +utils.o: ../sundials-cuda/utils.cpp ../sundials-cuda/sundials.h Makefile.aomp + $(CC) $(CFLAGS) -c $< -o $@ + +%.o: %.cpp ../sundials-cuda/sundials.h Makefile.aomp + $(CC) $(CFLAGS) -c $< -o $@ + +clean: + rm -rf $(program) $(obj) + +run: $(program) + $(LAUNCHER) ./$(program) 32 20000 6 1000 3 diff --git a/src/sundials-omp/Makefile.nvc b/src/sundials-omp/Makefile.nvc new file mode 100644 index 0000000000..e3b0cb6285 --- /dev/null +++ b/src/sundials-omp/Makefile.nvc @@ -0,0 +1,69 @@ +#=============================================================================== +# User Options +#=============================================================================== + +# Compiler can be set below, or via environment variable +CC = nvc++ +OPTIMIZE = yes +DEBUG = no +DEVICE = gpu +SM = cc70 +LAUNCHER = + +#=============================================================================== +# Program name & source code list +#=============================================================================== + +program = main + +source = main.cpp kernels.cpp utils.cpp + +obj = $(source:.cpp=.o) + +#=============================================================================== +# Sets Flags +#=============================================================================== + +# Standard Flags +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall -I../sundials-cuda + +# Linker Flags +LDFLAGS = + +# Debug Flags +ifeq ($(DEBUG),yes) + CFLAGS += -g + LDFLAGS += -g +endif + +# Optimization Flags +ifeq ($(OPTIMIZE),yes) + CFLAGS += -O3 +endif + +ifeq ($(DEVICE),gpu) + CFLAGS +=-Minfo -mp=gpu -gpu=$(SM)#,fastmath +else + CFLAGS += +endif +#=============================================================================== +# Targets to Build +#=============================================================================== + +$(program): $(obj) Makefile.nvc + $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) + +main.o: ../sundials-cuda/main.cpp ../sundials-cuda/sundials.h Makefile.nvc + $(CC) $(CFLAGS) -c $< -o $@ + +utils.o: ../sundials-cuda/utils.cpp ../sundials-cuda/sundials.h Makefile.nvc + $(CC) $(CFLAGS) -c $< -o $@ + +%.o: %.cpp ../sundials-cuda/sundials.h Makefile.nvc + $(CC) $(CFLAGS) -c $< -o $@ + +clean: + rm -rf $(program) $(obj) + +run: $(program) + $(LAUNCHER) ./$(program) 32 20000 6 1000 3 diff --git a/src/sundials-omp/kernels.cpp b/src/sundials-omp/kernels.cpp new file mode 100644 index 0000000000..80157402d9 --- /dev/null +++ b/src/sundials-omp/kernels.cpp @@ -0,0 +1,158 @@ +#include +#include +#include +#include +#include +#include +#include "sundials.h" + +/* + * ============================================================================= + * OpenMP target-offload port of the SUNDIALS implicit-integration miniapp. + * ============================================================================= + */ + +using clk = std::chrono::steady_clock; +static inline long ns_since(clk::time_point t0) { + return std::chrono::duration_cast(clk::now() - t0) + .count(); +} + +long sundials_miniapp(int steps, int newton, int /*bs*/, const sun_in *in, + sun_out *out) { + const IDXT m = in->m, nblocks = in->nblocks, blocknnz = in->blocknnz; + const size_t n = (size_t)m * nblocks; + const size_t nnz_total = (size_t)nblocks * blocknnz; + + const real reltol = in->reltol, Sabstol = in->Sabstol, rl1 = in->rl1, + ngamma = in->ngamma, h = in->h, r = in->r, fract = in->fract, + uround = in->uround, c = -in->gamma, damp = in->damp; + + // Inputs (resident on the device for the whole loop). + const IDXT *rowptr = in->rowptr, *colind = in->colind; + const REAL *Jvals = in->Jvals, *Vabstol = in->Vabstol, *cflag = in->cflag, + *mm = in->mm, *zn1 = in->zn1, *ycor = in->ycor, *fpred = in->fpred, + *ypred = in->ypred, *Min = in->Min; + + // Outputs (caller-allocated host buffers); copied back at region exit. + real *ewt = out->ewt, *ewtsv = out->ewtsv, *tempv = out->tempv, + *ftemp = out->ftemp, *Mdiag = out->Mdiag, *Mupd = out->Mupd, + *res = out->res, *Mbcsr = out->Mbcsr, *Ax = out->Ax, *ycur = out->ycur; + + // Device-only scratch (written but not verified on the host). + real *save = (real *)malloc(n * sizeof(real)); + real *yform = (real *)malloc(n * sizeof(real)); + real *ybuild = (real *)malloc(n * sizeof(real)); + + // ycur is the evolving state; seed it from the initial condition. + for (size_t i = 0; i < n; i++) ycur[i] = in->ycur0[i]; + + double dot = 0, wl2 = 0, mx = 0, l1 = 0, mn = 0; + long ns = 0; + + #pragma omp target data \ + map(to: rowptr[0:m + 1], colind[0:blocknnz], Jvals[0:nnz_total], \ + Vabstol[0:n], cflag[0:n], mm[0:n], zn1[0:n], ycor[0:n], \ + fpred[0:n], ypred[0:n], Min[0:n]) \ + map(tofrom: ycur[0:n]) \ + map(from: ewt[0:n], ewtsv[0:n], tempv[0:n], ftemp[0:n], Mdiag[0:n], \ + Mupd[0:n], res[0:n], Mbcsr[0:nnz_total], Ax[0:n]) \ + map(alloc: save[0:n], yform[0:n], ybuild[0:n]) + { + clk::time_point start = clk::now(); + + for (int s = 0; s < steps; s++) { + // ---- fused setup: ewtSS/SV + checkConstraints + CVDiag form/build/update + #pragma omp target teams distribute parallel for + for (size_t i = 0; i < n; i++) { + real yc = ycur[i]; + real w = 1.0 / (reltol * fabs(yc) + Sabstol); + ewt[i] = w; + ewtsv[i] = 1.0 / (reltol * fabs(yc) + Vabstol[i]); + + real tmp = (fabs(cflag[i]) >= 1.5) ? 1.0 : 0.0; + tmp = tmp * cflag[i] / w; + save[i] = -0.1 * tmp; + tempv[i] = (yc - 0.1 * tmp) * mm[i]; + + real ft = h * fpred[i] - zn1[i]; + ftemp[i] = ft; + yform[i] = r * ft + ypred[i]; + + real mval = fract * ft - h * (Min[i] - fpred[i]); + real yv = ft * w; + bool test = (fabs(yv) >= uround); + real bit = test ? 1.0 : 0.0; + real bitcomp = test ? 0.0 : -1.0; + yv = fract * ft * bit - bitcomp; + real md = mval / yv * bit - bitcomp; + Mdiag[i] = md; + ybuild[i] = yv; + + Mupd[i] = r * (1.0 / md - 1.0) + 1.0; + } + + // ---- scaleAddI: form Newton matrix M = c*J + I (BCSR) ---- + #pragma omp target teams distribute parallel for collapse(2) + for (IDXT block = 0; block < nblocks; block++) { + for (IDXT row = 0; row < m; row++) { + IDXT tmp = rowptr[row]; + IDXT rownnz = rowptr[row + 1] - tmp; + IDXT idxg = block * blocknnz + tmp; + for (IDXT j = 0; j < rownnz; j++) { + REAL v = c * Jvals[idxg + j]; + if (colind[tmp + j] == row) v += 1.0; + Mbcsr[idxg + j] = v; + } + } + } + + // ---- Newton iterations: fused matvec (M*ycur) + residual ---- + for (int it = 0; it < newton; it++) { + #pragma omp target teams distribute parallel for collapse(2) + for (IDXT block = 0; block < nblocks; block++) { + for (IDXT row = 0; row < m; row++) { + IDXT tmp = rowptr[row]; + IDXT rownnz = rowptr[row + 1] - tmp; + IDXT idxg = block * blocknnz + tmp; + IDXT rowg = block * m + row; + IDXT colg = block * m; + REAL sum = 0; + for (IDXT j = 0; j < rownnz; j++) + sum += Mbcsr[idxg + j] * ycur[colg + colind[tmp + j]]; + Ax[rowg] = sum; + res[rowg] = ngamma * sum + (rl1 * zn1[rowg] + ycor[rowg]); + } + } + } + + // ---- fused reduce + advance: 5 res reductions + state advance + renorm. + double sdot = 0, swl2 = 0, sl1 = 0; + double smx = -DBL_MAX, smn = DBL_MAX, symx = -DBL_MAX; + #pragma omp target teams distribute parallel for \ + reduction(+: sdot, swl2, sl1) reduction(max: smx, symx) \ + reduction(min: smn) + for (size_t i = 0; i < n; i++) { + real rv = res[i], p = rv * ewt[i], a = fabs(rv); + sdot += p; swl2 += p * p; sl1 += a; + smx = fmax(a, smx); smn = fmin(rv, smn); + real yn = ycur[i] - damp * rv; + ycur[i] = yn; + symx = fmax(fabs(yn), symx); + } + + // Renormalize the advanced state (power-iteration style). + real inv = 1.0 / symx; + #pragma omp target teams distribute parallel for + for (size_t i = 0; i < n; i++) ycur[i] = ycur[i] * inv; + + dot = sdot; wl2 = swl2; mx = smx; l1 = sl1; mn = smn; + } + + ns = ns_since(start); + } + + out->dot = dot; out->wl2 = wl2; out->mx = mx; out->l1 = l1; out->mn = mn; + free(save); free(yform); free(ybuild); + return ns; +} From 40f73db2474b001b1873d9d8abf99020383e9864 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Thu, 6 Aug 2026 06:13:35 -0700 Subject: [PATCH 2/7] [sundials-cuda] remove the CUDA in the printed message --- src/sundials-cuda/main.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/sundials-cuda/main.cpp b/src/sundials-cuda/main.cpp index b73958d83d..1d8b23e2bb 100644 --- a/src/sundials-cuda/main.cpp +++ b/src/sundials-cuda/main.cpp @@ -77,7 +77,7 @@ int main(int argc, char *argv[]) { in.h = 1e-3; in.r = 0.75; in.fract = 0.1; in.uround = 1e-13; in.gamma = 0.1; in.damp = 0.1; - printf("SUNDIALS implicit-integration miniapp (CUDA)\n"); + printf("SUNDIALS implicit-integration miniapp\n"); printf("Batch: %d systems of size %d (vector length n = %zu)\n", nblocks, m, n); printf("Block-CSR: %d nonzeros/block (%d per row), %zu total nonzeros\n", From 66364a07a9395e2c6c7e1bdbde55563dfe611c84 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Thu, 6 Aug 2026 07:40:30 -0700 Subject: [PATCH 3/7] [axpby-omp] add the example --- src/axpby-omp/CMakeLists.txt | 8 ++ src/axpby-omp/Makefile | 62 ++++++++++++++ src/axpby-omp/Makefile.aomp | 66 +++++++++++++++ src/axpby-omp/Makefile.nvc | 63 ++++++++++++++ src/axpby-omp/main.cpp | 158 +++++++++++++++++++++++++++++++++++ 5 files changed, 357 insertions(+) create mode 100644 src/axpby-omp/CMakeLists.txt create mode 100644 src/axpby-omp/Makefile create mode 100644 src/axpby-omp/Makefile.aomp create mode 100644 src/axpby-omp/Makefile.nvc create mode 100644 src/axpby-omp/main.cpp diff --git a/src/axpby-omp/CMakeLists.txt b/src/axpby-omp/CMakeLists.txt new file mode 100644 index 0000000000..a4b40632eb --- /dev/null +++ b/src/axpby-omp/CMakeLists.txt @@ -0,0 +1,8 @@ +# axpby-omp/CMakeLists.txt + +add_hecbench_benchmark( + NAME axpby + MODEL omp + SOURCES main.cpp + CATEGORIES algorithms +) diff --git a/src/axpby-omp/Makefile b/src/axpby-omp/Makefile new file mode 100644 index 0000000000..a05de4fe20 --- /dev/null +++ b/src/axpby-omp/Makefile @@ -0,0 +1,62 @@ +#=============================================================================== +# User Options +#=============================================================================== + +# Compiler can be set below, or via environment variable +CC = icpx +OPTIMIZE = yes +DEBUG = no +DEVICE = gpu +LAUNCHER = + +#=============================================================================== +# Program name & source code list +#=============================================================================== + +program = main + +source = main.cpp + +obj = $(source:.cpp=.o) + +#=============================================================================== +# Sets Flags +#=============================================================================== + +# Standard Flags +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall + +# Linker Flags +LDFLAGS = + +# Debug Flags +ifeq ($(DEBUG),yes) + CFLAGS += -g + LDFLAGS += -g +endif + +# Optimization Flags +ifeq ($(OPTIMIZE),yes) + CFLAGS += -O3 +endif + +ifeq ($(DEVICE),gpu) + CFLAGS +=-fiopenmp -fopenmp-targets=spir64 -D__STRICT_ANSI__ +else + CFLAGS +=-qopenmp +endif +#=============================================================================== +# Targets to Build +#=============================================================================== + +$(program): $(obj) Makefile + $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) + +%.o: %.cpp Makefile + $(CC) $(CFLAGS) -c $< -o $@ + +clean: + rm -rf $(program) $(obj) + +run: $(program) + $(LAUNCHER) ./$(program) 1000 diff --git a/src/axpby-omp/Makefile.aomp b/src/axpby-omp/Makefile.aomp new file mode 100644 index 0000000000..3e045c79b2 --- /dev/null +++ b/src/axpby-omp/Makefile.aomp @@ -0,0 +1,66 @@ +#=============================================================================== +# User Options +#=============================================================================== + +# Compiler can be set below, or via environment variable +CC = clang++ +OPTIMIZE = yes +DEBUG = no +DEVICE = gpu +ARCH = gfx906 +LAUNCHER = + +#=============================================================================== +# Program name & source code list +#=============================================================================== + +program = main + +source = main.cpp + +obj = $(source:.cpp=.o) + +#=============================================================================== +# Sets Flags +#=============================================================================== + +# Standard Flags +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall + +# Linker Flags +LDFLAGS = + +# Debug Flags +ifeq ($(DEBUG),yes) + CFLAGS += -g + LDFLAGS += -g +endif + +# Optimization Flags +ifeq ($(OPTIMIZE),yes) + CFLAGS += -O3 +endif + +ifeq ($(DEVICE),gpu) + CFLAGS += -target x86_64-pc-linux-gnu \ + -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa \ + -Xopenmp-target=amdgcn-amd-amdhsa \ + -march=$(ARCH) +else + CFLAGS +=-fopenmp +endif +#=============================================================================== +# Targets to Build +#=============================================================================== + +$(program): $(obj) Makefile.aomp + $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) + +%.o: %.cpp Makefile.aomp + $(CC) $(CFLAGS) -c $< -o $@ + +clean: + rm -rf $(program) $(obj) + +run: $(program) + $(LAUNCHER) ./$(program) 1000 diff --git a/src/axpby-omp/Makefile.nvc b/src/axpby-omp/Makefile.nvc new file mode 100644 index 0000000000..9217f67c9d --- /dev/null +++ b/src/axpby-omp/Makefile.nvc @@ -0,0 +1,63 @@ +#=============================================================================== +# User Options +#=============================================================================== + +# Compiler can be set below, or via environment variable +CC = nvc++ +OPTIMIZE = yes +DEBUG = no +DEVICE = gpu +SM = cc70 +LAUNCHER = + +#=============================================================================== +# Program name & source code list +#=============================================================================== + +program = main + +source = main.cpp + +obj = $(source:.cpp=.o) + +#=============================================================================== +# Sets Flags +#=============================================================================== + +# Standard Flags +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall + +# Linker Flags +LDFLAGS = + +# Debug Flags +ifeq ($(DEBUG),yes) + CFLAGS += -g + LDFLAGS += -g +endif + +# Optimization Flags +ifeq ($(OPTIMIZE),yes) + CFLAGS += -O3 +endif + +ifeq ($(DEVICE),gpu) + CFLAGS +=-Minfo -mp=gpu -gpu=$(SM)#,fastmath +else + CFLAGS += +endif +#=============================================================================== +# Targets to Build +#=============================================================================== + +$(program): $(obj) Makefile.nvc + $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) + +%.o: %.cpp Makefile.nvc + $(CC) $(CFLAGS) -c $< -o $@ + +clean: + rm -rf $(program) $(obj) + +run: $(program) + $(LAUNCHER) ./$(program) 1000 diff --git a/src/axpby-omp/main.cpp b/src/axpby-omp/main.cpp new file mode 100644 index 0000000000..e6d29f42ae --- /dev/null +++ b/src/axpby-omp/main.cpp @@ -0,0 +1,158 @@ +#include +#include +#include +#include +#include +#include +#include +#include + +// Multi-tensor AXPBY: out = a * x + b * y (elementwise, per tensor) +// The CUDA versions use a chunked "multi_tensor_apply" launcher: the work +// is split into fixed-size chunks and each thread block processes one chunk of +// one tensor, so a single kernel launch handles many (small) tensors at once. +// This OpenMP offload version mirrors that design: a flat list of +// (tensor, chunk) "blocks" is built on the host and processed in a single +// target region where each team handles one chunk. + +template +struct Tensor { + T* data_ptr; + int64_t numel; +}; + +template +void multi_tensor_axpby(int chunk_size, + std::vector>> &tensor_lists, + float a, float b) { + const int max_tensors = tensor_lists[0].size(); + const int dev = omp_get_default_device(); + + auto start = std::chrono::steady_clock::now(); + + // Build the flat chunk metadata (mirrors the CUDA multi_tensor_apply + // launcher, where each block maps to one chunk of one tensor). Store device + // addresses so the single target region can index straight into device + // memory without any per-tensor mapping. + std::vector chunk_x, chunk_y, chunk_out; + std::vector chunk_offset, chunk_len; + for (int n = 0; n < max_tensors; n++) { + scalar_t *x = (scalar_t*) omp_get_mapped_ptr(tensor_lists[0][n].data_ptr, dev); + scalar_t *y = (scalar_t*) omp_get_mapped_ptr(tensor_lists[1][n].data_ptr, dev); + scalar_t *out = (scalar_t*) omp_get_mapped_ptr(tensor_lists[2][n].data_ptr, dev); + const int64_t len = tensor_lists[0][n].numel; + const int64_t chunks_this_tensor = (len + chunk_size - 1) / chunk_size; + for (int64_t c = 0; c < chunks_this_tensor; c++) { + const int64_t offset = c * (int64_t)chunk_size; + chunk_x.push_back(x); + chunk_y.push_back(y); + chunk_out.push_back(out); + chunk_offset.push_back(offset); + chunk_len.push_back(std::min(chunk_size, len - offset)); + } + } + + const int nblocks = (int) chunk_x.size(); + scalar_t **xp = chunk_x.data(); + scalar_t **yp = chunk_y.data(); + scalar_t **op = chunk_out.data(); + int64_t *coff = chunk_offset.data(); + int64_t *clen = chunk_len.data(); + + // One team per chunk ("block"); threads within a team stride over the chunk. + #pragma omp target teams distribute \ + map(to: xp[0:nblocks], yp[0:nblocks], op[0:nblocks], \ + coff[0:nblocks], clen[0:nblocks]) + for (int blk = 0; blk < nblocks; blk++) { + scalar_t *x = xp[blk]; + scalar_t *y = yp[blk]; + scalar_t *out = op[blk]; + const int64_t off = coff[blk]; + const int64_t n = clen[blk]; + #pragma omp parallel for + for (int64_t i = 0; i < n; i++) { + out[off + i] = a * (float)x[off + i] + b * (float)y[off + i]; + } + } + + auto end = std::chrono::steady_clock::now(); + auto time = std::chrono::duration_cast(end - start).count(); + printf("Chunk size %8d | Total execution time of multi_tensor_axpby: %f (us)\n", + chunk_size, (time * 1e-3f)); +} + +int main(int argc, char* argv[]) +{ + if (argc != 2) { + printf("Usage: %s \n", argv[0]); + return 1; + } + const int max_tensors = atoi(argv[1]); + + std::vector>> tensor_lists_ref (3); + std::vector>> tensor_lists (3); + srand(123); + for (int n = 0; n < max_tensors; n++) { + int64_t length = rand() % (1024*1024) + 1024; + for (int d = 0; d < 3; d++) { + float *tensor = (float*) malloc (sizeof(float) * length); + float *tensor_ref = (float*) malloc (sizeof(float) * length); + if (d <= 1) { + for (int64_t i = 0; i < length; i++) + tensor[i] = tensor_ref[i] = rand() % length; + } + #pragma omp target enter data map(to: tensor[0:length]) + + Tensor t; + t.data_ptr = tensor; + t.numel = length; + tensor_lists[d].push_back(t); + + Tensor t_ref; + t_ref.data_ptr = tensor_ref; + t_ref.numel = length; + tensor_lists_ref[d].push_back(t_ref); + } + } + + const float a = 1.f; + const float b = 1.f; + + for (int chunk_size = 256; chunk_size <= 1024*1024; chunk_size = chunk_size * 2) { + + multi_tensor_axpby(chunk_size, tensor_lists, a, b); + + bool ok = true; + for (int n = 0; n < max_tensors; n++) { + auto x = tensor_lists_ref[0][n]; + auto y = tensor_lists_ref[1][n]; + auto z = tensor_lists_ref[2][n]; + for (int i = 0; i < x.numel; i++) { + z.data_ptr[i] = a * x.data_ptr[i] + b * y.data_ptr[i]; + } + float *zp = tensor_lists[2][n].data_ptr; + int64_t len = tensor_lists[2][n].numel; + #pragma omp target update from(zp[0:len]) + for (int i = 0; i < x.numel; i++) { + if (fabsf(zp[i] - z.data_ptr[i]) > 1e-3f) { + ok = false; + break; + } + } + if (!ok) break; + } + printf("%s\n", ok ? "PASS" : "FAIL"); + } + + for (int d = 0; d < 3; d++) { + for (int t = 0; t < max_tensors; t++) { + float *tensor = tensor_lists[d][t].data_ptr; + int64_t len = tensor_lists[d][t].numel; + #pragma omp target exit data map(delete: tensor[0:len]) + free(tensor_lists[d][t].data_ptr); + free(tensor_lists_ref[d][t].data_ptr); + } + } + + return 0; +} From aaec7ddfd005d261d51f94f1565491e05e89339a Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Thu, 6 Aug 2026 09:06:02 -0700 Subject: [PATCH 4/7] [dwconv-omp] add the example --- src/dwconv-omp/CMakeLists.txt | 8 ++ src/dwconv-omp/Makefile | 62 +++++++++ src/dwconv-omp/Makefile.aomp | 66 ++++++++++ src/dwconv-omp/Makefile.nvc | 63 ++++++++++ src/dwconv-omp/main.cpp | 228 ++++++++++++++++++++++++++++++++++ 5 files changed, 427 insertions(+) create mode 100644 src/dwconv-omp/CMakeLists.txt create mode 100644 src/dwconv-omp/Makefile create mode 100644 src/dwconv-omp/Makefile.aomp create mode 100644 src/dwconv-omp/Makefile.nvc create mode 100644 src/dwconv-omp/main.cpp diff --git a/src/dwconv-omp/CMakeLists.txt b/src/dwconv-omp/CMakeLists.txt new file mode 100644 index 0000000000..37333cdf7a --- /dev/null +++ b/src/dwconv-omp/CMakeLists.txt @@ -0,0 +1,8 @@ +# dwconv-omp/CMakeLists.txt + +add_hecbench_benchmark( + NAME dwconv + MODEL omp + SOURCES main.cpp + CATEGORIES algorithms +) diff --git a/src/dwconv-omp/Makefile b/src/dwconv-omp/Makefile new file mode 100644 index 0000000000..aac4a69113 --- /dev/null +++ b/src/dwconv-omp/Makefile @@ -0,0 +1,62 @@ +#=============================================================================== +# User Options +#=============================================================================== + +# Compiler can be set below, or via environment variable +CC = icpx +OPTIMIZE = yes +DEBUG = no +DEVICE = gpu +LAUNCHER = + +#=============================================================================== +# Program name & source code list +#=============================================================================== + +program = main + +source = main.cpp + +obj = $(source:.cpp=.o) + +#=============================================================================== +# Sets Flags +#=============================================================================== + +# Standard Flags +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall + +# Linker Flags +LDFLAGS = + +# Debug Flags +ifeq ($(DEBUG),yes) + CFLAGS += -g + LDFLAGS += -g +endif + +# Optimization Flags +ifeq ($(OPTIMIZE),yes) + CFLAGS += -O3 +endif + +ifeq ($(DEVICE),gpu) + CFLAGS +=-fiopenmp -fopenmp-targets=spir64 -D__STRICT_ANSI__ +else + CFLAGS +=-qopenmp +endif +#=============================================================================== +# Targets to Build +#=============================================================================== + +$(program): $(obj) Makefile + $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) + +%.o: %.cpp Makefile + $(CC) $(CFLAGS) -c $< -o $@ + +clean: + rm -rf $(program) $(obj) + +run: $(program) + $(LAUNCHER) ./$(program) 128 16 224 224 100 diff --git a/src/dwconv-omp/Makefile.aomp b/src/dwconv-omp/Makefile.aomp new file mode 100644 index 0000000000..51570d320b --- /dev/null +++ b/src/dwconv-omp/Makefile.aomp @@ -0,0 +1,66 @@ +#=============================================================================== +# User Options +#=============================================================================== + +# Compiler can be set below, or via environment variable +CC = clang++ +OPTIMIZE = yes +DEBUG = no +DEVICE = gpu +ARCH = gfx906 +LAUNCHER = + +#=============================================================================== +# Program name & source code list +#=============================================================================== + +program = main + +source = main.cpp + +obj = $(source:.cpp=.o) + +#=============================================================================== +# Sets Flags +#=============================================================================== + +# Standard Flags +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall + +# Linker Flags +LDFLAGS = + +# Debug Flags +ifeq ($(DEBUG),yes) + CFLAGS += -g + LDFLAGS += -g +endif + +# Optimization Flags +ifeq ($(OPTIMIZE),yes) + CFLAGS += -O3 +endif + +ifeq ($(DEVICE),gpu) + CFLAGS += -target x86_64-pc-linux-gnu \ + -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa \ + -Xopenmp-target=amdgcn-amd-amdhsa \ + -march=$(ARCH) +else + CFLAGS +=-fopenmp +endif +#=============================================================================== +# Targets to Build +#=============================================================================== + +$(program): $(obj) Makefile.aomp + $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) + +%.o: %.cpp Makefile.aomp + $(CC) $(CFLAGS) -c $< -o $@ + +clean: + rm -rf $(program) $(obj) + +run: $(program) + $(LAUNCHER) ./$(program) 128 16 224 224 100 diff --git a/src/dwconv-omp/Makefile.nvc b/src/dwconv-omp/Makefile.nvc new file mode 100644 index 0000000000..76eb389f94 --- /dev/null +++ b/src/dwconv-omp/Makefile.nvc @@ -0,0 +1,63 @@ +#=============================================================================== +# User Options +#=============================================================================== + +# Compiler can be set below, or via environment variable +CC = nvc++ +OPTIMIZE = yes +DEBUG = no +DEVICE = gpu +SM = cc70 +LAUNCHER = + +#=============================================================================== +# Program name & source code list +#=============================================================================== + +program = main + +source = main.cpp + +obj = $(source:.cpp=.o) + +#=============================================================================== +# Sets Flags +#=============================================================================== + +# Standard Flags +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall + +# Linker Flags +LDFLAGS = + +# Debug Flags +ifeq ($(DEBUG),yes) + CFLAGS += -g + LDFLAGS += -g +endif + +# Optimization Flags +ifeq ($(OPTIMIZE),yes) + CFLAGS += -O3 +endif + +ifeq ($(DEVICE),gpu) + CFLAGS +=-Minfo -mp=gpu -gpu=$(SM)#,fastmath +else + CFLAGS += +endif +#=============================================================================== +# Targets to Build +#=============================================================================== + +$(program): $(obj) Makefile.nvc + $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) + +%.o: %.cpp Makefile.nvc + $(CC) $(CFLAGS) -c $< -o $@ + +clean: + rm -rf $(program) $(obj) + +run: $(program) + $(LAUNCHER) ./$(program) 128 16 224 224 100 diff --git a/src/dwconv-omp/main.cpp b/src/dwconv-omp/main.cpp new file mode 100644 index 0000000000..8246099d70 --- /dev/null +++ b/src/dwconv-omp/main.cpp @@ -0,0 +1,228 @@ +#include +#include +#include +#include + +// Depthwise 2D convolution forward. + +// kSize is the kernel extent when it is known at compile time, and zero when +// the extent is only known at run time. +template +void conv_depthwise2d_forward_kernel( + scalar_t *__restrict__ input, + scalar_t *__restrict__ output, + scalar_t *__restrict__ weight, + scalar_t *__restrict__ bias, + bool biasEnabled, + int totalElements, + const int outputChannels, + const int depthwiseMultiplier, + const int inputWidth, const int inputHeight, + const int outputWidth, const int outputHeight, + const int kernelWidth, const int kernelHeight, + const int strideWidth, const int strideHeight, + const int padWidth, const int padHeight, + const int dilationWidth, const int dilationHeight) +{ + const int KW_LIMIT = (kSize != 0) ? kSize : kernelWidth; + const int KH_LIMIT = (kSize != 0) ? kSize : kernelHeight; + + #pragma omp target teams distribute parallel for + for (int linearIndex = 0; linearIndex < totalElements; linearIndex++) { + //calculate n,c,h,w indices, replacing modulos by divide and multiply add, + //result is same as would be in the code below + //const int n = linearIndex / batchStride; //batchStride = outputChannels * outputHeight * outputWidth + //const int c = (linearIndex / channelStride) % outputChannels; //channelStride = outputHeight * outputWidth + //const int h = (linearIndex / outputWidth) % outputHeight; + //const int w = linearIndex % outputWidth; + + int indtmp1 = linearIndex/outputWidth; + const int w = linearIndex - indtmp1 * outputWidth; + int indtmp2 = indtmp1/outputHeight; + const int h = indtmp1 - indtmp2 * outputHeight; + indtmp1 = indtmp2; + indtmp2 = indtmp1/outputChannels; + const int c = indtmp1 - indtmp2 * outputChannels; + const int n = indtmp2; + + int inputChannel = c; + int inputChannels = outputChannels; + if (depthwiseMultiplier != 1) { + inputChannel /= depthwiseMultiplier; + inputChannels /= depthwiseMultiplier; + } + + int weightOffset = c * kernelHeight * kernelWidth; + + scalar_t value = biasEnabled ? bias[c] : (scalar_t)0; + const int offset0 = (n * inputChannels + inputChannel) * inputHeight * inputWidth; + #pragma unroll + for (int kH = 0; kH < KH_LIMIT; ++kH) { + #pragma unroll + for (int kW = 0; kW < KW_LIMIT; ++kW) { + const int h_in = -padHeight + h * strideHeight + kH * dilationHeight; + const int w_in = -padWidth + w * strideWidth + kW * dilationWidth; + + if ((h_in >= 0) && (h_in < inputHeight) && (w_in >= 0) && (w_in < inputWidth)) { + const int offset = offset0 + h_in * inputWidth + w_in; + value += weight[weightOffset] * input[offset]; + } + ++weightOffset; + } + } + output[linearIndex] = value; + } +} + +template +void dwconv2d_forward (const int m, + const int n, + const int input_channels, + const int H, + const int W, + const int kH, + const int kW, + const int repeat, + const int padH = 1, + const int padW = 1, + const int strideH = 1, + const int strideW = 1, + const int dilateH = 1, + const int dilateW = 1) +{ + // The number of output_channels is a multiple of input_channels + const int output_channels = input_channels * m; + + // Weight Tensor is shape (output_channels, 1, kH, kW) + int weight_size = output_channels * kH * kW; + size_t weight_size_bytes = weight_size * sizeof(scalar_t); + + // Input Tensor is shape (N, input_channels, H, W) + int input_size = n * input_channels * H * W; + size_t input_size_bytes = input_size * sizeof(scalar_t); + int input_sizes[] = {n, input_channels, H, W}; + + // Bias has same number of channels as output + int bias_size = output_channels; + size_t bias_size_bytes = bias_size * sizeof(scalar_t); + + int pad_sizes[] = {padH, padW}; + int stride_sizes[] = {strideH, strideW}; + int dilation_sizes[] = {dilateH, dilateW}; + + int weight_sizes[] = {output_channels, 1, kH, kW}; + + int height = input_sizes[2]; // H + int width = input_sizes[3]; // W + int outputChannels = weight_sizes[0]; // output_channels + + int output_sizes[4]; + output_sizes[0] = input_sizes[0]; + output_sizes[1] = weight_sizes[0]; + + for (int d = 2; d < 4; d++) { + auto dilation_ = dilation_sizes[d - 2]; + auto kernel = dilation_ * (weight_sizes[d] - 1) + 1; + output_sizes[d] = (input_sizes[d] + (2 * pad_sizes[d - 2]) - kernel) / stride_sizes[d - 2] + 1; + } + + int outputHeight = output_sizes[2]; + int outputWidth = output_sizes[3]; + + int inputChannels = input_sizes[1]; + int depthwiseMultiplier = outputChannels / inputChannels; + + int output_size = 1; + for (int i = 0; i < 4; i++) + output_size *= output_sizes[i]; + + size_t output_size_bytes = output_size * sizeof(scalar_t); + + scalar_t *h_input = (scalar_t*) malloc (input_size_bytes); + scalar_t *h_weight = (scalar_t*) malloc (weight_size_bytes); + scalar_t *h_bias = (scalar_t*) malloc (bias_size_bytes); + scalar_t *h_output = (scalar_t*) malloc (output_size_bytes); + + srand(123); + for (int i = 0; i < input_size; i++) { + h_input[i] = rand() / (float)RAND_MAX; + } + for (int i = 0; i < weight_size; i++) { + h_weight[i] = rand() / (float)RAND_MAX; + } + for (int i = 0; i < bias_size; i++) { + h_bias[i] = rand() / (float)RAND_MAX; + } + + const bool has_bias = true; + + #pragma omp target data map(to: h_input[0:input_size], \ + h_weight[0:weight_size], \ + h_bias[0:bias_size]) \ + map(from: h_output[0:output_size]) + { + auto start = std::chrono::steady_clock::now(); + + for (int i = 0; i < repeat; i++) { + if (kW == 3 && kH == 3) { + conv_depthwise2d_forward_kernel<3, scalar_t>( + h_input, h_output, h_weight, h_bias, has_bias, output_size, outputChannels, depthwiseMultiplier, + width, height, outputWidth, outputHeight, + kW, kH, strideW, strideH, padW, padH, dilateW, dilateH); + } else if (kW == 1 && kH == 1) { + conv_depthwise2d_forward_kernel<1, scalar_t>( + h_input, h_output, h_weight, h_bias, has_bias, output_size, outputChannels, depthwiseMultiplier, + width, height, outputWidth, outputHeight, + kW, kH, strideW, strideH, padW, padH, dilateW, dilateH); + } else { + conv_depthwise2d_forward_kernel<0, scalar_t>( + h_input, h_output, h_weight, h_bias, has_bias, output_size, outputChannels, depthwiseMultiplier, + width, height, outputWidth, outputHeight, + kW, kH, strideW, strideH, padW, padH, dilateW, dilateH); + } + } + + auto end = std::chrono::steady_clock::now(); + auto time = std::chrono::duration_cast(end - start).count(); + printf("Average execution time of dwconv2d_forward kernel: %f (ms)\n", + time * 1e-6f / repeat); + } + + scalar_t sum = 0; + for (int i = 0; i < output_size; i++) { + sum += h_output[i]; + } + printf("Checksum = %f\n", sum / output_size); + + free(h_input); + free(h_output); + free(h_weight); + free(h_bias); +} + +int main(int argc, char* argv[]) +{ + if (argc != 6) { + printf("Usage: %s ", argv[0]); + printf(" \n"); + return 1; + } + + const int n = atoi(argv[1]); + const int c = atoi(argv[2]); + const int h = atoi(argv[3]); + const int w = atoi(argv[4]); + const int repeat = atoi(argv[5]); + + // sweep over depth multipliers and kernel sizes + for (int m = 1; m <= 4; m = m + 1) { + for (int k = 1; k <= 5; k = k + 2) { + printf("batch = %d, input channel = %d, height = %d, width = %d, ", + n, c, h, w); + printf("kernel size = %d, output channel = %d\n", k, m * c); + dwconv2d_forward(m, n, c, h, w, k, k, repeat); + } + } + + return 0; +} From 29d8bfc326e0fe90111be08a569fbd4da8cfa0ef Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Thu, 6 Aug 2026 09:20:25 -0700 Subject: [PATCH 5/7] [bincount-omp] add the example --- src/bincount-omp/CMakeLists.txt | 8 ++ src/bincount-omp/Makefile | 65 ++++++++++ src/bincount-omp/Makefile.aomp | 66 ++++++++++ src/bincount-omp/Makefile.nvc | 63 ++++++++++ src/bincount-omp/main.cpp | 216 ++++++++++++++++++++++++++++++++ 5 files changed, 418 insertions(+) create mode 100644 src/bincount-omp/CMakeLists.txt create mode 100644 src/bincount-omp/Makefile create mode 100644 src/bincount-omp/Makefile.aomp create mode 100644 src/bincount-omp/Makefile.nvc create mode 100644 src/bincount-omp/main.cpp diff --git a/src/bincount-omp/CMakeLists.txt b/src/bincount-omp/CMakeLists.txt new file mode 100644 index 0000000000..ef73144568 --- /dev/null +++ b/src/bincount-omp/CMakeLists.txt @@ -0,0 +1,8 @@ +# bincount-omp/CMakeLists.txt + +add_hecbench_benchmark( + NAME bincount + MODEL omp + SOURCES main.cpp + CATEGORIES algorithms +) diff --git a/src/bincount-omp/Makefile b/src/bincount-omp/Makefile new file mode 100644 index 0000000000..941ca68435 --- /dev/null +++ b/src/bincount-omp/Makefile @@ -0,0 +1,65 @@ +#=============================================================================== +# User Options +#=============================================================================== + +# Compiler can be set below, or via environment variable +CC = icpx +OPTIMIZE = yes +DEBUG = no +DEVICE = gpu +LAUNCHER = + +#=============================================================================== +# Program name & source code list +#=============================================================================== + +program = main + +source = main.cpp + +obj = $(source:.cpp=.o) + +#=============================================================================== +# Sets Flags +#=============================================================================== + +# Standard Flags +# The bin index is computed in float, so the host reference and the kernel must +# round identically for the exact comparison to hold. +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall -I../bincount-cuda -fp-model=precise + +# Linker Flags +LDFLAGS = + +# Debug Flags +ifeq ($(DEBUG),yes) + CFLAGS += -g + LDFLAGS += -g +endif + +# Optimization Flags +ifeq ($(OPTIMIZE),yes) + CFLAGS += -O3 +endif + +ifeq ($(DEVICE),gpu) + CFLAGS +=-fiopenmp -fopenmp-targets=spir64 -D__STRICT_ANSI__ -DHAVE_OMPX_DEVICE_INFO + LDFLAGS +=-Xopenmp-target-backend "-cl-fp32-correctly-rounded-divide-sqrt" +else + CFLAGS +=-qopenmp +endif +#=============================================================================== +# Targets to Build +#=============================================================================== + +$(program): $(obj) Makefile + $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) + +%.o: %.cpp ../bincount-cuda/reference.h Makefile + $(CC) $(CFLAGS) -c $< -o $@ + +clean: + rm -rf $(program) $(obj) + +run: $(program) + $(LAUNCHER) ./$(program) 10000000 100 diff --git a/src/bincount-omp/Makefile.aomp b/src/bincount-omp/Makefile.aomp new file mode 100644 index 0000000000..2ad07bd80d --- /dev/null +++ b/src/bincount-omp/Makefile.aomp @@ -0,0 +1,66 @@ +#=============================================================================== +# User Options +#=============================================================================== + +# Compiler can be set below, or via environment variable +CC = clang++ +OPTIMIZE = yes +DEBUG = no +DEVICE = gpu +ARCH = gfx906 +LAUNCHER = + +#=============================================================================== +# Program name & source code list +#=============================================================================== + +program = main + +source = main.cpp + +obj = $(source:.cpp=.o) + +#=============================================================================== +# Sets Flags +#=============================================================================== + +# Standard Flags +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall -I../bincount-cuda + +# Linker Flags +LDFLAGS = + +# Debug Flags +ifeq ($(DEBUG),yes) + CFLAGS += -g + LDFLAGS += -g +endif + +# Optimization Flags +ifeq ($(OPTIMIZE),yes) + CFLAGS += -O3 +endif + +ifeq ($(DEVICE),gpu) + CFLAGS += -target x86_64-pc-linux-gnu \ + -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa \ + -Xopenmp-target=amdgcn-amd-amdhsa \ + -march=$(ARCH) +else + CFLAGS +=-fopenmp +endif +#=============================================================================== +# Targets to Build +#=============================================================================== + +$(program): $(obj) Makefile.aomp + $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) + +%.o: %.cpp ../bincount-cuda/reference.h Makefile.aomp + $(CC) $(CFLAGS) -c $< -o $@ + +clean: + rm -rf $(program) $(obj) + +run: $(program) + $(LAUNCHER) ./$(program) 10000000 100 diff --git a/src/bincount-omp/Makefile.nvc b/src/bincount-omp/Makefile.nvc new file mode 100644 index 0000000000..6d5caf5979 --- /dev/null +++ b/src/bincount-omp/Makefile.nvc @@ -0,0 +1,63 @@ +#=============================================================================== +# User Options +#=============================================================================== + +# Compiler can be set below, or via environment variable +CC = nvc++ +OPTIMIZE = yes +DEBUG = no +DEVICE = gpu +SM = cc70 +LAUNCHER = + +#=============================================================================== +# Program name & source code list +#=============================================================================== + +program = main + +source = main.cpp + +obj = $(source:.cpp=.o) + +#=============================================================================== +# Sets Flags +#=============================================================================== + +# Standard Flags +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall -I../bincount-cuda + +# Linker Flags +LDFLAGS = + +# Debug Flags +ifeq ($(DEBUG),yes) + CFLAGS += -g + LDFLAGS += -g +endif + +# Optimization Flags +ifeq ($(OPTIMIZE),yes) + CFLAGS += -O3 +endif + +ifeq ($(DEVICE),gpu) + CFLAGS +=-Minfo -mp=gpu -gpu=$(SM)#,fastmath +else + CFLAGS += +endif +#=============================================================================== +# Targets to Build +#=============================================================================== + +$(program): $(obj) Makefile.nvc + $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) + +%.o: %.cpp ../bincount-cuda/reference.h Makefile.nvc + $(CC) $(CFLAGS) -c $< -o $@ + +clean: + rm -rf $(program) $(obj) + +run: $(program) + $(LAUNCHER) ./$(program) 10000000 100 diff --git a/src/bincount-omp/main.cpp b/src/bincount-omp/main.cpp new file mode 100644 index 0000000000..9eeb6c72ab --- /dev/null +++ b/src/bincount-omp/main.cpp @@ -0,0 +1,216 @@ +#include +#include +#include +#include +#include +#include +#include +#include "reference.h" + +#define threadsPerBlock 256 + +// Capacity of the per-team histogram. OpenMP has no way to size a team-local +// array at run time, so this is fixed when the kernel is built; the usable +// limit is the smaller of this and what the device actually offers. +// It is kept below a device's full 64 KB local memory so that the compiler's +// own per-kernel local-memory overhead still fits under the hardware limit. +#define sharedMemoryCapacity (63 * 1024) + +// Largest per-team local memory the device provides. OpenMP has no portable +// query for this; ompx_get_device_info is an Intel extension that also needs +// the offload runtime, so the Makefile only defines HAVE_OMPX_DEVICE_INFO for +// an offload build. Everything else falls back to the built-in capacity. +static int getDeviceLocalMemSize() +{ +#ifdef HAVE_OMPX_DEVICE_INFO + size_t localMemSize = 0, sizeRet = 0; + if (ompx_get_device_info(omp_get_default_device(), ompx_devinfo_local_mem_size, + sizeof(localMemSize), &localMemSize, &sizeRet) == 0) + return (int) localMemSize; +#endif + return sharedMemoryCapacity; +} + +#pragma omp declare target +template +static IndexType +getBin(input_t v, input_t minvalue, input_t maxvalue, IndexType nbins) +{ + IndexType bin = (v - minvalue) * nbins / (maxvalue - minvalue); + // while each bin is inclusive at the lower end and exclusive at the higher, + // i.e. [start, end) the last bin is inclusive at both, i.e. [start, end], in + // order to include maxvalue if exists therefore when bin == nbins, adjust bin + // to the last bin + if (bin == nbins) bin--; + return bin; +} +#pragma omp end declare target + +/* + Calculate the frequency of the input values. + The GPU offloaded kernel atomically updates the global histogram tensor. +*/ +template +void eval(IndexType input_size, int repeat) +{ + size_t input_size_bytes = sizeof(input_t) * input_size; + + input_t *input = (input_t*) malloc (input_size_bytes); + + // https://cplusplus.com/reference/random/normal_distribution/ + std::default_random_engine generator (123); + std::normal_distribution distribution(5.0,2.0); + for (int i = 0; i < input_size; i++) { + input[i] = distribution(generator); + } + + auto min_iter = std::min_element(input, input+input_size); + auto max_iter = std::max_element(input, input+input_size); + + input_t input_minvalue = *min_iter; + input_t input_maxvalue = *max_iter; + printf("Input min, max values: (%f %f)\n", (float)input_minvalue, (float)input_maxvalue); + + #pragma omp target enter data map(to: input[0:input_size]) + + const int maxSharedMemory = + std::min(getDeviceLocalMemSize(), (int)sharedMemoryCapacity); + printf("Maximum shared local memory size per block in bytes: %d\n", maxSharedMemory); + + for (IndexType nbins = 768; nbins <= 768 * 32; nbins = nbins * 2) { + + printf("\nNumber of bins: %d\n", nbins); + IndexType sharedMem = nbins * sizeof(output_t); + + IndexType output_size = nbins; + size_t output_size_bytes = sizeof(output_t) * output_size; + output_t *output = (output_t*) malloc (output_size_bytes); + + // reference + output_t *output_r = (output_t*) calloc (output_size, sizeof(output_t)); + reference( + output_r, input, nbins, input_minvalue, input_maxvalue, + input_size, output_size, repeat); + + #pragma omp target enter data map(alloc: output[0:output_size]) + + input_t minvalue = input_minvalue; + input_t maxvalue = input_maxvalue; + + // determine memory type to use in the kernel + printf("bincount using global atomics\n"); + + #pragma omp target teams distribute parallel for + for (IndexType i = 0; i < output_size; i++) output[i] = 0; + + auto start = std::chrono::steady_clock::now(); + for (int n = 0; n < repeat; n++) { + #pragma omp target teams distribute parallel for + for (IndexType linearIndex = 0; linearIndex < input_size; linearIndex++) { + const input_t v = input[linearIndex]; + if (v >= minvalue && v <= maxvalue) { + const IndexType bin = getBin( + v, minvalue, maxvalue, nbins); + #pragma omp atomic update + output[bin] += 1; + } + } + } + auto end = std::chrono::steady_clock::now(); + auto time = std::chrono::duration_cast(end - start).count(); + printf("Average execution time of bincount kernel: %f (us)\n", + (time * 1e-3f) / repeat); + + #pragma omp target update from(output[0:output_size]) + + int status = memcmp(output, output_r, output_size_bytes); + printf("%s\n", status ? "FAIL" : "PASS"); + + if (sharedMem <= maxSharedMemory) { + printf("\n"); + printf("bincount using global and local atomics\n"); + + // number of teams mirrors the CUDA grid dimension + const IndexType numTeams = + (input_size + threadsPerBlock - 1) / threadsPerBlock; + + #pragma omp target teams distribute parallel for + for (IndexType i = 0; i < output_size; i++) output[i] = 0; + + start = std::chrono::steady_clock::now(); + for (int n = 0; n < repeat; n++) { + #pragma omp target teams num_teams(numTeams) thread_limit(threadsPerBlock) + { + // Per-team histogram. Declaring it in the teams region makes it + // shared by the team's threads and lets the compiler place it in + // team-local memory (LDS/SLM). + output_t smem[sharedMemoryCapacity / sizeof(output_t)]; + + #pragma omp parallel + { + const int nthreads = omp_get_num_threads(); + const int tid = omp_get_thread_num(); + const int team = omp_get_team_num(); + const int nteams = omp_get_num_teams(); + + // zero the shared histogram + for (IndexType i = tid; i < nbins; i += nthreads) smem[i] = 0; + + #pragma omp barrier + + // atomically accumulate into the shared histogram + for (IndexType linearIndex = (IndexType)team * nthreads + tid; + linearIndex < input_size; + linearIndex += (IndexType)nteams * nthreads) { + const input_t v = input[linearIndex]; + if (v >= minvalue && v <= maxvalue) { + const IndexType bin = getBin( + v, minvalue, maxvalue, nbins); + #pragma omp atomic update + smem[bin] += 1; + } + } + + #pragma omp barrier + + // flush the shared histogram to the global output + for (IndexType i = tid; i < nbins; i += nthreads) { + #pragma omp atomic update + output[i] += smem[i]; + } + } + } + } + end = std::chrono::steady_clock::now(); + time = std::chrono::duration_cast(end - start).count(); + printf("Average execution time of bincount kernel: %f (us)\n", + (time * 1e-3f) / repeat); + + #pragma omp target update from(output[0:output_size]) + + int status = memcmp(output, output_r, output_size_bytes); + printf("%s\n", status ? "FAIL" : "PASS"); + } + + #pragma omp target exit data map(delete: output[0:output_size]) + free(output); + free(output_r); + } + + #pragma omp target exit data map(delete: input[0:input_size]) + free(input); +} + +int main(int argc, char* argv[]) +{ + if (argc != 3) { + printf("Usage: %s \n", argv[0]); + return 1; + } + const int n = atoi(argv[1]); + const int repeat = atoi(argv[2]); + + eval(n, repeat); + + return 0; +} From 155455ac032608fbe17976a0d8a3d3c4a74716cc Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Thu, 6 Aug 2026 15:33:46 -0700 Subject: [PATCH 6/7] [pointerchase-omp] add the example --- src/pointerchase-omp/CMakeLists.txt | 8 ++ src/pointerchase-omp/Makefile | 62 ++++++++++++ src/pointerchase-omp/Makefile.aomp | 66 +++++++++++++ src/pointerchase-omp/Makefile.nvc | 63 ++++++++++++ src/pointerchase-omp/main.cpp | 146 ++++++++++++++++++++++++++++ 5 files changed, 345 insertions(+) create mode 100644 src/pointerchase-omp/CMakeLists.txt create mode 100644 src/pointerchase-omp/Makefile create mode 100644 src/pointerchase-omp/Makefile.aomp create mode 100644 src/pointerchase-omp/Makefile.nvc create mode 100644 src/pointerchase-omp/main.cpp diff --git a/src/pointerchase-omp/CMakeLists.txt b/src/pointerchase-omp/CMakeLists.txt new file mode 100644 index 0000000000..03034a3bc1 --- /dev/null +++ b/src/pointerchase-omp/CMakeLists.txt @@ -0,0 +1,8 @@ +# pointerchase-omp/CMakeLists.txt + +add_hecbench_benchmark( + NAME pointerchase + MODEL omp + SOURCES main.cpp + CATEGORIES algorithms +) diff --git a/src/pointerchase-omp/Makefile b/src/pointerchase-omp/Makefile new file mode 100644 index 0000000000..e1557fc75b --- /dev/null +++ b/src/pointerchase-omp/Makefile @@ -0,0 +1,62 @@ +#=============================================================================== +# User Options +#=============================================================================== + +# Compiler can be set below, or via environment variable +CC = icpx +OPTIMIZE = yes +DEBUG = no +DEVICE = gpu +LAUNCHER = + +#=============================================================================== +# Program name & source code list +#=============================================================================== + +program = main + +source = main.cpp + +obj = $(source:.cpp=.o) + +#=============================================================================== +# Sets Flags +#=============================================================================== + +# Standard Flags +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall + +# Linker Flags +LDFLAGS = + +# Debug Flags +ifeq ($(DEBUG),yes) + CFLAGS += -g + LDFLAGS += -g +endif + +# Optimization Flags +ifeq ($(OPTIMIZE),yes) + CFLAGS += -O3 +endif + +ifeq ($(DEVICE),gpu) + CFLAGS +=-fiopenmp -fopenmp-targets=spir64 -D__STRICT_ANSI__ +else + CFLAGS +=-qopenmp +endif +#=============================================================================== +# Targets to Build +#=============================================================================== + +$(program): $(obj) Makefile + $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) + +%.o: %.cpp Makefile + $(CC) $(CFLAGS) -c $< -o $@ + +clean: + rm -rf $(program) $(obj) + +run: $(program) + $(LAUNCHER) ./$(program) diff --git a/src/pointerchase-omp/Makefile.aomp b/src/pointerchase-omp/Makefile.aomp new file mode 100644 index 0000000000..b8a8ba5f42 --- /dev/null +++ b/src/pointerchase-omp/Makefile.aomp @@ -0,0 +1,66 @@ +#=============================================================================== +# User Options +#=============================================================================== + +# Compiler can be set below, or via environment variable +CC = clang++ +OPTIMIZE = yes +DEBUG = no +DEVICE = gpu +ARCH = gfx906 +LAUNCHER = + +#=============================================================================== +# Program name & source code list +#=============================================================================== + +program = main + +source = main.cpp + +obj = $(source:.cpp=.o) + +#=============================================================================== +# Sets Flags +#=============================================================================== + +# Standard Flags +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall + +# Linker Flags +LDFLAGS = + +# Debug Flags +ifeq ($(DEBUG),yes) + CFLAGS += -g + LDFLAGS += -g +endif + +# Optimization Flags +ifeq ($(OPTIMIZE),yes) + CFLAGS += -O3 +endif + +ifeq ($(DEVICE),gpu) + CFLAGS += -target x86_64-pc-linux-gnu \ + -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa \ + -Xopenmp-target=amdgcn-amd-amdhsa \ + -march=$(ARCH) +else + CFLAGS +=-fopenmp +endif +#=============================================================================== +# Targets to Build +#=============================================================================== + +$(program): $(obj) Makefile.aomp + $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) + +%.o: %.cpp Makefile.aomp + $(CC) $(CFLAGS) -c $< -o $@ + +clean: + rm -rf $(program) $(obj) + +run: $(program) + $(LAUNCHER) ./$(program) diff --git a/src/pointerchase-omp/Makefile.nvc b/src/pointerchase-omp/Makefile.nvc new file mode 100644 index 0000000000..eadd3a6960 --- /dev/null +++ b/src/pointerchase-omp/Makefile.nvc @@ -0,0 +1,63 @@ +#=============================================================================== +# User Options +#=============================================================================== + +# Compiler can be set below, or via environment variable +CC = nvc++ +OPTIMIZE = yes +DEBUG = no +DEVICE = gpu +SM = cc70 +LAUNCHER = + +#=============================================================================== +# Program name & source code list +#=============================================================================== + +program = main + +source = main.cpp + +obj = $(source:.cpp=.o) + +#=============================================================================== +# Sets Flags +#=============================================================================== + +# Standard Flags +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall + +# Linker Flags +LDFLAGS = + +# Debug Flags +ifeq ($(DEBUG),yes) + CFLAGS += -g + LDFLAGS += -g +endif + +# Optimization Flags +ifeq ($(OPTIMIZE),yes) + CFLAGS += -O3 +endif + +ifeq ($(DEVICE),gpu) + CFLAGS +=-Minfo -mp=gpu -gpu=$(SM) +else + CFLAGS +=-mp +endif +#=============================================================================== +# Targets to Build +#=============================================================================== + +$(program): $(obj) Makefile.nvc + $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) + +%.o: %.cpp Makefile.nvc + $(CC) $(CFLAGS) -c $< -o $@ + +clean: + rm -rf $(program) $(obj) + +run: $(program) + $(LAUNCHER) ./$(program) diff --git a/src/pointerchase-omp/main.cpp b/src/pointerchase-omp/main.cpp new file mode 100644 index 0000000000..8b764af120 --- /dev/null +++ b/src/pointerchase-omp/main.cpp @@ -0,0 +1,146 @@ +#include +#include +#include +#include +#include +#include +#include + +const uint64_t latencyMemAccessCnt = 1000000; /* 1M total read accesses to gauge latency */ +const uint32_t _2MiB = 2 * 1024 * 1024; +const uint32_t strideLen = 16; /* cacheLine size 128 Bytes, 16 words */ + +/* Upper bound on the number of teams used */ +const uint32_t maxTeamCount = 256; + +struct LatencyNode { + struct LatencyNode *next; +}; + +void initBuffer(void* buffer, uint64_t buffer_size, bool measureDeviceToDeviceLatency) { + uint64_t n_ptrs = buffer_size / sizeof(struct LatencyNode); + + if (measureDeviceToDeviceLatency) { + // For device-to-device latency, create and initialize pattern on device + const int dev = omp_get_default_device(); + const int host = omp_get_initial_device(); + for (uint64_t i = 0; i < n_ptrs; i++) { + struct LatencyNode node; + uint64_t nextOffset = ((i + strideLen) % n_ptrs) * sizeof(struct LatencyNode); + // Set up pattern with device addresses + node.next = (struct LatencyNode*)((uint8_t*)buffer + nextOffset); + int status = omp_target_memcpy(buffer, &node, sizeof(struct LatencyNode), + i * sizeof(struct LatencyNode), 0, + dev, host); + if (status != 0) { + fprintf(stderr, "omp_target_memcpy failed with status %d\n", status); + exit(EXIT_FAILURE); + } + } + } else { + // For host-device latency, initialize pattern with host addresses + struct LatencyNode* hostMem = (struct LatencyNode*)buffer; + for (uint64_t i = 0; i < n_ptrs; i++) { + hostMem[i].next = &hostMem[(i + strideLen) % n_ptrs]; + } + } +} + + +double latencyPtrChaseKernel(void* data, uint64_t memAccessCnt, + uint32_t smCount) +{ + double latencySum = 0.0f; + uint32_t measuredTeamCount = 0; + struct LatencyNode *nodes = static_cast(data); + + // For smCount teams, each team has memAccessCnt pointer chases + for (uint32_t targetBlock = 0; targetBlock < smCount; ++targetBlock) { + int executed = 0; + auto start = std::chrono::steady_clock::now(); + + // The body of a teams region is executed by one thread per team + #pragma omp target teams num_teams(smCount) thread_limit(1) \ + is_device_ptr(nodes) map(tofrom: executed) + { + if ((uint32_t)omp_get_team_num() == targetBlock) { + executed = 1; + struct LatencyNode *p = nodes; + for (uint32_t i = 0; i < memAccessCnt; ++i) { + p = p->next; + } + + // Avoid compiler optimization: the store is never reached, but the + // compiler cannot prove it and therefore has to keep the pointer chase + // above. An assert() would not do, as it is compiled out when NDEBUG is + // defined. + if (p == nullptr) { + nodes[0].next = nullptr; + } + } + } + + auto end = std::chrono::steady_clock::now(); + if (executed) { + auto latency = + std::chrono::duration_cast(end - start) + .count(); + latencySum += latency; + ++measuredTeamCount; + } + } + if (measuredTeamCount == 0) { + fprintf(stderr, "No OpenMP target team executed the pointer chase\n"); + exit(EXIT_FAILURE); + } + return latencySum / + (memAccessCnt * measuredTeamCount); // finalLatencyPerAccessNs +} + +class MemPtrChaseOperation { + public: + MemPtrChaseOperation() { + // OpenMP has no portable query for the number of compute units + int teams = 0; + #pragma omp target teams num_teams(maxTeamCount) thread_limit(1) \ + map(tofrom: teams) + { + if (omp_get_team_num() == 0) teams = omp_get_num_teams(); + } + if (teams < 1) teams = 1; + smCount = (uint32_t)teams; + if (smCount > maxTeamCount) smCount = maxTeamCount; + } + ~MemPtrChaseOperation() = default; + double doPtrChase(void* peerBuffer) { + double lat = + latencyPtrChaseKernel(peerBuffer, latencyMemAccessCnt, smCount); + return lat; + } + private: + uint32_t smCount; +}; + +int main() { + const uint64_t buffer_size = _2MiB; + const bool measureDeviceToDeviceLatency = true; + + void *buffer = omp_target_alloc(buffer_size, omp_get_default_device()); + if (buffer == nullptr) { + fprintf(stderr, "Failed to allocate %lu bytes on the device\n", + (unsigned long) buffer_size); + return 1; + } + + // initialize the buffer + initBuffer(buffer, buffer_size, measureDeviceToDeviceLatency); + + // compute the latency of pointer chasing on the default device + MemPtrChaseOperation mpc; + + double lat = mpc.doPtrChase(buffer); + + printf("Latency per access on device: %lf (ns)\n", lat); + omp_target_free(buffer, omp_get_default_device()); + return 0; +} From e3f15f17ef0bd04653c08bb4639fe66907463b6b Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Thu, 6 Aug 2026 19:57:53 -0700 Subject: [PATCH 7/7] [dwconv] replace checksum reporting with reference-based validation --- src/dwconv-cuda/Makefile | 2 +- src/dwconv-cuda/main.cu | 17 +++++++-- src/dwconv-cuda/reference.h | 67 ++++++++++++++++++++++++++++++++++ src/dwconv-hip/CMakeLists.txt | 5 ++- src/dwconv-hip/Makefile | 5 ++- src/dwconv-hip/main.cu | 17 +++++++-- src/dwconv-omp/CMakeLists.txt | 3 ++ src/dwconv-omp/Makefile | 4 +- src/dwconv-omp/Makefile.aomp | 4 +- src/dwconv-omp/Makefile.nvc | 4 +- src/dwconv-omp/main.cpp | 17 +++++++-- src/dwconv-sycl/CMakeLists.txt | 5 ++- src/dwconv-sycl/Makefile | 4 +- src/dwconv-sycl/main.cpp | 17 +++++++-- 14 files changed, 146 insertions(+), 25 deletions(-) create mode 100644 src/dwconv-cuda/reference.h diff --git a/src/dwconv-cuda/Makefile b/src/dwconv-cuda/Makefile index b2e596e8bd..3c016b0c7f 100644 --- a/src/dwconv-cuda/Makefile +++ b/src/dwconv-cuda/Makefile @@ -47,7 +47,7 @@ endif $(program): $(obj) Makefile $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) -%.o: %.cu ../tensorAccessor-cuda/tensorAccessor.h Makefile +%.o: %.cu reference.h ../tensorAccessor-cuda/tensorAccessor.h Makefile $(CC) $(CFLAGS) -c $< -o $@ clean: diff --git a/src/dwconv-cuda/main.cu b/src/dwconv-cuda/main.cu index d4d83d9290..ce681ee60a 100644 --- a/src/dwconv-cuda/main.cu +++ b/src/dwconv-cuda/main.cu @@ -1,7 +1,9 @@ #include #include #include +#include #include +#include "reference.h" #include "tensorAccessor.h" template class PtrTraits = DefaultPtrTraits> @@ -154,6 +156,7 @@ void dwconv2d_forward (const int m, scalar_t *h_weight = (scalar_t*) malloc (weight_size_bytes); scalar_t *h_bias = (scalar_t*) malloc (bias_size_bytes); scalar_t *h_output = (scalar_t*) malloc (output_size_bytes); + scalar_t *h_reference = (scalar_t*) malloc (output_size_bytes); srand(123); for (int i = 0; i < input_size; i++) { @@ -216,11 +219,18 @@ void dwconv2d_forward (const int m, cudaMemcpy(h_output, d_output, output_size_bytes, cudaMemcpyDeviceToHost); - scalar_t sum = 0; + reference( + h_input, h_reference, h_weight, h_bias, has_bias, + n, inputChannels, height, width, outputChannels, outputHeight, outputWidth, + kH, kW, strideH, strideW, padH, padW, dilateH, dilateW); + + int errors = 0; for (int i = 0; i < output_size; i++) { - sum += h_output[i]; + const scalar_t tolerance = 1e-4f + 1e-4f * fabs(h_reference[i]); + if (fabs(h_output[i] - h_reference[i]) > tolerance) + errors++; } - printf("Checksum = %f\n", sum / output_size); + printf("%s\n", errors == 0 ? "PASS" : "FAIL"); cudaFree(d_input); cudaFree(d_output); @@ -229,6 +239,7 @@ void dwconv2d_forward (const int m, free(h_input); free(h_output); + free(h_reference); free(h_weight); free(h_bias); } diff --git a/src/dwconv-cuda/reference.h b/src/dwconv-cuda/reference.h new file mode 100644 index 0000000000..20324ea03f --- /dev/null +++ b/src/dwconv-cuda/reference.h @@ -0,0 +1,67 @@ +#ifndef DWCONV_REFERENCE_H +#define DWCONV_REFERENCE_H + +template +void reference( + const scalar_t *input, + scalar_t *output, + const scalar_t *weight, + const scalar_t *bias, + const bool bias_enabled, + const int batch_size, + const int input_channels, + const int input_height, + const int input_width, + const int output_channels, + const int output_height, + const int output_width, + const int kernel_height, + const int kernel_width, + const int stride_height, + const int stride_width, + const int pad_height, + const int pad_width, + const int dilation_height, + const int dilation_width) +{ + const int depthwise_multiplier = output_channels / input_channels; + + for (int n = 0; n < batch_size; ++n) { + for (int c = 0; c < output_channels; ++c) { + const int input_channel = c / depthwise_multiplier; + for (int h = 0; h < output_height; ++h) { + for (int w = 0; w < output_width; ++w) { + acc_t value = bias_enabled + ? static_cast(bias[c]) + : static_cast(0); + + for (int kh = 0; kh < kernel_height; ++kh) { + const int input_h = + -pad_height + h * stride_height + kh * dilation_height; + for (int kw = 0; kw < kernel_width; ++kw) { + const int input_w = + -pad_width + w * stride_width + kw * dilation_width; + if (input_h >= 0 && input_h < input_height && + input_w >= 0 && input_w < input_width) { + const int input_index = + ((n * input_channels + input_channel) * input_height + + input_h) * input_width + input_w; + const int weight_index = + (c * kernel_height + kh) * kernel_width + kw; + value += static_cast(weight[weight_index]) * + static_cast(input[input_index]); + } + } + } + + const int output_index = + ((n * output_channels + c) * output_height + h) * + output_width + w; + output[output_index] = static_cast(value); + } + } + } + } +} + +#endif diff --git a/src/dwconv-hip/CMakeLists.txt b/src/dwconv-hip/CMakeLists.txt index 84e7fd4d07..665ef2010e 100644 --- a/src/dwconv-hip/CMakeLists.txt +++ b/src/dwconv-hip/CMakeLists.txt @@ -1,6 +1,9 @@ # dwconv-hip/CMakeLists.txt -set(INC_DIR "${CMAKE_CURRENT_LIST_DIR}/../tensorAccessor-cuda") +set(INC_DIR + "${CMAKE_CURRENT_LIST_DIR}/../tensorAccessor-cuda" + "${CMAKE_CURRENT_LIST_DIR}/../dwconv-cuda" +) add_hecbench_benchmark( NAME dwconv diff --git a/src/dwconv-hip/Makefile b/src/dwconv-hip/Makefile index 206efbc7d6..e7912b82ae 100644 --- a/src/dwconv-hip/Makefile +++ b/src/dwconv-hip/Makefile @@ -23,7 +23,8 @@ obj = $(source:.cu=.o) #=============================================================================== # Standard Flags -CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall -I../tensorAccessor-cuda +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall -I../tensorAccessor-cuda \ + -I../dwconv-cuda # Linker Flags LDFLAGS = @@ -45,7 +46,7 @@ endif $(program): $(obj) Makefile $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) -%.o: %.cu ../tensorAccessor-cuda/tensorAccessor.h Makefile +%.o: %.cu ../dwconv-cuda/reference.h ../tensorAccessor-cuda/tensorAccessor.h Makefile $(CC) $(CFLAGS) -c $< -o $@ clean: diff --git a/src/dwconv-hip/main.cu b/src/dwconv-hip/main.cu index 8992f94bf0..c5a35e2434 100644 --- a/src/dwconv-hip/main.cu +++ b/src/dwconv-hip/main.cu @@ -1,7 +1,9 @@ #include #include #include +#include #include +#include "reference.h" #include "tensorAccessor.h" template class PtrTraits = DefaultPtrTraits> @@ -154,6 +156,7 @@ void dwconv2d_forward (const int m, scalar_t *h_weight = (scalar_t*) malloc (weight_size_bytes); scalar_t *h_bias = (scalar_t*) malloc (bias_size_bytes); scalar_t *h_output = (scalar_t*) malloc (output_size_bytes); + scalar_t *h_reference = (scalar_t*) malloc (output_size_bytes); srand(123); for (int i = 0; i < input_size; i++) { @@ -216,11 +219,18 @@ void dwconv2d_forward (const int m, hipMemcpy(h_output, d_output, output_size_bytes, hipMemcpyDeviceToHost); - scalar_t sum = 0; + reference( + h_input, h_reference, h_weight, h_bias, has_bias, + n, inputChannels, height, width, outputChannels, outputHeight, outputWidth, + kH, kW, strideH, strideW, padH, padW, dilateH, dilateW); + + int errors = 0; for (int i = 0; i < output_size; i++) { - sum += h_output[i]; + const scalar_t tolerance = 1e-4f + 1e-4f * fabs(h_reference[i]); + if (fabs(h_output[i] - h_reference[i]) > tolerance) + errors++; } - printf("Checksum = %f\n", sum / output_size); + printf("%s\n", errors == 0 ? "PASS" : "FAIL"); hipFree(d_input); hipFree(d_output); @@ -229,6 +239,7 @@ void dwconv2d_forward (const int m, free(h_input); free(h_output); + free(h_reference); free(h_weight); free(h_bias); } diff --git a/src/dwconv-omp/CMakeLists.txt b/src/dwconv-omp/CMakeLists.txt index 37333cdf7a..3b7d0a945d 100644 --- a/src/dwconv-omp/CMakeLists.txt +++ b/src/dwconv-omp/CMakeLists.txt @@ -1,8 +1,11 @@ # dwconv-omp/CMakeLists.txt +set(INC_DIR "${CMAKE_CURRENT_LIST_DIR}/../dwconv-cuda") + add_hecbench_benchmark( NAME dwconv MODEL omp SOURCES main.cpp CATEGORIES algorithms + INCLUDE_DIRS ${INC_DIR} ) diff --git a/src/dwconv-omp/Makefile b/src/dwconv-omp/Makefile index aac4a69113..ecca1e1c54 100644 --- a/src/dwconv-omp/Makefile +++ b/src/dwconv-omp/Makefile @@ -24,7 +24,7 @@ obj = $(source:.cpp=.o) #=============================================================================== # Standard Flags -CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall -I../dwconv-cuda # Linker Flags LDFLAGS = @@ -52,7 +52,7 @@ endif $(program): $(obj) Makefile $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) -%.o: %.cpp Makefile +%.o: %.cpp ../dwconv-cuda/reference.h Makefile $(CC) $(CFLAGS) -c $< -o $@ clean: diff --git a/src/dwconv-omp/Makefile.aomp b/src/dwconv-omp/Makefile.aomp index 51570d320b..a4a099acd6 100644 --- a/src/dwconv-omp/Makefile.aomp +++ b/src/dwconv-omp/Makefile.aomp @@ -25,7 +25,7 @@ obj = $(source:.cpp=.o) #=============================================================================== # Standard Flags -CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall -I../dwconv-cuda # Linker Flags LDFLAGS = @@ -56,7 +56,7 @@ endif $(program): $(obj) Makefile.aomp $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) -%.o: %.cpp Makefile.aomp +%.o: %.cpp ../dwconv-cuda/reference.h Makefile.aomp $(CC) $(CFLAGS) -c $< -o $@ clean: diff --git a/src/dwconv-omp/Makefile.nvc b/src/dwconv-omp/Makefile.nvc index 76eb389f94..962823714e 100644 --- a/src/dwconv-omp/Makefile.nvc +++ b/src/dwconv-omp/Makefile.nvc @@ -25,7 +25,7 @@ obj = $(source:.cpp=.o) #=============================================================================== # Standard Flags -CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall +CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall -I../dwconv-cuda # Linker Flags LDFLAGS = @@ -53,7 +53,7 @@ endif $(program): $(obj) Makefile.nvc $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) -%.o: %.cpp Makefile.nvc +%.o: %.cpp ../dwconv-cuda/reference.h Makefile.nvc $(CC) $(CFLAGS) -c $< -o $@ clean: diff --git a/src/dwconv-omp/main.cpp b/src/dwconv-omp/main.cpp index 8246099d70..01c4d4aa86 100644 --- a/src/dwconv-omp/main.cpp +++ b/src/dwconv-omp/main.cpp @@ -1,7 +1,9 @@ #include #include #include +#include #include +#include "reference.h" // Depthwise 2D convolution forward. @@ -142,6 +144,7 @@ void dwconv2d_forward (const int m, scalar_t *h_weight = (scalar_t*) malloc (weight_size_bytes); scalar_t *h_bias = (scalar_t*) malloc (bias_size_bytes); scalar_t *h_output = (scalar_t*) malloc (output_size_bytes); + scalar_t *h_reference = (scalar_t*) malloc (output_size_bytes); srand(123); for (int i = 0; i < input_size; i++) { @@ -188,14 +191,22 @@ void dwconv2d_forward (const int m, time * 1e-6f / repeat); } - scalar_t sum = 0; + reference( + h_input, h_reference, h_weight, h_bias, has_bias, + n, inputChannels, height, width, outputChannels, outputHeight, outputWidth, + kH, kW, strideH, strideW, padH, padW, dilateH, dilateW); + + int errors = 0; for (int i = 0; i < output_size; i++) { - sum += h_output[i]; + const scalar_t tolerance = 1e-4f + 1e-4f * fabs(h_reference[i]); + if (fabs(h_output[i] - h_reference[i]) > tolerance) + errors++; } - printf("Checksum = %f\n", sum / output_size); + printf("%s\n", errors == 0 ? "PASS" : "FAIL"); free(h_input); free(h_output); + free(h_reference); free(h_weight); free(h_bias); } diff --git a/src/dwconv-sycl/CMakeLists.txt b/src/dwconv-sycl/CMakeLists.txt index 830a21f390..c2dfd960dc 100644 --- a/src/dwconv-sycl/CMakeLists.txt +++ b/src/dwconv-sycl/CMakeLists.txt @@ -1,6 +1,9 @@ # dwconv-sycl/CMakeLists.txt -set(INC_DIR "${CMAKE_CURRENT_LIST_DIR}/../tensorAccessor-cuda") +set(INC_DIR + "${CMAKE_CURRENT_LIST_DIR}/../tensorAccessor-cuda" + "${CMAKE_CURRENT_LIST_DIR}/../dwconv-cuda" +) add_hecbench_benchmark( NAME dwconv diff --git a/src/dwconv-sycl/Makefile b/src/dwconv-sycl/Makefile index d808eaac6d..8db3094945 100644 --- a/src/dwconv-sycl/Makefile +++ b/src/dwconv-sycl/Makefile @@ -31,7 +31,7 @@ obj = $(source:.cpp=.o) # Standard Flags CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall \ - -fsycl -I../tensorAccessor-cuda \ + -fsycl -I../tensorAccessor-cuda -I../dwconv-cuda \ --gcc-toolchain=$(GCC_TOOLCHAIN) ifeq ($(VENDOR), AdaptiveCpp) @@ -73,7 +73,7 @@ endif $(program): $(obj) $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) -%.o: %.cpp ../tensorAccessor-cuda/tensorAccessor.h +%.o: %.cpp ../dwconv-cuda/reference.h ../tensorAccessor-cuda/tensorAccessor.h $(CC) $(CFLAGS) -c $< -o $@ clean: diff --git a/src/dwconv-sycl/main.cpp b/src/dwconv-sycl/main.cpp index 63f7b8a3d1..d633a950b5 100644 --- a/src/dwconv-sycl/main.cpp +++ b/src/dwconv-sycl/main.cpp @@ -1,7 +1,9 @@ #include #include #include +#include #include +#include "reference.h" #include "tensorAccessor.h" template class PtrTraits = DefaultPtrTraits> @@ -156,6 +158,7 @@ void dwconv2d_forward (sycl::queue &q, scalar_t *h_weight = (scalar_t*) malloc (weight_size_bytes); scalar_t *h_bias = (scalar_t*) malloc (bias_size_bytes); scalar_t *h_output = (scalar_t*) malloc (output_size_bytes); + scalar_t *h_reference = (scalar_t*) malloc (output_size_bytes); srand(123); for (int i = 0; i < input_size; i++) { @@ -235,11 +238,18 @@ void dwconv2d_forward (sycl::queue &q, q.memcpy(h_output, d_output, output_size_bytes).wait(); - scalar_t sum = 0; + reference( + h_input, h_reference, h_weight, h_bias, has_bias, + n, inputChannels, height, width, outputChannels, outputHeight, outputWidth, + kH, kW, strideH, strideW, padH, padW, dilateH, dilateW); + + int errors = 0; for (int i = 0; i < output_size; i++) { - sum += h_output[i]; + const scalar_t tolerance = 1e-4f + 1e-4f * fabs(h_reference[i]); + if (fabs(h_output[i] - h_reference[i]) > tolerance) + errors++; } - printf("Checksum = %f\n", sum / output_size); + printf("%s\n", errors == 0 ? "PASS" : "FAIL"); sycl::free(d_input, q); sycl::free(d_output, q); @@ -248,6 +258,7 @@ void dwconv2d_forward (sycl::queue &q, free(h_input); free(h_output); + free(h_reference); free(h_weight); free(h_bias); }