Publish

code/.gitignore

@@ -0,0 +1,6 @@
+program
+repl
+construct
+experiment_data
+generate_test_data
+benchmark

code/Makefile

@@ -0,0 +1,31 @@
+CC=gcc
+CPPC=g++
+CFLAGS=-I. -Wextra -Wall -g
+DEPS = fmindex.h util.h
+OBJ = fmindex.o util.o rapl.o
+EXES = program repl construct generate_test_data benchmark
+
+%.o: %.c $(DEPS)
+	$(CC) -c -o $@ $< $(CFLAGS)
+
+all: $(EXES)
+
+program: $(OBJ) program.o
+	$(CC) -o $@ $^ $(CFLAGS)
+
+repl: $(OBJ) repl.o
+	$(CC) -o $@ $^ $(CFLAGS)
+
+construct: $(OBJ) construct.o
+	$(CC) -o $@ $^ $(CFLAGS)
+
+generate_test_data: $(OBJ) generate_test_data.o
+	$(CPPC) -o $@ $^ $(CFLAGS)
+
+benchmark: $(OBJ) benchmark.o
+	$(CPPC) -o $@ $^ $(CFLAGS)
+
+.PHONY: clean all
+
+clean:
+	rm -f *.o *~ core $(EXES)

code/benchmark.c

@@ -0,0 +1,79 @@
+#define _GNU_SOURCE
+
+#include "fmindex.h"
+#include "rapl.h"
+#include "util.h"
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <time.h>
+#include <unistd.h>
+
+// Global variables which can be accessed in function reference.
+unsigned pattern_count, pattern_sz, max_match_count;
+char *patterns;
+fm_index *fm;
+float total_time;
+unsigned long total_matches = 0;
+unsigned long *match_indices;
+
+static void benchmark(void) {
+  float time1 = 0., time2 = 0.;
+  float start_time, end_time;
+  for (unsigned i = 0; i < pattern_count; ++i) {
+    ranges_t start, end;
+    start_time = (float)clock() / CLOCKS_PER_SEC;
+    FMIndexFindMatchRange(fm, &patterns[i * pattern_sz], pattern_sz, &start,
+                          &end);
+    end_time = (float)clock() / CLOCKS_PER_SEC;
+    time1 += end_time - start_time;
+
+    start_time = (float)clock() / CLOCKS_PER_SEC;
+    FMIndexFindRangeIndices(fm, start, end, &match_indices);
+    end_time = (float)clock() / CLOCKS_PER_SEC;
+    time2 += end_time - start_time;
+
+    total_matches += end - start;
+  }
+
+  total_time = time1 + time2;
+}
+
+int main(int argc, char **argv) {
+  if (argc < 3) {
+    fprintf(stderr, "Usage: $ %s <FMFILE> <TESTFILE>\n", argv[0]);
+    return 1;
+  }
+
+  fm = FMIndexReadFromFile(argv[1], 0);
+  if (!fm) {
+    fprintf(stderr, "Failed to read FM-index from file.\n");
+    return 1;
+  }
+
+  if (!(LoadTestData(argv[2], &patterns, &pattern_count, &pattern_sz,
+                     &max_match_count, 0))) {
+    fprintf(stderr, "Could not read test data file.\n");
+    return 1;
+  }
+
+  match_indices = calloc(max_match_count, sizeof(unsigned long));
+  if (!match_indices) {
+    fprintf(stderr, "Failed to allocate memory for match indices.\n");
+    return 1;
+  }
+
+  double total_joules;
+  if (rapl_sysfs(benchmark, &total_joules) != 0) {
+    fprintf(stderr, "Failed to get energy consumption\n");
+    return 1;
+  }
+
+  printf("%a %a %lu\n", total_time, total_joules, total_matches);
+
+  free(match_indices);
+  free(patterns);
+  FMIndexFree(fm);
+  return 0;
+}

code/benchmark_cpu.py

@@ -0,0 +1,64 @@
+import argparse
+import subprocess
+import os
+
+
+def main(repeats, count, maxmatches, lengths, dir, filenames):
+    for filename in filenames:
+        for length in lengths:
+            benchmark(repeats, count, maxmatches, length, dir, filename)
+
+
+def benchmark(repeats, count, maxmatches, length, dir, filename):
+    testfilename = f"{dir}/{filename}.cpu{length}.test"
+    fmfilename = f"{dir}/{filename}.fm"
+    textfilename = f"{dir}/{filename}"
+    resultfilename = f"{dir}/{filename}.cpu{length}.result"
+
+    gentestargs = ["./generate_test_data", textfilename, fmfilename, testfilename, str(count), str(length), str(maxmatches)]
+    benchmarkargs = ["./benchmark", fmfilename, testfilename]
+    print(" ".join(gentestargs))
+    print(" ".join(benchmarkargs))
+
+    # Remove result file if it already exists.
+    try:
+        os.remove(resultfilename)
+    except OSError:
+        pass
+
+    for n in range(repeats):
+        print(f"{n+1}/{repeats}")
+
+        # Create test file.
+        gentestproc = subprocess.Popen(gentestargs, universal_newlines=True, stdout=subprocess.PIPE, stderr=subprocess.PIPE)
+        stdout, stderr = gentestproc.communicate()
+        if stderr:
+            print(f">{stderr.strip()}")
+        ret = gentestproc.poll()
+        if ret != 0:
+            print(f"Error creating test data: {stdout.strip()}")
+            exit(1)
+
+        # Create result file.
+        with open(resultfilename, "a") as resultfile:
+            benchmarkproc = subprocess.Popen(benchmarkargs, stdout=resultfile, universal_newlines=True, stderr=subprocess.PIPE)
+            _, stderr = benchmarkproc.communicate()
+            if stderr:
+                print(f">{stderr.strip()}")
+            ret = benchmarkproc.poll()
+            if ret != 0:
+                print("Error benchmarking test data")
+                exit(1)
+
+
+if __name__ == "__main__":
+    parser = argparse.ArgumentParser()
+    parser.add_argument("-n", "--repeats", help="number of times to repeat each experiment", type=int, required=True)
+    parser.add_argument("-c", "--count", help="number of patterns", type=int, required=True)
+    parser.add_argument("-m", "--maxmatches", help="maximum number of matches per pattern", type=int, required=True)
+    parser.add_argument("-l", "--lengths", help="length of the patterns", type=int, nargs="+", default=[], required=True)
+    parser.add_argument("-d", "--dir", help="directory containing FM-indices and original texts (with the same name)", required=True)
+    parser.add_argument("-f", "--files", help="FM-index files to benchmark", nargs="+", default=[], required=True)
+    args = parser.parse_args()
+
+    main(args.repeats, args.count, args.maxmatches, args.lengths, args.dir, args.files)

code/construct.c

@@ -0,0 +1,30 @@
+#include "fmindex.h"
+#include "util.h"
+
+#include <stdio.h>
+
+int main(int argc, char *argv[]) {
+  if (argc < 3) {
+    printf("Usage: $ %s <INPUTFILE> <OUTPUTFILE>\n", argv[0]);
+    return 1;
+  }
+
+  char *s = ReadFile(argv[1]);
+  if (!s)
+    return 1;
+
+  fm_index *index = FMIndexConstruct(s);
+  if (!index) {
+    printf("Failed to construct index.\n");
+    return 1;
+  }
+
+  if (!FMIndexDumpToFile(index, argv[2])) {
+    printf("Failed to write FM-index to file.\n");
+    return 1;
+  }
+
+  FMIndexFree(index);
+  free(s);
+  return 0;
+}

code/cpu_results/dblp.xml.10MB.cpu4.result

@@ -0,0 +1,5 @@
+0x1.6477fp-2 0x1.22abebp+25 171360982
+0x1.5b44c8p-2 0x1.2b59f9p+25 174285929
+0x1.66d8p-2 0x1.28d3dcp+25 172040584
+0x1.59b5ep-2 0x1.29edabp+25 174000838
+0x1.5f7ep-2 0x1.1f099cp+25 168176711

