From ddac5a865b26ef4a78d4ce9fcc8fd92a49b2d1d1 Mon Sep 17 00:00:00 2001 From: minhhn2910 Date: Thu, 26 Oct 2023 15:52:18 +0800 Subject: [PATCH 1/4] add sample execution 1thread with stack and uint256 --- .gitignore | 1 + .vscode/c_cpp_properties.json | 16 +++ Makefile | 17 ++- include/cuevm_test.h | 9 ++ include/opcode.h | 11 ++ include/stack.cuh | 28 ++++ include/uint256.cuh | 60 ++++++++ src/{cu_evm.cu => cuevm.cu} | 94 ++++++++++++- src/cuevm_test.cu | 131 ++++++++++++++++++ src/stack.cu | 46 +++++++ src/uint256.cu | 253 ++++++++++++++++++++++++++++++++++ 11 files changed, 659 insertions(+), 7 deletions(-) create mode 100644 .vscode/c_cpp_properties.json create mode 100644 include/cuevm_test.h create mode 100644 include/opcode.h create mode 100644 include/stack.cuh create mode 100644 include/uint256.cuh rename src/{cu_evm.cu => cuevm.cu} (50%) create mode 100644 src/cuevm_test.cu create mode 100644 src/stack.cu create mode 100644 src/uint256.cu diff --git a/.gitignore b/.gitignore index cb385db..b4741ef 100644 --- a/.gitignore +++ b/.gitignore @@ -4,3 +4,4 @@ *.ptx *.cubin *.fatbin +*.o diff --git a/.vscode/c_cpp_properties.json b/.vscode/c_cpp_properties.json new file mode 100644 index 0000000..2489d95 --- /dev/null +++ b/.vscode/c_cpp_properties.json @@ -0,0 +1,16 @@ +{ + "configurations": [ + { + "name": "Linux", + "includePath": [ + "${workspaceFolder}/**" + ], + "defines": [], + "compilerPath": "/usr/bin/gcc", + "cStandard": "c11", + "cppStandard": "gnu++14", + "intelliSenseMode": "linux-gcc-x64" + } + ], + "version": 4 +} \ No newline at end of file diff --git a/Makefile b/Makefile index ac22b13..b218528 100644 --- a/Makefile +++ b/Makefile @@ -1,10 +1,17 @@ +# Object files +objects = cuevm.o cuevm_test.o stack.o uint256.o + +# Compiler NVCC = nvcc -NVCC_FLAGS = -I./include -lstdc++ -all: cuEVM +# Compiler Flags +NVCC_FLAGS = -I./include -lstdc++ -dc + +all: $(objects) + $(NVCC) $(objects) -o cuEVM -cuEVM: - $(NVCC) $(NVCC_FLAGS) -o cuEVM src/cu_evm.cu +%.o: src/%.cu + $(NVCC) $(NVCC_FLAGS) $< -o $@ clean: - rm -f cuEVM \ No newline at end of file + rm -f *.o cuEVM diff --git a/include/cuevm_test.h b/include/cuevm_test.h new file mode 100644 index 0000000..032c59f --- /dev/null +++ b/include/cuevm_test.h @@ -0,0 +1,9 @@ +#ifndef CUEVM_TEST_H +#define CUEVM_TEST_H +#include "stack.cuh" + + +void test_arithmetic_operations(); +void test_stack(); + +#endif // CUEVM_TEST_H diff --git a/include/opcode.h b/include/opcode.h new file mode 100644 index 0000000..41a726a --- /dev/null +++ b/include/opcode.h @@ -0,0 +1,11 @@ +#ifndef OPCODE_H +#define OPCODE_H + +#define ADD 0x01 +#define MUL 0x02 + +#define POP 0x50 +#define PUSH1 0x60 +// Add other opcode definitions here as your VM expands + +#endif // OPCODE_H diff --git a/include/stack.cuh b/include/stack.cuh new file mode 100644 index 0000000..1a6d0a5 --- /dev/null +++ b/include/stack.cuh @@ -0,0 +1,28 @@ +#ifndef STACK_CUH +#define STACK_CUH + +#include +#include +#include +#include +#include +#include "uint256.cuh" +#define STACK_SIZE 100 // For example, temporarily set the stack size of 100 + +typedef struct { + base_uint items[STACK_SIZE]; + int top; +} base_uint_stack; + +__host__ __device__ void init_stack(base_uint_stack* stack); + +__host__ __device__ bool push(base_uint_stack* stack, base_uint item); + +__host__ __device__ bool pop(base_uint_stack* stack, base_uint* item); + +__host__ __device__ bool swap_with_top(base_uint_stack* stack, int i); + +__host__ __device__ void print_stack(base_uint_stack* stack); + + +#endif // STACK_CUH diff --git a/include/uint256.cuh b/include/uint256.cuh new file mode 100644 index 0000000..eeff53f --- /dev/null +++ b/include/uint256.cuh @@ -0,0 +1,60 @@ +// base_uint256.cuh + +#ifndef BASE_UINT256_CUH +#define BASE_UINT256_CUH + +#include +#include +#include +#include +#include +#include + +#define BITS 256 +#define WIDTH (BITS / 32) + +typedef struct +{ + uint32_t pn[WIDTH]; +} base_uint; + +// utility functions +__host__ int hexToInt(const char *hex); + +__host__ void intToHex(int num, char *hex); + +__host__ bool hex_to_decimal(const char *hex_str, char *dec_str); + +__host__ __device__ void print_base_uint(const base_uint *val); + +// conversion operations +__host__ bool base_uint_set_hex(base_uint *val, const char *hex); + +__host__ void base_uint_to_string(const base_uint *val, char *out_str); + +__host__ bool int_to_base_uint(int int_val, base_uint *val); + +__host__ __device__ void base_uint_get_hex(const base_uint *val, char *hex); + + +// comparison operations +__host__ __device__ bool is_zero(const base_uint *num); + +// bitwise operations +__host__ __device__ base_uint bitwise_not(const base_uint *num); + +__host__ __device__ void base_uint_set_bit(base_uint *value, uint32_t bitpos); + +// arithmetic operations + +__host__ __device__ void base_uint_add(const base_uint *a, const base_uint *b, base_uint *result); + +__host__ __device__ bool base_uint_sub(const base_uint *a, const base_uint *b, base_uint *result); + +__host__ __device__ void base_uint_mul(const base_uint *a, const base_uint *b, base_uint *result); + +__host__ __device__ void base_uint_shift_left(base_uint *a, size_t bits); + +__host__ __device__ void base_uint_div(const base_uint *a, const base_uint *b, base_uint *quotient, base_uint *remainder); + +#endif // BASE_UINT256_CUH diff --git a/src/cu_evm.cu b/src/cuevm.cu similarity index 50% rename from src/cu_evm.cu rename to src/cuevm.cu index 32ef6ab..c85407c 100644 --- a/src/cu_evm.cu +++ b/src/cuevm.cu @@ -2,10 +2,14 @@ #include #include #include - +#include "uint256.cuh" +#include "stack.cuh" +#include "cuevm_test.h" +#include "opcode.h" #define NUMTHREAD 4096 - +#define DEBUG 1 // simple draft kernel for place holder +// simple testing opcodes and return the popped top of stack value __global__ void cuEVM(unsigned char *bytecode, unsigned char *input, size_t bytecode_len, size_t input_len, size_t num_threads) { int idx = threadIdx.x + blockIdx.x * blockDim.x; @@ -26,6 +30,85 @@ __global__ void cuEVM(unsigned char *bytecode, unsigned char *input, size_t byte printf("%02x ", input[i]); } printf("\n"); + + base_uint_stack stack; + init_stack(&stack); + + // push(&stack, a); + // pop(&stack, &b); + // debugging : print_stack(&stack); + // define 3 reusable temp uints for binary op + base_uint op1, op2, result; + for (size_t i = 0; i < bytecode_len; i++) + { + unsigned char opcode = bytecode[i]; + switch (opcode) + { + case ADD: // ADD + // TODO: check stack size + // future optimization : can override push pop ops and modify the stack directly + pop(&stack, &op1); + pop(&stack, &op2); + base_uint_add(&op1, &op2, &result); + + #if DEBUG + printf("ADD OPCODE: \n"); + printf("op1: "); + print_base_uint(&op1); + printf("op2: "); + print_base_uint(&op2); + printf("result: "); + print_base_uint(&result); + printf("\n***************\n"); + #endif + + push(&stack, result); + break; + + case MUL: // MUL + // TODO: check stack size + pop(&stack, &op1); + pop(&stack, &op2); + base_uint_mul(&op1, &op2, &result); + + #if DEBUG + printf("MUL OPCODE: \n"); + printf("op1: "); + print_base_uint(&op1); + printf("op2: "); + print_base_uint(&op2); + printf("result: "); + print_base_uint(&result); + printf("\n***************\n"); + #endif + + push(&stack, result); + break; + case PUSH1: + unsigned char push_val = bytecode[++i]; + result = { {push_val, 0, 0, 0, 0, 0, 0, 0} }; + push(&stack, result); + + #if DEBUG + printf("PUSH1 OPCODE: \n"); + printf("push_val: "); + print_base_uint(&result); + printf("\n***************\n"); + #endif + + break; + case POP: + pop(&stack, &result); + printf("Popped Stack value: "); + print_base_uint(&result); + printf("\n***************\n"); + break; + default: + printf("Unknown opcode 0x%02x at position %zu\n", opcode, i); + return; + } + } + } } } @@ -45,6 +128,8 @@ void hexStringToByteArray(const char *hexString, unsigned char *byteArray, int l } } + + int main(int argc, char *argv[]) { char *byte_code_hex = NULL; @@ -53,6 +138,7 @@ int main(int argc, char *argv[]) static struct option long_options[] = { {"bytecode", required_argument, 0, 'b'}, {"input", required_argument, 0, 'i'}, + {"test", no_argument, 0, 't'}, {0, 0, 0, 0}}; int opt; @@ -67,6 +153,10 @@ int main(int argc, char *argv[]) case 'i': input_hex = optarg; break; + case 't': + test_arithmetic_operations(); + test_stack(); + exit(0); default: fprintf(stderr, "Usage: %s --bytecode --input \n", argv[0]); exit(EXIT_FAILURE); diff --git a/src/cuevm_test.cu b/src/cuevm_test.cu new file mode 100644 index 0000000..e2dbe8e --- /dev/null +++ b/src/cuevm_test.cu @@ -0,0 +1,131 @@ +#include "cuevm_test.h" + +void test_arithmetic_operations() +{ + base_uint a, b, c, d; + + // Test addition + base_uint_set_hex(&a, "11111111111111111111111111111111"); + base_uint_set_hex(&b, "22222222222222222222222222222222"); + base_uint_add(&a, &b, &c); + printf("Addition Result: "); + + char buffer[BITS / 4 + 1] = {0}; + base_uint_get_hex(&c, buffer); + + printf("%s\n", buffer); + + if (strcmp(buffer, "0000000000000000000000000000000033333333333333333333333333333333") != 0) + { + printf("Addition failed!\n"); + } + // Test addition with carry + base_uint_set_hex(&a, "1"); + base_uint_set_hex(&b, "ffffffffffffffffffffffffffffffff"); + base_uint_add(&a, &b, &c); + printf("Addition Result: "); + + base_uint_get_hex(&c, buffer); + + printf("%s\n", buffer); + + if (strcmp(buffer, "0000000000000000000000000000000100000000000000000000000000000000") != 0) + { + printf("Addition failed!\n"); + } + // Test addition overflow carry + base_uint_set_hex(&a, "1234"); + base_uint_set_hex(&b, "ffffffffffffffffffffffffffffffffffffffffffffffffffffffffffffffff"); + base_uint_add(&a, &b, &c); + printf("Addition Result: "); + + base_uint_get_hex(&c, buffer); + + printf("%s\n", buffer); + + if (strcmp(buffer, "0000000000000000000000000000000000000000000000000000000000001233") != 0) + { + printf("Addition failed!\n"); + } + + // Test subtraction + base_uint_set_hex(&a, "ffffffffffffffffffffffffffffffff"); + base_uint_set_hex(&b, "fe"); + base_uint_sub(&a, &b, &c); + printf("Subtraction Result: "); + base_uint_get_hex(&c, buffer); + printf("%s\n", buffer); + if (strcmp(buffer, "00000000000000000000000000000000ffffffffffffffffffffffffffffff01") != 0) + { + printf("Subtraction failed!\n"); + } + + // Test subtraction underflow + base_uint_set_hex(&a, "01"); + base_uint_set_hex(&b, "ff"); + base_uint_sub(&a, &b, &c); + printf("Subtraction Result: "); + base_uint_get_hex(&c, buffer); + printf("%s\n", buffer); + if (strcmp(buffer, "ffffffffffffffffffffffffffffffffffffffffffffffffffffffffffffff02") != 0) + { + printf("Subtraction failed!\n"); + } + + // Test multiplication + base_uint_set_hex(&a, "ffffffffffffffffffffffffffffffff"); + base_uint_set_hex(&b, "ffffffffffffffffffffffffffffff"); + base_uint_mul(&a, &b, &c); + printf("Multiplication Result: "); + base_uint_get_hex(&c, buffer); + printf("%s\n", buffer); + if (strcmp(buffer, "00fffffffffffffffffffffffffffffeff000000000000000000000000000001") != 0) + { + printf("Multiplication failed!\n"); + } + // Test multiplication overflow + base_uint_set_hex(&a, "ffffffffffffffffffffffffffffffffffffffffffffffffffffffffffffffff"); + base_uint_set_hex(&b, "2"); + base_uint_mul(&a, &b, &c); + printf("Multiplication Result: "); + base_uint_get_hex(&c, buffer); + printf("%s\n", buffer); + if (strcmp(buffer, "0000000000000000000000000000000000000000000000000000000000000000") != 0) + { + printf("Multiplication overflow failed!\n"); + } + +} + + +void test_stack() { + base_uint_stack stack; + init_stack(&stack); + + // Test push and print + base_uint a = { {1, 2, 3, 4} }; + printf("Pushing: "); + for (int i = 0; i < WIDTH; i++) printf("%u ", a.pn[i]); + printf("\n"); + push(&stack, a); + print_stack(&stack); + + // Test pop + base_uint b; + if (pop(&stack, &b)) { + printf("Popped: "); + for (int i = 0; i < WIDTH; i++) printf("%u ", b.pn[i]); + printf("\n"); + } + print_stack(&stack); + + // Test swap with top + push(&stack, a); + base_uint c = { {5, 6, 7, 8} }; + push(&stack, c); + printf("Before swap with top:\n"); + print_stack(&stack); + swap_with_top(&stack, 0); + printf("After swap with top:\n"); + print_stack(&stack); +} diff --git a/src/stack.cu b/src/stack.cu new file mode 100644 index 0000000..17a4c72 --- /dev/null +++ b/src/stack.cu @@ -0,0 +1,46 @@ +#include "stack.cuh" + +__host__ __device__ void init_stack(base_uint_stack* stack) { + stack->top = -1; +} + +__host__ __device__ bool push(base_uint_stack* stack, base_uint item) { + if (stack->top >= STACK_SIZE - 1) { + return false; // Stack is full + } + stack->top++; + stack->items[stack->top] = item; + return true; +} + +__host__ __device__ bool pop(base_uint_stack* stack, base_uint* item) { + if (stack->top < 0) { + return false; // Stack is empty + } + *item = stack->items[stack->top]; + stack->top--; + return true; +} + +__host__ __device__ bool swap_with_top(base_uint_stack* stack, int i) { + if (stack->top < 0 || i > stack->top || i < 0) { + return false; // Stack is empty or index out of bounds + } + base_uint temp = stack->items[i]; + stack->items[i] = stack->items[stack->top]; + stack->items[stack->top] = temp; + return true; +} + +__host__ __device__ void print_stack(base_uint_stack* stack) { + printf("Stack: "); + for (int i = 0; i <= stack->top; i++) { + printf("["); + for (int j = 0; j < WIDTH; j++) { + printf("%u", stack->items[i].pn[j]); + if (j < WIDTH - 1) printf(","); + } + printf("] "); + } + printf("\n"); +} diff --git a/src/uint256.cu b/src/uint256.cu new file mode 100644 index 0000000..fd1854f --- /dev/null +++ b/src/uint256.cu @@ -0,0 +1,253 @@ + +#include "uint256.cuh" +// implementation +__host__ int hexToInt(const char *hex) +{ + int result = 0; + int len = strlen(hex); + + for (int i = 0; i < len; i++) + { + char c = tolower(hex[i]); + if (c >= '0' && c <= '9') + { + result = result * 16 + (c - '0'); + } + else if (c >= 'a' && c <= 'f') + { + result = result * 16 + (c - 'a' + 10); + } + else + { + // Invalid hexadecimal character + return -1; + } + } + return result; +} + +__host__ void intToHex(int num, char *hex) +{ + // Assuming hex has enough space + char *ptr = hex; + do + { + int remainder = num % 16; + if (remainder < 10) + { + *ptr++ = '0' + remainder; + } + else + { + *ptr++ = 'a' + (remainder - 10); + } + num /= 16; + } while (num != 0); + + *ptr-- = '\0'; // NULL-terminate the string and point to the last valid character + + // Reverse the string + char *start = hex; + while (start < ptr) + { + char t = *start; + *start = *ptr; + *ptr = t; + start++; + ptr--; + } +} + +__host__ bool hex_to_decimal(const char *hex_str, char *dec_str) +{ + unsigned long long result = 0; + unsigned long long place = 1; + + int len = strlen(hex_str); + for (int i = len - 1; i >= 0; i--) + { + char c = tolower(hex_str[i]); + int digit; + if (c >= '0' && c <= '9') + { + digit = c - '0'; + } + else if (c >= 'a' && c <= 'f') + { + digit = 10 + (c - 'a'); + } + else + { + return false; + } + + result += digit * place; + place *= 16; + } + + sprintf(dec_str, "%llu", result); + return true; +} + +__host__ bool base_uint_set_hex(base_uint *val, const char *hex) +{ + memset(val->pn, 0, sizeof(val->pn)); + + size_t len = strlen(hex); + if (len == 0 || len > BITS / 4) + return false; + + // Iterate through the string from end to start + for (size_t i = 0; i < len; i++) + { + char c = tolower(hex[len - 1 - i]); + uint32_t number = 0; + + if (c >= '0' && c <= '9') + { + number = c - '0'; + } + else if (c >= 'a' && c <= 'f') + { + number = c - 'a' + 10; + } + else + { + return false; // Invalid character + } + + // Determine which uint32_t element and position the hex character should be placed + val->pn[i / 8] |= (number << ((i % 8) * 4)); + } + return true; +} + +__host__ void base_uint_to_string(const base_uint *val, char *out_str) +{ + char hex_str[BITS / 4 + 1] = {0}; + base_uint_get_hex(val, hex_str); + + if (!hex_to_decimal(hex_str, out_str)) + { + strcpy(out_str, "Error"); + } +} + +__host__ bool int_to_base_uint(int int_val, base_uint *val) +{ + char *p; + sprintf(p, "%08x", int_val); + printf("%s\n", p); + return base_uint_set_hex(val, p); +} + +__host__ __device__ void base_uint_get_hex(const base_uint *val, char *hex) +{ + char *p = hex; + + for (int i = WIDTH - 1; i >= 0; i--) + { + // printf("%d ", val->pn[i]); + sprintf(p, "%08x", val->pn[i]); + p += 8; + } +} + +__host__ __device__ void print_base_uint(const base_uint *val) +{ + for (int i = 0; i < WIDTH; i++) + { + printf("%d ", val->pn[i]); + } +} + +__host__ __device__ bool is_zero(const base_uint *num) +{ + for (int i = 0; i < WIDTH; i++) + { + if (num->pn[i] != 0) + { + return false; + } + } + return true; +} + +__host__ __device__ base_uint bitwise_not(const base_uint *num) +{ + base_uint ret; + for (int i = 0; i < WIDTH; i++) + { + ret.pn[i] = ~num->pn[i]; + } + return ret; +} + +__host__ __device__ void base_uint_set_bit(base_uint *value, uint32_t bitpos) +{ + value->pn[bitpos / 32] |= (1 << (bitpos % 32)); +} + +__host__ __device__ void base_uint_add(const base_uint *a, const base_uint *b, base_uint *result) +{ + uint64_t carry = 0; + + for (size_t i = 0; i < WIDTH; i++) + { + uint64_t sum = (uint64_t)a->pn[i] + b->pn[i] + carry; + printf("%d %d = %d %d\n", a->pn[i], b->pn[i], sum, carry); + result->pn[i] = (uint32_t)sum; // Store lower 32 bits + carry = sum >> 32; // Take upper 32 bits as the next carry + } +} + +__host__ __device__ bool base_uint_sub(const base_uint *a, const base_uint *b, base_uint *result) +{ + uint64_t borrow = 0; + + for (size_t i = 0; i < WIDTH; i++) + { + uint64_t res = 0x100000000ULL + (uint64_t)a->pn[i] - b->pn[i] - borrow; + result->pn[i] = (uint32_t)res; + if (res >= 0x100000000ULL) + { + borrow = 0; + } + else + { + borrow = 1; + } + } + + // If borrow is still 1 after looping through all words, then a < b. + // Return false to indicate underflow + return borrow == 0; +} + +/* +Warming: +1. Not tested yet. +2. Overflow wraparound is not correctly implemented yet. +*/ +__host__ __device__ void base_uint_mul(const base_uint *a, const base_uint *b, base_uint *result) +{ + base_uint temp_result = {0}; + for (size_t i = 0; i < WIDTH; i++) + { + uint64_t carry = 0; + for (size_t j = 0; j < WIDTH; j++) + { + if (i + j < WIDTH) + { + uint64_t product = (uint64_t)a->pn[i] * b->pn[j] + temp_result.pn[i + j] + carry; + temp_result.pn[i + j] = (uint32_t)product; + carry = product >> 32; + } + } + } + + for (size_t i = 0; i < WIDTH; i++) + { + result->pn[i] = temp_result.pn[i]; + } +} From 5558f699cdd79dfb0273cd033f1a127cae55a4e3 Mon Sep 17 00:00:00 2001 From: minhhn2910 Date: Thu, 26 Oct 2023 16:58:57 +0800 Subject: [PATCH 2/4] Update README.md --- README.md | 27 +++++++++++++++++++++++++++ 1 file changed, 27 insertions(+) diff --git a/README.md b/README.md index a94e01a..e9bdc50 100644 --- a/README.md +++ b/README.md @@ -16,6 +16,33 @@ Or with custom nvcc path * `./cuEVM --bytecode [hex string of byte code] --input [hex string of input]` +Sample testing bytecode : + +``` +0x6006 PUSH1 0x06 +0x6007 PUSH1 0x07 +0x02 MUL +0x50 POP // Return 42 +0x600660070250 => POP 42 from the stack +``` +Usage : +``` +./cuEVM --bytecode 0x600660070250 --input 0x1234 +Bytecode: 60 06 60 07 02 50 +Input: 12 34 +PUSH1 OPCODE: +push_val: 6 0 0 0 0 0 0 0 +*************** +PUSH1 OPCODE: +push_val: 7 0 0 0 0 0 0 0 +*************** +MUL OPCODE: +op1: 7 0 0 0 0 0 0 0 op2: 6 0 0 0 0 0 0 0 result: 42 0 0 0 0 0 0 0 +*************** +Popped Stack value: 42 0 0 0 0 0 0 0 +*************** +``` + TODO: change options, configs, and how we use the tool in the future. ## Code structure From 98dc55a50834a4ce155dc40666f4386c45755c50 Mon Sep 17 00:00:00 2001 From: minhhn2910 Date: Thu, 26 Oct 2023 21:15:46 +0800 Subject: [PATCH 3/4] Add JUMPI and loop test case --- .gitignore | 1 + README.md | 47 +++++++++--- include/opcode.h | 13 ++++ include/processor.cuh | 42 +++++++++++ src/cuevm.cu | 170 +++++++++++++++++++++++++++++++++--------- 5 files changed, 227 insertions(+), 46 deletions(-) create mode 100644 include/processor.cuh diff --git a/.gitignore b/.gitignore index b4741ef..9835a29 100644 --- a/.gitignore +++ b/.gitignore @@ -5,3 +5,4 @@ *.cubin *.fatbin *.o +cuEVM diff --git a/README.md b/README.md index e9bdc50..df94d2b 100644 --- a/README.md +++ b/README.md @@ -12,7 +12,7 @@ Cuda implementation of EVM bytecode executor Or with custom nvcc path * `make NVCC=/usr/local/cuda-10.0/bin/nvcc` -## Usage +## Usage * `./cuEVM --bytecode [hex string of byte code] --input [hex string of input]` @@ -21,28 +21,53 @@ Sample testing bytecode : ``` 0x6006 PUSH1 0x06 0x6007 PUSH1 0x07 -0x02 MUL +0x02 MUL 0x50 POP // Return 42 0x600660070250 => POP 42 from the stack ``` Usage : ``` ./cuEVM --bytecode 0x600660070250 --input 0x1234 -Bytecode: 60 06 60 07 02 50 -Input: 12 34 -PUSH1 OPCODE: -push_val: 6 0 0 0 0 0 0 0 +Bytecode: 60 06 60 07 02 50 +Input: 12 34 +PUSH1 OPCODE: +push_val: 6 0 0 0 0 0 0 0 *************** -PUSH1 OPCODE: -push_val: 7 0 0 0 0 0 0 0 +PUSH1 OPCODE: +push_val: 7 0 0 0 0 0 0 0 *************** -MUL OPCODE: -op1: 7 0 0 0 0 0 0 0 op2: 6 0 0 0 0 0 0 0 result: 42 0 0 0 0 0 0 0 +MUL OPCODE: +op1: 7 0 0 0 0 0 0 0 op2: 6 0 0 0 0 0 0 0 result: 42 0 0 0 0 0 0 0 *************** -Popped Stack value: 42 0 0 0 0 0 0 0 +Popped Stack value: 42 0 0 0 0 0 0 0 *************** ``` +Loop with jumpi : +``` +LOGIC: +1. Perform 6 * 7 = res; +2. Loop: while(res!=0): res = res - 14 (loop 3 times) => STOP + +PC 0 : 0x6006 PUSH1 0x06 +PC 2 : 0x6007 PUSH1 0x07 +PC 4 : 0x02 MUL // TOP STACK 42 +PC 5 : 0x5b JUMPDEST // TAG1_JUMP +PC 7 : 0x600e PUSH 0x0e +PC 8 : 0x90 SWAP1 +PC 9 : 0x03 SUB // 42 - 14 // condition != 1 jump; +PC 10 : 0x6005 PUSH1 TAG1_JUMP // destination +PC 12 : 0x57 JUMPI +PC 13 : 0x50 POP // for testing +PC 14 : 0xf3 RETURN // for testing +Bytecode: +0x60066007025b600e900360055750f3 +``` +Run +``` +./cuEVM --bytecode 0x60066007025b600e900360055750f3 --input 0x1234 +``` + TODO: change options, configs, and how we use the tool in the future. ## Code structure diff --git a/include/opcode.h b/include/opcode.h index 41a726a..5f30bf5 100644 --- a/include/opcode.h +++ b/include/opcode.h @@ -3,9 +3,22 @@ #define ADD 0x01 #define MUL 0x02 +#define SUB 0x03 + #define POP 0x50 #define PUSH1 0x60 +#define PUSH2 0x61 + +#define SWAP1 0x90 + + +#define JUMP 0x56 +#define JUMPI 0x57 +#define JUMPDEST 0x5b + + +#define RETURN 0xf3 // Add other opcode definitions here as your VM expands #endif // OPCODE_H diff --git a/include/processor.cuh b/include/processor.cuh new file mode 100644 index 0000000..6512268 --- /dev/null +++ b/include/processor.cuh @@ -0,0 +1,42 @@ +#ifndef PROCESSOR_CUH +#define PROCESSOR_CUH + +#include +#include +#include +#include +#include +#include "uint256.cuh" +#include "stack.cuh" +#define EVM_VERSION "petersburg" // hardcode to only one version ! + +typedef struct { + base_uint origin; + base_uint block_numer; + base_uint block_difficulty; + // other fields +} environment; + +// Struct to represent a call frame +// Not implemented yet (cross contract) +// typedef struct { +// base_uint caller; +// base_uint callValue; +// base_uint gasLimit; // not supported +// uint8_t* inputData; +// size_t inputDataSize; +// } CallFrame; + +// Struct for the EVM processor +typedef struct { + // other fields as new functionality is added + // base_uint gasRemaining; not implemented + // CallFrame* callStack; // not implemented yet + base_uint_stack stack; + base_uint caller; + uint32_t programCounter; + // uint8_t* bytecode; temporarily not needed in single contract +} processor; +// util functions + +#endif // PROCESSOR_CUH diff --git a/src/cuevm.cu b/src/cuevm.cu index c85407c..0f4e3e0 100644 --- a/src/cuevm.cu +++ b/src/cuevm.cu @@ -6,6 +6,7 @@ #include "stack.cuh" #include "cuevm_test.h" #include "opcode.h" +#include "processor.cuh" #define NUMTHREAD 4096 #define DEBUG 1 // simple draft kernel for place holder @@ -31,27 +32,35 @@ __global__ void cuEVM(unsigned char *bytecode, unsigned char *input, size_t byte } printf("\n"); - base_uint_stack stack; - init_stack(&stack); + // todo refactor processor.execute to a function outside of this kernel + processor evm; + evm.programCounter = 0; // redundant but better safe. + init_stack(&evm.stack); - // push(&stack, a); - // pop(&stack, &b); - // debugging : print_stack(&stack); - // define 3 reusable temp uints for binary op - base_uint op1, op2, result; - for (size_t i = 0; i < bytecode_len; i++) - { - unsigned char opcode = bytecode[i]; - switch (opcode) + // push(&stack, a); + // pop(&stack, &b); + // debugging : print_stack(&stack); + // define 3 reusable temp uints for binary op + base_uint op1, op2, result; + uint safe_counter = 0; // safety counter prevent infinite loop + while (evm.programCounter < bytecode_len) { + unsigned char opcode = bytecode[evm.programCounter]; + safe_counter ++; + if (safe_counter > 100) { + printf("Safety counter exceeded, return from execution\n"); + return; + } + switch (opcode) + { case ADD: // ADD // TODO: check stack size // future optimization : can override push pop ops and modify the stack directly - pop(&stack, &op1); - pop(&stack, &op2); + pop(&evm.stack, &op1); + pop(&evm.stack, &op2); base_uint_add(&op1, &op2, &result); - #if DEBUG +#if DEBUG printf("ADD OPCODE: \n"); printf("op1: "); print_base_uint(&op1); @@ -60,18 +69,36 @@ __global__ void cuEVM(unsigned char *bytecode, unsigned char *input, size_t byte printf("result: "); print_base_uint(&result); printf("\n***************\n"); - #endif +#endif + + push(&evm.stack, result); + break; + case SUB: + pop(&evm.stack, &op1); + pop(&evm.stack, &op2); + base_uint_sub(&op1, &op2, &result); + +#if DEBUG + printf("SUB OPCODE: \n"); + printf("op1: "); + print_base_uint(&op1); + printf("op2: "); + print_base_uint(&op2); + printf("result: "); + print_base_uint(&result); + printf("\n***************\n"); +#endif - push(&stack, result); + push(&evm.stack, result); break; case MUL: // MUL // TODO: check stack size - pop(&stack, &op1); - pop(&stack, &op2); + pop(&evm.stack, &op1); + pop(&evm.stack, &op2); base_uint_mul(&op1, &op2, &result); - #if DEBUG +#if DEBUG printf("MUL OPCODE: \n"); printf("op1: "); print_base_uint(&op1); @@ -80,41 +107,116 @@ __global__ void cuEVM(unsigned char *bytecode, unsigned char *input, size_t byte printf("result: "); print_base_uint(&result); printf("\n***************\n"); - #endif +#endif - push(&stack, result); + push(&evm.stack, result); break; case PUSH1: - unsigned char push_val = bytecode[++i]; - result = { {push_val, 0, 0, 0, 0, 0, 0, 0} }; - push(&stack, result); + unsigned char push_val = bytecode[++evm.programCounter]; + result = {{push_val, 0, 0, 0, 0, 0, 0, 0}}; + push(&evm.stack, result); - #if DEBUG +#if DEBUG printf("PUSH1 OPCODE: \n"); printf("push_val: "); print_base_uint(&result); printf("\n***************\n"); - #endif +#endif + + break; + case PUSH2: + // Increment the program counter to point to the first byte of data + evm.programCounter++; + + // Read the two bytes from the bytecode + unsigned char byte1 = bytecode[evm.programCounter]; + unsigned char byte2 = bytecode[++evm.programCounter]; + + // Combine the two bytes into a single 16-bit value + uint16_t push_val_16 = (byte1 << 8) | byte2; + + // Convert the 16-bit value into your base_uint format + result = {{push_val_16, 0, 0, 0, 0, 0, 0, 0}}; + + // Push the value onto the stack + push(&evm.stack, result); + +#if DEBUG + printf("PUSH2 OPCODE: \n"); + printf("push_val: "); + print_base_uint(&result); + printf("\n***************\n"); +#endif break; + case POP: - pop(&stack, &result); + pop(&evm.stack, &result); printf("Popped Stack value: "); print_base_uint(&result); printf("\n***************\n"); break; + + case SWAP1: + printf("SWAP1 OPCODE BEFORE: \n"); + print_base_uint(&evm.stack.items[0]); + print_base_uint(&evm.stack.items[1]); + pop(&evm.stack, &op1); + pop(&evm.stack, &op2); + push(&evm.stack, op1); + push(&evm.stack, op2); + printf("\nSWAP1 OPCODE AFTER: \n"); + print_base_uint(&evm.stack.items[0]); + print_base_uint(&evm.stack.items[1]); + break; + case JUMPI: + pop(&evm.stack, &op1); + pop(&evm.stack, &op2); +#if DEBUG + printf("JUMPI OPCODE: \n"); + printf("Condition:\n"); + print_base_uint(&op2); + printf("Destination:\n"); + print_base_uint(&op1); + printf("is ZEROP op1: %d\n", is_zero(&op1)); +#endif + if (!is_zero(&op2)) + { + evm.programCounter = op1.pn[0]; + // TODO: check JUMPDEST in destiation ? + } + break; + + case JUMP: + pop(&evm.stack, &op1); + evm.programCounter = op1.pn[0];; + // TODO: check JUMPDEST in destiation ? + break; + + case JUMPDEST: + // do nothing + break; + case RETURN: + printf("RETURN OPCODE\n"); + evm.programCounter = bytecode_len; + break; default: - printf("Unknown opcode 0x%02x at position %zu\n", opcode, i); + printf("Unknown opcode %d at position %d\n", opcode, evm.programCounter); + printf("Return from execution\n"); return; - } - } + } + evm.programCounter++; + printf("Program counter: %d\n", evm.programCounter); + } } } } -int adjustedLength(char** hexString) { - if (strncmp(*hexString, "0x", 2) == 0 || strncmp(*hexString, "0X", 2) == 0) { - *hexString += 2; // Skip the "0x" prefix +int adjustedLength(char **hexString) +{ + if (strncmp(*hexString, "0x", 2) == 0 || strncmp(*hexString, "0X", 2) == 0) + { + *hexString += 2; // Skip the "0x" prefix return (strlen(*hexString) / 2); } return (strlen(*hexString) / 2); @@ -128,8 +230,6 @@ void hexStringToByteArray(const char *hexString, unsigned char *byteArray, int l } } - - int main(int argc, char *argv[]) { char *byte_code_hex = NULL; From 69a258486aa5352a4404a82684d3f4c0a69d265b Mon Sep 17 00:00:00 2001 From: minhhn2910 Date: Sat, 28 Oct 2023 16:54:44 +0800 Subject: [PATCH 4/4] Fix example + add DUP1 + add code formatting rule --- .clang-format | 5 + README.md | 17 ++- include/opcode.h | 1 + include/uint256.cuh | 17 ++- src/cuevm.cu | 344 +++++++++++++++++++++----------------------- src/cuevm_test.cu | 30 ++-- src/stack.cu | 4 +- src/uint256.cu | 151 +++++++------------ 8 files changed, 250 insertions(+), 319 deletions(-) create mode 100644 .clang-format diff --git a/.clang-format b/.clang-format new file mode 100644 index 0000000..de9397e --- /dev/null +++ b/.clang-format @@ -0,0 +1,5 @@ +BasedOnStyle: Google +IndentWidth: 4 +TabWidth: 4 +UseTab: Never +ColumnLimit: 120 \ No newline at end of file diff --git a/README.md b/README.md index df94d2b..068df25 100644 --- a/README.md +++ b/README.md @@ -53,21 +53,24 @@ PC 0 : 0x6006 PUSH1 0x06 PC 2 : 0x6007 PUSH1 0x07 PC 4 : 0x02 MUL // TOP STACK 42 PC 5 : 0x5b JUMPDEST // TAG1_JUMP -PC 7 : 0x600e PUSH 0x0e +PC 6 : 0x600e PUSH 0x0e PC 8 : 0x90 SWAP1 PC 9 : 0x03 SUB // 42 - 14 // condition != 1 jump; -PC 10 : 0x6005 PUSH1 TAG1_JUMP // destination -PC 12 : 0x57 JUMPI -PC 13 : 0x50 POP // for testing -PC 14 : 0xf3 RETURN // for testing +PC 10 : 0x80 DUP1 // DUP the result because JUMPI will remove it +PC 11 : 0x6005 PUSH1 TAG1_JUMP // destination +PC 13 : 0x57 JUMP I +PC 14 : 0x50 POP // for testing +PC 15 : 0xf3 RETURN // for testing Bytecode: -0x60066007025b600e900360055750f3 +0x60066007025b600e90038060055750f3 ``` Run ``` -./cuEVM --bytecode 0x60066007025b600e900360055750f3 --input 0x1234 +./cuEVM --bytecode 0x60066007025b600e90038060055750f3 --input 0x1234 ``` +Reference tools (www.evm.codes) for testing bytecode sequence : [Simulate test bytecode sequence](https://www.evm.codes/playground?fork=shanghai&unit=Wei&codeType=Bytecode&code=%27%7E6%7E7025b%7Ee900380%7E55750f3%27%7E600%01%7E_) + TODO: change options, configs, and how we use the tool in the future. ## Code structure diff --git a/include/opcode.h b/include/opcode.h index 5f30bf5..559f8c2 100644 --- a/include/opcode.h +++ b/include/opcode.h @@ -12,6 +12,7 @@ #define SWAP1 0x90 +#define DUP1 0x80 #define JUMP 0x56 #define JUMPI 0x57 diff --git a/include/uint256.cuh b/include/uint256.cuh index eeff53f..cd2d80d 100644 --- a/include/uint256.cuh +++ b/include/uint256.cuh @@ -3,18 +3,17 @@ #ifndef BASE_UINT256_CUH #define BASE_UINT256_CUH +#include +#include #include #include -#include #include -#include -#include +#include #define BITS 256 #define WIDTH (BITS / 32) -typedef struct -{ +typedef struct { uint32_t pn[WIDTH]; } base_uint; @@ -25,7 +24,7 @@ __host__ void intToHex(int num, char *hex); __host__ bool hex_to_decimal(const char *hex_str, char *dec_str); -__host__ __device__ void print_base_uint(const base_uint *val); +__host__ __device__ void print_base_uint(const base_uint *val); // conversion operations __host__ bool base_uint_set_hex(base_uint *val, const char *hex); @@ -36,7 +35,6 @@ __host__ bool int_to_base_uint(int int_val, base_uint *val); __host__ __device__ void base_uint_get_hex(const base_uint *val, char *hex); - // comparison operations __host__ __device__ bool is_zero(const base_uint *num); @@ -55,6 +53,7 @@ __host__ __device__ void base_uint_mul(const base_uint *a, const base_uint *b, b __host__ __device__ void base_uint_shift_left(base_uint *a, size_t bits); -__host__ __device__ void base_uint_div(const base_uint *a, const base_uint *b, base_uint *quotient, base_uint *remainder); +__host__ __device__ void base_uint_div(const base_uint *a, const base_uint *b, base_uint *quotient, + base_uint *remainder); -#endif // BASE_UINT256_CUH +#endif // BASE_UINT256_CUH diff --git a/src/cuevm.cu b/src/cuevm.cu index 0f4e3e0..e7d7384 100644 --- a/src/cuevm.cu +++ b/src/cuevm.cu @@ -1,40 +1,37 @@ +#include +#include #include #include -#include -#include -#include "uint256.cuh" -#include "stack.cuh" + #include "cuevm_test.h" #include "opcode.h" #include "processor.cuh" +#include "stack.cuh" +#include "uint256.cuh" #define NUMTHREAD 4096 #define DEBUG 1 // simple draft kernel for place holder // simple testing opcodes and return the popped top of stack value -__global__ void cuEVM(unsigned char *bytecode, unsigned char *input, size_t bytecode_len, size_t input_len, size_t num_threads) -{ +__global__ void cuEVM(unsigned char *bytecode, unsigned char *input, size_t bytecode_len, size_t input_len, + size_t num_threads) { int idx = threadIdx.x + blockIdx.x * blockDim.x; - if (idx < num_threads) - { - if (idx == 0) - { + if (idx < num_threads) { + if (idx == 0) { printf("Bytecode: "); - for (size_t i = 0; i < bytecode_len; i++) - { + for (size_t i = 0; i < bytecode_len; i++) { printf("%02x ", bytecode[i]); } printf("\n"); printf("Input: "); - for (size_t i = 0; i < input_len; i++) - { + for (size_t i = 0; i < input_len; i++) { printf("%02x ", input[i]); } printf("\n"); // todo refactor processor.execute to a function outside of this kernel processor evm; - evm.programCounter = 0; // redundant but better safe. + evm.programCounter = 0; // redundant but better safe. init_stack(&evm.stack); // push(&stack, a); @@ -42,229 +39,220 @@ __global__ void cuEVM(unsigned char *bytecode, unsigned char *input, size_t byte // debugging : print_stack(&stack); // define 3 reusable temp uints for binary op base_uint op1, op2, result; - uint safe_counter = 0; // safety counter prevent infinite loop - while (evm.programCounter < bytecode_len) - { + uint safe_counter = 0; // safety counter prevent infinite loop + while (evm.programCounter < bytecode_len) { unsigned char opcode = bytecode[evm.programCounter]; - safe_counter ++; + safe_counter++; if (safe_counter > 100) { printf("Safety counter exceeded, return from execution\n"); return; } - switch (opcode) - { - case ADD: // ADD - // TODO: check stack size - // future optimization : can override push pop ops and modify the stack directly - pop(&evm.stack, &op1); - pop(&evm.stack, &op2); - base_uint_add(&op1, &op2, &result); + switch (opcode) { + case ADD: // ADD + // TODO: check stack size + // future optimization : can override push pop ops and modify the stack directly + pop(&evm.stack, &op1); + pop(&evm.stack, &op2); + base_uint_add(&op1, &op2, &result); #if DEBUG - printf("ADD OPCODE: \n"); - printf("op1: "); - print_base_uint(&op1); - printf("op2: "); - print_base_uint(&op2); - printf("result: "); - print_base_uint(&result); - printf("\n***************\n"); + printf("ADD OPCODE: \n"); + printf("op1: "); + print_base_uint(&op1); + printf("op2: "); + print_base_uint(&op2); + printf("result: "); + print_base_uint(&result); + #endif - push(&evm.stack, result); - break; - case SUB: - pop(&evm.stack, &op1); - pop(&evm.stack, &op2); - base_uint_sub(&op1, &op2, &result); + push(&evm.stack, result); + break; + case SUB: + pop(&evm.stack, &op1); + pop(&evm.stack, &op2); + base_uint_sub(&op1, &op2, &result); #if DEBUG - printf("SUB OPCODE: \n"); - printf("op1: "); - print_base_uint(&op1); - printf("op2: "); - print_base_uint(&op2); - printf("result: "); - print_base_uint(&result); - printf("\n***************\n"); + printf("SUB OPCODE: \n"); + printf("op1: "); + print_base_uint(&op1); + printf("op2: "); + print_base_uint(&op2); + printf("result: "); + print_base_uint(&result); + #endif - push(&evm.stack, result); - break; + push(&evm.stack, result); + break; - case MUL: // MUL - // TODO: check stack size - pop(&evm.stack, &op1); - pop(&evm.stack, &op2); - base_uint_mul(&op1, &op2, &result); + case MUL: // MUL + // TODO: check stack size + pop(&evm.stack, &op1); + pop(&evm.stack, &op2); + base_uint_mul(&op1, &op2, &result); #if DEBUG - printf("MUL OPCODE: \n"); - printf("op1: "); - print_base_uint(&op1); - printf("op2: "); - print_base_uint(&op2); - printf("result: "); - print_base_uint(&result); - printf("\n***************\n"); + printf("MUL OPCODE: \n"); + printf("op1: "); + print_base_uint(&op1); + printf("op2: "); + print_base_uint(&op2); + printf("result: "); + print_base_uint(&result); + #endif - push(&evm.stack, result); - break; - case PUSH1: - unsigned char push_val = bytecode[++evm.programCounter]; - result = {{push_val, 0, 0, 0, 0, 0, 0, 0}}; - push(&evm.stack, result); + push(&evm.stack, result); + break; + case PUSH1: + unsigned char push_val = bytecode[++evm.programCounter]; + result = {{push_val, 0, 0, 0, 0, 0, 0, 0}}; + push(&evm.stack, result); #if DEBUG - printf("PUSH1 OPCODE: \n"); - printf("push_val: "); - print_base_uint(&result); - printf("\n***************\n"); + printf("PUSH1 OPCODE: \n"); + printf("push_val: "); + print_base_uint(&result); + #endif - break; - case PUSH2: - // Increment the program counter to point to the first byte of data - evm.programCounter++; + break; + case PUSH2: + // Increment the program counter to point to the first byte of data + evm.programCounter++; - // Read the two bytes from the bytecode - unsigned char byte1 = bytecode[evm.programCounter]; - unsigned char byte2 = bytecode[++evm.programCounter]; + // Read the two bytes from the bytecode + unsigned char byte1 = bytecode[evm.programCounter]; + unsigned char byte2 = bytecode[++evm.programCounter]; - // Combine the two bytes into a single 16-bit value - uint16_t push_val_16 = (byte1 << 8) | byte2; + // Combine the two bytes into a single 16-bit value + uint16_t push_val_16 = (byte1 << 8) | byte2; - // Convert the 16-bit value into your base_uint format - result = {{push_val_16, 0, 0, 0, 0, 0, 0, 0}}; + // Convert the 16-bit value into your base_uint format + result = {{push_val_16, 0, 0, 0, 0, 0, 0, 0}}; - // Push the value onto the stack - push(&evm.stack, result); + // Push the value onto the stack + push(&evm.stack, result); #if DEBUG - printf("PUSH2 OPCODE: \n"); - printf("push_val: "); - print_base_uint(&result); - printf("\n***************\n"); + printf("PUSH2 OPCODE: \n"); + printf("push_val: "); + print_base_uint(&result); + #endif - break; - - case POP: - pop(&evm.stack, &result); - printf("Popped Stack value: "); - print_base_uint(&result); - printf("\n***************\n"); - break; - - case SWAP1: - printf("SWAP1 OPCODE BEFORE: \n"); - print_base_uint(&evm.stack.items[0]); - print_base_uint(&evm.stack.items[1]); - pop(&evm.stack, &op1); - pop(&evm.stack, &op2); - push(&evm.stack, op1); - push(&evm.stack, op2); - printf("\nSWAP1 OPCODE AFTER: \n"); - print_base_uint(&evm.stack.items[0]); - print_base_uint(&evm.stack.items[1]); - break; - case JUMPI: - pop(&evm.stack, &op1); - pop(&evm.stack, &op2); + break; + + case POP: + pop(&evm.stack, &result); + printf("Popped Stack value: "); + print_base_uint(&result); + break; + + case SWAP1: + pop(&evm.stack, &op1); + pop(&evm.stack, &op2); + push(&evm.stack, op1); + push(&evm.stack, op2); + break; + + case DUP1: + printf("DUP1 OPCODE: \n"); + push(&evm.stack, evm.stack.items[evm.stack.top]); + break; + case JUMPI: + pop(&evm.stack, &op1); + pop(&evm.stack, &op2); #if DEBUG - printf("JUMPI OPCODE: \n"); - printf("Condition:\n"); - print_base_uint(&op2); - printf("Destination:\n"); - print_base_uint(&op1); - printf("is ZEROP op1: %d\n", is_zero(&op1)); + printf("JUMPI OPCODE: \n"); + printf("Condition:\n"); + print_base_uint(&op2); + printf("Destination:\n"); + print_base_uint(&op1); + printf("is ZEROP op1: %d\n", is_zero(&op1)); #endif - if (!is_zero(&op2)) - { + if (!is_zero(&op2)) { + evm.programCounter = op1.pn[0]; + // TODO: check JUMPDEST in destiation ? + } + break; + + case JUMP: + pop(&evm.stack, &op1); evm.programCounter = op1.pn[0]; // TODO: check JUMPDEST in destiation ? - } - break; - - case JUMP: - pop(&evm.stack, &op1); - evm.programCounter = op1.pn[0];; - // TODO: check JUMPDEST in destiation ? - break; - - case JUMPDEST: - // do nothing - break; - case RETURN: - printf("RETURN OPCODE\n"); - evm.programCounter = bytecode_len; - break; - default: - printf("Unknown opcode %d at position %d\n", opcode, evm.programCounter); - printf("Return from execution\n"); - return; + break; + + case JUMPDEST: + // do nothing + break; + case RETURN: + printf("RETURN OPCODE\n"); + evm.programCounter = bytecode_len; + break; + default: + printf("Unknown opcode %d at position %d\n", opcode, evm.programCounter); + printf("Return from execution\n"); + return; } evm.programCounter++; +#if DEBUG printf("Program counter: %d\n", evm.programCounter); + printf("Stack size: %d\n", evm.stack.top + 1); + print_stack(&evm.stack); + printf("\n***************\n"); +#endif } } } } -int adjustedLength(char **hexString) -{ - if (strncmp(*hexString, "0x", 2) == 0 || strncmp(*hexString, "0X", 2) == 0) - { - *hexString += 2; // Skip the "0x" prefix +int adjustedLength(char **hexString) { + if (strncmp(*hexString, "0x", 2) == 0 || strncmp(*hexString, "0X", 2) == 0) { + *hexString += 2; // Skip the "0x" prefix return (strlen(*hexString) / 2); } return (strlen(*hexString) / 2); } -void hexStringToByteArray(const char *hexString, unsigned char *byteArray, int length) -{ - for (int i = 0; i < length; i += 2) - { +void hexStringToByteArray(const char *hexString, unsigned char *byteArray, int length) { + for (int i = 0; i < length; i += 2) { sscanf(&hexString[i], "%2hhx", &byteArray[i / 2]); } } -int main(int argc, char *argv[]) -{ +int main(int argc, char *argv[]) { char *byte_code_hex = NULL; char *input_hex = NULL; - static struct option long_options[] = { - {"bytecode", required_argument, 0, 'b'}, - {"input", required_argument, 0, 'i'}, - {"test", no_argument, 0, 't'}, - {0, 0, 0, 0}}; + static struct option long_options[] = {{"bytecode", required_argument, 0, 'b'}, + {"input", required_argument, 0, 'i'}, + {"test", no_argument, 0, 't'}, + {0, 0, 0, 0}}; int opt; int option_index = 0; - while ((opt = getopt_long(argc, argv, "b:i:", long_options, &option_index)) != -1) - { - switch (opt) - { - case 'b': - byte_code_hex = optarg; - break; - case 'i': - input_hex = optarg; - break; - case 't': - test_arithmetic_operations(); - test_stack(); - exit(0); - default: - fprintf(stderr, "Usage: %s --bytecode --input \n", argv[0]); - exit(EXIT_FAILURE); + while ((opt = getopt_long(argc, argv, "b:i:", long_options, &option_index)) != -1) { + switch (opt) { + case 'b': + byte_code_hex = optarg; + break; + case 'i': + input_hex = optarg; + break; + case 't': + test_arithmetic_operations(); + test_stack(); + exit(0); + default: + fprintf(stderr, "Usage: %s --bytecode --input \n", argv[0]); + exit(EXIT_FAILURE); } } - if (!byte_code_hex || !input_hex) - { + if (!byte_code_hex || !input_hex) { fprintf(stderr, "Both --bytecode and --input flags are required\n"); exit(EXIT_FAILURE); } diff --git a/src/cuevm_test.cu b/src/cuevm_test.cu index e2dbe8e..6c6f638 100644 --- a/src/cuevm_test.cu +++ b/src/cuevm_test.cu @@ -1,7 +1,6 @@ #include "cuevm_test.h" -void test_arithmetic_operations() -{ +void test_arithmetic_operations() { base_uint a, b, c, d; // Test addition @@ -15,8 +14,7 @@ void test_arithmetic_operations() printf("%s\n", buffer); - if (strcmp(buffer, "0000000000000000000000000000000033333333333333333333333333333333") != 0) - { + if (strcmp(buffer, "0000000000000000000000000000000033333333333333333333333333333333") != 0) { printf("Addition failed!\n"); } // Test addition with carry @@ -29,8 +27,7 @@ void test_arithmetic_operations() printf("%s\n", buffer); - if (strcmp(buffer, "0000000000000000000000000000000100000000000000000000000000000000") != 0) - { + if (strcmp(buffer, "0000000000000000000000000000000100000000000000000000000000000000") != 0) { printf("Addition failed!\n"); } // Test addition overflow carry @@ -43,8 +40,7 @@ void test_arithmetic_operations() printf("%s\n", buffer); - if (strcmp(buffer, "0000000000000000000000000000000000000000000000000000000000001233") != 0) - { + if (strcmp(buffer, "0000000000000000000000000000000000000000000000000000000000001233") != 0) { printf("Addition failed!\n"); } @@ -55,8 +51,7 @@ void test_arithmetic_operations() printf("Subtraction Result: "); base_uint_get_hex(&c, buffer); printf("%s\n", buffer); - if (strcmp(buffer, "00000000000000000000000000000000ffffffffffffffffffffffffffffff01") != 0) - { + if (strcmp(buffer, "00000000000000000000000000000000ffffffffffffffffffffffffffffff01") != 0) { printf("Subtraction failed!\n"); } @@ -67,8 +62,7 @@ void test_arithmetic_operations() printf("Subtraction Result: "); base_uint_get_hex(&c, buffer); printf("%s\n", buffer); - if (strcmp(buffer, "ffffffffffffffffffffffffffffffffffffffffffffffffffffffffffffff02") != 0) - { + if (strcmp(buffer, "ffffffffffffffffffffffffffffffffffffffffffffffffffffffffffffff02") != 0) { printf("Subtraction failed!\n"); } @@ -79,8 +73,7 @@ void test_arithmetic_operations() printf("Multiplication Result: "); base_uint_get_hex(&c, buffer); printf("%s\n", buffer); - if (strcmp(buffer, "00fffffffffffffffffffffffffffffeff000000000000000000000000000001") != 0) - { + if (strcmp(buffer, "00fffffffffffffffffffffffffffffeff000000000000000000000000000001") != 0) { printf("Multiplication failed!\n"); } // Test multiplication overflow @@ -90,20 +83,17 @@ void test_arithmetic_operations() printf("Multiplication Result: "); base_uint_get_hex(&c, buffer); printf("%s\n", buffer); - if (strcmp(buffer, "0000000000000000000000000000000000000000000000000000000000000000") != 0) - { + if (strcmp(buffer, "0000000000000000000000000000000000000000000000000000000000000000") != 0) { printf("Multiplication overflow failed!\n"); } - } - void test_stack() { base_uint_stack stack; init_stack(&stack); // Test push and print - base_uint a = { {1, 2, 3, 4} }; + base_uint a = {{1, 2, 3, 4}}; printf("Pushing: "); for (int i = 0; i < WIDTH; i++) printf("%u ", a.pn[i]); printf("\n"); @@ -121,7 +111,7 @@ void test_stack() { // Test swap with top push(&stack, a); - base_uint c = { {5, 6, 7, 8} }; + base_uint c = {{5, 6, 7, 8}}; push(&stack, c); printf("Before swap with top:\n"); print_stack(&stack); diff --git a/src/stack.cu b/src/stack.cu index 17a4c72..516dff6 100644 --- a/src/stack.cu +++ b/src/stack.cu @@ -1,8 +1,6 @@ #include "stack.cuh" -__host__ __device__ void init_stack(base_uint_stack* stack) { - stack->top = -1; -} +__host__ __device__ void init_stack(base_uint_stack* stack) { stack->top = -1; } __host__ __device__ bool push(base_uint_stack* stack, base_uint item) { if (stack->top >= STACK_SIZE - 1) { diff --git a/src/uint256.cu b/src/uint256.cu index fd1854f..19ab93c 100644 --- a/src/uint256.cu +++ b/src/uint256.cu @@ -1,24 +1,17 @@ #include "uint256.cuh" // implementation -__host__ int hexToInt(const char *hex) -{ +__host__ int hexToInt(const char *hex) { int result = 0; int len = strlen(hex); - for (int i = 0; i < len; i++) - { + for (int i = 0; i < len; i++) { char c = tolower(hex[i]); - if (c >= '0' && c <= '9') - { + if (c >= '0' && c <= '9') { result = result * 16 + (c - '0'); - } - else if (c >= 'a' && c <= 'f') - { + } else if (c >= 'a' && c <= 'f') { result = result * 16 + (c - 'a' + 10); - } - else - { + } else { // Invalid hexadecimal character return -1; } @@ -26,30 +19,24 @@ __host__ int hexToInt(const char *hex) return result; } -__host__ void intToHex(int num, char *hex) -{ +__host__ void intToHex(int num, char *hex) { // Assuming hex has enough space char *ptr = hex; - do - { + do { int remainder = num % 16; - if (remainder < 10) - { + if (remainder < 10) { *ptr++ = '0' + remainder; - } - else - { + } else { *ptr++ = 'a' + (remainder - 10); } num /= 16; } while (num != 0); - *ptr-- = '\0'; // NULL-terminate the string and point to the last valid character + *ptr-- = '\0'; // NULL-terminate the string and point to the last valid character // Reverse the string char *start = hex; - while (start < ptr) - { + while (start < ptr) { char t = *start; *start = *ptr; *ptr = t; @@ -58,26 +45,19 @@ __host__ void intToHex(int num, char *hex) } } -__host__ bool hex_to_decimal(const char *hex_str, char *dec_str) -{ +__host__ bool hex_to_decimal(const char *hex_str, char *dec_str) { unsigned long long result = 0; unsigned long long place = 1; int len = strlen(hex_str); - for (int i = len - 1; i >= 0; i--) - { + for (int i = len - 1; i >= 0; i--) { char c = tolower(hex_str[i]); int digit; - if (c >= '0' && c <= '9') - { + if (c >= '0' && c <= '9') { digit = c - '0'; - } - else if (c >= 'a' && c <= 'f') - { + } else if (c >= 'a' && c <= 'f') { digit = 10 + (c - 'a'); - } - else - { + } else { return false; } @@ -89,31 +69,23 @@ __host__ bool hex_to_decimal(const char *hex_str, char *dec_str) return true; } -__host__ bool base_uint_set_hex(base_uint *val, const char *hex) -{ +__host__ bool base_uint_set_hex(base_uint *val, const char *hex) { memset(val->pn, 0, sizeof(val->pn)); size_t len = strlen(hex); - if (len == 0 || len > BITS / 4) - return false; + if (len == 0 || len > BITS / 4) return false; // Iterate through the string from end to start - for (size_t i = 0; i < len; i++) - { + for (size_t i = 0; i < len; i++) { char c = tolower(hex[len - 1 - i]); uint32_t number = 0; - if (c >= '0' && c <= '9') - { + if (c >= '0' && c <= '9') { number = c - '0'; - } - else if (c >= 'a' && c <= 'f') - { + } else if (c >= 'a' && c <= 'f') { number = c - 'a' + 10; - } - else - { - return false; // Invalid character + } else { + return false; // Invalid character } // Determine which uint32_t element and position the hex character should be placed @@ -122,99 +94,79 @@ __host__ bool base_uint_set_hex(base_uint *val, const char *hex) return true; } -__host__ void base_uint_to_string(const base_uint *val, char *out_str) -{ +__host__ void base_uint_to_string(const base_uint *val, char *out_str) { char hex_str[BITS / 4 + 1] = {0}; base_uint_get_hex(val, hex_str); - if (!hex_to_decimal(hex_str, out_str)) - { + if (!hex_to_decimal(hex_str, out_str)) { strcpy(out_str, "Error"); } } -__host__ bool int_to_base_uint(int int_val, base_uint *val) -{ +__host__ bool int_to_base_uint(int int_val, base_uint *val) { char *p; sprintf(p, "%08x", int_val); printf("%s\n", p); return base_uint_set_hex(val, p); } -__host__ __device__ void base_uint_get_hex(const base_uint *val, char *hex) -{ +__host__ __device__ void base_uint_get_hex(const base_uint *val, char *hex) { char *p = hex; - for (int i = WIDTH - 1; i >= 0; i--) - { + for (int i = WIDTH - 1; i >= 0; i--) { // printf("%d ", val->pn[i]); sprintf(p, "%08x", val->pn[i]); p += 8; } } -__host__ __device__ void print_base_uint(const base_uint *val) -{ - for (int i = 0; i < WIDTH; i++) - { +__host__ __device__ void print_base_uint(const base_uint *val) { + for (int i = 0; i < WIDTH; i++) { printf("%d ", val->pn[i]); } } -__host__ __device__ bool is_zero(const base_uint *num) -{ - for (int i = 0; i < WIDTH; i++) - { - if (num->pn[i] != 0) - { +__host__ __device__ bool is_zero(const base_uint *num) { + for (int i = 0; i < WIDTH; i++) { + if (num->pn[i] != 0) { return false; } } return true; } -__host__ __device__ base_uint bitwise_not(const base_uint *num) -{ +__host__ __device__ base_uint bitwise_not(const base_uint *num) { base_uint ret; - for (int i = 0; i < WIDTH; i++) - { + for (int i = 0; i < WIDTH; i++) { ret.pn[i] = ~num->pn[i]; } return ret; } -__host__ __device__ void base_uint_set_bit(base_uint *value, uint32_t bitpos) -{ +__host__ __device__ void base_uint_set_bit(base_uint *value, uint32_t bitpos) { value->pn[bitpos / 32] |= (1 << (bitpos % 32)); } -__host__ __device__ void base_uint_add(const base_uint *a, const base_uint *b, base_uint *result) -{ +__host__ __device__ void base_uint_add(const base_uint *a, const base_uint *b, base_uint *result) { uint64_t carry = 0; - for (size_t i = 0; i < WIDTH; i++) - { + for (size_t i = 0; i < WIDTH; i++) { uint64_t sum = (uint64_t)a->pn[i] + b->pn[i] + carry; printf("%d %d = %d %d\n", a->pn[i], b->pn[i], sum, carry); - result->pn[i] = (uint32_t)sum; // Store lower 32 bits - carry = sum >> 32; // Take upper 32 bits as the next carry + result->pn[i] = (uint32_t)sum; // Store lower 32 bits + carry = sum >> 32; // Take upper 32 bits as the next carry } } -__host__ __device__ bool base_uint_sub(const base_uint *a, const base_uint *b, base_uint *result) -{ +__host__ __device__ bool base_uint_sub(const base_uint *a, const base_uint *b, base_uint *result) { uint64_t borrow = 0; - for (size_t i = 0; i < WIDTH; i++) - { + for (size_t i = 0; i < WIDTH; i++) { uint64_t res = 0x100000000ULL + (uint64_t)a->pn[i] - b->pn[i] - borrow; result->pn[i] = (uint32_t)res; - if (res >= 0x100000000ULL) - { + if (res >= 0x100000000ULL) { borrow = 0; - } - else - { + } else { borrow = 1; } } @@ -229,16 +181,12 @@ Warming: 1. Not tested yet. 2. Overflow wraparound is not correctly implemented yet. */ -__host__ __device__ void base_uint_mul(const base_uint *a, const base_uint *b, base_uint *result) -{ +__host__ __device__ void base_uint_mul(const base_uint *a, const base_uint *b, base_uint *result) { base_uint temp_result = {0}; - for (size_t i = 0; i < WIDTH; i++) - { + for (size_t i = 0; i < WIDTH; i++) { uint64_t carry = 0; - for (size_t j = 0; j < WIDTH; j++) - { - if (i + j < WIDTH) - { + for (size_t j = 0; j < WIDTH; j++) { + if (i + j < WIDTH) { uint64_t product = (uint64_t)a->pn[i] * b->pn[j] + temp_result.pn[i + j] + carry; temp_result.pn[i + j] = (uint32_t)product; carry = product >> 32; @@ -246,8 +194,7 @@ __host__ __device__ void base_uint_mul(const base_uint *a, const base_uint *b, b } } - for (size_t i = 0; i < WIDTH; i++) - { + for (size_t i = 0; i < WIDTH; i++) { result->pn[i] = temp_result.pn[i]; } }