Light Field Super-Resolution (LFSR)

• Light Field Super-Resolution (LFSR) is a computational technique that enhances light field images by reconstructing high-resolution spatial details while retaining angular information.
• LFSR methods integrate multi-view data through interpolation, depth estimation, and deep learning to improve image quality for applications like virtual reality and digital refocusing.
• Recent advances in LFSR focus on balancing computational efficiency with performance, using metrics such as PSNR and SSIM to quantify improvements over conventional methods.

GEMMbench is a community-oriented framework and methodology for evaluating the performance of matrix-matrix multiplication (GEMM) implementations in a reproducible, comparable, and extensible way. It is built on top of the Collective Knowledge (CK) framework, which provides the plumbing for packaging experiments, capturing metadata, automating runs, and sharing results and artifacts.

Motivation and problem statement

• Despite GEMM’s ubiquity in benchmarking architectures, compilers, and libraries, the field has long lacked a common methodology. Papers often vary in subtle but impactful ways: specialization to C = A × B (i.e., α = 1, β = 0) vs general α,β; data type (SGEMM/DGEMM/complex); layout (N vs T) and whether transposition overheads are included; rectangular vs square shapes and size regimes; cache effects and heterogeneous data transfer overheads; whether power/energy is measured; and the extent of human effort vs generator sophistication.
• These variations complicate interpretation, comparison, and scientific progress. GEMMbench addresses this by standardizing configuration, recording metadata, automating experiments, validating results, and enabling collaborative crowdsourcing of kernels and data.

Mathematical definition and performance metrics

• GEMM definition (real-valued): $C \leftarrow \alpha A B + \beta C$, with $A \in \mathbb{R}^{m\times k}$, $B \in \mathbb{R}^{k\times n}$, $C \in \mathbb{R}^{m\times n}$, and scalars $\alpha, \beta$.
• FLOP count (real GEMM): $\text{FLOPs} = 2 m n k$.
• Performance metric: $\text{GFLOPs/s} = \frac{2 m n k}{t} \times 10^{-9}$, where $t$ is execution time in seconds.

Overview of the framework and methodology

• Built on CK: GEMMbench leverages CK’s “pipelines” and metadata to run experiments under controlled conditions and record all relevant parameters (platform, device, kernel variant, data type, layout, sizes, work-group sizes, etc.).
• Reproducibility: CK captures the complete experiment context (software versions, OpenCL driver and platform selection, DVFS settings, command-line parameters, kernel source and JSON metadata, statistical repetitions). Experiments can be replayed exactly using a single command.
• Collaboration and crowdsourcing: The GEMMbench repository defines datasets describing kernels via JSON, while the community can pull/push experiment repositories, share result archives, and contribute new kernels and datasets. CK’s web UI lets users browse experiments, copy replay commands, and compare runs.
• Experiment structure: GEMMbench provides scripts, parameter sweeps, and standard views. Results are aggregated with basic statistics (mean, std) and, where available, energy measurements.

Supported implementations and data types

• Initial focus: handwritten OpenCL kernels for single-precision (SGEMM) and double-precision (DGEMM) real-valued matrices. Layouts include N (non-transposed) and T (transposed), with explicit metadata capturing transA/transB.
• Thread coarsening and vectorization: Work-items can produce single ($1\times 1$) or multiple ($d_i \times d_j$) output elements. GEMMbench includes variants where each work-item computes a vector of four adjacent outputs in the j-dimension (“4×1”), as well as a barrier-optimized “4×1” kernel to improve cache utilization. In practice, such kernels often leverage OpenCL vector types (e.g., float4 for SGEMM), but GEMMbench’s metadata is agnostic and focuses on what each work-item computes (di,dj) rather than mandating a specific type.
• Work-item behavior examples:
  • SGEMM_NT_1x1: each work-item computes a single element C[i,j] using get_global_id(0)=j and get_global_id(1)=i over k in a simple accumulation loop.
  • SGEMM_NT_4x1: each work-item computes four adjacent elements along j (coarsening and vectorization in the output dimension). A barrier-optimized variant synchronizes work-items to improve cache behavior.