code/cpu_results/dblp.xml.10MB.cpu6.result

@@ -0,0 +1,5 @@
+0x1.dc032p-3 0x1.885c8ap+24 108982070
+0x1.e32bp-3 0x1.8e6becp+24 110444451
+0x1.e262ap-3 0x1.8d2bf8p+24 110958122
+0x1.f334p-3 0x1.9b5deep+24 115123603
+0x1.ce874p-3 0x1.81151p+24 113346544

code/cpu_results/dblp.xml.10MB.cpu8.result

@@ -0,0 +1,5 @@
+0x1.d3c5p-4 0x1.903b2ap+23 53234288
+0x1.c491cp-4 0x1.86f7dcp+23 51182855
+0x1.cfe28p-4 0x1.8f3564p+23 52348451
+0x1.c04p-4 0x1.800afep+23 50877186
+0x1.f90c8p-4 0x1.b10df8p+23 50214738

code/cpu_results/dblp.xml.15MB.cpu4.result

@@ -0,0 +1,5 @@
+0x1.fb7a8p-2 0x1.a2b59e8p+25 256535338
+0x1.f103ap-2 0x1.9a685cp+25 250410896
+0x1.f7b0fp-2 0x1.a13b028p+25 248846051
+0x1.f6ec6p-2 0x1.a16cb88p+25 253974661
+0x1.ff299p-2 0x1.a58ed2p+25 258322531

code/cpu_results/dblp.xml.15MB.cpu6.result

@@ -0,0 +1,5 @@
+0x1.4d16ap-2 0x1.1620e78p+25 165252248
+0x1.53f52p-2 0x1.1b1bb98p+25 168443910
+0x1.438b4p-2 0x1.0d4fcd8p+25 160215390
+0x1.4b71ep-2 0x1.12ba34p+25 163605068
+0x1.4fafcp-2 0x1.1698d4p+25 166128629

code/cpu_results/dblp.xml.15MB.cpu8.result

@@ -0,0 +1,5 @@
+0x1.41d8ap-3 0x1.11bfap+24 74552832
+0x1.3eefep-3 0x1.0daaa3p+24 74737159
+0x1.3c48ap-3 0x1.0aea5ap+24 74006790
+0x1.4209ap-3 0x1.12066fp+24 75652221
+0x1.3e1a2p-3 0x1.0f36e4p+24 74332277

code/cpu_results/dblp.xml.20MB.cpu4.result

@@ -0,0 +1,5 @@
+0x1.4045dp-1 0x1.09ff364p+26 319604955
+0x1.45e9cp-1 0x1.0e48e38p+26 325973389
+0x1.4ec6cp-1 0x1.167383cp+26 334037324
+0x1.44e6cp-1 0x1.0b42dd4p+26 324111950
+0x1.494b4p-1 0x1.0ec5e1p+26 329018290

code/cpu_results/dblp.xml.20MB.cpu6.result

@@ -0,0 +1,5 @@
+0x1.ba04ep-2 0x1.6f1d43p+25 215259281
+0x1.b1094p-2 0x1.67563b8p+25 211398266
+0x1.a9b66p-2 0x1.63b7078p+25 211580469
+0x1.a923ap-2 0x1.62cc108p+25 211496190
+0x1.a2bd4p-2 0x1.5db353p+25 209139670

code/cpu_results/dblp.xml.20MB.cpu8.result

@@ -0,0 +1,5 @@
+0x1.9d7fcp-3 0x1.5d2c61p+24 97664968
+0x1.a1b28p-3 0x1.61a48cp+24 99282005
+0x1.8a2ecp-3 0x1.4df333p+24 92818122
+0x1.8d93p-3 0x1.4f55f7p+24 94123638
+0x1.9f80cp-3 0x1.5ca19ep+24 99177125

code/cpu_results/dna.10MB.cpu4.result

@@ -0,0 +1,5 @@
+0x1.845828p-1 0x1.35f82dp+26 398786955
+0x1.83880cp-1 0x1.3729878p+26 398482564
+0x1.80a4aap-1 0x1.3b15634p+26 394461926
+0x1.82cccep-1 0x1.3dccf9p+26 395082462
+0x1.7d2d2p-1 0x1.3768a68p+26 389967743

code/cpu_results/dna.10MB.cpu6.result

@@ -0,0 +1,5 @@
+0x1.125d58p-4 0x1.e6b25cp+22 30743684
+0x1.1521ecp-4 0x1.ea8d6cp+22 30733306
+0x1.11be7p-4 0x1.e61bbp+22 30733306
+0x1.13663p-4 0x1.eda22p+22 30733306
+0x1.11cec8p-4 0x1.ee1d24p+22 30417546

code/cpu_results/dna.10MB.cpu8.result

@@ -0,0 +1,5 @@
+0x1.117a6p-6 0x1.3a3a3p+21 2970948
+0x1.17e02p-6 0x1.3a5c88p+21 3124607
+0x1.18c4cp-6 0x1.3b61ep+21 3124607
+0x1.11e68p-6 0x1.344a18p+21 2892252
+0x1.0cb7p-6 0x1.37acp+21 2892252

code/cpu_results/dna.15MB.cpu4.result

@@ -0,0 +1,5 @@
+0x1.220476p+0 0x1.dccf43cp+26 593315165
+0x1.22aa62p+0 0x1.dcfa4c8p+26 592263458
+0x1.22da9p+0 0x1.dd50e78p+26 593423513
+0x1.1ff9cp+0 0x1.d815f2cp+26 590120821
+0x1.23b0fcp+0 0x1.ddc499p+26 596923038

code/cpu_results/dna.15MB.cpu6.result

@@ -0,0 +1,5 @@
+0x1.8e7eep-4 0x1.5d36a2p+23 45834923
+0x1.90183p-4 0x1.5d6cfep+23 46403238
+0x1.8ed448p-4 0x1.5d5062p+23 46403238
+0x1.8d61c8p-4 0x1.5b89p+23 46113391
+0x1.8d3038p-4 0x1.5b6894p+23 46485028

code/cpu_results/dna.15MB.cpu8.result

@@ -0,0 +1,5 @@
+0x1.6c464p-6 0x1.82f5b8p+21 4495841
+0x1.4d9eep-6 0x1.6b5b4p+21 4544987
+0x1.4460ep-6 0x1.63f73p+21 4407471
+0x1.4afa4p-6 0x1.69936p+21 4407471
+0x1.4609ep-6 0x1.64212p+21 4295384

code/cpu_results/dna.20MB.cpu4.result

@@ -0,0 +1,5 @@
+0x1.6d8ecp+0 0x1.2f4f214p+27 746010747
+0x1.71c79cp+0 0x1.2f301bp+27 749721648
+0x1.6ec31cp+0 0x1.2f3a726p+27 747612561
+0x1.70d87p+0 0x1.2edfc9ep+27 750618917
+0x1.6ec052p+0 0x1.2f4e6a2p+27 748739834

code/cpu_results/dna.20MB.cpu6.result

@@ -0,0 +1,5 @@
+0x1.0299ep-3 0x1.be91e4p+23 61438240
+0x1.00d65p-3 0x1.b9f4c6p+23 61287159
+0x1.05935p-3 0x1.c33dccp+23 62472530
+0x1.065c2p-3 0x1.bc9bc8p+23 62363451
+0x1.045de8p-3 0x1.b96f42p+23 61764726

code/cpu_results/dna.20MB.cpu8.result

@@ -0,0 +1,5 @@
+0x1.7bfacp-6 0x1.8d7cdp+21 5743186
+0x1.86b58p-6 0x1.928bd8p+21 6072658
+0x1.9694p-6 0x1.9fdc58p+21 6469630
+0x1.8e7ccp-6 0x1.998cdp+21 6232401
+0x1.8481cp-6 0x1.94466p+21 5901980

code/cpu_results/proteins.10MB.cpu4.result

