Merge remote-tracking branch 'origin/main'

sbip-sg · Dec 5, 2024 · 907199c · 907199c
2 parents 789b8e7 + 3cb3af2
commit 907199c
Show file tree

Hide file tree

Showing 96 changed files with 5,751 additions and 2,255 deletions.
diff --git a/.gitignore b/.gitignore
@@ -23,4 +23,7 @@ dtmp
 cuevm.out
 geth.out
 *.out
-*.jsonl
+*.jsonl
+
+# fuzzing
+**/__pycache__/
diff --git a/.gitmodules b/.gitmodules
@@ -3,7 +3,7 @@
 	url = https://github.com/sdcioc/CuBigInt.git
 [submodule "CuEVM/libraries/CGBN"]
 	path = CuEVM/libraries/CGBN
-	url = https://github.com/sdcioc/CGBN.git
+	url = https://github.com/sbip-sg/CGBN.git
 [submodule "CuEVM/libraries/CuCrypto"]
 	path = CuEVM/libraries/CuCrypto
 	url = https://github.com/sdcioc/CuCrypto.git
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -13,8 +13,9 @@ option(ENABLE_PAIRING_CODE "Enable ec pairing precompile code" OFF)
 option(ENABLE_EIP_3155_OPTIONAL "Enable EIP_3155 optional tracer" ON)
 option(TESTS "Enable tests" ON)
 option(TESTS_GPU "Enable GPU tests" OFF)
+option(BUILD_LIBRARY "Build Python library" OFF)
 set(EVM_VERSION "SHANGHAI" CACHE STRING "EVM version to use")
-
+set(CGBN_TPI "32" CACHE STRING "Thread per instance for CGBN")
 project("cuevm"
         VERSION 0.1.0
         DESCRIPTION "A CUDA EVM Interpreter"
@@ -35,8 +36,11 @@ if(DOCS)
         add_subdirectory("${PROJECT_SOURCE_DIR}/docs")
 endif()
 
-if (CPU OR GPU OR TESTS)
+if (CPU OR GPU OR TESTS OR BUILD_LIBRARY)
     add_subdirectory("${PROJECT_SOURCE_DIR}/CuEVM")
+    if(BUILD_LIBRARY)
+        target_compile_definitions(CuEVM PRIVATE BUILD_LIBRARY)
+    endif()
     if(ENABLE_EIP_3155)
         target_compile_definitions(CuEVM PRIVATE EIP_3155)
     endif()
@@ -58,20 +62,20 @@ if(CPU)
         $<TARGET_OBJECTS:CGBN>
         $<TARGET_OBJECTS:CuBigInt>
     )
-    
+
     set_target_properties(${CPU_INTERPRETER}
         PROPERTIES
             CUDA_SEPARABLE_COMPILATION ON
             CUDA_RESOLVE_DEVICE_SYMBOLS ON
             POSITION_INDEPENDENT_CODE ON
     )
-    
+
     # Link the library to its dependencies
     # our own libraries first
     target_link_libraries(${CPU_INTERPRETER} PRIVATE CGBN CuCrypto CuBigInt CuEVM)
     # then the external ones
     target_link_libraries(${CPU_INTERPRETER} PUBLIC gmp cjson ${CUDA_LIBRARIES})
-    
+
     target_compile_options(${CPU_INTERPRETER}  PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-lineinfo --std=c++20 -rdc=true --expt-relaxed-constexpr>)
     if(ENABLE_EIP_3155)
         target_compile_definitions(${CPU_INTERPRETER} PRIVATE EIP_3155)
@@ -80,43 +84,99 @@ if(CPU)
         target_compile_definitions(${CPU_INTERPRETER} PRIVATE EIP_3155_OPTIONAL)
     endif()
     target_compile_definitions(${CPU_INTERPRETER} PRIVATE EVM_VERSION=${EVM_VERSION})
+
+    target_compile_definitions(${CPU_INTERPRETER} PRIVATE CGBN_TPI=${CGBN_TPI})
 endif()
 
 if(GPU)
-    set(GPU_INTERPRETER ${PROJECT_NAME}_GPU)
 