Kernel metadata and configuration

• Kernel metadata is stored as JSON to capture:
  • type: “S” (single) or “D” (double),
  • transA/transB: “N” or “T”,
  • di, dj: thread coarsening (block of rows/columns per work-item),
  • name, file, and other attributes.
• Example metadata (SGEMM_NT_1x1): name=“SGEMM_NT_1x1”, file=“SGEMM_NT_1x1.cl”, type=“S”, transA=“N”, transB=“T”, di=1, dj=1.
• CK captures and varies parameters such as:
  • local work size (LWS; e.g., (8,8), (1,16), (2,32), etc.),
  • global sizes implicitly derived from problem order and kernel coarsening,
  • matrix order n (for square cases in initial experiments),
  • platform/device selection (OpenCL platform/device IDs),
  • scalars α, β and initial C (for generality, though many experiments use α=1, β=0),
  • number of statistical repetitions.

Experimental methodology: measurement, platforms, and tools

• Measurement approach:
  • Runtime is captured via CK pipelines and converted to GFLOPs/s using the standard formula.
  • Statistical rigor: multiple repetitions per experiment point; mean and standard deviation reported (computed using pandas).
  • Energy: on supported platforms (Odroid XU3), GPU and memory energy is estimated using integrated power sensors. The method averages power at entry and exit of the kernel region and multiplies by the region’s execution time (a deliberate, simple approximation).
• Platform used in the initial evaluation:
  • Hardware: Odroid XU3 (Samsung Exynos 5422), ARM Mali-T628 GPU (cluster 0: 4 cores at up to 600 MHz; cluster 1: 2 cores), CPU clusters (Cortex-A7 x4 up to 1.4 GHz and Cortex-A15 x4 up to 2.0 GHz), 2 GiB LPDDR3 RAM (≤ 14.9 GB/s), and integrated power sensors for RAM, CPU clusters, and GPU clusters.
  • Software: Ubuntu 14.04.x, Linux kernel 3.10.69, OpenCL 1.1 Full Profile (driver version 4.0), host compiler Clang++ 3.6.
• Controlled conditions:
  • GPU DVFS disabled; GPU frequency set to 600 MHz.
  • CPU governors set to performance mode; CPU frequencies at max.
  • Mali-T628 concurrency considerations: default LWS (8,8) typically supports up to four concurrently executing work-groups per core, depending on register usage.

Results and insights

• Varying matrix order (64 to 1024), fixed LWS = (8,8), SGEMM NT kernels on Mali-T628 GPU cluster 0:
  • SGEMM_NT_1x1 (naïve 1×1):
  • Low but stable performance up to ≈3 GFLOPs/s; e.g., around 2.96 GFLOPs/s at n=64; ≈3.01 GFLOPs/s at n=128; ≈2.62 GFLOPs/s at n=1024.
  • SGEMM_NT_4x1 (vectorized 4×1 without barrier):
  • Peaks over 12 GFLOPs/s (≈12.16 GFLOPs/s at n=128), but exhibits dramatic drops at orders that are multiples of 256 (n=256, 512, 768, 1024), falling to ≈3.86 GFLOPs/s at n=256, ≈1.84 at n=512, ≈2.17 at n=768, and ≈1.55 at n=1024. High run-to-run variation at some sizes (e.g., n=640 mean ≈4.77 GFLOPs/s with std ≈3.74).
  • SGEMM_NT_4x1_barrier (vectorized 4×1 with barrier):
  • Achieves up to ≈11.39 GFLOPs/s (n=128) and, critically, maintains high performance of ≈9–10 GFLOPs/s for larger orders (≥256), e.g., ≈9.40 at n=256, ≈9.24 at n=512, ≈9.85 at n=640, ≈9.84 at n=1024, with low variation.
  • Key insight: simple vectorization can be fragile without cache-aware coordination; adding a barrier to improve cache utilization stabilizes and sustains high performance at larger sizes.
