Merge branch 'main' into yellow_paper

.clang-format

@@ -0,0 +1,5 @@
+BasedOnStyle: Google
+IndentWidth: 4
+TabWidth: 4
+UseTab: Never
+ColumnLimit: 120

.gitignore

@@ -15,4 +15,6 @@ build/
 _build/
 _static/
 _templates/
-docs/api/
+docs/api/
+*.o
+cuEVM

.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
+}

Makefile

@@ -1,3 +1,4 @@
+# Compiler
 NVCC = nvcc
 NVCC_FLAGS = -I./CGBN/include -lstdc++ -lm -lgmp -lcjson -rdc=true --std c++20 -lcudadevrt -lineinfo
 GCC = gcc
@@ -6,10 +7,6 @@ GPP = g++
 GPP_FLAGS = -I./CGBN/include -lm -lgmp -lcjson 
 OUT_DIRECTORY = ./out
 
-all: cuEVM
-
-cuEVM:
-	$(NVCC) $(NVCC_FLAGS) -o $(OUT_DIRECTORY)/cuEVM src/cu_evm.cu
 
 test_gmp: src/test/test_gmp.c
 	$(GCC) -o $(OUT_DIRECTORY)/test_gmp src/test/test_gmp.c $(GCC_FLAGS) 
@@ -39,4 +36,4 @@ cpu_debug_interpreter: src/interpreter.cu
 	$(NVCC) $(NVCC_FLAGS) -g -G -o $(OUT_DIRECTORY)/$@ $<
 
 clean:
-	rm -f $(OUT_DIRECTORY)/*
+	rm -f $(OUT_DIRECTORY)/*

README.md

@@ -12,7 +12,7 @@ Cuda implementation of EVM bytecode executor
 Or for running with cpu
 * `make cpu_interpreter`
 
-## Usage 
+## Usage
 
 * `clear && compute-sanitizer --tool memcheck --leak-check=full ./out/interpreter --input [inpot_json_file] --output [output_json_file]`
 * `clear && valgrind --leak-check=full --show-leak-kinds=all ./out/*`
@@ -21,6 +21,61 @@ Or for running with cpu
 Easy test:
 * `./out/cpu_interpreter --input ./input/evm_arith.json --output ./output/evm_test.json`
 
+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
+***************
+```
+
+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 6 : 0x600e PUSH 0x0e
+PC 8 : 0x90 SWAP1
+PC 9 : 0x03 SUB  // 42 - 14 // condition != 1 jump;
+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:
+0x60066007025b600e90038060055750f3
+```
+Run
+```
+./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

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

include/opcode.h

@@ -0,0 +1,25 @@
+#ifndef OPCODE_H
+#define OPCODE_H
+
+#define ADD 0x01
+#define MUL 0x02
+#define SUB 0x03
+
+
+#define POP 0x50
+#define PUSH1 0x60
+#define PUSH2 0x61
+
+#define SWAP1 0x90
+
+#define DUP1 0x80
+
+#define JUMP 0x56
+#define JUMPI 0x57
+#define JUMPDEST 0x5b
+
+
+#define RETURN 0xf3
+// Add other opcode definitions here as your VM expands
+
+#endif // OPCODE_H

include/processor.cuh

@@ -0,0 +1,42 @@
+#ifndef PROCESSOR_CUH
+#define PROCESSOR_CUH
+
+#include <stdbool.h>
+#include <stdint.h>
+#include <string.h>
+#include <stdio.h>
+#include <ctype.h>
+#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

include/stack.cuh

@@ -0,0 +1,28 @@
+#ifndef STACK_CUH
+#define STACK_CUH
+
+#include <stdbool.h>
+#include <stdint.h>
+#include <string.h>
+#include <stdio.h>
+#include <ctype.h>
+#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

include/uint256.cuh

@@ -0,0 +1,59 @@
+// base_uint256.cuh
+
+#ifndef BASE_UINT256_CUH
+#define BASE_UINT256_CUH
+
+#include <ctype.h>
+#include <cuda_runtime.h>
+#include <stdbool.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <string.h>
+
+#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