hls-performance-thesis/code/fpga/final.cl
2021-07-03 17:59:32 +02:00

72 lines
2.3 KiB
Common Lisp

#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];
}
}