• Varying local work size (LWS) for SGEMM_NT_4x1_barrier across selected orders:
  • Systematic exploration of LWS from 16 to 128 work-items per work-group revealed up to ~20% improvement over the baseline (8,8).
  • Notable patterns:
  • Shapes with smaller first dimension (j) and larger second (i) performed better than their transposes; e.g., (2,16) outperformed (16,2) by 1.5–3×.
  • Best observed performance was 11.9 GFLOPs/s with LWS (1,16).
  • For small LWS (16, 32), s_l = 1 (smaller j) tended to be best; for larger LWS (64, 128), s_l = 4 was often preferable.
• Energy consumption (GPU and memory) on Odroid XU3:
  • At n=128 and n=384, when the two vectorized kernels matched in performance, they also matched in energy.
  • At n=1024, the non-cache-optimized 4×1 kernel was dramatically worse:
  • GPU energy: ≈1.1766 J (4×1) vs ≈0.0296 J (4×1_barrier) — the barrier kernel is ~40× more energy efficient; the 1×1 kernel ≈1.0883 J, i.e., even the naïve kernel is ~75% faster and ~10% more GPU-energy efficient than the non-cache-optimized vectorized one.
  • Memory energy: ≈0.4976 J (4×1) vs ≈0.01117 J (4×1_barrier) — ~45× improvement; 1×1 ≈0.1133 J (≈5× better than 4×1).
  • Insight: cache-aware synchronization not only stabilizes performance but massively improves energy efficiency at large sizes.

Kernel optimization techniques covered

• Thread coarsening and vectorization: increasing di,dj so each work-item computes multiple outputs reduces global memory traffic per output and can leverage hardware SIMD lanes.
• Work-group sizing: careful choice of local work size (LWS) influences occupancy, cache behavior, and scheduling; asymmetric shapes with s_l < s_h in j×i often outperform their transpose on Mali-T628.
• Barrier-based coordination: synchronizing work-items in a work-group can improve cache reuse and reduce memory system contention; in the SGEMM_NT_4x1_barrier kernel this translated to sustained performance at large problem sizes.

Validation and correctness

• GEMMbench includes a reference CPU implementation to validate OpenCL results via element-wise comparison with an epsilon threshold.
• Handling floating-point discrepancies:
  • Initial epsilon ε = 1e−5 proved too strict (and an early bug used abs() instead of fabs()); empirical tuning found ε = 0.1 worked well with inputs drawn from U(−0.5, +0.5).
  • Failures were observed for certain kernels at small sizes (e.g., SGEMM_NT_4x1_barrier at n=64, 96, 192). Some failures were intermittent (hard to reproduce), others repeatable (suggestive of specific issues).
  • GEMMbench’s reproducibility (ck replay experiment:… --point=…) streamlines debugging and community triage.

CK integration details: installing, running, reproducing, sharing

• Installation outline:
  • Install CK, add its bin to PATH, and verify version (e.g., 1.6.12).
  • Pull GEMMbench and its CK dependencies (ck-autotuning, ck-env).
  • Register an OpenCL driver (setup soft:lib.opencl.*), pointing CK to your libOpenCL.
  • Optionally customize platform init scripts to disable DVFS and pin frequencies; add them to PATH.
  • Compile the launcher program (program:gemmbench-cl-launcher-1.0); CK auto-fetches cJSON and xOpenME as needed.
• Running:
  • Use CK’s run interface to select flavour (SGEMM NN, SGEMM NT, DGEMM NN, DGEMM NT) and kernel variant.
  • Override parameters with extra_run_cmd; e.g., “-p 1 -d 1 -n 512 -lws 4,16”.
  • Provided scripts:
  • SGEMM_NT/_setup_program_pipeline.sh; explore-f-n (vary order); explore-n-lws (vary LWS).