@@ -0,0 +1,5 @@
+0x1.39b24p-6 0x1.5654cp+21 4815936
+0x1.371e4p-6 0x1.556278p+21 4696974
+0x1.3561p-6 0x1.50c5d8p+21 4712326
+0x1.3b6ccp-6 0x1.5829e8p+21 4915309
+0x1.3af18p-6 0x1.503c8p+21 4856245

code/cpu_results/proteins.10MB.cpu6.result

@@ -0,0 +1,5 @@
+0x1.118a8p-6 0x1.38894p+21 2572993
+0x1.15a98p-6 0x1.31103p+21 2748401
+0x1.15968p-6 0x1.33adbp+21 2761448
+0x1.18ca8p-6 0x1.364d08p+21 2771358
+0x1.19dbcp-6 0x1.361f5p+21 2807302

code/cpu_results/proteins.10MB.cpu8.result

@@ -0,0 +1,5 @@
+0x1.248ccp-6 0x1.309438p+21 2028498
+0x1.1ba54p-6 0x1.3654bp+21 1966015
+0x1.2766p-6 0x1.43da8p+21 2049837
+0x1.238fcp-6 0x1.4395d8p+21 2008116
+0x1.25804p-6 0x1.36f11p+21 2061600

code/cpu_results/proteins.15MB.cpu4.result

@@ -0,0 +1,5 @@
+0x1.5bc68p-6 0x1.711bcp+21 5783326
+0x1.6099p-6 0x1.7d223p+21 5898449
+0x1.66854p-6 0x1.7739ap+21 6078789
+0x1.6189cp-6 0x1.787828p+21 6050872
+0x1.6b70cp-6 0x1.7dd76p+21 6253011

code/cpu_results/proteins.15MB.cpu6.result

@@ -0,0 +1,5 @@
+0x1.263dp-6 0x1.490a08p+21 2963082
+0x1.2396p-6 0x1.4304e8p+21 2984599
+0x1.298f8p-6 0x1.4e9adp+21 3180703
+0x1.27a28p-6 0x1.4a6158p+21 3016497
+0x1.1fd18p-6 0x1.3c4108p+21 2913659

code/cpu_results/proteins.15MB.cpu8.result

@@ -0,0 +1,5 @@
+0x1.36784p-6 0x1.635e9p+21 2393193
+0x1.3a3a4p-6 0x1.51c57p+21 2397301
+0x1.38688p-6 0x1.572878p+21 2449818
+0x1.35684p-6 0x1.44aa68p+21 2267396
+0x1.3395cp-6 0x1.53db88p+21 2294787

code/cpu_results/proteins.20MB.cpu4.result

@@ -0,0 +1,5 @@
+0x1.6376p-6 0x1.79c41p+21 5720934
+0x1.620cp-6 0x1.748edp+21 5745857
+0x1.9e548p-6 0x1.a69888p+21 6040126
+0x1.6031p-6 0x1.7a4d68p+21 5780235
+0x1.6ee28p-6 0x1.888e2p+21 6174243

code/cpu_results/proteins.20MB.cpu6.result

@@ -0,0 +1,5 @@
+0x1.1d6a8p-6 0x1.378d78p+21 2639463
+0x1.16ddp-6 0x1.37a828p+21 2403897
+0x1.1bec8p-6 0x1.379518p+21 2522317
+0x1.1828p-6 0x1.37952p+21 2494930
+0x1.1b8dp-6 0x1.3dbc9p+21 2490305

code/cpu_results/proteins.20MB.cpu8.result

@@ -0,0 +1,5 @@
+0x1.2e9d8p-6 0x1.4bb4dp+21 1915812
+0x1.2c43p-6 0x1.442aap+21 1825102
+0x1.2931p-6 0x1.440ffp+21 1824534
+0x1.31928p-6 0x1.440ep+21 1962021
+0x1.2cef8p-6 0x1.483648p+21 1863384

code/fmindex.c