-    # GPU
-    add_executable(${GPU_INTERPRETER}
-        "${PROJECT_SOURCE_DIR}/interpreter/interpreter.cu"
-        $<TARGET_OBJECTS:CuEVM>
-        $<TARGET_OBJECTS:CuCrypto>
-        $<TARGET_OBJECTS:CGBN>
-        $<TARGET_OBJECTS:CuBigInt>
-    )
-
-    set_target_properties(${GPU_INTERPRETER}
+    if(BUILD_LIBRARY)
+        # Find Python libraries
+        find_package(PythonLibs REQUIRED)
+        include_directories(${PYTHON_INCLUDE_DIRS})
+        # Build as a shared library
+        add_library(${PROJECT_NAME} SHARED
+            "${PROJECT_SOURCE_DIR}/interpreter/libcuevm.cu"
+            "${PROJECT_SOURCE_DIR}/interpreter/interpreter.cu"
+            "${PROJECT_SOURCE_DIR}/interpreter/python_utils.cu"
+            $<TARGET_OBJECTS:CuEVM>
+            $<TARGET_OBJECTS:CuCrypto>
+            $<TARGET_OBJECTS:CGBN>
+            $<TARGET_OBJECTS:CuBigInt>
+        )
+
+        # Apply flags for library
+        target_compile_options(${PROJECT_NAME} PRIVATE  $<$<COMPILE_LANGUAGE:CUDA>:--shared -Xcompiler '-fPIC'>)
+        if(CMAKE_BUILD_TYPE MATCHES Debug)
+            target_compile_options(${PROJECT_NAME} PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-g>)
+        endif()
+        # set_target_properties(${PROJECT_NAME} PROPERTIES OUTPUT_NAME "cuevm")
+        set_target_properties(${PROJECT_NAME}
         PROPERTIES
+            OUTPUT_NAME "cuevm"
             CUDA_SEPARABLE_COMPILATION ON
             CUDA_RESOLVE_DEVICE_SYMBOLS ON
-            POSITION_INDEPENDENT_CODE ON
-    )
-
-    # Link the library to its dependencies
-    # our own libraries first
-    target_link_libraries(${GPU_INTERPRETER} PRIVATE CGBN CuCrypto CuBigInt CuEVM)
-    # then the external ones
-    target_link_libraries(${GPU_INTERPRETER} PUBLIC gmp cjson ${CUDA_LIBRARIES})
-
-    target_compile_options(${GPU_INTERPRETER}  PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-lineinfo --std=c++20 -rdc=true --expt-relaxed-constexpr>)
-    if(ENABLE_EIP_3155)
-        target_compile_definitions(${GPU_INTERPRETER} PRIVATE EIP_3155)
-    endif()
-    if(ENABLE_EIP_3155_OPTIONAL)
-        target_compile_definitions(${GPU_INTERPRETER} PRIVATE EIP_3155_OPTIONAL)
-    endif()
-    target_compile_definitions(${GPU_INTERPRETER} PRIVATE EVM_VERSION=${EVM_VERSION})
+            CUDA_RUNTIME_LIBRARY Shared  # Ensure CUDA runtime library is shared
+        )
+
+        # Link the library to its dependencies
+        # our own libraries first
+        target_link_libraries(${PROJECT_NAME} PRIVATE CGBN CuCrypto CuBigInt CuEVM)
+        # then the external ones
+        target_link_libraries(${PROJECT_NAME} PUBLIC gmp cjson ${PYTHON_LIBRARIES} ${CUDA_LIBRARIES})
+
+        target_compile_options(${PROJECT_NAME} PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-lineinfo --std=c++20 -rdc=true --expt-relaxed-constexpr>)
+        if(ENABLE_EIP_3155)
+            target_compile_definitions(${PROJECT_NAME} PRIVATE EIP_3155)
+        endif()
+        if(ENABLE_EIP_3155_OPTIONAL)
+            target_compile_definitions(${PROJECT_NAME} PRIVATE EIP_3155_OPTIONAL)
+        endif()
+        target_compile_definitions(${PROJECT_NAME} PRIVATE EVM_VERSION=${EVM_VERSION})
+
+        target_compile_definitions(${PROJECT_NAME} PRIVATE GPU)
+
+        target_compile_definitions(${PROJECT_NAME} PRIVATE CGBN_TPI=${CGBN_TPI})
+
+    else()
+
+        set(GPU_INTERPRETER ${PROJECT_NAME}_GPU)
+        # GPU
+        add_executable(${GPU_INTERPRETER}
+            "${PROJECT_SOURCE_DIR}/interpreter/interpreter.cu"
+            $<TARGET_OBJECTS:CuEVM>
+            $<TARGET_OBJECTS:CuCrypto>
+            $<TARGET_OBJECTS:CGBN>
+            $<TARGET_OBJECTS:CuBigInt>
+        )
+
+        set_target_properties(${GPU_INTERPRETER}
+            PROPERTIES
+                CUDA_SEPARABLE_COMPILATION ON
+                CUDA_RESOLVE_DEVICE_SYMBOLS ON
+                POSITION_INDEPENDENT_CODE ON
+        )
+
+        # Link the library to its dependencies
+        # our own libraries first
+        target_link_libraries(${GPU_INTERPRETER} PRIVATE CGBN CuCrypto CuBigInt CuEVM)
+        # then the external ones
+        target_link_libraries(${GPU_INTERPRETER} PUBLIC gmp cjson ${CUDA_LIBRARIES})
+
+        target_compile_options(${GPU_INTERPRETER} PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-lineinfo --std=c++20 -rdc=true --expt-relaxed-constexpr>)
+        if(ENABLE_EIP_3155)
+            target_compile_definitions(${GPU_INTERPRETER} PRIVATE EIP_3155)
+        endif()
+        if(ENABLE_EIP_3155_OPTIONAL)
+            target_compile_definitions(${GPU_INTERPRETER} PRIVATE EIP_3155_OPTIONAL)
+        endif()
+        target_compile_definitions(${GPU_INTERPRETER} PRIVATE EVM_VERSION=${EVM_VERSION})
+
+        target_compile_definitions(${GPU_INTERPRETER} PRIVATE GPU)
+
+        target_compile_definitions(${GPU_INTERPRETER} PRIVATE CGBN_TPI=${CGBN_TPI})
+    endif() 
 
