Program Listing for File arith.cuh

↰ Return to documentation for file (src/arith.cuh)

// cuEVM: CUDA Ethereum Virtual Machine implementation
// Copyright 2023 Stefan-Dan Ciocirlan (SBIP - Singapore Blockchain Innovation Programme)
// Author: Stefan-Dan Ciocirlan
// Data: 2023-11-30
// SPDX-License-Identifier: MIT

#ifndef _ARITH_H_
#define _ARITH_H_

#include <stdio.h>
#include <stdint.h>
#include <stdlib.h>
#include <cuda.h>
#include <gmp.h>
#ifndef __CGBN_H__
#define __CGBN_H__
#include <cgbn/cgbn.h>
#endif
#include "data_content.h"

template <class params>
class arith_env_t
{
public:
  typedef cgbn_context_t<params::TPI, params> context_t;
  typedef cgbn_env_t<context_t, params::BITS> env_t;
  typedef typename env_t::cgbn_t bn_t;
  typedef typename env_t::cgbn_wide_t bn_wide_t;
  typedef cgbn_mem_t<params::BITS> evm_word_t;
  static const uint32_t BITS = params::BITS;
  static const uint32_t BYTES = params::BITS / 8;
  static const uint32_t LIMBS = params::BITS / 32;
  static const uint32_t ADDRESS_BYTES = 20;
  context_t _context;
  env_t _env;
  uint32_t _instance;
  __device__ __forceinline__ arith_env_t(
    cgbn_monitor_t monitor,
    cgbn_error_report_t *report,
    uint32_t instance
  ) : _context(monitor, report, instance),
      _env(_context),
      _instance(instance)
  {
  }

  __device__ __forceinline__ arith_env_t(
    cgbn_monitor_t monitor
  ) : _context(monitor),
      _env(_context),
      _instance(0)
  {
  }

  __host__ __forceinline__ arith_env_t(
    cgbn_monitor_t monitor,
    uint32_t instance
  ) : _context(monitor),
      _env(_context),
      _instance(instance)
  {
  }

  __host__ __device__ __forceinline__ arith_env_t(
    const arith_env_t &env
  ) : _context(env._context),
      _env(_context),
      _instance(env._instance)
  {
  }

  __host__ __device__ __forceinline__ void address_conversion(
    bn_t &address
  )
  {
    int32_t address_bits = int32_t(ADDRESS_BYTES) * 8;
    cgbn_bitwise_mask_and(_env, address, address, address_bits);
  }
  __host__ __device__ __forceinline__ size_t memory_from_cgbn(
    uint8_t *dst,
    bn_t &src
  )
  {
    for (uint32_t idx = 0; idx < BYTES; idx++)
    {
      dst[idx] = cgbn_extract_bits_ui32(_env, src, BITS - (idx + 1) * 8, 8);
    }
    return BYTES;
  }

  __host__ __device__ __forceinline__ void cgbn_from_memory(
    bn_t &dst,
    uint8_t *src
  )
  {
    for (uint32_t idx = 0; idx < BYTES; idx++)
    {
      cgbn_insert_bits_ui32(_env, dst, dst, BITS - (idx + 1) * 8, 8, src[idx]);
    }
  }

  __host__ __device__ __forceinline__ void cgbn_from_size_t(
    bn_t &dst,
    size_t src
  )
  {
    cgbn_set_ui32(_env, dst, 0);
    cgbn_insert_bits_ui32(_env, dst, dst, 32, 32, (src >> 32));
    cgbn_insert_bits_ui32(_env, dst, dst, 0, 32, src);
  }

  __host__ __device__ __forceinline__ int32_t size_t_from_cgbn(
    size_t &dst,
    bn_t &src
  )
  {
    bn_t MAX_SIZE_T;
    cgbn_set_ui32(_env, MAX_SIZE_T, 1);
    cgbn_shift_left(_env, MAX_SIZE_T, MAX_SIZE_T, 64);
    dst = 0;
    dst = cgbn_extract_bits_ui32(_env, src, 0, 32);
    dst |= ((size_t)cgbn_extract_bits_ui32(_env, src, 32, 32)) << 32;
    if (cgbn_compare(_env, src, MAX_SIZE_T) >= 0)
    {
      return 1;
    }
    else
    {
      return 0;
    }
  }

  __host__ __device__ __forceinline__ int32_t uint64_t_from_cgbn(
    uint64_t &dst,
    bn_t &src
  )
  {
    bn_t MAX_uint64_t;
    cgbn_set_ui32(_env, MAX_uint64_t, 1);
    cgbn_shift_left(_env, MAX_uint64_t, MAX_uint64_t, 64);
    dst = 0;
    dst = cgbn_extract_bits_ui32(_env, src, 0, 32);
    dst |= ((uint64_t)cgbn_extract_bits_ui32(_env, src, 32, 32)) << 32;
    if (cgbn_compare(_env, src, MAX_uint64_t) >= 0)
    {
      return 1;
    }
    else
    {
      return 0;
    }
  }
  __host__ void hex_string_from_cgbn_memory(
    char *dst_hex_string,
    evm_word_t &src_cgbn_mem,
    uint32_t count = LIMBS
  )
  {
    dst_hex_string[0] = '0';
    dst_hex_string[1] = 'x';
    for (uint32_t idx = 0; idx < count; idx++)
    {
      sprintf(
        dst_hex_string + 2 + idx * 8,
        "%08x",
        src_cgbn_mem._limbs[count - 1 - idx]
      );
    }
    dst_hex_string[count * 8 + 2] = '\0';
  }