@@ -0,0 +1,345 @@
+#define _GNU_SOURCE
+
+#include "fmindex.h"
+#include "util.h"
+
+#include <limits.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+inline static int string_index(char *s, char c) { return strchr(s, c) - s; }
+
+static int CompareChar(const void *a, const void *b) {
+  char i = *(char *)a;
+  char j = *(char *)b;
+
+  // Dollar sign is always sorted first.
+  if (i == '$')
+    return -1;
+  if (j == '$')
+    return 1;
+
+  return i > j;
+}
+
+/* Converts the given string to a newly allocated string
+ *  of all distinct characters, sorted in lexicographical order.
+ * The dollar sign is always sorted first.
+ * Return NULL on memory allocation error.
+ */
+char *TextToAlphabet(char *text, size_t sz) {
+  char *alphabet = calloc(2, sizeof(char));
+  if (!alphabet)
+    return NULL;
+
+  alphabet[0] = '$';
+  alphabet[1] = '\0';
+  int len = 2;
+  for (size_t i = 0; i < sz; ++i) {
+    if (!strchr(alphabet, text[i])) {
+      len += 1;
+      alphabet = realloc(alphabet, len);
+      alphabet[len - 2] = text[i];
+      alphabet[len - 1] = '\0';
+    }
+  }
+
+  qsort(alphabet, len - 1, sizeof(char), &CompareChar);
+
+  return alphabet;
+}
+
+static int CompareSuffixArray(const void *a, const void *b, void *arg) {
+  sa_t i = *(sa_t *)a;
+  sa_t j = *(sa_t *)b;
+  char *s = arg;
+
+  while (s[i] != '\0' && s[j] != '\0' && s[i] == s[j]) {
+    ++i;
+    ++j;
+  }
+
+  if (s[i] == '\0')
+    return -1;
+  if (s[j] == '\0')
+    return 1;
+
+  return s[i] > s[j];
+}
+
+/* Construct a newly allocated suffix array for the given string.
+ * The suffix array holds the index for each suffix in the given string.
+ * Then the indices are sorted by the lexicographical ordering of the suffices.
+ * Return NULL on memory allocation error.
+ */
+sa_t *ConstructSuffixArray(char *s, size_t sz) {
+  sa_t *suffix_array = calloc(sz + 1, sizeof(sa_t));
+  if (!suffix_array)
+    return NULL;
+
+  for (size_t i = 0; i < sz + 1; ++i)
+    suffix_array[i] = i;
+
+  qsort_r(suffix_array, sz + 1, sizeof(sa_t), &CompareSuffixArray, s);
+
+  return suffix_array;
+}
+
+/* Construct the Burrows-Wheeler transform of the given string,
+ *  using the corresponding suffix array.
+ * Return NULL on memory allocation error.
+ */
+char *ConstructBWT(char *s, size_t sz, sa_t *suffix_array) {
+  char *bwt = calloc(sz + 2, sizeof(char));
+  if (!bwt)
+    return NULL;
+
+  for (size_t i = 0; i < sz + 1; ++i) {
+    sa_t n = suffix_array[i];
+    // Index 0 is always the dollar sign.
+    bwt[i] = (n) ? s[suffix_array[i] - 1] : '$';
+  }
+  bwt[sz + 1] = '\0';
+
+  return bwt;
+}
+
+/* Construct the rank matrix for the given Burrows-Wheeler transformed string.
+ * The rank matrix is an (alphabet X bwt) sized array that holds the accumulated
+ *  counts of encountered characters in the BWT.
+ * A newly allocated rank matrix is returned, or NULL on memory error.
+ */
+ranks_t *ConstructRankMatrix(char *bwt, size_t sz, char *alphabet) {
+  size_t alphabet_sz = strlen(alphabet);
+
+  ranks_t *rank_matrix = calloc(sz * alphabet_sz, sizeof(ranks_t));
+  if (!rank_matrix) {
+    printf("Failed to allocate %lu bytes.\n", sz * alphabet_sz);
+    return NULL;
+  }
+
+  ranks_t acc[alphabet_sz];
+
+  // Initialize counts with zero.
+  for (size_t i = 0; i < alphabet_sz; ++i)
+    acc[i] = 0;
+
+  for (size_t i = 0; i < sz; ++i) {
+    char c = bwt[i];
+
+    // Update accumulator.
+    for (unsigned j = 0; j < alphabet_sz; ++j)
+      if (c == alphabet[j]) {
+        ++acc[j];
+        break;
+      }
+
+    for (unsigned j = 0; j < alphabet_sz; ++j)
+      rank_matrix[i * alphabet_sz + j] = acc[j];
+  }
+
+  return rank_matrix;
+}
+
+/* Calculate the ranges in the "F column" where each character
+ *  in the alphabet appears.
+ * Return a (2 X alphabet) matrix where each row means the starting
+ *  and ending range for a character.
+ * Return NULL on memory error.
+ */
+ranges_t *ConstructCharacterRanges(char *bwt, size_t sz, char *alphabet) {
+  size_t alphabet_sz = strlen(alphabet);
+  unsigned long counts[alphabet_sz];
+
+  // Set counts to zero.
+  for (size_t i = 0; i < alphabet_sz; ++i)
+    counts[i] = 0;
+
+  // Count total amounts of characters.
+  for (size_t i = 0; i < sz; ++i)
+    counts[string_index(alphabet, bwt[i])]++;
+
+  // Accumulate counts.
+  size_t acc = 0;
+  for (size_t i = 0; i < alphabet_sz; ++i) {
+    counts[i] += acc;
+    acc = counts[i];
+  }
+
+  // Calculate ranges.
+  ranges_t *ranges = calloc(2 * alphabet_sz, sizeof(ranges_t));
+  if (!ranges)
+    return NULL;
+
+  unsigned long cur_idx = 0;
+  for (size_t i = 0; i < alphabet_sz; ++i) {
+    ranges[2 * i] = cur_idx;
+    ranges[2 * i + 1] = counts[i];
+    cur_idx = counts[i];
+  }
+
+  return ranges;
+}
+
+/* Find the range of matches for the given pattern in the F column of the
+ *  given FM-index.
+ */
+void FMIndexFindMatchRange(fm_index *fm, char *pattern, size_t pattern_sz,
+                           ranges_t *start, ranges_t *end) {
+  int p_idx = pattern_sz - 1;
+  char c = pattern[p_idx];
+  // Initial range is all instances of the last character in pattern.
+  *start = fm->ranges[2 * string_index(fm->alphabet, c)];
+  *end = fm->ranges[2 * string_index(fm->alphabet, c) + 1];
+
+  p_idx -= 1;
+  while (p_idx >= 0 && *end > 1) {
+    c = pattern[p_idx];
+    ranges_t range_start = fm->ranges[2 * string_index(fm->alphabet, c)];
+    int alphabet_idx = string_index(fm->alphabet, c);
+    *start =
+        range_start + fm->ranks[fm->alphabet_sz * (*start - 1) + alphabet_idx];
+    *end = range_start + fm->ranks[fm->alphabet_sz * (*end - 1) + alphabet_idx];
+    p_idx -= 1;
+  }
+}
+
+/* Find the matching indices in the original text for the given
+ *  range in the "F column" of the Burrows-Wheeler matrix.
+ */
+void FMIndexFindRangeIndices(fm_index *fm, ranges_t start, ranges_t end,
+                             unsigned long **match_indices) {
+  for (unsigned long i = 0; i < end - start; ++i)
+    (*match_indices)[i] = fm->sa[start + i];
+}
+
+fm_index *FMIndexConstruct(char *s) {
+  fm_index *index = malloc(sizeof(fm_index));
+  if (!index)
+    return NULL;
+  memset(index, 0, sizeof(fm_index));
+
+  size_t sz = strlen(s);
+  if (!(index->alphabet = TextToAlphabet(s, sz)))
+    goto error;
+  index->alphabet_sz = strlen(index->alphabet);
+  if (!(index->sa = ConstructSuffixArray(s, sz)))
+    goto error;
+  if (!(index->bwt = ConstructBWT(s, sz, index->sa)))
+    goto error;
+  ++sz; // Because of the added dollar sign.
+  index->bwt_sz = sz;
+  if (!(index->ranks = ConstructRankMatrix(index->bwt, sz, index->alphabet)))
+    goto error;
+  if (!(index->ranges =
+            ConstructCharacterRanges(index->bwt, sz, index->alphabet)))
+    goto error;
+
+  return index;
+
+error:
+  if (index->alphabet)
+    free(index->alphabet);
+  if (index->sa)
+    free(index->sa);
+  if (index->bwt)
+    free(index->bwt);
+  if (index->ranks)
+    free(index->ranks);
+  if (index->ranges)
+    free(index->ranges);
+
+  free(index);
+  return NULL;
+}
+
+void FMIndexFree(fm_index *index) {
+  free(index->alphabet);
+  free(index->sa);
+  free(index->bwt);
+  free(index->ranks);
+  free(index->ranges);
+  free(index);
+}
+
+int FMIndexDumpToFile(fm_index *index, char *filename) {
+  FILE *f = fopen(filename, "w");
+  if (!f)
+    return 0;
+
+  fwrite(&index->bwt_sz, sizeof(index->bwt_sz), 1, f);
+  fwrite(index->bwt, sizeof(char), index->bwt_sz, f);
+  fwrite(&index->alphabet_sz, sizeof(index->alphabet_sz), 1, f);
+  fwrite(index->alphabet, sizeof(char), index->alphabet_sz, f);
+  fwrite(index->ranges, sizeof(ranges_t), 2 * index->alphabet_sz, f);
+  fwrite(index->ranks, sizeof(ranks_t), index->bwt_sz * index->alphabet_sz, f);
+  fwrite(index->sa, sizeof(sa_t), index->bwt_sz, f);
+
+  fclose(f);
+  return 1;
+}
+
+fm_index *FMIndexReadFromFile(char *filename, int aligned) {
+  FILE *f = fopen(filename, "r");
+  if (!f)
+    return NULL;
+
+  fm_index *index;
+
+  if (!(index = calloc(1, sizeof(fm_index))))
+    goto error;
+
+  fread(&index->bwt_sz, sizeof(index->bwt_sz), 1, f);
+
+  if (!MaybeMallocAligned((void **)&index->bwt,
+                          (index->bwt_sz + 1) * sizeof(char), aligned))
+    goto error;
+  fread(index->bwt, sizeof(char), index->bwt_sz, f);
+  index->bwt[index->bwt_sz] = '\0';
+
+  fread(&index->alphabet_sz, sizeof(index->alphabet_sz), 1, f);
+
+  if (!MaybeMallocAligned((void **)&index->alphabet,
+                          (index->alphabet_sz + 1) * sizeof(char), aligned))
+    goto error;
+  fread(index->alphabet, sizeof(char), index->alphabet_sz, f);
+  index->alphabet[index->alphabet_sz] = '\0';
+
+  if (!MaybeMallocAligned((void **)&index->ranges,
+                          2 * index->alphabet_sz * sizeof(ranges_t), aligned))
+    goto error;
+  fread(index->ranges, sizeof(ranges_t), 2 * index->alphabet_sz, f);
+
+  if (!MaybeMallocAligned((void **)&index->ranks,
+                          index->bwt_sz * index->alphabet_sz * sizeof(ranks_t),
+                          aligned))
+    goto error;
+  fread(index->ranks, sizeof(ranks_t), index->bwt_sz * index->alphabet_sz, f);
+
+  if (!MaybeMallocAligned((void **)&index->sa, index->bwt_sz * sizeof(sa_t),
+                          aligned))
+    goto error;
+  fread(index->sa, sizeof(sa_t), index->bwt_sz, f);
+
+  fclose(f);
+  return index;
+
+error:
+  if (index) {
+    if (index->bwt)
+      free(index->bwt);
+    if (index->alphabet)
+      free(index->alphabet);
+    if (index->ranges)
+      free(index->ranges);
+    if (index->ranks)
+      free(index->ranks);
+    if (index->sa)
+      free(index->sa);
+    free(index);
+  }
+
+  fclose(f);
+  return NULL;
+}