-    target_compile_definitions(${GPU_INTERPRETER} PRIVATE GPU)
 endif()
 
 if(TESTS)

diff --git a/CuEVM/CMakeLists.txt b/CuEVM/CMakeLists.txt
@@ -29,7 +29,7 @@ if (NOT CMAKE_CUDA_ARCHITECTURES)
     set(CMAKE_CUDA_ARCHITECTURES ${CUDA_COMPUTE_CAPABILITY})
 endif()
 set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR})
-
+set(CGBN_TPI "32" CACHE STRING "Thread per instance for CGBN")
 
 find_package(GMP REQUIRED)
 find_package(cJSON REQUIRED) # Assuming cJSON is available as a find package
@@ -74,5 +74,5 @@ target_link_libraries(${PROJECT_NAME} PUBLIC gmp cjson ${CUDA_LIBRARIES})
 
 # Add specific NVCC flags using target_compile_options (if necessary)
 target_compile_options(${PROJECT_NAME}  PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-lineinfo --std=c++20 -rdc=true --expt-relaxed-constexpr>)
-
+target_compile_definitions(${PROJECT_NAME} PRIVATE CGBN_TPI=${CGBN_TPI})
 
diff --git a/CuEVM/include/CuEVM/core/byte_array.cuh b/CuEVM/include/CuEVM/core/byte_array.cuh
@@ -11,6 +11,7 @@
 #include <stdio.h>
 #include <stdlib.h>
 