  __host__ int32_t cgbn_memory_from_hex_string(
    evm_word_t &dst_cgbn_memory,
    char *src_hex_string
  )
  {
    mpz_t value;
    size_t written;
    mpz_init(value);
    if (
      (src_hex_string[0] == '0') &&
      ((src_hex_string[1] == 'x') || (src_hex_string[1] == 'X'))
    )
    {
      mpz_set_str(value, src_hex_string + 2, 16);
    }
    else
    {
      mpz_set_str(value, src_hex_string, 16);
    }
    if (mpz_sizeinbase(value, 2) > BITS)
    {
      return 1;
    }
    mpz_export(
      dst_cgbn_memory._limbs,
      &written,
      -1,
      sizeof(uint32_t),
      0,
      0,
      value
    );
    while (written < LIMBS)
    {
      dst_cgbn_memory._limbs[written++] = 0;
    }
    mpz_clear(value);
    return 0;
  }

  __host__ __device__ __forceinline__ void cgbn_memory_from_size_t(
    evm_word_t &dst_cgbn_memory,
    size_t src
  )
  {
    bn_t src_cgbn;
    cgbn_from_size_t(src_cgbn, src);
    cgbn_store(_env, &dst_cgbn_memory, src_cgbn);
  }

  __host__ __device__ __forceinline__ void print_cgbn_memory(
    evm_word_t &src_cgbn_memory
  )
  {
    for (uint32_t idx = 0; idx < LIMBS; idx++)
      printf("%08x ", src_cgbn_memory._limbs[LIMBS - 1 - idx]);
    printf("\n");
  }

  __host__ __device__ __forceinline__ int32_t has_gas(
    bn_t &gas_limit,
    bn_t &gas_used,
    uint32_t &error_code
  )
  {
    int32_t gas_sign = cgbn_compare(_env, gas_limit, gas_used);
    error_code = (gas_sign < 0) ? ERROR_GAS_LIMIT_EXCEEDED : error_code;
    return (gas_sign >= 0) && (error_code == ERR_NONE);
  }

  __host__ __device__ __forceinline__ void max_gas_call(
    bn_t &gas_capped,
    bn_t &gas_limit,
    bn_t &gas_used
  )
  {
      // compute the remaining gas
      bn_t gas_left;
      cgbn_sub(_env, gas_left, gas_limit, gas_used);

      // gas capped = (63/64) * gas_left
      cgbn_div_ui32(_env, gas_capped, gas_left, 64);
      cgbn_sub(_env, gas_capped, gas_left, gas_capped);
  }

  __host__ __device__ __forceinline__ uint8_t *get_data(
    data_content_t &data_content,
    bn_t &index,
    bn_t &length,
    size_t &available_size
  )
  {
    available_size = 0;
    size_t index_s;
    int32_t overflow = size_t_from_cgbn(index_s, index);
    if (
        (overflow != 0) ||
        (index_s >= data_content.size))
    {
        return NULL;
    }
    else
    {
        size_t length_s;
        overflow = size_t_from_cgbn(length_s, length);
        if (
            (overflow != 0) ||
            (length_s > data_content.size - index_s))
        {
            available_size = data_content.size - index_s;
            return data_content.data + index_s;
        }
        else
        {
            available_size = length_s;
            return data_content.data + index_s;
        }
    }
  }

  __host__ __device__ __forceinline__ void evm_words_gas_cost
  (
    bn_t &gas_used,
    bn_t &length,
    uint32_t gas_per_word
  )
  {
    // gas_used += gas_per_word * emv word count of length
    // length = (length + 31) / 32
    bn_t evm_words_gas;
    cgbn_add_ui32(_env, evm_words_gas, length, BYTES -1);
    cgbn_div_ui32(_env, evm_words_gas, evm_words_gas, BYTES);
    cgbn_mul_ui32(_env, evm_words_gas, evm_words_gas, gas_per_word);
    cgbn_add(_env, gas_used, gas_used, evm_words_gas);
  }

  __host__ __device__ __forceinline__ void initcode_cost(
    bn_t &gas_used,
    bn_t &initcode_length
  )
  {
    // gas_used += GAS_INITCODE_WORD_COST * emv word count of initcode
    // length = (initcode_length + 31) / 32
    evm_words_gas_cost(gas_used, initcode_length, GAS_INITCODE_WORD_COST);
  }

  __host__ __device__ __forceinline__ void keccak_cost(
    bn_t &gas_used,
    bn_t &length
  )
  {
    evm_words_gas_cost(gas_used, length, GAS_KECCAK256_WORD);
  }



};

#endif