code/fmindex.h

@@ -0,0 +1,35 @@
+#pragma once
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#include <stdlib.h>
+
+typedef unsigned ranges_t;
+typedef unsigned ranks_t;
+typedef unsigned sa_t;
+typedef struct fm_index {
+  char *bwt;
+  size_t bwt_sz;
+  char *alphabet;
+  size_t alphabet_sz;
+  ranks_t *ranks;
+  sa_t *sa;
+  ranges_t *ranges;
+} fm_index;
+
+fm_index *FMIndexConstruct(char *s);
+void FMIndexFree(fm_index *index);
+
+fm_index *FMIndexReadFromFile(char *filename, int aligned);
+int FMIndexDumpToFile(fm_index *index, char *filename);
+
+void FMIndexFindMatchRange(fm_index *fm, char *pattern, size_t pattern_sz,
+                           ranges_t *start, ranges_t *end);
+void FMIndexFindRangeIndices(fm_index *fm, ranges_t start, ranges_t end,
+                             unsigned long **match_indices);
+
+#ifdef __cplusplus
+}
+#endif

code/fpga/.gitignore

@@ -0,0 +1,13 @@
+verify
+sw_emu
+hw_emu
+hw
+.Xil
+.run
+emconfig.json
+vitis_env.sh
+.ipcache
+example_output
+test.txt
+*.jou
+*.str

code/fpga/Makefile

@@ -0,0 +1,106 @@
+ifndef PLATFORM
+    $(error PLATFORM is not set.)
+endif
+
+ifndef TARGET
+    $(info TARGET is not set; using sw_emu.)
+endif
+
+TARGET ?= sw_emu
+
+TARGETS := sw_emu hw_emu hw
+
+ifeq ($(filter $(TARGET),$(TARGETS)),)
+    $(error TARGET can only be one of: "$(TARGETS)")
+endif
+
+VXXFLAGS := -t ${TARGET} --log_dir $(TARGET) --report_dir $(TARGET) --temp_dir $(TARGET) -I/usr/include/x86_64-linux-gnu -Wno-unused-label
+GXXFLAGS := -Wall -g -std=c++11 -I${XILINX_XRT}/include/ -L${XILINX_XRT}/lib/ -lOpenCL -lpthread -lrt -lstdc++ -I..
+PROJ_HEADERS := ../fmindex.h ../util.h
+PROJ_OBJS := ../fmindex.o ../util.o
+
+ifeq ($(TARGET), hw)
+	EMULATION_FLAG :=
+else
+	EMULATION_FLAG := XCL_EMULATION_MODE=$(TARGET)
+endif
+
+.PHONY: all run-verify clean cleanall unopt opt
+
+all: $(TARGET) verify benchmark unopt opt
+
+unopt: $(TARGET)/unopt.xclbin
+
+ndrange: $(TARGET)/ndrange.xclbin
+
+ndrange2: $(TARGET)/ndrange2.xclbin
+
+memory: $(TARGET)/memory.xclbin
+
+final: $(TARGET)/final.xclbin
+
+run-verify: verify $(TARGET)/$(KERNEL).xclbin
+	@test -n "$(FMFILE)" || (echo "FMFILE undefined" ; exit 1)
+	@test -n "$(TESTFILE)" || (echo "TESTFILE undefined" ; exit 1)
+	@test -n "$(KERNEL)" || (echo "KERNEL undefined" ; exit 1)
+	@test -n "$(NDRANGE)" || (echo "NDRANGE undefined" ; exit 1)
+	@test -n "$(LOCALSIZE)" || (echo "LOCALSIZE undefined" ; exit 1)
+	cd $(TARGET) && $(EMULATION_FLAG) ../verify ../$(FMFILE) ../$(TARGET)/$(KERNEL).xclbin ../$(TESTFILE) $(NDRANGE) $(LOCALSIZE)
+
+$(TARGET):
+	mkdir $(TARGET)
+
+%.o: %.cpp $(PROJ_HEADERS)
+	g++ -c -o $@ $< $(GXXFLAGS)
+
+verify: verify.o
+	g++ -o $@ $(PROJ_OBJS) $< $(GXXFLAGS)
+
+benchmark: benchmark.o
+	g++ -o $@ $(PROJ_OBJS) $< $(GXXFLAGS)
+
+$(TARGET)/unopt.xo: unopt.cl no_ndrange.cfg
+	v++ -c -k fmindex $< $(VXXFLAGS) --config no_ndrange.cfg -o $@
+
+$(TARGET)/unopt.xclbin: $(TARGET)/unopt.xo emconfig.json
+	v++ -l $< $(VXXFLAGS) --config no_ndrange.cfg -o $@
+	mv -t $(TARGET) xrc.log xcd.log
+
+$(TARGET)/memory.xo: memory.cl no_ndrange.cfg
+	v++ -c -k fmindex $< $(VXXFLAGS) --config no_ndrange.cfg -o $@
+
+$(TARGET)/memory.xclbin: $(TARGET)/memory.xo emconfig.json
+	v++ -l $< $(VXXFLAGS) --config no_ndrange.cfg -o $@
+	mv -t $(TARGET) xrc.log xcd.log
+
+$(TARGET)/ndrange.xo: ndrange.cl ndrange.cfg
+	v++ -c -k fmindex $< $(VXXFLAGS) --config ndrange.cfg -o $@
+
+$(TARGET)/ndrange.xclbin: $(TARGET)/ndrange.xo emconfig.json
+	v++ -l $< $(VXXFLAGS) --config ndrange.cfg -o $@
+	mv -t $(TARGET) xrc.log xcd.log
+
+$(TARGET)/ndrange2.xo: ndrange2.cl ndrange2.cfg
+	v++ -c -k fmindex $< $(VXXFLAGS) --config ndrange2.cfg -o $@
+
+$(TARGET)/ndrange2.xclbin: $(TARGET)/ndrange2.xo emconfig.json
+	v++ -l $< $(VXXFLAGS) --config ndrange2.cfg -o $@
+	mv -t $(TARGET) xrc.log xcd.log
+
+$(TARGET)/final.xo: final.cl ndrange.cfg
+	v++ -c -k fmindex $< $(VXXFLAGS) --config ndrange.cfg -o $@
+
+$(TARGET)/final.xclbin: $(TARGET)/final.xo emconfig.json
+	v++ -l $< $(VXXFLAGS) --config ndrange.cfg -o $@
+	mv -t $(TARGET) xrc.log xcd.log
+
+emconfig.json:
+	emconfigutil --platform $(PLATFORM) --nd 1
+
+clean:
+	rm -rf emconfig.json *.info *.link_summary *.compile_summary \
+	    *.xclbin *.xo *.log _x $(TARGET) verify TempConfig *.csv *.run_summary \
+	    *.o .Xil .run .ipcache
+
+cleanall: clean
+	rm -rf $(TARGETS)

code/fpga/benchmark.cpp

