Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
13 changes: 11 additions & 2 deletions .gitignore
Original file line number Diff line number Diff line change
@@ -1,6 +1,15 @@
.vscode/
benchmarks/
grammars/
simple_benchmarks/
test_benchmarks/
b2
main
*.txt
main
gpu/mba
gpu/mba_alt
gpu/mba_grammar
gpu/mba_grammar_gen
gpu/mbalight
gpu/mbalight_alt
gpu/mba_nocache
gpu/mba_rpn
40 changes: 40 additions & 0 deletions gpu/dsl.cuh
Original file line number Diff line number Diff line change
@@ -0,0 +1,40 @@
#pragma once

#include <stdint.h>

typedef uint32_t(*Op1)(uint32_t);
typedef uint32_t(*Op2)(uint32_t, uint32_t);
typedef uint32_t(*Op3)(uint32_t, uint32_t, uint32_t);

// Device functions for MBA operators
__device__ uint32_t op_and(uint32_t a, uint32_t b) { return a & b; }
__device__ uint32_t op_or(uint32_t a, uint32_t b) { return a | b; }
__device__ uint32_t op_xor(uint32_t a, uint32_t b) { return a ^ b; }
__device__ uint32_t op_add(uint32_t a, uint32_t b) { return a + b; }
__device__ uint32_t op_sub(uint32_t a, uint32_t b) { return a - b; }
__device__ uint32_t op_mul(uint32_t a, uint32_t b) { return a * b; }
__device__ uint32_t op_not(uint32_t a) { return ~a; }

// Number of operators (maximum is 16)
static const int opCount = 7;

// Symbols
static const char* DSLSymbols[] = {
"&", "|", "^", "+", "-", "*", "~"
};

// Arities (maximum per operator is 3)
static const int DSLArities[] = {
2, 2, 2, 2, 2, 2, 1
};

// Pointers to device functions
__device__ void* DSLPointers[] = {
(void*)op_and,
(void*)op_or,
(void*)op_xor,
(void*)op_add,
(void*)op_sub,
(void*)op_mul,
(void*)op_not
};
Loading