• Reproducibility:
  • Pull paper’s experiment repo (gemmbench-adapt16) to replay published results.
  • Launch CK web (ck start web), browse experiments and copy “ck replay … --point=…” commands for exact replays.
• Sharing results:
  • Move locally created experiments to a new CK-backed Git repo (ck mv experiment:… repo:new:experiment:…).
  • Commit, publish, and share the URL; GEMMbench maintainers will list community repos.

Limitations and future work

• Initial scope: square matrices, real single/double precision, OpenCL kernels; complex numbers, rectangular shapes, and non-OpenCL implementations are planned.
• Kernel set is intentionally small and optimized for ARM Mali-T600; more kernels targeting diverse architectures are welcome.
• Some validation failures and intermittent bugs remain; community participation and tooling (GPUVerify, Oclgrind) are encouraged to uncover and fix issues.
• Energy measurement method (start/end averaging) is simplistic; more precise instrumentation would be beneficial.
• Legal constraints can restrict sharing of benchmark results on some vendor platforms; the Odroid BSP is free of such limitations.
• Broader CK strengths (multi-OS/compiler integration, data mining/predictive analytics) will be leveraged for future auto-tuning and performance portability studies.

Practical guidance: using GEMMbench to evaluate your own GEMM kernel

• Prepare your kernel and metadata:
  • Implement your OpenCL kernel and place it in the GEMMbench repository structure.
  • Create a JSON metadata entry specifying type (“S”/“D”), transA/transB (“N”/“T”), di/dj (thread coarsening), file, and name.
• Configure experimental conditions:
  • Disable DVFS and pin CPU/GPU frequencies to reduce run-to-run variability.
  • Choose a baseline LWS (e.g., (8,8)) and a set of matrix orders (e.g., 64, 128, 256, 384, 512, 640, 768, 896, 1024).
  • Define the number of repetitions for statistical confidence (≥4).
• Parameter sweeps:
  • Vary matrix order to identify cache thresholds and stability issues; watch for patterns at powers of two or multiples of cache-sensitive sizes (e.g., 256).
  • Sweep LWS with asymmetric shapes (prefer j×i where smaller j < larger i); try (1,16), (2,16), (2,32), (4,16), (4,32), (8,8), (16,4), etc. The Mali-T628 data suggests (1,16) and (2,32) are strong baselines, and (4,16) can be competitive at higher LWS.
  • Explore di/dj coarsening factors (e.g., 1×1, 4×1, 1×4) to balance memory traffic and compute utilization.
• Measure and validate:
  • Record runtime; compute GFLOPs/s via $A \in \mathbb{R}^{m\times k}$0.
  • If energy sensors exist, capture GPU and memory energy for the kernel region (entry/exit averaging).
  • Validate results against the CPU reference with ε appropriate to your input distribution; start with ε=0.1 for U(−0.5, +0.5) and adjust as needed.
• Reproduce and share:
  • Use CK’s web UI to inspect runs and copy replay commands.
  • Publish your experiments via a CK repo and share the link with the GEMMbench community.
• Interpret results:
  • Look for sustained performance at large orders, not just peaks; cache-aware synchronization can transform both performance and energy efficiency.
  • Examine variation (std dev) to identify intermittency or instability; use ck replay to isolate and debug.
  • Use insights from LWS sweeps to guide non-exhaustive tuning; prioritize asymmetric LWS and barrier usage where beneficial.

In summary, GEMMbench provides the scaffolding to turn GEMM benchmarking from ad hoc, apples-to-oranges comparisons into a reproducible, community-driven methodology. Through explicit kernel metadata, parameterized pipelines, statistical reporting, and shared repositories, it enables fair evaluation across architectures, data sizes, and data types, and helps surface both optimization techniques (vectorization, coarsening, barriers, work-group sizing) and pitfalls (cache-sensitive regressions, numerical tolerance issues). The initial results on ARM Mali-T628 highlight the outsized impact of cache-aware synchronization on both performance and energy, and exemplify how GEMMbench’s structure makes such insights discoverable, confirmable, and sharable.