@@ -0,0 +1,171 @@
+#define CL_HPP_CL_1_2_DEFAULT_BUILD
+#define CL_HPP_TARGET_OPENCL_VERSION 120
+#define CL_HPP_MINIMUM_OPENCL_VERSION 120
+#define CL_HPP_ENABLE_PROGRAM_CONSTRUCTION_FROM_ARRAY_COMPATIBILITY 1
+#define CL_USE_DEPRECATED_OPENCL_1_2_APIS
+
+#include <CL/cl2.hpp>
+#include <fstream>
+#include <iostream>
+#include <stdlib.h>
+#include <string.h>
+#include <time.h>
+#include <unistd.h>
+#include <vector>
+
+#include "../fmindex.h"
+#include "../util.h"
+
+std::vector<cl::Device> get_xilinx_devices();
+char *read_binary_file(const std::string &xclbin_file_name, unsigned &nb);
+
+int main(int argc, char **argv) {
+  if (argc < 6) {
+    fprintf(
+        stderr,
+        "Usage: %s <FMINDEXFILE> <XCLBIN> <TESTFILE> <NDRANGE> <LOCALSIZE>\n",
+        argv[0]);
+    return 1;
+  }
+
+  int use_ndrange = atoi(argv[4]);
+  int local_size = atoi(argv[5]);
+
+  // Load FM-index.
+  fm_index *index = FMIndexReadFromFile(argv[1], 1);
+  if (!index) {
+    fprintf(stderr, "Could not read FM-index from file.\n");
+    return 1;
+  }
+
+  // Load test file.
+  unsigned pattern_count, pattern_sz, max_match_count;
+  char *patterns;
+
+  if (!(LoadTestData(argv[3], &patterns, &pattern_count, &pattern_sz,
+                     &max_match_count, 1))) {
+    fprintf(stderr, "Could not read test data file.\n");
+    return 1;
+  }
+
+  // Initialize OpenCL.
+  cl_int err;
+  unsigned fileBufSize;
+  std::vector<cl::Device> devices = get_xilinx_devices();
+  devices.resize(1);
+  cl::Device device = devices[0];
+  cl::Context context(device, NULL, NULL, NULL, &err);
+  char *fileBuf = read_binary_file(argv[2], fileBufSize);
+  cl::Program::Binaries bins{{fileBuf, fileBufSize}};
+  cl::Program program(context, devices, bins, NULL, &err);
+  cl::CommandQueue q(context, device, CL_QUEUE_PROFILING_ENABLE, &err);
+  cl::Kernel kernel(program, "fmindex", &err);
+
+  // Create OpenCL buffers for host data.
+  cl::Buffer bwt_buf(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
+                     sizeof(char) * index->bwt_sz, index->bwt, &err);
+  cl::Buffer alphabet_buf(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
+                          sizeof(char) * index->alphabet_sz, index->alphabet,
+                          &err);
+  cl::Buffer ranks_buf(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
+                       sizeof(ranks_t) * index->bwt_sz * index->alphabet_sz,
+                       index->ranks, &err);
+  cl::Buffer sa_buf(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
+                    sizeof(sa_t) * index->bwt_sz, index->sa, &err);
+  cl::Buffer ranges_buf(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
+                        sizeof(ranges_t) * 2 * index->alphabet_sz,
+                        index->ranges, &err);
+  cl::Buffer patterns_buf(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
+                          sizeof(char) * pattern_count * pattern_sz, patterns,
+                          &err);
+  cl::Buffer out_buf(context, CL_MEM_WRITE_ONLY,
+                     sizeof(unsigned long) *
+                         (pattern_count * (max_match_count + 1)),
+                     NULL, &err);
+
+  // Map OpenCL buffers to kernel arguments.
+  kernel.setArg(0, bwt_buf);
+  kernel.setArg(1, alphabet_buf);
+  kernel.setArg(2, ranks_buf);
+  kernel.setArg(3, sa_buf);
+  kernel.setArg(4, ranges_buf);
+  kernel.setArg(5, patterns_buf);
+  kernel.setArg(6, out_buf);
+
+  // Map OpenCL buffers to kernel arguments including scalars.
+  kernel.setArg(0, bwt_buf);
+  kernel.setArg(1, alphabet_buf);
+  kernel.setArg(2, ranks_buf);
+  kernel.setArg(3, sa_buf);
+  kernel.setArg(4, ranges_buf);
+  kernel.setArg(5, patterns_buf);
+  kernel.setArg(6, out_buf);
+  kernel.setArg(7, index->bwt_sz);
+  kernel.setArg(8, index->alphabet_sz);
+  kernel.setArg(9, pattern_count);
+  kernel.setArg(10, pattern_sz);
+  kernel.setArg(11, max_match_count + 1);
+
+  // Schedule transfer of inputs and output.
+  q.enqueueMigrateMemObjects(
+      {bwt_buf, alphabet_buf, ranks_buf, sa_buf, ranges_buf, patterns_buf}, 0);
+
+  if (use_ndrange)
+    q.enqueueNDRangeKernel(kernel, 0, pattern_count, local_size);
+  else
+    q.enqueueTask(kernel);
+
+  q.enqueueMigrateMemObjects({out_buf}, CL_MIGRATE_MEM_OBJECT_HOST);
+
+  q.finish();
+
+  // Release resources.
+  delete[] fileBuf;
+  FMIndexFree(index);
+  return EXIT_SUCCESS;
+}
+
+// ------------------------------------------------------------------------------------
+// Utility functions
+// ------------------------------------------------------------------------------------
+std::vector<cl::Device> get_xilinx_devices() {
+  size_t i;
+  cl_int err;
+  std::vector<cl::Platform> platforms;
+  err = cl::Platform::get(&platforms);
+  cl::Platform platform;
+  for (i = 0; i < platforms.size(); i++) {
+    platform = platforms[i];
+    std::string platformName = platform.getInfo<CL_PLATFORM_NAME>(&err);
+    if (platformName == "Xilinx") {
+      std::cerr << "INFO: Found Xilinx Platform" << std::endl;
+      break;
+    }
+  }
+  if (i == platforms.size()) {
+    std::cerr << "ERROR: Failed to find Xilinx platform" << std::endl;
+    exit(EXIT_FAILURE);
+  }
+
+  // Getting ACCELERATOR Devices and selecting 1st such device
+  std::vector<cl::Device> devices;
+  err = platform.getDevices(CL_DEVICE_TYPE_ACCELERATOR, &devices);
+  return devices;
+}
+
+char *read_binary_file(const std::string &xclbin_file_name, unsigned &nb) {
+  if (access(xclbin_file_name.c_str(), R_OK) != 0) {
+    printf("ERROR: %s xclbin not available please build\n",
+           xclbin_file_name.c_str());
+    exit(EXIT_FAILURE);
+  }
+  // Loading XCL Bin into char buffer
+  std::cerr << "INFO: Loading '" << xclbin_file_name << "'\n";
+  std::ifstream bin_file(xclbin_file_name.c_str(), std::ifstream::binary);
+  bin_file.seekg(0, bin_file.end);
+  nb = bin_file.tellg();
+  bin_file.seekg(0, bin_file.beg);
+  char *buf = new char[nb];
+  bin_file.read(buf, nb);
+  return buf;
+}

code/fpga/benchmark_fpga.py