+#include <CuEVM/utils/cuda_utils.cuh>
 namespace CuEVM {
 enum PaddingDirection { NO_PADDING = 0, LEFT_PADDING = 1, RIGHT_PADDING = 2 };
 
@@ -200,8 +201,13 @@ struct byte_array_t {
      */
 
     __host__ __device__ static void reset_return_data(byte_array_t *&return_data_ptr) {
-        if (return_data_ptr != nullptr) delete return_data_ptr;
-        return_data_ptr = new byte_array_t();
+        return_data_ptr->free();
+        // printf("reset return data thread %d", THREADIDX);
+        // __ONE_GPU_THREAD_WOSYNC_BEGIN__
+        // if (return_data_ptr->data != nullptr) delete[] return_data_ptr->data;
+        // return_data_ptr->data = nullptr;
+        // __ONE_GPU_THREAD_WOSYNC_END__
+        // return_data_ptr->size = 0;
     }
 
    private:

diff --git a/CuEVM/include/CuEVM/core/evm_word.cuh b/CuEVM/include/CuEVM/core/evm_word.cuh
@@ -4,7 +4,6 @@
 // SPDX-License-Identifier: MIT
 #pragma once
 
-#include <CGBN/cgbn.h>
 #include <cuda.h>
 #include <stdint.h>
 
@@ -30,6 +29,7 @@ struct evm_word_t : cgbn_mem_t<CuEVM::word_bits> {
      */
     __host__ __device__ evm_word_t(uint32_t value);
 
+    __host__ __device__ void set_zero();
     /**
      * The assignment operator.
      * @param[in] src The source evm_word_t
@@ -54,6 +54,7 @@ struct evm_word_t : cgbn_mem_t<CuEVM::word_bits> {
      * @return 1 for equal, 0 otherwise
      */
     __host__ __device__ int32_t operator==(const uint32_t &value) const;
+    __host__ __device__ int32_t operator<(const uint32_t &value) const;
     /**
      * Set the evm_word_t from a hex string.
      * The hex string is in Big Endian format.
@@ -62,13 +63,21 @@ struct evm_word_t : cgbn_mem_t<CuEVM::word_bits> {
      */
     __host__ int32_t from_hex(const char *hex_string);
     /**
-     * Set the evm_word_t from a byte array.
+     * Set the evm_word_t from a byte array, optimized for parallel threads.
      * The byte array is in Big Endian format.
      * @param[in] byte_array The source byte array
      * @param[in] endian The endian format
      * @return 0 for success, 1 otherwise
      */
     __host__ __device__ int32_t from_byte_array_t(byte_array_t &byte_array, int32_t endian = LITTLE_ENDIAN);
+    /**
+     * Set the evm_word_t from a byte array. Compatible with CPU version.
+     * The byte array is in Big Endian format.
+     * @param[in] byte_array The source byte array
+     * @param[in] endian The endian format
+     * @return 0 for success, 1 otherwise
+     */
+    __host__ int32_t from_byte_array_t_loop(byte_array_t &byte_array, int32_t endian = LITTLE_ENDIAN);
     /**
      * Set the evm_word_t from a size_t.
      * @param[in] value The source size_t
@@ -115,6 +124,7 @@ struct evm_word_t : cgbn_mem_t<CuEVM::word_bits> {
     __host__ __device__ char *to_hex(char *hex_string = nullptr, int32_t pretty = 0,
                                      uint32_t count = CuEVM::cgbn_limbs) const;
 
+    __host__ __device__ char *address_to_hex(char *hex_string = nullptr, uint32_t count = CuEVM::cgbn_limbs) const;
     __host__ __device__ void print_as_compact_hex() const;
 
     /**

diff --git a/CuEVM/include/CuEVM/core/jump_destinations.cuh b/CuEVM/include/CuEVM/core/jump_destinations.cuh
@@ -19,29 +19,33 @@ namespace CuEVM {
  */
 struct jump_destinations_t {
    private:
-    byte_array_t
-        destinations; /**< The array of valid JUMPDESTs pc \f$D(c)\f$ */
-
+    uint16_t* destinations; /**< The array of valid JUMPDESTs pc \f$D(c)\f$ */
+    uint32_t capacity;
+    uint32_t size; /**< The size of the jump dest, smaller than capacity*/
    public:
     /**
      * Cosntructor from the given code
      * @param[in] byte_code The code
      */
-    __host__ __device__ jump_destinations_t(CuEVM::byte_array_t &byte_code);
+    __host__ __device__ jump_destinations_t(CuEVM::byte_array_t& byte_code);
 
     /**
      * Destructor of the class
      * free the alocated memory
      */
     __host__ __device__ ~jump_destinations_t();
 
+    __host__ __device__ void set_bytecode(CuEVM::byte_array_t& byte_code);
+
     /**
      * Find out if a given pc is a valid JUMPDEST
      * @param[in] pc The pc to check
      * @return 0 if is a valid JUMPDEST, 1 otherwise
      */
     __host__ __device__ uint32_t has(uint32_t pc);
 
+    __host__ __device__ int32_t grow_capacity(uint32_t new_capacity);
+
     /**
      * Print the destinations
      */