@@ -0,0 +1,232 @@
+import argparse
+import subprocess
+import os
+import shutil
+import json
+from pathlib import Path
+
+
+# Thanks Tristan Laan for data extraction functions.
+
+
+def _calculate_wattage(data: list) -> float:
+    return (int(data[1]) / 1000) * (int(data[2]) / 1000) \
+        + (int(data[3]) / 1000) * (int(data[4]) / 1000)
+
+
+def _get_power_profile_data(directory: Path) -> list:
+    file = next(directory.glob('power_profile_*.csv'))
+    data = []
+
+    with file.open() as f:
+        for i, line in enumerate(f):
+            if i < 2:
+                continue
+            csv_data = line.split(',')
+            data.append({'timestamp': float(csv_data[0]),
+                         'power': _calculate_wattage(csv_data)})
+
+    return data
+
+
+def _get_timeline_data(directory: Path) -> dict:
+    file = next(directory.glob('timeline_trace.csv'))
+    data = {}
+
+    with file.open() as f:
+        for i, line in enumerate(f):
+            cells = line.split(',')
+            if i < 13:
+                continue
+            if cells[0] == 'Footer':
+                break
+
+            if cells[1].startswith("KERNEL") and cells[1].endswith("all") and cells[2] in ["START", "END"]:
+                data[cells[2]] = cells[0]
+
+    return data
+
+
+def _is_mode(line: str):
+    if line == 'OpenCL API Calls':
+        return True
+    if line == 'Kernel Execution':
+        return True
+    if line == 'Compute Unit Utilization':
+        return True
+    if line == 'Data Transfer: Host to Global Memory':
+        return True
+    if line == 'Data Transfer: Kernels to Global Memory':
+        return True
+
+
+def _get_profile_summary_data(file: Path) -> dict:
+    data = dict()
+    mode = None
+    skip_line = False
+    with file.open() as f:
+        for line in f:
+            line = line.strip()
+            if skip_line:
+                skip_line = False
+                continue
+
+            if _is_mode(line):
+                mode = line
+                data[mode] = []
+                skip_line = True
+                continue
+
+            if line == '':
+                mode = None
+
+            if not mode:
+                continue
+
+            csv_data = line.split(',')
+
+            if mode == 'OpenCL API Calls':
+                data[mode].append({
+                        'name': csv_data[0],
+                        'calls': int(csv_data[1]),
+                        'time': float(csv_data[2])
+                    })
+
+            if mode == 'Kernel Execution':
+                data[mode].append({
+                        'kernel': csv_data[0],
+                        'enqueues': int(csv_data[1]),
+                        'time': float(csv_data[2])
+                    })
+
+            if mode == 'Compute Unit Utilization':
+                data[mode].append({
+                        'cu': csv_data[1],
+                        'kernel': csv_data[2],
+                        'time': float(csv_data[9])
+                    })
+
+            if mode == 'Data Transfer: Host to Global Memory':
+                data[mode].append({
+                        'type': csv_data[1],
+                        'transfers': int(csv_data[2]),
+                        'speed': float(csv_data[3]),
+                        'utilization': float(csv_data[4]),
+                        'size': float(csv_data[5]),
+                        'time': float(csv_data[6])
+                    })
+
+
+            if mode == 'Data Transfer: Kernels to Global Memory':
+                data[mode].append({
+                        'interface': csv_data[3],
+                        'type': csv_data[4],
+                        'transfers': int(csv_data[5]),
+                        'speed': float(csv_data[6]),
+                        'utilization': float(csv_data[7]),
+                        'size': float(csv_data[8])
+                    })
+
+    return data
+
+
+def read_data(directory: Path) -> dict:
+    data = dict()
+    profile = directory / 'profile_summary.csv'
+    try:
+        data['power'] = _get_power_profile_data(directory)
+        data['timeline'] = _get_timeline_data(directory)
+    except StopIteration:
+        pass
+    if profile.exists():
+        data.update(_get_profile_summary_data(profile))
+    return data
+
+
+def write_data(data: dict, file: Path):
+    with file.open('w') as f:
+        json.dump(data, f)
+
+
+class cd:
+    """Context manager for changing the current working directory"""
+    def __init__(self, newPath):
+        self.newPath = os.path.expanduser(newPath)
+
+    def __enter__(self):
+        self.savedPath = os.getcwd()
+        os.chdir(self.newPath)
+
+    def __exit__(self, etype, value, traceback):
+        os.chdir(self.savedPath)
+
+
+def main(repeats, count, maxmatches, lengths, dir, filenames, kernel, target, localsize, ndrange):
+    for filename in filenames:
+        for length in lengths:
+            benchmark(repeats, count, maxmatches, length, dir, filename, kernel, target, localsize, ndrange)
+
+
+def benchmark(repeats, count, maxmatches, length, dir, filename, kernel, target, localsize, ndrange):
+    textfilename = f"../{dir}/{filename}"
+    fmfilename = f"{textfilename}.fm"
+    resultdir = f"../{kernel}_results/{filename}.len{length}"
+    testfilename = f"{resultdir}/test.txt"
+
+    if target != "hw":
+        env = dict(os.environ, XCL_EMULATION_MODE=target)
+    else:
+        env = os.environ
+
+    gentestargs = ["../../generate_test_data", textfilename, fmfilename, testfilename, str(count), str(length), str(maxmatches)]
+    benchmarkargs = ["../benchmark", fmfilename, f"{kernel}.xclbin", testfilename, str(ndrange), str(localsize)]
+    print(" ".join(gentestargs))
+    print(" ".join(benchmarkargs))
+
+    with cd(target):
+        shutil.rmtree(resultdir, ignore_errors=True)
+        Path(resultdir).mkdir(parents=True, exist_ok=True)
+
+        for n in range(repeats):
+            print(f"{n+1}/{repeats}")
+
+            # Create test file.
+            gentestproc = subprocess.Popen(gentestargs, universal_newlines=True, stdout=subprocess.PIPE, stderr=subprocess.PIPE)
+            stdout, stderr = gentestproc.communicate()
+            if stderr:
+                print(f">{stderr.strip()}")
+            ret = gentestproc.poll()
+            if ret != 0:
+                print(f"Error creating test data: {stdout.strip()}")
+                exit(1)
+
+            # Execute kernel
+            benchmarkproc = subprocess.Popen(benchmarkargs, stdout=subprocess.PIPE, universal_newlines=True, stderr=subprocess.PIPE, env=env)
+            _, stderr = benchmarkproc.communicate()
+            if stderr:
+                print(f">{stderr.strip()}")
+            ret = benchmarkproc.poll()
+            if ret != 0:
+                print("Error benchmarking test data")
+                exit(1)
+
+            # Move data to result directory.
+            data = read_data(Path("."))
+            write_data(data, Path(resultdir) / f"run{n}.json")
+
+
+if __name__ == "__main__":
+    parser = argparse.ArgumentParser()
+    parser.add_argument("-n", "--repeats", help="number of times to repeat each experiment", type=int, required=True)
+    parser.add_argument("-c", "--count", help="number of patterns", type=int, required=True)
+    parser.add_argument("-m", "--maxmatches", help="maximum number of matches per pattern", type=int, required=True)
+    parser.add_argument("-l", "--lengths", help="length of the patterns", type=int, nargs="+", default=[], required=True)
+    parser.add_argument("-d", "--dir", help="directory containing FM-indices and original texts (with the same name)", required=True)
+    parser.add_argument("-f", "--files", help="FM-index files to benchmark", nargs="+", default=[], required=True)
+    parser.add_argument("-k", "--kernel", help="FPGA kernel", required=True)
+    parser.add_argument("-t", "--target", help="compilation target", default="sw_emu", required=False)
+    parser.add_argument("-s", "--localsize", help="NDRange local size", type=int, required=True)
+    parser.add_argument("-r", "--ndrange", help="use NDRange", type=int, required=True)
+    args = parser.parse_args()
+
+    main(args.repeats, args.count, args.maxmatches, args.lengths, args.dir, args.files, args.kernel, args.target, args.localsize, args.ndrange)

code/fpga/final.cl

@@ -0,0 +1,72 @@
+#define LOCAL_SIZE 300
+#define MAX_PATTERN_SZ 8
+#define MAX_ALPHABET_SZ 97
+#define MAX_RANGES_SZ (2 * MAX_ALPHABET_SZ)
+#define PATTERNS_SZ (MAX_PATTERN_SZ * LOCAL_SIZE)
+
+inline static int string_index(__private char *s, char c) {
+  int i = 0;
+  while (1) {
+    if (s[i] == c)
+      return i;
+    ++i;
+  }
+}
+
+kernel
+__attribute__((reqd_work_group_size(LOCAL_SIZE, 1, 1)))
+__attribute__((xcl_zero_global_work_offset))
+void fmindex(__global char *bwt,
+             __global char *alphabet,
+             __global unsigned *ranks,
+             __global unsigned *sa,
+             __global unsigned *ranges,
+             __global char *patterns,
+             __global unsigned long *out,
+             size_t bwt_sz, size_t alphabet_sz, unsigned pattern_count,
+             unsigned pattern_sz, unsigned out_sz) {
+  int group_id = get_group_id(0);
+
+  __local char _patterns[PATTERNS_SZ];
+  __attribute__((xcl_pipeline_loop(1)))
+  for (unsigned i = 0; i < PATTERNS_SZ; ++i)
+    _patterns[i] = patterns[group_id * LOCAL_SIZE * pattern_sz + i];
+
+  __attribute__((xcl_pipeline_workitems)) {
+    int work_id = get_global_id(0);
+    int local_id = get_local_id(0);
+
+    __private char _alphabet[MAX_ALPHABET_SZ];
+    __attribute__((xcl_pipeline_loop(1)))
+    for (unsigned i = 0; i < alphabet_sz; ++i)
+      _alphabet[i] = alphabet[i];
+
+    __private unsigned _ranges[MAX_RANGES_SZ];
+    __attribute__((xcl_pipeline_loop(1)))
+    for (unsigned i = 0; i < 2 * alphabet_sz; ++i)
+      _ranges[i] = ranges[i];
+
+    int p_idx = pattern_sz - 1;
+    char c = _patterns[local_id * pattern_sz + p_idx];
+    int alphabet_idx = string_index(_alphabet, c);
+    unsigned start = _ranges[2 * alphabet_idx];
+    unsigned end = _ranges[2 * alphabet_idx + 1];
+  
+    p_idx -= 1;
+    while (p_idx >= 0 && end > 1) {
+      c = _patterns[local_id * pattern_sz + p_idx];
+      alphabet_idx = string_index(_alphabet, c);
+      unsigned range_start = _ranges[2 * alphabet_idx];
+      start = range_start + ranks[alphabet_sz * (start - 1) + alphabet_idx];
+      end = range_start + ranks[alphabet_sz * (end - 1) + alphabet_idx];
+      p_idx -= 1;
+    }
+  
+    unsigned long match_count = end - start;
+    out[work_id * out_sz] = match_count;
+
+    __attribute__((xcl_pipeline_loop(1)))
+    for (unsigned i = 0; i < match_count; ++i)
+      out[work_id * out_sz + i + 1] = sa[start + i];
+  }
+}

code/fpga/final_results/dna.10MB.len8/run2.json

@@ -0,0 +1 @@
+{"power": [{"timestamp": 37.8581, "power": 35.3288}, {"timestamp": 58.5703, "power": 35.3288}, {"timestamp": 79.0738, "power": 35.3288}, {"timestamp": 99.609, "power": 35.3288}, {"timestamp": 120.151, "power": 35.3288}, {"timestamp": 140.789, "power": 35.3288}, {"timestamp": 161.382, "power": 35.3288}, {"timestamp": 181.94, "power": 35.3288}, {"timestamp": 202.498, "power": 35.3288}, {"timestamp": 223.0, "power": 35.3288}, {"timestamp": 243.546, "power": 35.3288}, {"timestamp": 264.094, "power": 35.3288}, {"timestamp": 284.638, "power": 35.3288}, {"timestamp": 305.188, "power": 35.3288}, {"timestamp": 325.74, "power": 35.3288}, {"timestamp": 346.294, "power": 35.3288}, {"timestamp": 366.809, "power": 35.3288}, {"timestamp": 387.352, "power": 35.3288}, {"timestamp": 407.899, "power": 35.3288}, {"timestamp": 428.445, "power": 35.3288}, {"timestamp": 449.024, "power": 35.3288}, {"timestamp": 469.623, "power": 35.3288}, {"timestamp": 490.116, "power": 35.3288}, {"timestamp": 510.633, "power": 35.3288}, {"timestamp": 531.131, "power": 35.3288}, {"timestamp": 551.647, "power": 35.3288}, {"timestamp": 572.344, "power": 35.3288}, {"timestamp": 592.961, "power": 35.3288}, {"timestamp": 613.414, "power": 35.3288}, {"timestamp": 633.856, "power": 35.3288}, {"timestamp": 654.298, "power": 35.3288}, {"timestamp": 674.764, "power": 35.3288}, {"timestamp": 695.258, "power": 35.3288}, {"timestamp": 715.801, "power": 35.3288}, {"timestamp": 736.339, "power": 35.3288}, {"timestamp": 756.873, "power": 35.3288}, {"timestamp": 777.787, "power": 35.3288}, {"timestamp": 798.354, "power": 35.3288}], "timeline": {"START": "470.05948", "END": "523.05116"}, "OpenCL API Calls": [{"name": "clSetKernelArg", "calls": 19, "time": 284.273}, {"name": "clFinish", "calls": 1, "time": 229.365}, {"name": "clReleaseKernel", "calls": 1, "time": 86.4594}, {"name": "clCreateProgramWithBinary", "calls": 1, "time": 63.9265}, {"name": "clReleaseContext", "calls": 1, "time": 51.3423}, {"name": "clCreateContext", "calls": 1, "time": 35.5788}, {"name": "clReleaseProgram", "calls": 1, "time": 11.7884}, {"name": "clEnqueueMigrateMemObjects", "calls": 2, "time": 0.367593}, {"name": "clCreateKernel", "calls": 1, "time": 0.34037}, {"name": "clEnqueueNDRangeKernel", "calls": 1, "time": 0.146976}, {"name": "clReleaseMemObject", "calls": 21, "time": 0.077769}, {"name": "clRetainMemObject", "calls": 14, "time": 0.037016}, {"name": "clGetExtensionFunctionAddress", "calls": 1, "time": 0.025618}, {"name": "clCreateBuffer", "calls": 7, "time": 0.024742}, {"name": "clReleaseDevice", "calls": 2, "time": 0.023307}, {"name": "clGetPlatformInfo", "calls": 4, "time": 0.010228}, {"name": "clGetExtensionFunctionAddressForPlatform", "calls": 1, "time": 0.009001}, {"name": "clReleaseCommandQueue", "calls": 1, "time": 0.0084}, {"name": "clCreateCommandQueue", "calls": 1, "time": 0.007222}, {"name": "clGetDeviceIDs", "calls": 2, "time": 0.006521}, {"name": "clRetainDevice", "calls": 2, "time": 0.004685}], "Kernel Execution": [{"kernel": "fmindex", "enqueues": 1, "time": 52.9917}], "Compute Unit Utilization": [{"cu": "fmindex_1", "kernel": "fmindex", "time": 47.4359}, {"cu": "fmindex_2", "kernel": "fmindex", "time": 49.6304}, {"cu": "fmindex_3", "kernel": "fmindex", "time": 48.7461}, {"cu": "fmindex_4", "kernel": "fmindex", "time": 52.5406}, {"cu": "fmindex_5", "kernel": "fmindex", "time": 52.4821}], "Data Transfer: Host to Global Memory": [{"type": "READ", "transfers": 1, "speed": 7632.600307, "utilization": 79.506253, "size": 969240.0, "time": 126.986867}, {"type": "WRITE", "transfers": 1, "speed": 7027.815445, "utilization": 73.206411, "size": 346090.0, "time": 49.245769}], "Data Transfer: Kernels to Global Memory": [{"interface": "DDR[1:3]", "type": "READ", "transfers": 564951, "speed": 102.623, "utilization": 0.890823, "size": 0.00812744}, {"interface": "DDR[1:3]", "type": "WRITE", "transfers": 35684, "speed": 371.67, "utilization": 3.2263, "size": 0.118922}, {"interface": "DDR[1:3]", "type": "READ", "transfers": 604431, "speed": 104.444, "utilization": 0.906631, "size": 0.00811912}, {"interface": "DDR[1:3]", "type": "WRITE", "transfers": 38166, "speed": 372.547, "utilization": 3.23391, "size": 0.119464}, {"interface": "DDR[1:3]", "type": "READ", "transfers": 587746, "speed": 103.73, "utilization": 0.900437, "size": 0.0081225}, {"interface": "DDR[1:3]", "type": "WRITE", "transfers": 37061, "speed": 372.351, "utilization": 3.23221, "size": 0.119424}, {"interface": "DDR[1:3]", "type": "READ", "transfers": 670506, "speed": 109.093, "utilization": 0.946991, "size": 0.00810738}, {"interface": "DDR[1:3]", "type": "WRITE", "transfers": 42277, "speed": 376.434, "utilization": 3.26766, "size": 0.12035}, {"interface": "DDR[1:3]", "type": "READ", "transfers": 669545, "speed": 108.988, "utilization": 0.946073, "size": 0.00810754}, {"interface": "DDR[1:3]", "type": "WRITE", "transfers": 42234, "speed": 376.111, "utilization": 3.26485, "size": 0.120291}]}