Changeset 5260


Ignore:
Timestamp:
Jan 14, 2017, 3:49:56 PM (10 months ago)
Author:
nmedfort
Message:

Changes working towards simplifying accessing stream elements + some modifications to simplify include / forward declarations within the CodeGen? library.

Location:
icGREP/icgrep-devel/icgrep
Files:
50 edited

Legend:

Unmodified
Added
Removed
  • icGREP/icgrep-devel/icgrep/CMakeLists.txt

    r5252 r5260  
    6060ENDIF()
    6161
    62 
    6362add_library(CodeGen ${KERNEL_SRC} ${IDISA_SRC} object_cache.cpp)
    6463add_library(PabloADT ${PABLO_SRC})
     
    6867add_library(UCDlib UCD/unicode_set.cpp UCD/ucd_compiler.cpp UCD/PropertyObjects.cpp UCD/resolve_properties.cpp UCD/UnicodeNameData.cpp)
    6968
    70 
    7169# add the executable
    7270target_link_libraries (PabloADT CodeGen ${REQ_LLVM_LIBRARIES})
     
    8280add_executable(editd editd/editd.cpp editd/pattern_compiler.cpp toolchain.cpp editd/editdscan_kernel.cpp editd/editd_gpu_kernel.cpp editd/editd_cpu_kernel.cpp)
    8381add_executable(array-test array-test.cpp toolchain.cpp)
     82
     83## IWYU detects superfluous includes and when the include can be replaced with a forward declaration.
     84## It can be obtained using "apt-get install iwyu" or from "github.com/include-what-you-use".
     85
     86#find_program(IWYU_PATH NAMES include-what-you-use iwyu)
     87#if(IWYU_PATH)
     88#cmake_minimum_required(VERSION 3.3 FATAL_ERROR)
     89#execute_process(COMMAND ${CMAKE_CXX_COMPILER} -print-libgcc-file-name OUTPUT_VARIABLE LIBGCC_FILE)
     90#get_filename_component(LIBGCC_PATH ${LIBGCC_FILE} DIRECTORY)
     91#include_directories("${LIBGCC_PATH}/include")
     92#set_property(TARGET CodeGen PROPERTY CXX_INCLUDE_WHAT_YOU_USE ${IWYU_PATH})
     93#set_property(TARGET PabloADT PROPERTY CXX_INCLUDE_WHAT_YOU_USE ${IWYU_PATH})
     94#set_property(TARGET RegExpADT PROPERTY CXX_INCLUDE_WHAT_YOU_USE ${IWYU_PATH})
     95#set_property(TARGET RegExpCompiler PROPERTY CXX_INCLUDE_WHAT_YOU_USE ${IWYU_PATH})
     96#set_property(TARGET CCADT PROPERTY CXX_INCLUDE_WHAT_YOU_USE ${IWYU_PATH})
     97#set_property(TARGET UCDlib PROPERTY CXX_INCLUDE_WHAT_YOU_USE ${IWYU_PATH})
     98#set_property(TARGET icgrep PROPERTY CXX_INCLUDE_WHAT_YOU_USE ${IWYU_PATH})
     99#set_property(TARGET u8u16 PROPERTY CXX_INCLUDE_WHAT_YOU_USE ${IWYU_PATH})
     100#set_property(TARGET wc PROPERTY CXX_INCLUDE_WHAT_YOU_USE ${IWYU_PATH})
     101#set_property(TARGET editd PROPERTY CXX_INCLUDE_WHAT_YOU_USE ${IWYU_PATH})
     102#set_property(TARGET array-test PROPERTY CXX_INCLUDE_WHAT_YOU_USE ${IWYU_PATH})
     103#endif()
    84104
    85105IF (PRINT_TIMING_INFORMATION)
  • icGREP/icgrep-devel/icgrep/IR_Gen/CBuilder.cpp

    r5245 r5260  
    77#include "CBuilder.h"
    88#include <llvm/IR/IRBuilder.h>
     9#include <llvm/IR/Module.h>
    910#include <llvm/IR/Constants.h>
    1011#include <llvm/IR/Intrinsics.h>
     
    1314#include <llvm/IR/TypeBuilder.h>
    1415
    15 
     16using namespace llvm;
    1617
    1718// ssize_t write(int fildes, const void *buf, size_t nbyte);
     
    290291    return CreateCall(pthreadJoinFunc, {thread, value_ptr});
    291292}
     293
     294CBuilder::CBuilder(llvm::Module * m, unsigned archBitWidth, unsigned CacheAlignment)
     295: IRBuilder<>(m->getContext())
     296, mMod(m)
     297, mCacheLineAlignment(CacheAlignment)
     298, mSizeType(getIntNTy(archBitWidth)) {
     299}
  • icGREP/icgrep-devel/icgrep/IR_Gen/CBuilder.h

    r5245 r5260  
    66#define CBUILDER_H
    77
    8 #include <llvm/IR/Module.h>
    9 #include <llvm/IR/Constant.h>
    10 #include <llvm/IR/Type.h>
    11 #include <llvm/IR/Value.h>
     8#include <string>
    129#include <llvm/IR/IRBuilder.h>
    13 #include <llvm/Support/Host.h>
    14 #include <llvm/ADT/Triple.h>
    15 #include <IR_Gen/types/streamtype.h>
    16 #include <boost/container/flat_map.hpp>
     10#include <llvm/IR/Constants.h>
     11namespace llvm { class Function; }
     12namespace llvm { class IntegerType; }
     13namespace llvm { class Module; }
     14namespace llvm { class PointerType; }
     15namespace llvm { class Type; }
     16namespace llvm { class Value; }
    1717
    18 using namespace llvm;
    19 
    20 
    21 class CBuilder : public IRBuilder<> {
     18class CBuilder : public llvm::IRBuilder<> {
    2219   
    2320public:
    2421   
    25     CBuilder(Module * m, unsigned archBitWidth, unsigned CacheAlignment=64)
    26     : IRBuilder<>(m->getContext())
    27     , mMod(m)
    28     , mCacheLineAlignment(CacheAlignment)
    29     , mSizeType(getIntNTy(archBitWidth)) {
    30     }
     22    CBuilder(llvm::Module * m, unsigned archBitWidth, unsigned CacheAlignment=64);
    3123   
    3224    virtual ~CBuilder() {}
    3325
    34     Module * getModule() const {
     26    llvm::Module * getModule() const {
    3527        return mMod;
    3628    }
    3729   
    38     void setModule(Module * m)  {
     30    void setModule(llvm::Module * m)  {
    3931        mMod = m;
    4032    }
    4133   
    42     Function * GetPrintf();
    43     Value * CreateMalloc(Type * type, Value * size);
    44     Value * CreateAlignedMalloc(Type * type, Value * size, const unsigned alignment);
    45     void CreateFree(Value * const ptr);
    46     void CreateAlignedFree(Value * const ptr, const bool ptrMayBeNull = false);
    47     Value * CreateRealloc(Value * ptr, Value * size);
    48     void CreateMemZero(Value * ptr, Value * size, const unsigned alignment = 1);
     34    llvm::Function * GetPrintf();
     35    llvm::Value * CreateMalloc(llvm::Type * type, llvm::Value * size);
     36    llvm::Value * CreateAlignedMalloc(llvm::Type * type, llvm::Value * size, const unsigned alignment);
     37    void CreateFree(llvm::Value * const ptr);
     38    void CreateAlignedFree(llvm::Value * const ptr, const bool ptrMayBeNull = false);
     39    llvm::Value * CreateRealloc(llvm::Value * ptr, llvm::Value * size);
     40    void CreateMemZero(llvm::Value * ptr, llvm::Value * size, const unsigned alignment = 1);
    4941
    5042    inline llvm::AllocaInst * CreateCacheAlignedAlloca(llvm::Type * Ty, llvm::Value * ArraySize = nullptr) {
     
    5749    //
    5850    // ssize_t write(int fildes, const void *buf, size_t nbyte);
    59     Value * CreateWriteCall(Value * fildes, Value * buf, Value * nbyte);
     51    llvm::Value * CreateWriteCall(llvm::Value * fildes, llvm::Value * buf, llvm::Value * nbyte);
    6052   
    6153    // Create calls to Posix thread (pthread.h) functions.
     
    6355    //  int pthread_create(pthread_t *thread, const pthread_attr_t *attr,
    6456    //                    void *(*start_routine)(void*), void *arg);
    65     Value * CreatePThreadCreateCall(Value * thread, Value * attr, Function * start_routine, Value * arg);
     57    llvm::Value * CreatePThreadCreateCall(llvm::Value * thread, llvm::Value * attr, llvm::Function * start_routine, llvm::Value * arg);
    6658    //  void pthread_exit(void *value_ptr);
    67     Value * CreatePThreadExitCall(Value * value_ptr);
     59    llvm::Value * CreatePThreadExitCall(llvm::Value * value_ptr);
    6860    //  int pthread_join(pthread_t thread, void **value_ptr);
    69     Value * CreatePThreadJoinCall(Value * thread, Value * value_ptr);
     61    llvm::Value * CreatePThreadJoinCall(llvm::Value * thread, llvm::Value * value_ptr);
    7062   
    71     void CallPrintInt(const std::string & name, Value * const value);
     63    void CallPrintInt(const std::string & name, llvm::Value * const value);
    7264   
    7365    inline llvm::IntegerType * getSizeTy() const {
     
    7668   
    7769    inline llvm::ConstantInt * getSize(const size_t value) const {
    78         return ConstantInt::get(getSizeTy(), value);
     70        return llvm::ConstantInt::get(getSizeTy(), value);
    7971    }
    8072   
    81     PointerType * getVoidPtrTy() const;
     73    llvm::PointerType * getVoidPtrTy() const;
    8274   
    8375    inline unsigned getCacheAlignment() const {
     
    8577    }
    8678   
    87     virtual llvm::LoadInst* CreateAtomicLoadAcquire(Value * ptr);
    88     virtual llvm::StoreInst *  CreateAtomicStoreRelease(Value * val, Value * ptr);
     79    virtual llvm::LoadInst* CreateAtomicLoadAcquire(llvm::Value * ptr);
     80
     81    virtual llvm::StoreInst *  CreateAtomicStoreRelease(llvm::Value * val, llvm::Value * ptr);
    8982   
    9083protected:
    91     Module *            mMod;
     84    llvm::Module *      mMod;
    9285    unsigned            mCacheLineAlignment;
    93     IntegerType *      mSizeType;
     86    llvm::IntegerType * mSizeType;
    9487};
    9588
  • icGREP/icgrep-devel/icgrep/IR_Gen/idisa_avx_builder.cpp

    r5115 r5260  
    1010#include <llvm/IR/Intrinsics.h>
    1111#include <llvm/IR/Function.h>
    12 #include <iostream>
     12#include <llvm/IR/Module.h>
    1313
    1414namespace IDISA {
  • icGREP/icgrep-devel/icgrep/IR_Gen/idisa_builder.cpp

    r5245 r5260  
    1010#include <llvm/IR/Intrinsics.h>
    1111#include <llvm/IR/Function.h>
     12#include <llvm/IR/Module.h>
    1213#include <llvm/Support/raw_ostream.h>
    1314#include <llvm/IR/TypeBuilder.h>
     15
     16using namespace llvm;
    1417
    1518namespace IDISA {
     
    378381}
    379382
    380 Type * IDISA_Builder::getStreamTy(const unsigned FieldWidth) {
     383StreamType * IDISA_Builder::getStreamTy(const unsigned FieldWidth) {
    381384    const auto f = mStreamTypes.find(FieldWidth);
    382385    if (LLVM_LIKELY(f != mStreamTypes.end())) {
  • icGREP/icgrep-devel/icgrep/IR_Gen/idisa_builder.h

    r5245 r5260  
    77 *  icgrep is a trademark of International Characters.
    88 */
    9 #include <llvm/IR/Module.h>
    10 #include <llvm/IR/Constant.h>
    11 #include <llvm/IR/Type.h>
    12 #include <llvm/IR/Value.h>
    13 #include <llvm/IR/IRBuilder.h>
    14 #include <llvm/Support/Host.h>
    15 #include <llvm/ADT/Triple.h>
    16 #include <IR_Gen/CBuilder.h>
     9#include "CBuilder.h"
     10#include <llvm/IR/DerivedTypes.h>
    1711#include <IR_Gen/types/streamtype.h>
    1812#include <boost/container/flat_map.hpp>
    19 
    20 using namespace llvm;
     13namespace llvm { class Constant; }
     14namespace llvm { class LoadInst; }
     15namespace llvm { class Module; }
     16namespace llvm { class Type; }
     17namespace llvm { class Value; }
    2118
    2219namespace IDISA {
     
    2825public:
    2926
    30     IDISA_Builder(Module * m, unsigned archBitWidth, unsigned bitBlockWidth, unsigned stride, unsigned CacheAlignment=64);
     27    IDISA_Builder(llvm::Module * m, unsigned archBitWidth, unsigned bitBlockWidth, unsigned stride, unsigned CacheAlignment=64);
    3128
    3229    virtual ~IDISA_Builder();
     
    3431    std::string getBitBlockTypeName() const;  // A short string such as v4i64 or i256.
    3532
    36     Value * bitCast(Value * a) {
     33    llvm::Value * bitCast(llvm::Value * a) {
    3734        return (a->getType() == mBitBlockType) ? a : CreateBitCast(a, mBitBlockType);
    3835    }
     
    4643    }
    4744
    48     Constant * allZeroes() const {
     45    llvm::Constant * allZeroes() const {
    4946        return mZeroInitializer;
    5047    }
    5148
    52     Constant * allOnes() const {
     49    llvm::Constant * allOnes() const {
    5350        return mOneInitializer;
    5451    }
    5552   
    5653
    57     LoadInst * CreateBlockAlignedLoad(Value * const ptr);
    58     LoadInst * CreateBlockAlignedLoad(Value * const ptr, Value * const index);
    59     LoadInst * CreateBlockAlignedLoad(Value * const ptr, std::initializer_list<Value *> indices);
     54    llvm::LoadInst * CreateBlockAlignedLoad(llvm::Value * const ptr);
     55    llvm::LoadInst * CreateBlockAlignedLoad(llvm::Value * const ptr, llvm::Value * const index);
     56    llvm::LoadInst * CreateBlockAlignedLoad(llvm::Value * const ptr, std::initializer_list<llvm::Value *> indices);
    6057
    61     void CreateBlockAlignedStore(Value * const value, Value * const ptr);
    62     void CreateBlockAlignedStore(Value * const value, Value * const ptr, Value * const index);
    63     void CreateBlockAlignedStore(Value * const value, Value * const ptr, std::initializer_list<Value *> indices);
     58    void CreateBlockAlignedStore(llvm::Value * const value, llvm::Value * const ptr);
     59    void CreateBlockAlignedStore(llvm::Value * const value, llvm::Value * const ptr, llvm::Value * const index);
     60    void CreateBlockAlignedStore(llvm::Value * const value, llvm::Value * const ptr, std::initializer_list<llvm::Value *> indices);
    6461
    65     VectorType * fwVectorType(unsigned fw);
     62    llvm::VectorType * fwVectorType(unsigned fw);
    6663
    67     Constant * simd_himask(unsigned fw);
    68     Constant * simd_lomask(unsigned fw);
     64    llvm::Constant * simd_himask(unsigned fw);
     65    llvm::Constant * simd_lomask(unsigned fw);
    6966   
    70     virtual Value * simd_fill(unsigned fw, Value * a);
     67    virtual llvm::Value * simd_fill(unsigned fw, llvm::Value * a);
    7168
    72     virtual Value * simd_add(unsigned fw, Value * a, Value * b);
    73     virtual Value * simd_sub(unsigned fw, Value * a, Value * b);
    74     virtual Value * simd_mult(unsigned fw, Value * a, Value * b);
    75     virtual Value * simd_eq(unsigned fw, Value * a, Value * b);
    76     virtual Value * simd_gt(unsigned fw, Value * a, Value * b);
    77     virtual Value * simd_ugt(unsigned fw, Value * a, Value * b);
    78     virtual Value * simd_lt(unsigned fw, Value * a, Value * b);
    79     virtual Value * simd_ult(unsigned fw, Value * a, Value * b);
    80     virtual Value * simd_max(unsigned fw, Value * a, Value * b);
    81     virtual Value * simd_umax(unsigned fw, Value * a, Value * b);
    82     virtual Value * simd_min(unsigned fw, Value * a, Value * b);
    83     virtual Value * simd_umin(unsigned fw, Value * a, Value * b);
    84     virtual Value * simd_if(unsigned fw, Value * cond, Value * a, Value * b);
     69    virtual llvm::Value * simd_add(unsigned fw, llvm::Value * a, llvm::Value * b);
     70    virtual llvm::Value * simd_sub(unsigned fw, llvm::Value * a, llvm::Value * b);
     71    virtual llvm::Value * simd_mult(unsigned fw, llvm::Value * a, llvm::Value * b);
     72    virtual llvm::Value * simd_eq(unsigned fw, llvm::Value * a, llvm::Value * b);
     73    virtual llvm::Value * simd_gt(unsigned fw, llvm::Value * a, llvm::Value * b);
     74    virtual llvm::Value * simd_ugt(unsigned fw, llvm::Value * a, llvm::Value * b);
     75    virtual llvm::Value * simd_lt(unsigned fw, llvm::Value * a, llvm::Value * b);
     76    virtual llvm::Value * simd_ult(unsigned fw, llvm::Value * a, llvm::Value * b);
     77    virtual llvm::Value * simd_max(unsigned fw, llvm::Value * a, llvm::Value * b);
     78    virtual llvm::Value * simd_umax(unsigned fw, llvm::Value * a, llvm::Value * b);
     79    virtual llvm::Value * simd_min(unsigned fw, llvm::Value * a, llvm::Value * b);
     80    virtual llvm::Value * simd_umin(unsigned fw, llvm::Value * a, llvm::Value * b);
     81    virtual llvm::Value * simd_if(unsigned fw, llvm::Value * cond, llvm::Value * a, llvm::Value * b);
    8582   
    86     virtual Value * simd_slli(unsigned fw, Value * a, unsigned shift);
    87     virtual Value * simd_srli(unsigned fw, Value * a, unsigned shift);
    88     virtual Value * simd_srai(unsigned fw, Value * a, unsigned shift);
     83    virtual llvm::Value * simd_slli(unsigned fw, llvm::Value * a, unsigned shift);
     84    virtual llvm::Value * simd_srli(unsigned fw, llvm::Value * a, unsigned shift);
     85    virtual llvm::Value * simd_srai(unsigned fw, llvm::Value * a, unsigned shift);
    8986   
    90     virtual Value * simd_cttz(unsigned fw, Value * a);
    91     virtual Value * simd_popcount(unsigned fw, Value * a);
     87    virtual llvm::Value * simd_cttz(unsigned fw, llvm::Value * a);
     88    virtual llvm::Value * simd_popcount(unsigned fw, llvm::Value * a);
    9289   
    93     virtual Value * esimd_mergeh(unsigned fw, Value * a, Value * b);
    94     virtual Value * esimd_mergel(unsigned fw, Value * a, Value * b);
    95     virtual Value * esimd_bitspread(unsigned fw, Value * bitmask);
     90    virtual llvm::Value * esimd_mergeh(unsigned fw, llvm::Value * a, llvm::Value * b);
     91    virtual llvm::Value * esimd_mergel(unsigned fw, llvm::Value * a, llvm::Value * b);
     92    virtual llvm::Value * esimd_bitspread(unsigned fw, llvm::Value * bitmask);
    9693   
    97     virtual Value * hsimd_packh(unsigned fw, Value * a, Value * b);
    98     virtual Value * hsimd_packl(unsigned fw, Value * a, Value * b);
    99     virtual Value * hsimd_packh_in_lanes(unsigned lanes, unsigned fw, Value * a, Value * b);
    100     virtual Value * hsimd_packl_in_lanes(unsigned lanes, unsigned fw, Value * a, Value * b);
     94    virtual llvm::Value * hsimd_packh(unsigned fw, llvm::Value * a, llvm::Value * b);
     95    virtual llvm::Value * hsimd_packl(unsigned fw, llvm::Value * a, llvm::Value * b);
     96    virtual llvm::Value * hsimd_packh_in_lanes(unsigned lanes, unsigned fw, llvm::Value * a, llvm::Value * b);
     97    virtual llvm::Value * hsimd_packl_in_lanes(unsigned lanes, unsigned fw, llvm::Value * a, llvm::Value * b);
    10198
    102     virtual Value * hsimd_signmask(unsigned fw, Value * a);
     99    virtual llvm::Value * hsimd_signmask(unsigned fw, llvm::Value * a);
    103100   
    104     virtual Value * mvmd_extract(unsigned fw, Value * a, unsigned fieldIndex);
    105     virtual Value * mvmd_insert(unsigned fw, Value * blk, Value * elt, unsigned fieldIndex);
    106     virtual Value * mvmd_slli(unsigned fw, Value * a, unsigned shift);
    107     virtual Value * mvmd_srli(unsigned fw, Value * a, unsigned shift);
    108     virtual Value * mvmd_dslli(unsigned fw, Value * a, Value * b, unsigned shift);
     101    virtual llvm::Value * mvmd_extract(unsigned fw, llvm::Value * a, unsigned fieldIndex);
     102    virtual llvm::Value * mvmd_insert(unsigned fw, llvm::Value * blk, llvm::Value * elt, unsigned fieldIndex);
     103    virtual llvm::Value * mvmd_slli(unsigned fw, llvm::Value * a, unsigned shift);
     104    virtual llvm::Value * mvmd_srli(unsigned fw, llvm::Value * a, unsigned shift);
     105    virtual llvm::Value * mvmd_dslli(unsigned fw, llvm::Value * a, llvm::Value * b, unsigned shift);
    109106   
    110107   
    111     virtual Value * bitblock_any(Value * a);
     108    virtual llvm::Value * bitblock_any(llvm::Value * a);
    112109    // full add producing {carryout, sum}
    113     virtual std::pair<Value *, Value *> bitblock_add_with_carry(Value * a, Value * b, Value * carryin);
     110    virtual std::pair<llvm::Value *, llvm::Value *> bitblock_add_with_carry(llvm::Value * a, llvm::Value * b, llvm::Value * carryin);
    114111    // full shift producing {shiftout, shifted}
    115     virtual std::pair<Value *, Value *> bitblock_advance(Value * a, Value * shiftin, unsigned shift);
    116     virtual Value * bitblock_mask_from(Value * pos);
    117     virtual Value * bitblock_set_bit(Value * pos);
     112    virtual std::pair<llvm::Value *, llvm::Value *> bitblock_advance(llvm::Value * a, llvm::Value * shiftin, unsigned shift);
     113    virtual llvm::Value * bitblock_mask_from(llvm::Value * pos);
     114    virtual llvm::Value * bitblock_set_bit(llvm::Value * pos);
    118115   
    119     Value * simd_and(Value * a, Value * b);
    120     Value * simd_or(Value * a, Value * b);
    121     Value * simd_xor(Value * a, Value * b);
    122     Value * simd_not(Value * a);
    123     Value * fwCast(unsigned fw, Value * a);
     116    llvm::Value * simd_and(llvm::Value * a, llvm::Value * b);
     117    llvm::Value * simd_or(llvm::Value * a, llvm::Value * b);
     118    llvm::Value * simd_xor(llvm::Value * a, llvm::Value * b);
     119    llvm::Value * simd_not(llvm::Value * a);
     120    llvm::Value * fwCast(unsigned fw, llvm::Value * a);
    124121   
    125     inline VectorType * getBitBlockType() const {
     122    inline llvm::VectorType * getBitBlockType() const {
    126123        return mBitBlockType;
    127124    }
    128125
    129     inline Type * getStreamSetTy(const unsigned NumElements = 1, const unsigned FieldWidth = 1) {
    130         return ArrayType::get(getStreamTy(FieldWidth), NumElements);
     126    inline llvm::ArrayType * getStreamSetTy(const unsigned NumElements = 1, const unsigned FieldWidth = 1) {
     127        return llvm::ArrayType::get(getStreamTy(FieldWidth), NumElements);
    131128    }
    132129   
    133     Type * getStreamTy(const unsigned FieldWidth = 1);
     130    StreamType * getStreamTy(const unsigned FieldWidth = 1);
    134131
    135     void CallPrintRegister(const std::string & regName, Value * const value);
     132    void CallPrintRegister(const std::string & regName, llvm::Value * const value);
    136133   
    137134protected:
    138135    unsigned            mBitBlockWidth;
    139136    unsigned            mStride;
    140     VectorType *        mBitBlockType;
     137    llvm::VectorType *  mBitBlockType;
    141138
    142     Constant *          mZeroInitializer;
    143     Constant *          mOneInitializer;
    144     Constant *          mPrintRegisterFunction;
     139    llvm::Constant *    mZeroInitializer;
     140    llvm::Constant *    mOneInitializer;
     141    llvm::Constant *    mPrintRegisterFunction;
    145142    StreamTypes         mStreamTypes;
    146143};
    147144
    148 inline LoadInst * IDISA_Builder::CreateBlockAlignedLoad(Value * const ptr) {
     145inline llvm::LoadInst * IDISA_Builder::CreateBlockAlignedLoad(llvm::Value * const ptr) {
    149146    return CreateAlignedLoad(ptr, mBitBlockWidth / 8);
    150147}
    151148
    152 inline LoadInst * IDISA_Builder::CreateBlockAlignedLoad(Value * const ptr, Value * const index) {
     149inline llvm::LoadInst * IDISA_Builder::CreateBlockAlignedLoad(llvm::Value * const ptr, llvm::Value * const index) {
    153150    return CreateBlockAlignedLoad(CreateGEP(ptr, index));
    154151}
    155152
    156 inline LoadInst * IDISA_Builder::CreateBlockAlignedLoad(Value * const ptr, std::initializer_list<Value *> indices) {
     153inline llvm::LoadInst * IDISA_Builder::CreateBlockAlignedLoad(llvm::Value * const ptr, std::initializer_list<llvm::Value *> indices) {
    157154    return CreateBlockAlignedLoad(CreateGEP(ptr, indices));
    158155}
    159156
    160 inline void IDISA_Builder::CreateBlockAlignedStore(Value * const value, Value * const ptr) {
     157inline void IDISA_Builder::CreateBlockAlignedStore(llvm::Value * const value, llvm::Value * const ptr) {
    161158    CreateAlignedStore(value, ptr, mBitBlockWidth / 8);
    162159}
    163160
    164 inline void IDISA_Builder::CreateBlockAlignedStore(Value * const value, Value * const ptr, Value * const index) {
     161inline void IDISA_Builder::CreateBlockAlignedStore(llvm::Value * const value, llvm::Value * const ptr, llvm::Value * const index) {
    165162    CreateBlockAlignedStore(value, CreateGEP(ptr, index));
    166163}
    167164
    168 inline void IDISA_Builder::CreateBlockAlignedStore(Value * const value, Value * const ptr, std::initializer_list<Value *> indices) {
     165inline void IDISA_Builder::CreateBlockAlignedStore(llvm::Value * const value, llvm::Value * const ptr, std::initializer_list<llvm::Value *> indices) {
    169166    CreateBlockAlignedStore(value, CreateGEP(ptr, indices));
    170167}
  • icGREP/icgrep-devel/icgrep/IR_Gen/idisa_i64_builder.cpp

    r4977 r5260  
    1010#include <llvm/IR/Intrinsics.h>
    1111#include <llvm/IR/Function.h>
     12#include <llvm/IR/Module.h>
    1213
    1314namespace IDISA {
  • icGREP/icgrep-devel/icgrep/IR_Gen/idisa_i64_builder.h

    r5238 r5260  
    77 *  icgrep is a trademark of International Characters.
    88 */
    9 #include <llvm/IR/Module.h>
    10 #include <llvm/IR/Constant.h>
    11 #include <llvm/IR/Type.h>
    12 #include <llvm/IR/Value.h>
    139#include <IR_Gen/idisa_builder.h>
    1410
  • icGREP/icgrep-devel/icgrep/IR_Gen/idisa_nvptx_builder.cpp

    r5230 r5260  
    1111#include <llvm/IR/Function.h>
    1212#include <llvm/IR/InlineAsm.h>
    13 #include <sstream>
     13#include <llvm/IR/Module.h>
    1414
    1515namespace IDISA {
     
    256256    Value * conv = CreateZExt(input, int32ty);
    257257
    258     std::ostringstream AsmStream;
    259     AsmStream << "{.reg .pred %p1; ";
    260     AsmStream << "setp.ne.u32 %p1, $1, 0; ";
    261     AsmStream << "vote.ballot.b32  $0, %p1;}";
     258    const char * AsmStream = "{.reg .pred %p1;"
     259                             "setp.ne.u32 %p1, $1, 0;"
     260                             "vote.ballot.b32  $0, %p1;}";
    262261    FunctionType * AsmFnTy = FunctionType::get(int32ty, int32ty, false);
    263     llvm::InlineAsm *IA = llvm::InlineAsm::get(AsmFnTy, AsmStream.str(), "=r,r", true, false);
     262    llvm::InlineAsm *IA = llvm::InlineAsm::get(AsmFnTy, AsmStream, "=r,r", true, false);
    264263    llvm::CallInst * result = CreateCall(IA, conv);
    265264    result->addAttribute(llvm::AttributeSet::FunctionIndex, llvm::Attribute::NoUnwind);
  • icGREP/icgrep-devel/icgrep/IR_Gen/idisa_nvptx_builder.h

    r5238 r5260  
    77*/
    88
    9 #include <IR_Gen/idisa_builder.h>
    109#include <IR_Gen/idisa_i64_builder.h>
    11 #include <iostream>
     10
    1211using namespace llvm;
    1312
  • icGREP/icgrep-devel/icgrep/IR_Gen/idisa_sse_builder.cpp

    r5118 r5260  
    1010#include <llvm/IR/Intrinsics.h>
    1111#include <llvm/IR/Function.h>
    12 #include <iostream>
     12#include <llvm/IR/Module.h>
    1313
    1414namespace IDISA {
  • icGREP/icgrep-devel/icgrep/IR_Gen/idisa_target.cpp

    r5238 r5260  
    44 */
    55
     6#include "idisa_target.h"
    67#include <toolchain.h>
    78#include <IR_Gen/idisa_avx_builder.h>
  • icGREP/icgrep-devel/icgrep/IR_Gen/idisa_target.h

    r5238 r5260  
    77#define IDISA_TARGET_H
    88
    9 #include <llvm/IR/Module.h>
    10 #include <llvm/IR/Type.h>
    11 #include <IR_Gen/idisa_builder.h>
     9namespace llvm { class Module; }
     10namespace IDISA { class IDISA_Builder; }
    1211
    1312namespace IDISA {
    1413   
    15 IDISA::IDISA_Builder * GetIDISA_Builder(Module * m);
     14IDISA::IDISA_Builder * GetIDISA_Builder(llvm::Module * m);
    1615
    17 IDISA::IDISA_Builder * GetIDISA_GPU_Builder(Module * m);
     16IDISA::IDISA_Builder * GetIDISA_GPU_Builder(llvm::Module * m);
    1817
    1918}
  • icGREP/icgrep-devel/icgrep/IR_Gen/types/streamtype.cpp

    r5238 r5260  
    11#include "streamtype.h"
    2 #include "IR_Gen/idisa_builder.h"
     2#include <IR_Gen/idisa_builder.h>  // for IDISA_Builder
     3#include <llvm/IR/DerivedTypes.h>  // for ArrayType, VectorType
    34
    45namespace IDISA {
     
    67llvm::Type * StreamType::resolveType(IDISA_Builder * const builder) {
    78    if (mFieldWidth == 1) return builder->getBitBlockType();
    8     return ArrayType::get(builder->getBitBlockType(), mFieldWidth);
     9    return llvm::ArrayType::get(builder->getBitBlockType(), mFieldWidth);
    910//    return llvm::VectorType::get(builder->getIntNTy(mFieldWidth), builder->getBitBlockWidth() / mFieldWidth);
    1011}
  • icGREP/icgrep-devel/icgrep/IR_Gen/types/streamtype.h

    r5230 r5260  
    22#define STREAMTYPE_H
    33
    4 #include <llvm/IR/DerivedTypes.h>
    5 #include <llvm/IR/LLVMContext.h>
     4#include <llvm/IR/Type.h>
     5namespace IDISA { class IDISA_Builder; }
     6namespace llvm { class LLVMContext; }
    67
    78namespace IDISA {
  • icGREP/icgrep-devel/icgrep/base64.cpp

    r5255 r5260  
    3838#include <boost/interprocess/mapped_region.hpp>
    3939#include <fcntl.h>
     40
     41
     42using namespace llvm;
     43
    4044static cl::OptionCategory base64Options("base64 Options",
    4145                                            "Transcoding control options.");
  • icGREP/icgrep-devel/icgrep/editd/editd.cpp

    r5257 r5260  
    6868using namespace kernel;
    6969using namespace pablo;
    70 
    71 std::string IRFilename = "editd.ll";
    72 std::string PTXFilename = "editd.ptx";
     70using namespace parabix;
     71
     72const static std::string IRFilename = "editd.ll";
     73const static std::string PTXFilename = "editd.ptx";
    7374
    7475struct matchPosition
  • icGREP/icgrep-devel/icgrep/editd/editd_cpu_kernel.cpp

    r5247 r5260  
    66#include <kernels/kernel.h>
    77#include <IR_Gen/idisa_builder.h>
     8#include <llvm/IR/Module.h>
    89#include <llvm/Support/raw_ostream.h>
    910#include <iostream>
    1011
     12using namespace llvm;
     13
    1114namespace kernel {
    12 using namespace llvm;
    1315
    1416void editdCPUKernel::bitblock_advance_ci_co(Value * val, unsigned shift, Value * stideCarryArr, unsigned carryIdx, std::vector<std::vector<Value *>> & adv, std::vector<std::vector<int>> & calculated, int i, int j) const {
     
    6466    Value * stideCarryArr = getScalarField(kernelStuctParam, "srideCarry");
    6567    Value * blockNo = getScalarField(kernelStuctParam, blockNoScalar);
    66     Value * ccStreamPtr = getStreamSetBlockPtr(kernelStuctParam, "CCStream", blockNo);
    67     Value * resultStreamPtr = getStreamSetBlockPtr(kernelStuctParam, "ResultStream", blockNo);
    68 
    6968   
    7069    unsigned carryIdx = 0;
     
    7776    Value * pattCh = iBuilder->CreateLoad(pattPtr);
    7877    Value * pattIdx = iBuilder->CreateAnd(iBuilder->CreateLShr(pattCh, 1), ConstantInt::get(int8ty, 3));
    79     Value * pattStreamPtr = iBuilder->CreateGEP(ccStreamPtr, {iBuilder->getInt32(0), iBuilder->CreateZExt(pattIdx, int32ty)});
     78    Value * pattStreamPtr = getStream(kernelStuctParam, "CCStream", blockNo, iBuilder->CreateZExt(pattIdx, int32ty));
    8079    Value * pattStream = iBuilder->CreateLoad(pattStreamPtr);
    8180    pattPos = iBuilder->CreateAdd(pattPos, ConstantInt::get(int32ty, 1));
     
    8685    }
    8786
    88     for(unsigned i = 1; i<mPatternLen; i++){     
     87    for(unsigned i = 1; i < mPatternLen; i++){
    8988        pattPtr = iBuilder->CreateGEP(pattStartPtr, pattPos);
    9089        pattCh = iBuilder->CreateLoad(pattPtr);
    9190        pattIdx = iBuilder->CreateAnd(iBuilder->CreateLShr(pattCh, 1), ConstantInt::get(int8ty, 3));
    92         pattStreamPtr = iBuilder->CreateGEP(ccStreamPtr, {iBuilder->getInt32(0), iBuilder->CreateZExt(pattIdx, int32ty)});
     91        pattStreamPtr = getStream(kernelStuctParam, "CCStream", blockNo, iBuilder->CreateZExt(pattIdx, int32ty));
    9392        pattStream = iBuilder->CreateLoad(pattStreamPtr);
    9493
     
    108107    }
    109108   
    110     Value * ptr = iBuilder->CreateGEP(resultStreamPtr, {iBuilder->getInt32(0), iBuilder->getInt32(0)});
    111     iBuilder->CreateStore(e[mPatternLen-1][0], ptr);
     109    Value * ptr = getStream(kernelStuctParam, "ResultStream", blockNo, iBuilder->getInt32(0));
     110    iBuilder->CreateStore(e[mPatternLen - 1][0], ptr);
    112111    for(unsigned j = 1; j<= mEditDistance; j++){
    113         ptr = iBuilder->CreateGEP(resultStreamPtr, {iBuilder->getInt32(0), iBuilder->getInt32(j)});
     112        ptr = getStream(kernelStuctParam, "ResultStream", blockNo, iBuilder->getInt32(j));
    114113        iBuilder->CreateStore(iBuilder->CreateAnd(e[mPatternLen-1][j], iBuilder->CreateNot(e[mPatternLen-1][j-1])), ptr);
    115114    }
     
    123122}
    124123
     124editdCPUKernel::editdCPUKernel(IDISA::IDISA_Builder * b, unsigned dist, unsigned pattLen) :
     125KernelBuilder(b, "editd_cpu",
     126             {Binding{b->getStreamSetTy(4), "CCStream"}},
     127             {Binding{b->getStreamSetTy(dist + 1), "ResultStream"}},
     128             {Binding{PointerType::get(b->getInt8Ty(), 1), "pattStream"},
     129             Binding{PointerType::get(ArrayType::get(b->getBitBlockType(), pattLen * (dist + 1) * 4), 0), "srideCarry"}},
     130             {},
     131             {Binding{b->getBitBlockType(), "EOFmask"}}),
     132mEditDistance(dist),
     133mPatternLen(pattLen){
     134
     135}
     136
    125137}
    126138
  • icGREP/icgrep-devel/icgrep/editd/editd_cpu_kernel.h

    r5246 r5260  
    1919public:
    2020   
    21     editdCPUKernel(IDISA::IDISA_Builder * b, unsigned dist, unsigned pattLen) :
    22     KernelBuilder(b, "editd_cpu",
    23                   {Binding{b->getStreamSetTy(4), "CCStream"}},
    24                   {Binding{b->getStreamSetTy(dist + 1), "ResultStream"}},
    25                   {Binding{PointerType::get(b->getInt8Ty(), 1), "pattStream"},
    26                   Binding{PointerType::get(ArrayType::get(b->getBitBlockType(), pattLen * (dist + 1) * 4), 0), "srideCarry"}},
    27                   {},
    28                   {Binding{b->getBitBlockType(), "EOFmask"}}),
    29     mEditDistance(dist),
    30     mPatternLen(pattLen){}
     21    editdCPUKernel(IDISA::IDISA_Builder * b, unsigned dist, unsigned pattLen);
    3122   
    3223   
     
    3425    void generateDoBlockMethod() const override;
    3526    void generateFinalBlockMethod() const override;
    36     void bitblock_advance_ci_co(Value * val, unsigned shift, Value * stideCarryArr, unsigned carryIdx, std::vector<std::vector<Value *>> & adv, std::vector<std::vector<int>> & calculated, int i, int j) const;
     27    void bitblock_advance_ci_co(llvm::Value * val, unsigned shift, llvm::Value * stideCarryArr, unsigned carryIdx, std::vector<std::vector<llvm::Value *>> & adv, std::vector<std::vector<int>> & calculated, int i, int j) const;
    3728    unsigned mEditDistance;
    3829    unsigned mPatternLen;
  • icGREP/icgrep-devel/icgrep/editd/editd_gpu_kernel.cpp

    r5247 r5260  
    66#include <kernels/kernel.h>
    77#include <IR_Gen/idisa_builder.h>
     8#include <llvm/IR/Module.h>
    89#include <llvm/Support/raw_ostream.h>
    910#include <iostream>
    1011
    11 namespace kernel {
    1212using namespace llvm;
    1313
     14namespace kernel {
    1415
    1516void bitblock_advance_ci_co(IDISA::IDISA_Builder * iBuilder, Value * val, unsigned shift, Value * stideCarryArr, unsigned carryIdx, std::vector<std::vector<Value *>> & adv, std::vector<std::vector<int>> & calculated, int i, int j){   
     
    6667    Value * stideCarryArr = getScalarField(kernelStuctParam, "srideCarry");
    6768    Value * blockNo = getScalarField(kernelStuctParam, blockNoScalar);
    68     Value * ccStreamPtr = getStreamSetBlockPtr(kernelStuctParam, "CCStream", blockNo);
    69     Value * resultStreamPtr = getStreamSetBlockPtr(kernelStuctParam, "ResultStream", blockNo);
    7069    Value * pattLen = ConstantInt::get(int32ty, mPatternLen+1);
    7170    Value * pattPos = ConstantInt::get(int32ty, 0);
     
    7372    unsigned carryIdx = 0;
    7473
    75     std::vector<std::vector<Value *>> e(mPatternLen, std::vector<Value *>(mEditDistance+1));
    76     std::vector<std::vector<Value *>> adv(mPatternLen, std::vector<Value *>(mEditDistance+1));
    77     std::vector<std::vector<int>> calculated(mPatternLen, std::vector<int>(mEditDistance+1));
    78     for(unsigned i=0; i<mPatternLen; i++)
    79         for(unsigned j=0; j<=mEditDistance; j++)
    80             calculated[i][j] = 0;
     74    std::vector<std::vector<Value *>> e(mPatternLen, std::vector<Value *>(mEditDistance + 1));
     75    std::vector<std::vector<Value *>> adv(mPatternLen, std::vector<Value *>(mEditDistance + 1));
     76    std::vector<std::vector<int>> calculated(mPatternLen, std::vector<int>(mEditDistance + 1, 0));
     77
    8178    Function * bidFunc = cast<Function>(m->getOrInsertFunction("llvm.nvvm.read.ptx.sreg.ctaid.x", int32ty, nullptr));
    8279    Value * bid = iBuilder->CreateCall(bidFunc);
     
    8582    Value * pattCh = iBuilder->CreateLoad(pattPtr);
    8683    Value * pattIdx = iBuilder->CreateAnd(iBuilder->CreateLShr(pattCh, 1), ConstantInt::get(int8ty, 3));
    87     Value * pattStreamPtr = iBuilder->CreateGEP(ccStreamPtr, {iBuilder->getInt32(0), iBuilder->CreateZExt(pattIdx, int32ty)});
     84    Value * pattStreamPtr = getStream(kernelStuctParam, "CCStream", blockNo, iBuilder->CreateZExt(pattIdx, int32ty));
    8885    Value * pattStream = iBuilder->CreateLoad(pattStreamPtr);
    8986    pattPos = iBuilder->CreateAdd(pattPos, ConstantInt::get(int32ty, 1));
     
    9895        pattCh = iBuilder->CreateLoad(pattPtr);
    9996        pattIdx = iBuilder->CreateAnd(iBuilder->CreateLShr(pattCh, 1), ConstantInt::get(int8ty, 3));
    100         pattStreamPtr = iBuilder->CreateGEP(ccStreamPtr, {iBuilder->getInt32(0), iBuilder->CreateZExt(pattIdx, int32ty)});
     97        pattStreamPtr = getStream(kernelStuctParam, "CCStream", blockNo, iBuilder->CreateZExt(pattIdx, int32ty));
    10198        pattStream = iBuilder->CreateLoad(pattStreamPtr);
    10299
     
    116113    }
    117114
    118     Value * ptr = iBuilder->CreateGEP(resultStreamPtr, {iBuilder->getInt32(0), iBuilder->getInt32(0)});
     115    Value * ptr = getStream(kernelStuctParam, "ResultStream", blockNo, iBuilder->getInt32(0));
    119116    iBuilder->CreateStore(e[mPatternLen-1][0], ptr);
    120117    for(unsigned j = 1; j<= mEditDistance; j++){
    121         ptr = iBuilder->CreateGEP(resultStreamPtr, {iBuilder->getInt32(0), iBuilder->getInt32(j)});
    122         iBuilder->CreateStore(iBuilder->CreateAnd(e[mPatternLen-1][j], iBuilder->CreateNot(e[mPatternLen-1][j-1])), ptr);
     118        ptr = getStream(kernelStuctParam, "ResultStream", blockNo, iBuilder->getInt32(j));
     119        iBuilder->CreateStore(iBuilder->CreateAnd(e[mPatternLen - 1][j], iBuilder->CreateNot(e[mPatternLen - 1][j - 1])), ptr);
    123120    }
    124121
     
    131128}
    132129
     130editdGPUKernel::editdGPUKernel(IDISA::IDISA_Builder * b, unsigned dist, unsigned pattLen) :
     131KernelBuilder(b, "editd_gpu",
     132              {Binding{b->getStreamSetTy(4), "CCStream"}},
     133              {Binding{b->getStreamSetTy(dist + 1), "ResultStream"}},
     134              {Binding{PointerType::get(b->getInt8Ty(), 1), "pattStream"},
     135              Binding{PointerType::get(ArrayType::get(b->getBitBlockType(), pattLen * (dist + 1) * 4), 0), "srideCarry"}},
     136              {},
     137              {Binding{b->getBitBlockType(), "EOFmask"}}),
     138mEditDistance(dist),
     139mPatternLen(pattLen) {
     140
     141}
     142
    133143}
    134144
  • icGREP/icgrep-devel/icgrep/editd/editd_gpu_kernel.h

    r5246 r5260  
    1919public:
    2020   
    21     editdGPUKernel(IDISA::IDISA_Builder * b, unsigned dist, unsigned pattLen) :
    22     KernelBuilder(b, "editd_gpu",
    23                   {Binding{b->getStreamSetTy(4), "CCStream"}},
    24                   {Binding{b->getStreamSetTy(dist + 1), "ResultStream"}},
    25                   {Binding{PointerType::get(b->getInt8Ty(), 1), "pattStream"},
    26                   Binding{PointerType::get(ArrayType::get(b->getBitBlockType(), pattLen * (dist + 1) * 4), 0), "srideCarry"}},
    27                   {},
    28                   {Binding{b->getBitBlockType(), "EOFmask"}}),
    29     mEditDistance(dist),
    30     mPatternLen(pattLen){}
     21    editdGPUKernel(IDISA::IDISA_Builder * b, unsigned dist, unsigned pattLen);
    3122   
    3223   
  • icGREP/icgrep-devel/icgrep/editd/editdscan_kernel.cpp

    r5246 r5260  
    88#include <llvm/IR/Intrinsics.h>
    99#include <IR_Gen/idisa_builder.h>
     10#include <llvm/IR/Module.h>
    1011#include <llvm/Support/raw_os_ostream.h>
    1112#include <iostream>
     
    3536    Value * scanwordPos = iBuilder->CreateMul(blockNo, ConstantInt::get(blockNo->getType(), iBuilder->getBitBlockWidth()));
    3637   
    37     Value * matchResultsPtr = getStreamSetBlockPtr(kernelStuctParam, "matchResults", blockNo);
    38 
    3938    std::vector<Value * > matchWordVectors;
    4039    for(unsigned d = 0; d <= mEditDistance; d++){
    41         Value * matches = iBuilder->CreateBlockAlignedLoad(matchResultsPtr, {iBuilder->getInt32(0), iBuilder->getInt32(d)});
     40        Value * ptr = getStream(kernelStuctParam, "matchResults", blockNo, iBuilder->getInt32(d));
     41        Value * matches = iBuilder->CreateBlockAlignedLoad(ptr);
    4242        matchWordVectors.push_back(iBuilder->CreateBitCast(matches, scanwordVectorType));
    4343    }
     
    101101}
    102102
     103editdScanKernel::editdScanKernel(IDISA::IDISA_Builder * iBuilder, unsigned dist) :
     104KernelBuilder(iBuilder, "scanMatch",
     105              {Binding{iBuilder->getStreamSetTy(dist + 1), "matchResults"}},
     106              {}, {}, {}, {}),
     107mEditDistance(dist),
     108mScanwordBitWidth(iBuilder->getSizeTy()->getBitWidth()) {
     109
    103110}
     111
     112}
  • icGREP/icgrep-devel/icgrep/editd/editdscan_kernel.h

    r5246 r5260  
    1919class editdScanKernel : public KernelBuilder {
    2020public:
    21     editdScanKernel(IDISA::IDISA_Builder * iBuilder, unsigned dist) :
    22     KernelBuilder(iBuilder, "scanMatch",
    23                   {Binding{iBuilder->getStreamSetTy(dist + 1), "matchResults"}},
    24                   {}, {}, {}, {}),
    25     mEditDistance(dist),
    26     mScanwordBitWidth(iBuilder->getSizeTy()->getBitWidth()) {}
     21    editdScanKernel(IDISA::IDISA_Builder * iBuilder, unsigned dist);
    2722       
    2823private:
  • icGREP/icgrep-devel/icgrep/icgrep-devel.files

    r5240 r5260  
    797797utf8_encoder.h
    798798wc.cpp
     799cc/cc_compiler.cpp
     800cc/cc_compiler.h
     801editd/editd.cpp
     802editd/editd_cpu_kernel.cpp
     803editd/editd_cpu_kernel.h
     804editd/editd_gpu_kernel.cpp
     805editd/editd_gpu_kernel.h
     806editd/EditdCudaDriver.h
     807editd/editdscan_kernel.cpp
     808editd/editdscan_kernel.h
     809editd/pattern_compiler.cpp
     810editd/pattern_compiler.h
     811IR_Gen/types/streamtype.cpp
     812IR_Gen/types/streamtype.h
     813IR_Gen/CBuilder.cpp
     814IR_Gen/CBuilder.h
     815IR_Gen/CudaDriver.h
     816IR_Gen/idisa_avx_builder.cpp
     817IR_Gen/idisa_avx_builder.h
     818IR_Gen/idisa_builder.cpp
     819IR_Gen/idisa_builder.h
     820IR_Gen/idisa_i64_builder.cpp
     821IR_Gen/idisa_i64_builder.h
     822IR_Gen/idisa_nvptx_builder.cpp
     823IR_Gen/idisa_nvptx_builder.h
     824IR_Gen/idisa_sse_builder.cpp
     825IR_Gen/idisa_sse_builder.h
     826IR_Gen/idisa_target.cpp
     827IR_Gen/idisa_target.h
     828IR_Gen/llvm2ptx.h
     829kernels/cc_kernel.cpp
     830kernels/cc_kernel.h
     831kernels/deletion.cpp
     832kernels/deletion.h
     833kernels/interface.cpp
     834kernels/interface.h
     835kernels/kernel.cpp
     836kernels/kernel.h
     837kernels/mmap_kernel.cpp
     838kernels/mmap_kernel.h
     839kernels/p2s_kernel.cpp
     840kernels/p2s_kernel.h
     841kernels/pipeline.cpp
     842kernels/pipeline.h
     843kernels/radix64.cpp
     844kernels/radix64.h
     845kernels/s2p_kernel.cpp
     846kernels/s2p_kernel.h
     847kernels/scanmatchgen.cpp
     848kernels/scanmatchgen.h
     849kernels/stdout_kernel.cpp
     850kernels/stdout_kernel.h
     851kernels/streamset.cpp
     852kernels/streamset.h
     853kernels/symboltablepipeline.cpp
     854kernels/symboltablepipeline.h
     855pablo/analysis/pabloverifier.cpp
     856pablo/analysis/pabloverifier.hpp
     857pablo/optimizers/booleanreassociationpass.cpp
     858pablo/optimizers/booleanreassociationpass.h
     859pablo/optimizers/codemotionpass.cpp
     860pablo/optimizers/codemotionpass.h
     861pablo/optimizers/distributivepass.cpp
     862pablo/optimizers/distributivepass.h
     863pablo/optimizers/graph-facade.hpp
     864pablo/optimizers/maxsat.hpp
     865pablo/optimizers/pablo_automultiplexing.cpp
     866pablo/optimizers/pablo_automultiplexing.hpp
     867pablo/optimizers/pablo_bddminimization.cpp
     868pablo/optimizers/pablo_bddminimization.h
     869pablo/optimizers/pablo_simplifier.cpp
     870pablo/optimizers/pablo_simplifier.hpp
     871pablo/optimizers/schedulingprepass.cpp
     872pablo/optimizers/schedulingprepass.h
     873pablo/passes/factorizedfg.cpp
     874pablo/passes/factorizedfg.h
     875pablo/passes/flattenassociativedfg.cpp
     876pablo/passes/flattenassociativedfg.h
     877pablo/passes/flattenif.cpp
     878pablo/passes/flattenif.hpp
     879pablo/arithmetic.h
     880pablo/boolean.h
     881pablo/branch.cpp
     882pablo/branch.h
     883pablo/builder.cpp
     884pablo/builder.hpp
     885pablo/carry_data.h
     886pablo/carry_manager.cpp
     887pablo/carry_manager.h
     888pablo/codegenstate.cpp
     889pablo/codegenstate.h
     890pablo/expression_map.hpp
     891pablo/pablo_compiler.cpp
     892pablo/pablo_compiler.h
     893pablo/pablo_kernel.cpp
     894pablo/pablo_kernel.h
     895pablo/pablo_toolchain.cpp
     896pablo/pablo_toolchain.h
     897pablo/pabloAST.cpp
     898pablo/pabloAST.h
     899pablo/pe_advance.h
     900pablo/pe_call.h
     901pablo/pe_constant.h
     902pablo/pe_count.h
     903pablo/pe_infile.h
     904pablo/pe_integer.h
     905pablo/pe_lookahead.h
     906pablo/pe_matchstar.h
     907pablo/pe_ones.h
     908pablo/pe_scanthru.h
     909pablo/pe_setithbit.h
     910pablo/pe_string.h
     911pablo/pe_var.h
     912pablo/pe_zeroes.h
     913pablo/printer_pablos.cpp
     914pablo/printer_pablos.h
     915pablo/prototype.cpp
     916pablo/prototype.h
     917pablo/ps_assign.h
     918pablo/symbol_generator.cpp
     919pablo/symbol_generator.h
     920re/printer_re.cpp
     921re/printer_re.h
     922re/re_alt.h
     923re/re_analysis.cpp
     924re/re_analysis.h
     925re/re_any.h
     926re/re_assertion.h
     927re/re_cc.cpp
     928re/re_cc.h
     929re/re_compiler.cpp
     930re/re_compiler.h
     931re/re_diff.cpp
     932re/re_diff.h
     933re/re_end.h
     934re/re_intersect.cpp
     935re/re_intersect.h
     936re/re_memoizer.hpp
     937re/re_name.h
     938re/re_name_resolve.cpp
     939re/re_name_resolve.h
     940re/re_nullable.cpp
     941re/re_nullable.h
     942re/re_parser.cpp
     943re/re_parser.h
     944re/re_parser_bre.cpp
     945re/re_parser_bre.h
     946re/re_parser_ere.cpp
     947re/re_parser_ere.h
     948re/re_parser_helper.h
     949re/re_parser_pcre.cpp
     950re/re_parser_pcre.h
     951re/re_parser_prosite.cpp
     952re/re_parser_prosite.h
     953re/re_re.cpp
     954re/re_re.h
     955re/re_rep.cpp
     956re/re_rep.h
     957re/re_seq.h
     958re/re_simplifier.cpp
     959re/re_simplifier.h
     960re/re_start.h
     961re/re_toolchain.cpp
     962re/re_toolchain.h
     963re/re_utility.cpp
     964re/re_utility.h
     965UCD/Blocks.h
     966UCD/CaseFolding_txt.cpp
     967UCD/CaseFolding_txt.h
     968UCD/DerivedAge.h
     969UCD/DerivedBidiClass.h
     970UCD/DerivedBinaryProperties.h
     971UCD/DerivedCombiningClass.h
     972UCD/DerivedCoreProperties.h
     973UCD/DerivedDecompositionType.h
     974UCD/DerivedGeneralCategory.h
     975UCD/DerivedJoiningGroup.h
     976UCD/DerivedJoiningType.h
     977UCD/DerivedNormalizationProps.h
     978UCD/DerivedNumericType.h
     979UCD/EastAsianWidth.h
     980UCD/GraphemeBreakProperty.h
     981UCD/HangulSyllableType.h
     982UCD/LineBreak.h
     983UCD/PropertyAliases.h
     984UCD/PropertyObjects.cpp
     985UCD/PropertyObjects.h
     986UCD/PropertyObjectTable.h
     987UCD/PropertyValueAliases.h
     988UCD/PropList.h
     989UCD/resolve_properties.cpp
     990UCD/resolve_properties.h
     991UCD/ScriptExtensions.h
     992UCD/Scripts.h
     993UCD/SentenceBreakProperty.h
     994UCD/ucd_compiler.cpp
     995UCD/ucd_compiler.hpp
     996UCD/unicode_set.cpp
     997UCD/unicode_set.h
     998UCD/UnicodeNameData.cpp
     999UCD/UnicodeNameData.h
     1000UCD/WordBreakProperty.h
     1001util/aligned_allocator.h
     1002util/papi_helper.hpp
     1003util/slab_allocator.h
     1004array-test.cpp
     1005base64.cpp
     1006generate_predefined_ucd_functions.cpp
     1007grep_engine.cpp
     1008grep_engine.h
     1009grep_type.h
     1010hrtime.h
     1011icgrep.cpp
     1012object_cache.cpp
     1013object_cache.h
     1014symboltable.cpp
     1015toolchain.cpp
     1016toolchain.h
     1017u8u16.cpp
     1018utf16_encoder.cpp
     1019utf16_encoder.h
     1020utf8_encoder.cpp
     1021utf8_encoder.h
     1022wc.cpp
     1023cc/cc_compiler.cpp
     1024cc/cc_compiler.h
     1025editd/editd.cpp
     1026editd/editd_cpu_kernel.cpp
     1027editd/editd_cpu_kernel.h
     1028editd/editd_gpu_kernel.cpp
     1029editd/editd_gpu_kernel.h
     1030editd/EditdCudaDriver.h
     1031editd/editdscan_kernel.cpp
     1032editd/editdscan_kernel.h
     1033editd/pattern_compiler.cpp
     1034editd/pattern_compiler.h
     1035IR_Gen/types/streamtype.cpp
     1036IR_Gen/types/streamtype.h
     1037IR_Gen/CBuilder.cpp
     1038IR_Gen/CBuilder.h
     1039IR_Gen/CudaDriver.h
     1040IR_Gen/idisa_avx_builder.cpp
     1041IR_Gen/idisa_avx_builder.h
     1042IR_Gen/idisa_builder.cpp
     1043IR_Gen/idisa_builder.h
     1044IR_Gen/idisa_i64_builder.cpp
     1045IR_Gen/idisa_i64_builder.h
     1046IR_Gen/idisa_nvptx_builder.cpp
     1047IR_Gen/idisa_nvptx_builder.h
     1048IR_Gen/idisa_sse_builder.cpp
     1049IR_Gen/idisa_sse_builder.h
     1050IR_Gen/idisa_target.cpp
     1051IR_Gen/idisa_target.h
     1052IR_Gen/llvm2ptx.h
     1053kernels/cc_kernel.cpp
     1054kernels/cc_kernel.h
     1055kernels/deletion.cpp
     1056kernels/deletion.h
     1057kernels/interface.cpp
     1058kernels/interface.h
     1059kernels/kernel.cpp
     1060kernels/kernel.h
     1061kernels/mmap_kernel.cpp
     1062kernels/mmap_kernel.h
     1063kernels/p2s_kernel.cpp
     1064kernels/p2s_kernel.h
     1065kernels/pipeline.cpp
     1066kernels/pipeline.h
     1067kernels/radix64.cpp
     1068kernels/radix64.h
     1069kernels/s2p_kernel.cpp
     1070kernels/s2p_kernel.h
     1071kernels/scanmatchgen.cpp
     1072kernels/scanmatchgen.h
     1073kernels/stdout_kernel.cpp
     1074kernels/stdout_kernel.h
     1075kernels/streamset.cpp
     1076kernels/streamset.h
     1077kernels/symboltablepipeline.cpp
     1078kernels/symboltablepipeline.h
     1079pablo/analysis/pabloverifier.cpp
     1080pablo/analysis/pabloverifier.hpp
     1081pablo/optimizers/booleanreassociationpass.cpp
     1082pablo/optimizers/booleanreassociationpass.h
     1083pablo/optimizers/codemotionpass.cpp
     1084pablo/optimizers/codemotionpass.h
     1085pablo/optimizers/distributivepass.cpp
     1086pablo/optimizers/distributivepass.h
     1087pablo/optimizers/graph-facade.hpp
     1088pablo/optimizers/maxsat.hpp
     1089pablo/optimizers/pablo_automultiplexing.cpp
     1090pablo/optimizers/pablo_automultiplexing.hpp
     1091pablo/optimizers/pablo_bddminimization.cpp
     1092pablo/optimizers/pablo_bddminimization.h
     1093pablo/optimizers/pablo_simplifier.cpp
     1094pablo/optimizers/pablo_simplifier.hpp
     1095pablo/optimizers/schedulingprepass.cpp
     1096pablo/optimizers/schedulingprepass.h
     1097pablo/passes/factorizedfg.cpp
     1098pablo/passes/factorizedfg.h
     1099pablo/passes/flattenassociativedfg.cpp
     1100pablo/passes/flattenassociativedfg.h
     1101pablo/passes/flattenif.cpp
     1102pablo/passes/flattenif.hpp
     1103pablo/arithmetic.h
     1104pablo/boolean.h
     1105pablo/branch.cpp
     1106pablo/branch.h
     1107pablo/builder.cpp
     1108pablo/builder.hpp
     1109pablo/carry_data.h
     1110pablo/carry_manager.cpp
     1111pablo/carry_manager.h
     1112pablo/codegenstate.cpp
     1113pablo/codegenstate.h
     1114pablo/expression_map.hpp
     1115pablo/pablo_compiler.cpp
     1116pablo/pablo_compiler.h
     1117pablo/pablo_kernel.cpp
     1118pablo/pablo_kernel.h
     1119pablo/pablo_toolchain.cpp
     1120pablo/pablo_toolchain.h
     1121pablo/pabloAST.cpp
     1122pablo/pabloAST.h
     1123pablo/pe_advance.h
     1124pablo/pe_call.h
     1125pablo/pe_constant.h
     1126pablo/pe_count.h
     1127pablo/pe_infile.h
     1128pablo/pe_integer.h
     1129pablo/pe_lookahead.h
     1130pablo/pe_matchstar.h
     1131pablo/pe_ones.h
     1132pablo/pe_scanthru.h
     1133pablo/pe_setithbit.h
     1134pablo/pe_string.h
     1135pablo/pe_var.h
     1136pablo/pe_zeroes.h
     1137pablo/printer_pablos.cpp
     1138pablo/printer_pablos.h
     1139pablo/prototype.cpp
     1140pablo/prototype.h
     1141pablo/ps_assign.h
     1142pablo/symbol_generator.cpp
     1143pablo/symbol_generator.h
     1144re/printer_re.cpp
     1145re/printer_re.h
     1146re/re_alt.h
     1147re/re_analysis.cpp
     1148re/re_analysis.h
     1149re/re_any.h
     1150re/re_assertion.h
     1151re/re_cc.cpp
     1152re/re_cc.h
     1153re/re_compiler.cpp
     1154re/re_compiler.h
     1155re/re_diff.cpp
     1156re/re_diff.h
     1157re/re_end.h
     1158re/re_intersect.cpp
     1159re/re_intersect.h
     1160re/re_memoizer.hpp
     1161re/re_name.h
     1162re/re_name_resolve.cpp
     1163re/re_name_resolve.h
     1164re/re_nullable.cpp
     1165re/re_nullable.h
     1166re/re_parser.cpp
     1167re/re_parser.h
     1168re/re_parser_bre.cpp
     1169re/re_parser_bre.h
     1170re/re_parser_ere.cpp
     1171re/re_parser_ere.h
     1172re/re_parser_helper.h
     1173re/re_parser_pcre.cpp
     1174re/re_parser_pcre.h
     1175re/re_parser_prosite.cpp
     1176re/re_parser_prosite.h
     1177re/re_re.cpp
     1178re/re_re.h
     1179re/re_rep.cpp
     1180re/re_rep.h
     1181re/re_seq.h
     1182re/re_simplifier.cpp
     1183re/re_simplifier.h
     1184re/re_start.h
     1185re/re_toolchain.cpp
     1186re/re_toolchain.h
     1187re/re_utility.cpp
     1188re/re_utility.h
     1189UCD/Blocks.h
     1190UCD/CaseFolding_txt.cpp
     1191UCD/CaseFolding_txt.h
     1192UCD/DerivedAge.h
     1193UCD/DerivedBidiClass.h
     1194UCD/DerivedBinaryProperties.h
     1195UCD/DerivedCombiningClass.h
     1196UCD/DerivedCoreProperties.h
     1197UCD/DerivedDecompositionType.h
     1198UCD/DerivedGeneralCategory.h
     1199UCD/DerivedJoiningGroup.h
     1200UCD/DerivedJoiningType.h
     1201UCD/DerivedNormalizationProps.h
     1202UCD/DerivedNumericType.h
     1203UCD/EastAsianWidth.h
     1204UCD/GraphemeBreakProperty.h
     1205UCD/HangulSyllableType.h
     1206UCD/LineBreak.h
     1207UCD/PropertyAliases.h
     1208UCD/PropertyObjects.cpp
     1209UCD/PropertyObjects.h
     1210UCD/PropertyObjectTable.h
     1211UCD/PropertyValueAliases.h
     1212UCD/PropList.h
     1213UCD/resolve_properties.cpp
     1214UCD/resolve_properties.h
     1215UCD/ScriptExtensions.h
     1216UCD/Scripts.h
     1217UCD/SentenceBreakProperty.h
     1218UCD/ucd_compiler.cpp
     1219UCD/ucd_compiler.hpp
     1220UCD/unicode_set.cpp
     1221UCD/unicode_set.h
     1222UCD/UnicodeNameData.cpp
     1223UCD/UnicodeNameData.h
     1224UCD/WordBreakProperty.h
     1225util/aligned_allocator.h
     1226util/papi_helper.hpp
     1227util/slab_allocator.h
     1228array-test.cpp
     1229base64.cpp
     1230generate_predefined_ucd_functions.cpp
     1231grep_engine.cpp
     1232grep_engine.h
     1233grep_type.h
     1234hrtime.h
     1235icgrep.cpp
     1236object_cache.cpp
     1237object_cache.h
     1238symboltable.cpp
     1239toolchain.cpp
     1240toolchain.h
     1241u8u16.cpp
     1242utf16_encoder.cpp
     1243utf16_encoder.h
     1244utf8_encoder.cpp
     1245utf8_encoder.h
     1246wc.cpp
  • icGREP/icgrep-devel/icgrep/kernels/cc_kernel.cpp

    r5246 r5260  
    88#include <cc/cc_compiler.h>
    99#include <pablo/builder.hpp>
     10#include <llvm/IR/Module.h>
    1011
    1112using namespace cc;
     
    2526    Value * blockNo = getScalarField(self, blockNoScalar);
    2627   
    27     Value * codeUnitStreamBlock_ptr = getStreamSetBlockPtr(self, "codeUnitStream", blockNo);
    28     Value * ccStreamBlock_ptr = getStreamSetBlockPtr(self, "ccStream", blockNo);
    29 
    3028    unsigned packCount = 8 * mCodeUnitSize; 
    3129    unsigned codeUnitWidth = 8 * mCodeUnitSize;
    3230    Value * codeUnitPack[packCount];
    3331    for (unsigned i = 0; i < packCount; i++) {
    34         codeUnitPack[i] = iBuilder->CreateBlockAlignedLoad(codeUnitStreamBlock_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(0), iBuilder->getInt32(i)});
     32        Value * ptr = getStream(self, "codeUnitStream", blockNo, iBuilder->getInt32(0), iBuilder->getInt32(i));
     33        codeUnitPack[i] = iBuilder->CreateBlockAlignedLoad(ptr);
    3534    }
    3635
    37     std::vector<Value *> ccStreams;
    3836    for (unsigned j = 0; j < mCharClasses.size();  j++) {
    3937        Value * theCCstream = iBuilder->allZeroes();
     
    6866            theCCstream = iBuilder->simd_or(theCCstream, pack);
    6967        }
    70         iBuilder->CreateBlockAlignedStore(theCCstream, ccStreamBlock_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(j)});
     68        Value * ptr = getStream(self, "ccStream", blockNo, iBuilder->getInt32(j));
     69        iBuilder->CreateBlockAlignedStore(theCCstream, ptr);
    7170    }
    7271 
  • icGREP/icgrep-devel/icgrep/kernels/deletion.cpp

    r5247 r5260  
    44 */
    55
    6 #include <kernels/kernel.h>
    7 #include <kernels/deletion.h>
     6#include "deletion.h"
    87#include <IR_Gen/idisa_builder.h>
    98#include <llvm/IR/Value.h>
     9#include <llvm/IR/Module.h>
    1010
    11 std::vector<Value *> parallel_prefix_deletion_masks(IDISA::IDISA_Builder * iBuilder, unsigned fw, Value * del_mask) {
     11using namespace llvm;
     12
     13namespace kernel {
     14
     15inline std::vector<Value *> parallel_prefix_deletion_masks(IDISA::IDISA_Builder * iBuilder, const unsigned fw, Value * del_mask) {
    1216    Value * m = iBuilder->simd_not(del_mask);
    1317    Value * mk = iBuilder->simd_slli(fw, del_mask, 1);
     
    2630}
    2731
    28 Value * apply_parallel_prefix_deletion(IDISA::IDISA_Builder * iBuilder, unsigned fw, Value * del_mask, std::vector<Value *> mv, Value * strm) {
     32inline Value * apply_parallel_prefix_deletion(IDISA::IDISA_Builder * iBuilder, const unsigned fw, Value * del_mask, const std::vector<Value *> & mv, Value * strm) {
    2933    Value * s = iBuilder->simd_and(strm, iBuilder->simd_not(del_mask));
    3034    for (unsigned i = 0; i < mv.size(); i++) {
     
    3640}
    3741
    38 Value * partial_sum_popcount(IDISA::IDISA_Builder * iBuilder, unsigned fw, Value * mask) {
    39     Value * per_field = iBuilder->simd_popcount(fw, mask);
    40     for (unsigned move = 1; move < iBuilder->getBitBlockWidth()/fw; move *= 2) {
    41         per_field = iBuilder->simd_add(fw, per_field, iBuilder->mvmd_slli(fw, per_field, move));
     42inline Value * partial_sum_popcount(IDISA::IDISA_Builder * iBuilder, const unsigned fw, Value * mask) {
     43    Value * field = iBuilder->simd_popcount(fw, mask);
     44    const auto count = iBuilder->getBitBlockWidth() / fw;
     45    for (unsigned move = 1; move < count; move *= 2) {
     46        field = iBuilder->simd_add(fw, field, iBuilder->mvmd_slli(fw, field, move));
    4247    }
    43     return per_field;
     48    return field;
    4449}
    4550
     
    4853// Outputs: the deleted streams, plus a partial sum popcount
    4954
    50 namespace kernel {
     55void DeletionKernel::generateDoBlockMethod() const {
    5156
    52 void DeletionKernel::generateDoBlockMethod() const {
    5357    auto savePoint = iBuilder->saveIP();
    5458    Module * m = iBuilder->getModule();
     
    6266    Value * blockNo = getScalarField(self, blockNoScalar);
    6367
    64     Value * inputStreamPtr = getStreamSetBlockPtr(self, "inputStreamSet", blockNo);
    65 
    66     Value * delMaskPtr = getStreamSetBlockPtr(self, "delMaskSet", blockNo);
    67 
    68     Value * outputStreamPtr = getStreamSetBlockPtr(self, "outputStreamSet", blockNo);
    69 
    70     Value * delCountPtr = getStreamSetBlockPtr(self, "deletionCounts", blockNo);
    71 
    72     Value * delMask = iBuilder->CreateBlockAlignedLoad(delMaskPtr, {iBuilder->getInt32(0), iBuilder->getInt32(0)});
     68    Value * delMaskPtr = getStream(self, "delMaskSet", blockNo, iBuilder->getInt32(0));
     69    Value * delMask = iBuilder->CreateBlockAlignedLoad(delMaskPtr);
    7370
    7471    std::vector<Value *> move_masks = parallel_prefix_deletion_masks(iBuilder, mDeletionFieldWidth, delMask);
    7572
    7673    for (unsigned j = 0; j < mStreamCount; ++j) {
    77         Value * input = iBuilder->CreateBlockAlignedLoad(inputStreamPtr, {iBuilder->getInt32(0), iBuilder->getInt32(j)});
     74        Value * inputStreamPtr = getStream(self, "inputStreamSet", blockNo, iBuilder->getInt32(j));
     75        Value * input = iBuilder->CreateBlockAlignedLoad(inputStreamPtr);
     76
    7877        Value * output = apply_parallel_prefix_deletion(iBuilder, mDeletionFieldWidth, delMask, move_masks, input);
    79         iBuilder->CreateBlockAlignedStore(output, outputStreamPtr, {iBuilder->getInt32(0), iBuilder->getInt32(j)});
     78
     79        Value * outputStreamPtr = getStream(self, "outputStreamSet", blockNo, iBuilder->getInt32(j));
     80        iBuilder->CreateBlockAlignedStore(output, outputStreamPtr);
    8081    }
    81     Value * counts = partial_sum_popcount(iBuilder, mDeletionFieldWidth, iBuilder->simd_not(delMask));
    82     iBuilder->CreateBlockAlignedStore(iBuilder->bitCast(counts), delCountPtr, {iBuilder->getInt32(0), iBuilder->getInt32(0)});
     82    Value * delCount = partial_sum_popcount(iBuilder, mDeletionFieldWidth, iBuilder->simd_not(delMask));
     83    Value * delCountPtr = getStream(self, "deletionCounts", blockNo, iBuilder->getInt32(0));
     84    iBuilder->CreateBlockAlignedStore(iBuilder->bitCast(delCount), delCountPtr);
    8385    /* Stream deletion has only been applied within fields; the actual number of data items
    8486     * has not yet changed.   */
     
    103105    Value * self = getParameter(finalBlockFunction, "self");
    104106    Value * blockNo = getScalarField(self, blockNoScalar);
    105     Value * delMaskBlock = getStreamSetBlockPtr(self, "delMaskSet", blockNo);
    106107    Value * remaining = iBuilder->CreateZExt(remainingBytes, iBuilder->getIntNTy(blockSize));
    107108    Value * EOF_del = iBuilder->bitCast(iBuilder->CreateShl(Constant::getAllOnesValue(iBuilder->getIntNTy(blockSize)), remaining));
    108     Value * const delmaskPtr = iBuilder->CreateGEP(delMaskBlock, {iBuilder->getInt32(0), iBuilder->getInt32(0)});
     109    Value * const delmaskPtr = getStream(self, "delMaskSet", blockNo, iBuilder->getInt32(0));
    109110    Value * const delmaskVal = iBuilder->CreateBlockAlignedLoad(delmaskPtr);
    110111    iBuilder->CreateBlockAlignedStore(iBuilder->CreateOr(EOF_del, delmaskVal), delmaskPtr);
     
    121122}
    122123
     124DeletionKernel::DeletionKernel(IDISA::IDISA_Builder * iBuilder, unsigned fw, unsigned streamCount) :
     125KernelBuilder(iBuilder, "del",
     126              {Binding{iBuilder->getStreamSetTy(streamCount), "inputStreamSet"},
     127               Binding{iBuilder->getStreamSetTy(), "delMaskSet"}},
     128              {Binding{iBuilder->getStreamSetTy(streamCount), "outputStreamSet"},
     129               Binding{iBuilder->getStreamSetTy(), "deletionCounts"}},
     130              {}, {}, {}),
     131mDeletionFieldWidth(fw),
     132mStreamCount(streamCount) {
     133
    123134}
     135
     136}
  • icGREP/icgrep-devel/icgrep/kernels/deletion.h

    r5246 r5260  
    2525// algorithm.
    2626//
    27 // Deletion Mask Calculation
    28 
    29 std::vector<llvm::Value *> parallel_prefix_deletion_masks(IDISA::IDISA_Builder * iBuilder, unsigned fw, llvm::Value * del_mask);
    30 
    31 // Applying Deletion Masks to a Stream
    32 
    33 llvm::Value * apply_parallel_prefix_deletion(IDISA::IDISA_Builder * iBuilder, unsigned fw, llvm::Value * del_mask, std::vector<llvm::Value *> mv, llvm::Value * strm);
    3427
    3528using namespace parabix;
     
    3932class DeletionKernel : public kernel::KernelBuilder {
    4033public:
    41     DeletionKernel(IDISA::IDISA_Builder * iBuilder, unsigned fw, unsigned streamCount) :
    42     KernelBuilder(iBuilder, "del",
    43                   {Binding{iBuilder->getStreamSetTy(streamCount), "inputStreamSet"},
    44                    Binding{iBuilder->getStreamSetTy(), "delMaskSet"}},
    45                   {Binding{iBuilder->getStreamSetTy(streamCount), "outputStreamSet"},
    46                    Binding{iBuilder->getStreamSetTy(), "deletionCounts"}},
    47                   {}, {}, {}),
    48     mDeletionFieldWidth(fw),
    49     mStreamCount(streamCount) {}
     34
     35    DeletionKernel(IDISA::IDISA_Builder * iBuilder, unsigned fw, unsigned streamCount);
    5036   
     37protected:
     38
     39    void generateDoBlockMethod() const override;
     40
     41    void generateFinalBlockMethod() const override;
     42
    5143private:
    52     void generateDoBlockMethod() const override;
    53     void generateFinalBlockMethod() const override;
    5444    unsigned mDeletionFieldWidth;
    5545    unsigned mStreamCount;
  • icGREP/icgrep-devel/icgrep/kernels/interface.cpp

    r5257 r5260  
    55
    66#include "interface.h"
     7#include <llvm/IR/Value.h>         // for Value
     8#include <llvm/IR/CallingConv.h>   // for ::C
     9#include <llvm/IR/DerivedTypes.h>  // for FunctionType (ptr only), PointerType
     10#include <llvm/IR/Function.h>      // for Function, Function::arg_iterator
    711#include <llvm/IR/Module.h>
    8 #include <llvm/IR/Type.h>
    9 #include <llvm/IR/Value.h>
    10 #include <llvm/Support/raw_ostream.h>
     12#include <IR_Gen/idisa_builder.h>
     13namespace llvm { class Module; }
     14namespace llvm { class Type; }
    1115
    1216using namespace llvm;
    13 using namespace parabix;
    1417
    1518void KernelInterface::addKernelDeclarations(Module * client) {
  • icGREP/icgrep-devel/icgrep/kernels/interface.h

    r5257 r5260  
    77#define KERNEL_INTERFACE_H
    88
    9 #include <string>
    10 #include <vector>
    11 #include <llvm/IR/Type.h>
    12 #include <IR_Gen/idisa_builder.h>
    13 #include <kernels/streamset.h>
     9#include <string>  // for string
     10#include <vector>  // for vector
     11namespace IDISA { class IDISA_Builder; }
     12namespace llvm { class Module; }
     13namespace llvm { class StructType; }
     14namespace llvm { class Type; }
     15namespace llvm { class Value; }
    1416
    1517struct Binding {
     
    4850   
    4951    // Add ExternalLinkage method declarations for the kernel to a given client module.
    50     void addKernelDeclarations(Module * client);
     52    void addKernelDeclarations(llvm::Module * client);
    5153    virtual void createInstance() = 0;
    52     void setInitialArguments(std::vector<Value *> initialParameters);
     54    void setInitialArguments(std::vector<llvm::Value *> args);
    5355    llvm::Value * getInstance() const { return mKernelInstance; }
    5456
     
    100102    IDISA::IDISA_Builder * const iBuilder;
    101103    std::string mKernelName;
    102     std::vector<Value *> mInitialArguments;
     104    std::vector<llvm::Value *> mInitialArguments;
    103105    std::vector<Binding> mStreamSetInputs;
    104106    std::vector<Binding> mStreamSetOutputs;
  • icGREP/icgrep-devel/icgrep/kernels/kernel.cpp

    r5257 r5260  
    55
    66#include "kernel.h"
    7 #include <llvm/IR/Module.h>
    8 #include <llvm/IR/Type.h>
    9 #include <llvm/IR/Value.h>
    10 #include <llvm/Support/raw_ostream.h>
    11 #include <llvm/Support/ErrorHandling.h>
    12 #include <toolchain.h>
     7#include <llvm/IR/Value.h>               // for Value
     8#include <llvm/Support/ErrorHandling.h>  // for report_fatal_error
     9#include <toolchain.h>                   // for BufferSegments, SegmentSize
     10#include "IR_Gen/idisa_builder.h"        // for IDISA_Builder
     11#include "kernels/streamset.h"           // for StreamSetBuffer
     12#include "llvm/ADT/StringRef.h"          // for StringRef, operator==
     13#include "llvm/IR/CallingConv.h"         // for ::C
     14#include "llvm/IR/Constant.h"            // for Constant
     15#include "llvm/IR/Constants.h"           // for ConstantInt
     16#include "llvm/IR/Function.h"            // for Function, Function::arg_iter...
     17#include "llvm/IR/Instructions.h"        // for LoadInst (ptr only), PHINode
     18#include "llvm/Support/Compiler.h"       // for LLVM_UNLIKELY
     19namespace llvm { class BasicBlock; }
     20namespace llvm { class Module; }
     21namespace llvm { class Type; }
    1322
    1423using namespace llvm;
    1524using namespace kernel;
     25using namespace parabix;
    1626
    1727KernelBuilder::KernelBuilder(IDISA::IDISA_Builder * builder,
     
    213223    Value * processed = getProcessedItemCount(self, mStreamSetInputs[0].name);
    214224    Value * itemsAvail = iBuilder->CreateSub(availablePos, processed);
    215 #ifndef NDEBUG
    216     iBuilder->CallPrintInt(mKernelName + "_itemsAvail", itemsAvail);
    217 #endif
     225//#ifndef NDEBUG
     226//    iBuilder->CallPrintInt(mKernelName + "_itemsAvail", itemsAvail);
     227//#endif
    218228    Value * stridesToDo = iBuilder->CreateUDiv(blocksToDo, strideBlocks);
    219229    Value * stridesAvail = iBuilder->CreateUDiv(itemsAvail, stride);
     
    243253    iBuilder->CreateBr(segmentDone);
    244254    iBuilder->SetInsertPoint(segmentDone);
    245 #ifndef NDEBUG
    246     iBuilder->CallPrintInt(mKernelName + "_processed", processed);
    247 #endif
     255//#ifndef NDEBUG
     256//    iBuilder->CallPrintInt(mKernelName + "_processed", processed);
     257//#endif
    248258    for (unsigned i = 0; i < mStreamSetOutputs.size(); i++) {
    249259        Value * produced = getProducedItemCount(self, mStreamSetOutputs[i].name);
     
    293303    Value * processed = getProcessedItemCount(self, mStreamSetInputs[0].name);
    294304    Value * itemsAvail = iBuilder->CreateSub(availablePos, processed);
    295 #ifndef NDEBUG
    296     iBuilder->CallPrintInt(mKernelName + "_itemsAvail final", itemsAvail);
    297 #endif
     305//#ifndef NDEBUG
     306//    iBuilder->CallPrintInt(mKernelName + "_itemsAvail final", itemsAvail);
     307//#endif
    298308    Value * stridesToDo = iBuilder->CreateUDiv(blocksToDo, strideBlocks);
    299309    Value * stridesAvail = iBuilder->CreateUDiv(itemsAvail, stride);
     
    317327    createFinalBlockCall(self, remainingItems);
    318328    processed = iBuilder->CreateAdd(processed, remainingItems);
    319     setProcessedItemCount(self, mStreamSetInputs[0].name, processed);
    320        
    321 #ifndef NDEBUG
    322     iBuilder->CallPrintInt(mKernelName + "_processed final", processed);
    323 #endif
     329    setProcessedItemCount(self, mStreamSetInputs[0].name, processed);       
     330//#ifndef NDEBUG
     331//    iBuilder->CallPrintInt(mKernelName + "_processed final", processed);
     332//#endif
    324333    for (unsigned i = 0; i < mStreamSetOutputs.size(); i++) {
    325334        Value * produced = getProducedItemCount(self, mStreamSetOutputs[i].name);
     
    402411}
    403412
    404 void KernelBuilder::setBlockNo(Value * self, Value * newFieldVal) const {
     413void KernelBuilder::setBlockNo(Value * self, Value * value) const {
    405414    Value * ptr = iBuilder->CreateGEP(self, {iBuilder->getInt32(0), getScalarIndex(blockNoScalar)});
    406     iBuilder->CreateStore(newFieldVal, ptr);
     415    iBuilder->CreateStore(value, ptr);
    407416}
    408417
     
    424433}
    425434
    426 size_t KernelBuilder::getStreamSetBufferSize(Value * /* self */, const std::string & name) const {
    427     const unsigned index = getStreamSetIndex(name);
    428     StreamSetBuffer * buf = nullptr;
    429     if (index < mStreamSetInputs.size()) {
    430         buf = mStreamSetInputBuffers[index];
    431     } else {
    432         buf = mStreamSetOutputBuffers[index - mStreamSetInputs.size()];
    433     }
    434     return buf->getBufferSize();
    435 }
    436 
    437435Value * KernelBuilder::getStreamSetStructPtr(Value * self, const std::string & name) const {
    438436    return getScalarField(self, name + structPtrSuffix);
    439437}
    440438
    441 Value * KernelBuilder::getStreamSetBlockPtr(Value * self, const std::string &name, Value * blockNo) const {
    442     Value * const structPtr = getStreamSetStructPtr(self, name);
    443     const unsigned index = getStreamSetIndex(name);
    444     StreamSetBuffer * buf = nullptr;
    445     if (index < mStreamSetInputs.size()) {
    446         buf = mStreamSetInputBuffers[index];
     439inline const StreamSetBuffer * KernelBuilder::getStreamSetBuffer(const std::string & name) const {
     440    const unsigned structIdx = getStreamSetIndex(name);
     441    if (structIdx < mStreamSetInputs.size()) {
     442        return mStreamSetInputBuffers[structIdx];
    447443    } else {
    448         buf = mStreamSetOutputBuffers[index - mStreamSetInputs.size()];
    449     }   
    450     return buf->getStreamSetBlockPointer(structPtr, blockNo);
    451 }
    452 
    453 Value * KernelBuilder::getStream(Value * self, const std::string & name, Value * blockNo, Value * index) {
    454     return iBuilder->CreateGEP(getStreamSetBlockPtr(self, name, blockNo), {iBuilder->getInt32(0), index});
     444        return mStreamSetOutputBuffers[structIdx - mStreamSetInputs.size()];
     445    }
     446}
     447
     448Value * KernelBuilder::getStreamSetPtr(Value * self, const std::string & name, Value * blockNo) const {
     449    return getStreamSetBuffer(name)->getStreamSetPtr(getStreamSetStructPtr(self, name), blockNo);
     450}
     451
     452Value * KernelBuilder::getStream(Value * self, const std::string & name, Value * blockNo, Value * index) const {
     453    return getStreamSetBuffer(name)->getStream(getStreamSetStructPtr(self, name), blockNo, index);
     454}
     455
     456Value * KernelBuilder::getStream(Value * self, const std::string & name, Value * blockNo, Value * index1, Value * index2) const {
     457    assert (index1->getType() == index2->getType());
     458    return getStreamSetBuffer(name)->getStream(getStreamSetStructPtr(self, name), blockNo, index1, index2);
     459}
     460
     461Value * KernelBuilder::getStreamView(Value * self, const std::string & name, Value * blockNo, Value * index) const {
     462    return getStreamSetBuffer(name)->getStreamView(getStreamSetStructPtr(self, name), blockNo, index);
     463}
     464
     465Value * KernelBuilder::getStreamView(llvm::Type * type, Value * self, const std::string & name, Value * blockNo, Value * index) const {
     466    return getStreamSetBuffer(name)->getStreamView(type, getStreamSetStructPtr(self, name), blockNo, index);
    455467}
    456468
  • icGREP/icgrep-devel/icgrep/kernels/kernel.h

    r5257 r5260  
    77#define KERNEL_BUILDER_H
    88
    9 #include "streamset.h"
    10 #include "interface.h"
    11 #include <vector>
    12 #include <llvm/IR/Type.h>
     9#include <string>           // for string
     10#include <memory>           // for unique_ptr
     11#include "interface.h"      // for KernelInterface
     12#include <boost/container/flat_map.hpp>
    1313#include <IR_Gen/idisa_builder.h>
    14 #include <boost/container/flat_map.hpp>
     14namespace llvm { class ConstantInt; }
     15namespace llvm { class Function; }
     16namespace llvm { class IntegerType; }
     17namespace llvm { class LoadInst; }
     18namespace llvm { class Type; }
     19namespace llvm { class Value; }
     20namespace parabix { class StreamSetBuffer; }
    1521
    1622const std::string blockNoScalar = "blockNo";
     
    2228const std::string blkMaskSuffix = "_blkMask";
    2329
    24 using namespace parabix;
    2530namespace kernel {
    2631   
     
    3237    // the full implementation of all required methods.     
    3338    //
    34     std::unique_ptr<Module> createKernelModule(const std::vector<StreamSetBuffer *> & inputs, const std::vector<StreamSetBuffer *> & outputs);
     39    std::unique_ptr<llvm::Module> createKernelModule(const std::vector<parabix::StreamSetBuffer *> & inputs, const std::vector<parabix::StreamSetBuffer *> & outputs);
    3540   
    3641    // Generate the Kernel to the current module (iBuilder->getModule()).
    37     void generateKernel(const std::vector<StreamSetBuffer *> & inputs, const std::vector<StreamSetBuffer *> & outputs);
     42    void generateKernel(const std::vector<parabix::StreamSetBuffer *> & inputs, const std::vector<parabix::StreamSetBuffer *> & outputs);
    3843   
    3944    void createInstance() override;
    4045
    41     Function * generateThreadFunction(const std::string & name) const;
     46    llvm::Function * generateThreadFunction(const std::string & name) const;
    4247
    43     Value * getBlockNo(Value * self) const;
    44     virtual Value * getProcessedItemCount(Value * self, const std::string & ssName) const override;
    45     virtual Value * getProducedItemCount(Value * self, const std::string & ssName) const override;
     48    llvm::Value * getBlockNo(llvm::Value * self) const;
     49    virtual llvm::Value * getProcessedItemCount(llvm::Value * self, const std::string & ssName) const override;
     50    virtual llvm::Value * getProducedItemCount(llvm::Value * self, const std::string & ssName) const override;
    4651   
    4752    bool hasNoTerminateAttribute() { return mNoTerminateAttribute;}
    4853   
    49     Value * getTerminationSignal(Value * self) const override;
     54    llvm::Value * getTerminationSignal(llvm::Value * self) const override;
    5055   
    51     inline IntegerType * getSizeTy() const {
     56    inline llvm::IntegerType * getSizeTy() const {
    5257        return getBuilder()->getSizeTy();
    5358    }
    5459
    55     inline Type * getStreamTy(const unsigned FieldWidth = 1) {
     60    inline llvm::Type * getStreamTy(const unsigned FieldWidth = 1) {
    5661        return getBuilder()->getStreamTy(FieldWidth);
    5762    }
    5863   
    59     inline Type * getStreamSetTy(const unsigned NumElements = 1, const unsigned FieldWidth = 1) {
     64    inline llvm::Type * getStreamSetTy(const unsigned NumElements = 1, const unsigned FieldWidth = 1) {
    6065        return getBuilder()->getStreamSetTy(NumElements, FieldWidth);
    6166    }
     
    7075    // data has been extracted from the kernel for further pipeline processing, the
    7176    // segment number must be incremented and stored using releaseLogicalSegmentNo.
    72     LoadInst * acquireLogicalSegmentNo(Value * self) const;
     77    llvm::LoadInst * acquireLogicalSegmentNo(llvm::Value * self) const;
    7378
    74     void releaseLogicalSegmentNo(Value * self, Value * newFieldVal) const;
     79    void releaseLogicalSegmentNo(llvm::Value * self, llvm::Value * newFieldVal) const;
    7580
    7681    virtual ~KernelBuilder() = 0;
    7782   
    78     std::vector<StreamSetBuffer *> getStreamSetInputBuffers() {return mStreamSetInputBuffers;}
    79     std::vector<StreamSetBuffer *> getStreamSetOutputBuffers() {return mStreamSetOutputBuffers;}
     83    const std::vector<const parabix::StreamSetBuffer *> & getStreamSetInputBuffers() const { return mStreamSetInputBuffers; }
     84
     85    const std::vector<const parabix::StreamSetBuffer *> & getStreamSetOutputBuffers() const { return mStreamSetOutputBuffers; }
    8086
    8187
     
    110116    virtual void generateDoBlockMethod() const = 0;
    111117
    112     virtual void generateDoBlockLogic(Value * self, Value * blockNo) const;
     118    virtual void generateDoBlockLogic(llvm::Value * self, llvm::Value * blockNo) const;
    113119
    114120    // Each kernel builder subtypre must also specify the logic for processing the
     
    129135    // Add an additional scalar field to the KernelState struct.
    130136    // Must occur before any call to addKernelDeclarations or createKernelModule.
    131     unsigned addScalar(Type * type, const std::string & name);
     137    unsigned addScalar(llvm::Type * type, const std::string & name);
    132138
    133139    unsigned getScalarCount() const;
     
    137143   
    138144    // Get the index of a named scalar field within the kernel state struct.
    139     ConstantInt * getScalarIndex(const std::string & name) const;
     145    llvm::ConstantInt * getScalarIndex(const std::string & name) const;
    140146   
    141147    // Get the value of a scalar field for a given instance.
    142     Value * getScalarField(Value * self, const std::string & fieldName) const;
     148    llvm::Value * getScalarField(llvm::Value * self, const std::string & fieldName) const;
    143149
    144150    // Set the value of a scalar field for a given instance.
    145     void setScalarField(Value * self, const std::string & fieldName, Value * newFieldVal) const;
     151    void setScalarField(llvm::Value * self, const std::string & fieldName, llvm::Value * newFieldVal) const;
    146152   
    147153    // Get a parameter by name.
    148     Value * getParameter(Function * f, const std::string & paramName) const;
     154    llvm::Value * getParameter(llvm::Function * f, const std::string & paramName) const;
    149155
    150     Value * getStream(Value * self, const std::string & name, Value * blockNo, const unsigned index) {
    151         return getStream(self, name, blockNo, iBuilder->getInt32(index));
    152     }
     156    llvm::Value * getStream(llvm::Value * self, const std::string & name, llvm::Value * blockNo, llvm::Value * index) const;
    153157
    154     Value * getStream(Value * self, const std::string & name, Value * blockNo, Value * index);
    155    
     158    llvm::Value * getStream(llvm::Value * self, const std::string & name, llvm::Value * blockNo, llvm::Value * index1, llvm::Value * index2) const;
     159
     160    llvm::Value * getStreamView(llvm::Value * self, const std::string & name, llvm::Value * blockNo, llvm::Value * index) const;
     161
     162    llvm::Value * getStreamView(llvm::Type * type, llvm::Value * self, const std::string & name, llvm::Value * blockNo, llvm::Value * index) const;
     163
    156164    // Stream set helpers.
    157165    unsigned getStreamSetIndex(const std::string & name) const;
    158166   
    159     Value * getScalarFieldPtr(Value * self, const std::string & name) const;
     167    llvm::Value * getScalarFieldPtr(llvm::Value * self, const std::string & name) const;
    160168
    161     Value * getStreamSetStructPtr(Value * self, const std::string & name) const;
     169    llvm::Value * getStreamSetStructPtr(llvm::Value * self, const std::string & name) const;
    162170
    163     size_t getStreamSetBufferSize(Value * self, const std::string & name) const;
     171    llvm::Value * getStreamSetPtr(llvm::Value * self, const std::string & name, llvm::Value * blockNo) const;
     172   
     173    void setBlockNo(llvm::Value * self, llvm::Value * value) const;
    164174
    165     Value * getStreamSetBlockPtr(Value * self, const std::string & name, Value * blockNo) const;
    166    
    167     void setBlockNo(Value * self, Value * newFieldVal) const;
     175    virtual void setProcessedItemCount(llvm::Value * self, const std::string & ssName, llvm::Value * newFieldVal) const;
    168176
    169     virtual void setProcessedItemCount(Value * self, const std::string & ssName, Value * newFieldVal) const;
     177    virtual void setProducedItemCount(llvm::Value * self, const std::string & ssName, llvm::Value * newFieldVal) const;
    170178
    171     virtual void setProducedItemCount(Value * self, const std::string & ssName, Value * newFieldVal) const;
     179    void setTerminationSignal(llvm::Value * self) const;
    172180
    173     void setTerminationSignal(Value * self) const;
     181private:
     182
     183    const parabix::StreamSetBuffer * getStreamSetBuffer(const std::string & name) const;
    174184
    175185protected:
    176186
    177     std::vector<Type *>             mKernelFields;
    178     NameMap                         mKernelMap;
    179     NameMap                         mStreamSetNameMap;
    180     std::vector<StreamSetBuffer *>  mStreamSetInputBuffers;
    181     std::vector<StreamSetBuffer *>  mStreamSetOutputBuffers;
    182     bool                            mNoTerminateAttribute;
     187    std::vector<llvm::Type *>                       mKernelFields;
     188    NameMap                                         mKernelMap;
     189    NameMap                                         mStreamSetNameMap;
     190    std::vector<const parabix::StreamSetBuffer *>   mStreamSetInputBuffers;
     191    std::vector<const parabix::StreamSetBuffer *>   mStreamSetOutputBuffers;
     192    bool                                            mNoTerminateAttribute;
    183193
    184194};
  • icGREP/icgrep-devel/icgrep/kernels/mmap_kernel.cpp

    r5259 r5260  
    33 *  This software is licensed to the public under the Open Software License 3.0.
    44 */
    5 #include <kernels/mmap_kernel.h>
     5#include "mmap_kernel.h"
     6#include <llvm/IR/Function.h>  // for Function, Function::arg_iterator
     7#include <llvm/IR/Module.h>
    68#include <IR_Gen/idisa_builder.h>
     9#include <kernels/streamset.h>
     10namespace llvm { class BasicBlock; }
     11namespace llvm { class Constant; }
     12namespace llvm { class Module; }
     13namespace llvm { class Value; }
     14
     15using namespace llvm;
    716
    817namespace kernel {
     
    7685    iBuilder->restoreIP(savePoint);
    7786}
    78    
     87
     88MMapSourceKernel::MMapSourceKernel(IDISA::IDISA_Builder * iBuilder, unsigned blocksPerSegment, unsigned codeUnitWidth) :
     89KernelBuilder(iBuilder, "mmap_source",
     90              {}, {Binding{iBuilder->getStreamSetTy(1, codeUnitWidth), "sourceBuffer"}},
     91              {Binding{iBuilder->getSizeTy(), "fileSize"}}, {}, {})
     92, mSegmentBlocks(blocksPerSegment)
     93, mCodeUnitWidth(codeUnitWidth) {
     94
    7995}
     96
     97}
  • icGREP/icgrep-devel/icgrep/kernels/mmap_kernel.h

    r5257 r5260  
    66#define MMAP_KERNEL_H
    77
    8 #include "streamset.h"
    98#include "kernel.h"
    10 #include <llvm/IR/Type.h>
    11 
    129namespace IDISA { class IDISA_Builder; }
    1310
     
    2017class MMapSourceKernel : public KernelBuilder {
    2118public:
    22     MMapSourceKernel(IDISA::IDISA_Builder * iBuilder, unsigned blocksPerSegment = 1, unsigned codeUnitWidth = 8) :
    23     KernelBuilder(iBuilder, "mmap_source",
    24                   {}, {Binding{iBuilder->getStreamSetTy(1, codeUnitWidth), "sourceBuffer"}},
    25                   {Binding{iBuilder->getSizeTy(), "fileSize"}}, {}, {}),
    26     mSegmentBlocks(blocksPerSegment),
    27     mCodeUnitWidth(codeUnitWidth) {}
     19    MMapSourceKernel(IDISA::IDISA_Builder * iBuilder, unsigned blocksPerSegment = 1, unsigned codeUnitWidth = 8);
    2820   
    2921private:
  • icGREP/icgrep-devel/icgrep/kernels/p2s_kernel.cpp

    r5247 r5260  
    33#include "IR_Gen/idisa_builder.h"
    44#include <llvm/IR/Type.h>
     5#include <llvm/IR/Module.h>
    56#include <iostream>
    67#include <stdint.h>
     
    89#include <llvm/Support/raw_ostream.h>
    910
    10 
     11using namespace llvm;
    1112
    1213namespace kernel{
     
    4041}
    4142               
    42 void p2sKernel::generateDoBlockMethod() const {
     43void P2SKernel::generateDoBlockMethod() const {
    4344    auto savePoint = iBuilder->saveIP();
    4445    Module * m = iBuilder->getModule();
     
    5051    Value * self = getParameter(doBlockFunction, "self");
    5152    Value * blockNo = getScalarField(self, blockNoScalar);
    52     Value * basisBitsBlock_ptr = getStreamSetBlockPtr(self, "basisBits", blockNo);
    53     Value * byteStreamBlock_ptr = getStreamSetBlockPtr(self, "byteStream", blockNo);
    54 
    5553    Value * p_bitblock[8];
    5654    for (unsigned i = 0; i < 8; i++) {
    57         p_bitblock[i] = iBuilder->CreateBlockAlignedLoad(basisBitsBlock_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(i)});
     55        Value * ptr = getStream(self, "basisBits", blockNo, iBuilder->getInt32(i));
     56        p_bitblock[i] = iBuilder->CreateBlockAlignedLoad(ptr);
    5857    }
    5958    Value * s_bytepack[8];
    6059    p2s(iBuilder, p_bitblock, s_bytepack);
    6160    for (unsigned j = 0; j < 8; ++j) {
    62         iBuilder->CreateBlockAlignedStore(s_bytepack[j], byteStreamBlock_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(0), iBuilder->getInt32(j)});
    63     }
    64     iBuilder->CreateRetVoid();
    65     iBuilder->restoreIP(savePoint);
    66 }
    67        
    68 void p2sKernel_withCompressedOutput::generateDoBlockMethod() const {
    69     auto savePoint = iBuilder->saveIP();
    70     Module * m = iBuilder->getModule();
    71     Type * i8PtrTy = iBuilder->getInt8PtrTy();
    72     Type * i32 = iBuilder->getIntNTy(32);
    73     Type * bitBlockPtrTy = llvm::PointerType::get(iBuilder->getBitBlockType(), 0);
    74    
    75     Function * doBlockFunction = m->getFunction(mKernelName + doBlock_suffix);
    76    
     61        Value * ptr = getStream(self, "byteStream", blockNo, iBuilder->getInt32(0), iBuilder->getInt32(j));
     62        iBuilder->CreateBlockAlignedStore(s_bytepack[j], ptr);
     63    }
     64    iBuilder->CreateRetVoid();
     65    iBuilder->restoreIP(savePoint);
     66}
     67
     68void P2SKernelWithCompressedOutput::generateDoBlockMethod() const {
     69    auto savePoint = iBuilder->saveIP();
     70    Module * m = iBuilder->getModule();
     71    Type * i8PtrTy = iBuilder->getInt8PtrTy();
     72    Type * i32 = iBuilder->getIntNTy(32);
     73    Type * bitBlockPtrTy = llvm::PointerType::get(iBuilder->getBitBlockType(), 0);
     74
     75    Function * doBlockFunction = m->getFunction(mKernelName + doBlock_suffix);
     76
    7777    iBuilder->SetInsertPoint(BasicBlock::Create(iBuilder->getContext(), "entry", doBlockFunction, 0));
    7878    Value * self = getParameter(doBlockFunction, "self");
    7979    Value * blockNo = getScalarField(self, blockNoScalar);
    80     Value * basisBitsBlock_ptr = getStreamSetBlockPtr(self, "basisBits", blockNo);
    81     Value * delCountBlock_ptr = getStreamSetBlockPtr(self, "deletionCounts", blockNo);
    82     Value * byteStreamBlock_ptr = getStreamSetBlockPtr(self, "byteStream", blockNo);
    83    
     80
     81
     82
    8483    Value * basisBits[8];
    8584    for (unsigned i = 0; i < 8; i++) {
    86         basisBits[i] = iBuilder->CreateBlockAlignedLoad(basisBitsBlock_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(i)});
     85        Value * basisBitsBlock_ptr = getStream(self, "basisBits", blockNo, iBuilder->getInt32(i));
     86        basisBits[i] = iBuilder->CreateBlockAlignedLoad(basisBitsBlock_ptr);
    8787    }
    8888    Value * bytePack[8];
    8989    p2s(iBuilder, basisBits, bytePack);
    90    
     90
    9191    unsigned units_per_register = iBuilder->getBitBlockWidth()/8;
    92    
    93     Value * unit_counts = iBuilder->fwCast(units_per_register, iBuilder->CreateBlockAlignedLoad(delCountBlock_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(0)}));
    94    
    95     Value * output_ptr = iBuilder->CreateBitCast(byteStreamBlock_ptr, i8PtrTy);
    96     Value * offset = ConstantInt::get(i32, 0);
    97    
     92    Value * delCountBlock_ptr = getStream(self, "deletionCounts", blockNo, iBuilder->getInt32(0));
     93    Value * unit_counts = iBuilder->fwCast(units_per_register, iBuilder->CreateBlockAlignedLoad(delCountBlock_ptr));
     94
     95    Value * output_ptr = getStreamView(i8PtrTy, self, "byteStream", blockNo, iBuilder->getInt32(0));
     96    Value * offset = iBuilder->getInt32(0);
    9897    for (unsigned j = 0; j < 8; ++j) {
    9998        iBuilder->CreateAlignedStore(bytePack[j], iBuilder->CreateBitCast(iBuilder->CreateGEP(output_ptr, offset), bitBlockPtrTy), 1);
     
    104103}
    105104
    106 void p2s_16Kernel::generateDoBlockMethod() const {
    107     auto savePoint = iBuilder->saveIP();
    108     Module * m = iBuilder->getModule();
    109    
    110     Function * doBlockFunction = m->getFunction(mKernelName + doBlock_suffix);
    111    
    112     iBuilder->SetInsertPoint(BasicBlock::Create(iBuilder->getContext(), "entry", doBlockFunction, 0));
    113     Value * self = getParameter(doBlockFunction, "self");
    114     Value * blockNo = getScalarField(self, blockNoScalar);
    115     Value * basisBitsBlock_ptr = getStreamSetBlockPtr(self, "basisBits", blockNo);
    116     Value * i16StreamBlock_ptr = getStreamSetBlockPtr(self, "i16Stream", blockNo);
     105void P2S16Kernel::generateDoBlockMethod() const {
     106    auto savePoint = iBuilder->saveIP();
     107    Module * m = iBuilder->getModule();
     108   
     109    Function * doBlockFunction = m->getFunction(mKernelName + doBlock_suffix);
     110   
     111    iBuilder->SetInsertPoint(BasicBlock::Create(iBuilder->getContext(), "entry", doBlockFunction, 0));
     112    Value * self = getParameter(doBlockFunction, "self");
     113    Value * blockNo = getScalarField(self, blockNoScalar);   
    117114   
    118115    Value * hi_input[8];
    119116    for (unsigned j = 0; j < 8; ++j) {
    120         hi_input[j] = iBuilder->CreateBlockAlignedLoad(basisBitsBlock_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(0), iBuilder->getInt32(j)});
     117        Value * ptr = getStream(self, "basisBits", blockNo, iBuilder->getInt32(0), iBuilder->getInt32(j));
     118        hi_input[j] = iBuilder->CreateBlockAlignedLoad(ptr);
    121119    }
    122120    Value * hi_bytes[8];
     
    125123    Value * lo_input[8];
    126124    for (unsigned j = 0; j < 8; ++j) {
    127         lo_input[j] = iBuilder->CreateBlockAlignedLoad(basisBitsBlock_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(0), iBuilder->getInt32(j+8)});
     125        Value * ptr = getStream(self, "basisBits", blockNo, iBuilder->getInt32(0), iBuilder->getInt32(j + 8));
     126        lo_input[j] = iBuilder->CreateBlockAlignedLoad(ptr);
    128127    }
    129128    Value * lo_bytes[8];
     
    133132        Value * merge0 = iBuilder->bitCast(iBuilder->esimd_mergel(8, hi_bytes[j], lo_bytes[j]));
    134133        Value * merge1 = iBuilder->bitCast(iBuilder->esimd_mergeh(8, hi_bytes[j], lo_bytes[j]));
    135         // iBuilder->getInt32(0),
    136         iBuilder->CreateBlockAlignedStore(merge0, i16StreamBlock_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(2*j)});
    137         iBuilder->CreateBlockAlignedStore(merge1, i16StreamBlock_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(2*j+1)});
    138     }
    139     iBuilder->CreateRetVoid();
    140     iBuilder->restoreIP(savePoint);
    141 }
    142 
    143 void p2s_16Kernel_withCompressedOutput::generateDoBlockMethod() const {
     134        Value * ptr0 = getStream(self, "i16Stream", blockNo, iBuilder->getInt32(2 * j));
     135        iBuilder->CreateBlockAlignedStore(merge0, ptr0);
     136        Value * ptr1 = getStream(self, "i16Stream", blockNo, iBuilder->getInt32(2 * j + 1));
     137        iBuilder->CreateBlockAlignedStore(merge1, ptr1);
     138    }
     139    iBuilder->CreateRetVoid();
     140    iBuilder->restoreIP(savePoint);
     141}
     142
     143void P2S16KernelWithCompressedOutput::generateDoBlockMethod() const {
    144144    auto savePoint = iBuilder->saveIP();
    145145    Module * m = iBuilder->getModule();
     
    150150
    151151    iBuilder->SetInsertPoint(BasicBlock::Create(iBuilder->getContext(), "entry", doBlockFunction, 0));
    152     Constant * stride = iBuilder->getSize(iBuilder->getStride());
    153152
    154153    Value * self = getParameter(doBlockFunction, "self");
    155154    Value * blockNo = getScalarField(self, blockNoScalar);
    156     Value * basisBitsBlock_ptr = getStreamSetBlockPtr(self, "basisBits", blockNo);
    157     Value * delCountBlock_ptr = getStreamSetBlockPtr(self, "deletionCounts", blockNo);
     155
     156    Value * hi_input[8];
     157    for (unsigned j = 0; j < 8; ++j) {
     158        Value * ptr = getStream(self, "basisBits", blockNo, iBuilder->getInt32(j));
     159        hi_input[j] = iBuilder->CreateBlockAlignedLoad(ptr);
     160    }
     161    Value * hi_bytes[8];
     162    p2s(iBuilder, hi_input, hi_bytes);
     163
     164    Value * lo_input[8];
     165    for (unsigned j = 0; j < 8; ++j) {
     166        Value * ptr = getStream(self, "basisBits", blockNo, iBuilder->getInt32(j + 8));
     167        lo_input[j] = iBuilder->CreateBlockAlignedLoad(ptr);
     168    }
     169    Value * lo_bytes[8];
     170    p2s(iBuilder, lo_input, lo_bytes);
     171
     172    Value * delCountBlock_ptr = getStream(self, "deletionCounts", blockNo, iBuilder->getInt32(0));
     173    Value * unit_counts = iBuilder->fwCast(iBuilder->getBitBlockWidth() / 16, iBuilder->CreateBlockAlignedLoad(delCountBlock_ptr));
     174
     175    PointerType * int16PtrTy = PointerType::get(iBuilder->getInt16Ty(), 0);
     176    ConstantInt * stride = iBuilder->getSize(iBuilder->getStride());
    158177    Value * i16UnitsGenerated = getProducedItemCount(self, "i16Stream"); // units generated to buffer
    159178    Value * i16BlockNo = iBuilder->CreateUDiv(i16UnitsGenerated, stride);
    160 
    161     Value * i16StreamBase_ptr = iBuilder->CreateBitCast(getStreamSetBlockPtr(self, "i16Stream", i16BlockNo), PointerType::get(iBuilder->getInt16Ty(), 0));
    162 
    163     Value * u16_output_ptr = iBuilder->CreateGEP(i16StreamBase_ptr, iBuilder->CreateURem(i16UnitsGenerated, stride));
    164 
    165 
    166     Value * hi_input[8];
    167     for (unsigned j = 0; j < 8; ++j) {
    168         hi_input[j] = iBuilder->CreateBlockAlignedLoad(basisBitsBlock_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(j)});
    169     }
    170     Value * hi_bytes[8];
    171     p2s(iBuilder, hi_input, hi_bytes);
    172 
    173     Value * lo_input[8];
    174     for (unsigned j = 0; j < 8; ++j) {
    175         lo_input[j] = iBuilder->CreateBlockAlignedLoad(basisBitsBlock_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(j+8)});
    176     }
    177     Value * lo_bytes[8];
    178     p2s(iBuilder, lo_input, lo_bytes);
    179 
    180     const auto UTF_16_units_per_register = iBuilder->getBitBlockWidth() / 16;
    181 
    182     Value * unit_counts = iBuilder->fwCast(UTF_16_units_per_register, iBuilder->CreateBlockAlignedLoad(delCountBlock_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(0)}));
    183 
     179    Value * u16_output_ptr = getStreamView(int16PtrTy, self, "i16Stream", i16BlockNo, iBuilder->CreateURem(i16UnitsGenerated, stride));
    184180    Value * offset = ConstantInt::get(i32, 0);
    185 
    186181    for (unsigned j = 0; j < 8; ++j) {
    187182        Value * merge0 = iBuilder->bitCast(iBuilder->esimd_mergel(8, hi_bytes[j], lo_bytes[j]));
    188183        Value * merge1 = iBuilder->bitCast(iBuilder->esimd_mergeh(8, hi_bytes[j], lo_bytes[j]));
    189184        iBuilder->CreateAlignedStore(merge0, iBuilder->CreateBitCast(iBuilder->CreateGEP(u16_output_ptr, offset), bitBlockPtrTy), 1);
    190         offset = iBuilder->CreateZExt(iBuilder->CreateExtractElement(unit_counts, iBuilder->getInt32(2*j)), i32);
     185        offset = iBuilder->CreateZExt(iBuilder->CreateExtractElement(unit_counts, iBuilder->getInt32(2 * j)), i32);
    191186        iBuilder->CreateAlignedStore(merge1, iBuilder->CreateBitCast(iBuilder->CreateGEP(u16_output_ptr, offset), bitBlockPtrTy), 1);
    192         offset = iBuilder->CreateZExt(iBuilder->CreateExtractElement(unit_counts, iBuilder->getInt32(2*j+1)), i32);
    193     }
    194 
     187        offset = iBuilder->CreateZExt(iBuilder->CreateExtractElement(unit_counts, iBuilder->getInt32(2 * j + 1)), i32);
     188    }
    195189    i16UnitsGenerated = iBuilder->CreateAdd(i16UnitsGenerated, iBuilder->CreateZExt(offset, iBuilder->getSizeTy()));
    196190    setProducedItemCount(self, "i16Stream", i16UnitsGenerated);
     
    199193}
    200194
    201 void p2s_16Kernel_withCompressedOutput::generateFinalBlockMethod() const {
     195void P2S16KernelWithCompressedOutput::generateFinalBlockMethod() const {
    202196    auto savePoint = iBuilder->saveIP();
    203197    Module * m = iBuilder->getModule();
     
    214208    }
    215209    Value * i16UnitsGenerated = getProducedItemCount(self, "i16Stream"); // units generated to buffer
    216 
    217210    iBuilder->CreateCall(doBlockFunction, doBlockArgs);
    218211    i16UnitsGenerated = getProducedItemCount(self, "i16Stream"); // units generated to buffer
     
    226219}
    227220   
    228    
    229 }
     221}
  • icGREP/icgrep-devel/icgrep/kernels/p2s_kernel.h

    r5246 r5260  
    1717
    1818   
    19 class p2sKernel : public KernelBuilder {
     19class P2SKernel : public KernelBuilder {
    2020public:
    21     p2sKernel(IDISA::IDISA_Builder * iBuilder) :
     21    P2SKernel(IDISA::IDISA_Builder * iBuilder) :
    2222    KernelBuilder(iBuilder, "p2s",
    2323                  {Binding{iBuilder->getStreamSetTy(8, 1), "basisBits"}},
     
    3030};
    3131
    32 class p2sKernel_withCompressedOutput : public KernelBuilder {
     32class P2SKernelWithCompressedOutput : public KernelBuilder {
    3333public:
    34     p2sKernel_withCompressedOutput(IDISA::IDISA_Builder * iBuilder) :
     34    P2SKernelWithCompressedOutput(IDISA::IDISA_Builder * iBuilder) :
    3535    KernelBuilder(iBuilder, "p2s_compress",
    3636                  {Binding{iBuilder->getStreamSetTy(8, 1), "basisBits"}, Binding{iBuilder->getStreamSetTy(1, 1), "deletionCounts"}},
     
    4343   
    4444
    45 class p2s_16Kernel : public KernelBuilder {
     45class P2S16Kernel : public KernelBuilder {
    4646public:
    47     p2s_16Kernel(IDISA::IDISA_Builder * iBuilder) :
     47    P2S16Kernel(IDISA::IDISA_Builder * iBuilder) :
    4848    KernelBuilder(iBuilder, "p2s_16",
    4949                  {Binding{iBuilder->getStreamSetTy(16, 1), "basisBits"}},
     
    5757
    5858   
    59 class p2s_16Kernel_withCompressedOutput : public KernelBuilder {
     59class P2S16KernelWithCompressedOutput : public KernelBuilder {
    6060public:
    61     p2s_16Kernel_withCompressedOutput(IDISA::IDISA_Builder * iBuilder) :
     61    P2S16KernelWithCompressedOutput(IDISA::IDISA_Builder * iBuilder) :
    6262    KernelBuilder(iBuilder, "p2s_16_compress",
    6363                  {Binding{iBuilder->getStreamSetTy(16, 1), "basisBits"}, Binding{iBuilder->getStreamSetTy(1, 1), "deletionCounts"}},
  • icGREP/icgrep-devel/icgrep/kernels/pipeline.cpp

    r5258 r5260  
    1414
    1515using namespace kernel;
     16using namespace parabix;
     17using namespace llvm;
     18
     19#if 0
    1620
    1721using BufferMap = std::unordered_map<StreamSetBuffer *, std::pair<KernelBuilder *, unsigned>>;
    1822
    19 
    20 static void createStreamBufferMap(BufferMap & bufferMap, std::vector<KernelBuilder *> kernels) {
     23static void createStreamBufferMap(BufferMap & bufferMap, const std::vector<KernelBuilder *> & kernels) {
    2124    for (auto k: kernels) {
    2225        auto outputSets = k->getStreamSetOutputBuffers();
     
    5760    return iBuilder->CreateUDiv(itemsToDo, iBuilder->getSize(iBuilder->getStride()));
    5861}
    59                                    
    60 
    61 
    62 Function * generateSegmentParallelPipelineThreadFunction(std::string name, IDISA::IDISA_Builder * iBuilder, std::vector<KernelBuilder *> kernels, Type * sharedStructType, int id) {
     62
     63#endif
     64
     65Function * generateSegmentParallelPipelineThreadFunction(std::string name, IDISA::IDISA_Builder * iBuilder, const std::vector<KernelBuilder *> & kernels, Type * sharedStructType, int id) {
    6366
    6467    Module * m = iBuilder->getModule();
     
    176179
    177180
    178 void generateSegmentParallelPipeline(IDISA::IDISA_Builder * iBuilder, std::vector<KernelBuilder *> kernels) {
     181void generateSegmentParallelPipeline(IDISA::IDISA_Builder * iBuilder, const std::vector<KernelBuilder *> & kernels) {
    179182   
    180183    unsigned threadNum = codegen::ThreadNum;
     
    231234}
    232235
    233 void generatePipelineParallel(IDISA::IDISA_Builder * iBuilder, std::vector<KernelBuilder *> kernels) {
     236void generatePipelineParallel(IDISA::IDISA_Builder * iBuilder, const std::vector<KernelBuilder *> & kernels) {
    234237 
    235238    Type * pthreadTy = iBuilder->getSizeTy();
     
    271274
    272275
    273 void generatePipelineLoop(IDISA::IDISA_Builder * iBuilder, std::vector<KernelBuilder *> kernels) {
     276void generatePipelineLoop(IDISA::IDISA_Builder * iBuilder, const std::vector<KernelBuilder *> & kernels) {
    274277    for (auto k : kernels) k->createInstance();
    275278    //BufferMap bufferMap;
  • icGREP/icgrep-devel/icgrep/kernels/pipeline.h

    r5238 r5260  
    66#define PIPELINE_H
    77
    8 #include <IR_Gen/idisa_builder.h>
    9 #include <kernels/interface.h>
    10 #include <kernels/kernel.h>
     8#include <vector>
     9namespace IDISA { class IDISA_Builder; }
     10namespace kernel { class KernelBuilder; }
    1111
    12 void generateSegmentParallelPipeline(IDISA::IDISA_Builder * iBuilder, std::vector<kernel::KernelBuilder *> kernels);
     12void generateSegmentParallelPipeline(IDISA::IDISA_Builder * iBuilder, const std::vector<kernel::KernelBuilder *> & kernels);
    1313
    14 void generatePipelineLoop(IDISA::IDISA_Builder * iBuilder, std::vector<kernel::KernelBuilder *> kernels);
     14void generatePipelineLoop(IDISA::IDISA_Builder * iBuilder, const std::vector<kernel::KernelBuilder *> & kernels);
    1515
    16 void generatePipelineParallel(IDISA::IDISA_Builder * iBuilder, std::vector<kernel::KernelBuilder *> kernels);
     16void generatePipelineParallel(IDISA::IDISA_Builder * iBuilder, const std::vector<kernel::KernelBuilder *> & kernels);
    1717
    1818#endif // PIPELINE_H
  • icGREP/icgrep-devel/icgrep/kernels/radix64.cpp

    r5247 r5260  
    77#include <kernels/kernel.h>
    88#include <IR_Gen/idisa_builder.h>
     9#include <llvm/IR/Module.h>
    910#include <llvm/Support/raw_ostream.h>
    1011
     12using namespace llvm;
     13
    1114namespace kernel {
    12 using namespace llvm;
    1315
    1416// This kernel produces an expanded input stream by duplicating every third byte.
     
    8688    Constant * packSize = iBuilder->getSize(PACK_SIZE);
    8789    Constant * loopItemCount = iBuilder->getSize(3 * PACK_SIZE); // 3 packs per loop.
    88     UndefValue * undefPack = UndefValue::get(iBuilder->fwVectorType(parabix::i8));
     90    UndefValue * undefPack = UndefValue::get(iBuilder->fwVectorType(8));
    8991   
    9092    const unsigned packAlign = iBuilder->getBitBlockWidth()/8;
     
    114116    Value * blockNo = getScalarField(self, blockNoScalar);
    115117
    116     Value * sourceBlockPtr = getStreamSetBlockPtr(self, "sourceStream", blockNo);
     118    // A block is made up of 8 packs.  Get the pointer to the first pack (changes the type of the pointer only).
     119    Value * sourcePackPtr = getStream(self, "sourceStream", blockNo, iBuilder->getInt32(0), iBuilder->getInt32(0));
    117120
    118121    Value * outputGenerated = getProducedItemCount(self, "expandedStream"); // bytes previously generated to output
    119122    Value * outputBlockNo = iBuilder->CreateUDiv(outputGenerated, stride);
    120 
    121     Value * outputBlockPtr = getStreamSetBlockPtr(self, "expandedStream", outputBlockNo);
    122 
    123     // A block is made up of 8 packs.  Get the pointer to the first pack (changes the type of the pointer only).
    124     Value * sourcePackPtr = iBuilder->CreateGEP(sourceBlockPtr, {iBuilder->getInt32(0), iBuilder->getInt32(0), iBuilder->getInt32(0)});
    125     Value * outputPackPtr = iBuilder->CreateGEP(outputBlockPtr, {iBuilder->getInt32(0), iBuilder->getInt32(0), iBuilder->getInt32(0)});
     123    Value * outputPackPtr = getStream(self, "expandedStream", outputBlockNo, iBuilder->getInt32(0), iBuilder->getInt32(0));
     124
    126125    Value * hasFullLoop = iBuilder->CreateICmpUGE(itemsToDo, loopItemCount);
    127126
     
    307306//    xwvuts|  nlkjzy|  barqpm|  hgfedc    Target
    308307void radix64Kernel::generateDoBlockLogic(Value * self, Value * blockNo) const {
    309     Value * expandedStream = getStreamSetBlockPtr(self, "expandedStream", blockNo);
    310     Value * radix64stream = getStreamSetBlockPtr(self, "radix64stream",blockNo);
    311308
    312309    Value * step_right_6 = iBuilder->simd_fill(32, ConstantInt::get(iBuilder->getInt32Ty(), 0x00C00000));
     
    318315   
    319316    for (unsigned i = 0; i < 8; i++) {
    320         Value * bytepack = iBuilder->CreateBlockAlignedLoad(expandedStream, {iBuilder->getInt32(0), iBuilder->getInt32(0), iBuilder->getInt32(i)});
     317        Value * expandedStream = getStream(self, "expandedStream", blockNo, iBuilder->getInt32(0), iBuilder->getInt32(i));
     318        Value * bytepack = iBuilder->CreateBlockAlignedLoad(expandedStream);
    321319
    322320        Value * right_6_result = iBuilder->simd_srli(32, iBuilder->simd_and(bytepack, step_right_6), 6);
     
    335333        Value * radix64pack = iBuilder->bitCast(mid);
    336334
    337         iBuilder->CreateBlockAlignedStore(radix64pack, radix64stream, {iBuilder->getInt32(0), iBuilder->getInt32(0), iBuilder->getInt32(i)});
     335        Value * radix64stream = getStream(self, "radix64stream",blockNo, iBuilder->getInt32(0), iBuilder->getInt32(i));
     336        iBuilder->CreateBlockAlignedStore(radix64pack, radix64stream);
    338337    }
    339338    Value * produced = getProducedItemCount(self, "radix64stream");
     
    363362    Constant * packSize = iBuilder->getSize(PACK_SIZE);
    364363    Value * blockNo = getScalarField(self, blockNoScalar);
    365     Value * expandedstream_ptr = getStreamSetBlockPtr(self, "expandedStream", blockNo);
    366     Value * radix64stream_ptr = getStreamSetBlockPtr(self, "radix64stream", blockNo);
    367364
    368365    Value * step_right_6 = iBuilder->simd_fill(32, iBuilder->getInt32(0x00C00000));
     
    383380    loopRemain->addIncoming(remainingBytes, radix64_fb_entry);
    384381
    385     Value * bytepack = iBuilder->CreateBlockAlignedLoad(expandedstream_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(0), idx});
     382    Value * expandedStreamLoopPtr = getStream(self, "expandedStream", blockNo, iBuilder->getInt32(0), idx);
     383    Value * bytepack = iBuilder->CreateBlockAlignedLoad(expandedStreamLoopPtr);
    386384    Value * right_6_result = iBuilder->simd_srli(32, iBuilder->simd_and(bytepack, step_right_6), 6);
    387385    Value * right_4_result = iBuilder->simd_srli(32, iBuilder->simd_and(bytepack, step_right_4), 4);
     
    399397    Value * radix64pack = iBuilder->bitCast(mid);
    400398
    401     iBuilder->CreateBlockAlignedStore(radix64pack, radix64stream_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(0), idx});
     399    Value * radix64streamPtr = getStream(self, "radix64stream", blockNo, iBuilder->getInt32(0), idx);
     400    iBuilder->CreateBlockAlignedStore(radix64pack, radix64streamPtr);
    402401
    403402    Value* nextIdx = iBuilder->CreateAdd(idx, ConstantInt::get(iBuilder->getInt32Ty(), 1));
     
    414413    iBuilder->SetInsertPoint(handleRemainFirstByte);
    415414    // At least one padding byte required.
    416 
    417     Value * i8output_ptr = iBuilder->CreatePointerCast(radix64stream_ptr, iBuilder->getInt8PtrTy());
    418     Value * i8input_ptr = iBuilder->CreatePointerCast(expandedstream_ptr, iBuilder->getInt8PtrTy());
     415    Value * i8input_ptr = getStreamView(iBuilder->getInt8PtrTy(), self, "expandedStream", blockNo, iBuilder->getInt32(0));
    419416    Value * remainOutputStart = iBuilder->CreateSub(remainingBytes, remainMod4);
    420417
    421     Value * firstRemainByte = iBuilder->CreateLoad(iBuilder->CreateGEP(i8input_ptr, iBuilder->getInt32(0)));
     418    Value * firstRemainByte = iBuilder->CreateLoad(i8input_ptr);
    422419
    423420    Value * first_move_right_2_mask = ConstantInt::get(iBuilder->getInt8Ty(), 0xFC);
     
    427424    Value * first_move_left_4_byte = iBuilder->CreateShl(iBuilder->CreateAnd(firstRemainByte, first_move_left_4_mask), 4);
    428425
    429     iBuilder->CreateStore(first_output_byte, iBuilder->CreateGEP(i8output_ptr, iBuilder->CreateAdd(remainOutputStart, iBuilder->getInt64(0))));
    430 
     426
     427    Value * i8OutPtr0 = getStreamView(iBuilder->getInt8PtrTy(), self, "radix64stream", blockNo, remainOutputStart);
     428
     429    iBuilder->CreateStore(first_output_byte, i8OutPtr0);
    431430
    432431    iBuilder->CreateCondBr(iBuilder->CreateICmpEQ(remainMod4, iBuilder->getSize(1)), handleNoRemainSecondByte, handleRemainSecondByte);
     
    437436    Value * second_move_right_4_byte = iBuilder->CreateLShr(iBuilder->CreateAnd(secondRemainByte, second_move_right_4_mask), 4);
    438437    Value * second_output_byte = iBuilder->CreateOr(first_move_left_4_byte, second_move_right_4_byte);
    439     iBuilder->CreateStore(second_output_byte, iBuilder->CreateGEP(i8output_ptr, iBuilder->CreateAdd(remainOutputStart, iBuilder->getInt64(1))));
     438
     439    Value * i8OutPtr1 = getStreamView(iBuilder->getInt8PtrTy(), self, "radix64stream", blockNo, iBuilder->CreateAdd(remainOutputStart, iBuilder->getInt64(1)));
     440
     441    iBuilder->CreateStore(second_output_byte, i8OutPtr1);
    440442
    441443    Value * second_move_left_2_mask = ConstantInt::get(iBuilder->getInt8Ty(), 0x0F);
    442444    Value * second_move_left_2_byte = iBuilder->CreateShl(iBuilder->CreateAnd(secondRemainByte, second_move_left_2_mask), 2);
    443     iBuilder->CreateStore(second_move_left_2_byte, iBuilder->CreateGEP(i8output_ptr, iBuilder->CreateAdd(remainOutputStart, iBuilder->getInt64(2))));
     445
     446    Value * i8OutPtr2 = getStreamView(iBuilder->getInt8PtrTy(), self, "radix64stream", blockNo, iBuilder->CreateAdd(remainOutputStart, iBuilder->getInt64(2)));
     447
     448    iBuilder->CreateStore(second_move_left_2_byte, i8OutPtr2);
    444449    iBuilder->CreateBr(fbExit);
    445450
    446451    iBuilder->SetInsertPoint(handleNoRemainSecondByte);
    447     iBuilder->CreateStore(first_move_left_4_byte, iBuilder->CreateGEP(i8output_ptr, iBuilder->CreateAdd(remainOutputStart, iBuilder->getInt64(1))));
     452
     453    i8OutPtr1 = getStreamView(iBuilder->getInt8PtrTy(), self, "radix64stream", blockNo, iBuilder->CreateAdd(remainOutputStart, iBuilder->getInt64(1)));
     454
     455    iBuilder->CreateStore(first_move_left_4_byte, i8OutPtr1);
    448456    iBuilder->CreateBr(fbExit);
    449457
     
    474482
    475483
    476 void base64Kernel::generateDoBlockLogic(Value * self, Value * blockNo) const {
    477     Value * radix64stream_ptr = getStreamSetBlockPtr(self, "radix64stream", blockNo);
    478     Value * base64stream_ptr = getStreamSetBlockPtr(self, "base64stream", blockNo);   
     484void base64Kernel::generateDoBlockLogic(Value * self, Value * blockNo) const {       
    479485    for (unsigned i = 0; i < 8; i++) {
    480         Value * bytepack = iBuilder->CreateBlockAlignedLoad(radix64stream_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(0), iBuilder->getInt32(i)});
     486        Value * radix64stream_ptr = getStream(self, "radix64stream", blockNo, iBuilder->getInt32(0), iBuilder->getInt32(i));
     487        Value * bytepack = iBuilder->CreateBlockAlignedLoad(radix64stream_ptr);
    481488        Value * mask_gt_25 = iBuilder->simd_ugt(8, bytepack, iBuilder->simd_fill(8, iBuilder->getInt8(25)));
    482489        Value * mask_gt_51 = iBuilder->simd_ugt(8, bytepack, iBuilder->simd_fill(8, iBuilder->getInt8(51)));
     
    494501        Value * t0_62 = iBuilder->simd_sub(8, t0_61, iBuilder->simd_and(mask_eq_62, iBuilder->simd_fill(8, iBuilder->getInt8(15))));
    495502        Value * base64pack = iBuilder->simd_sub(8, t0_62, iBuilder->simd_and(mask_eq_63, iBuilder->simd_fill(8, iBuilder->getInt8(2))));
    496         iBuilder->CreateBlockAlignedStore(iBuilder->bitCast(base64pack), base64stream_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(0), iBuilder->getInt32(i)});
     503        Value * base64stream_ptr = getStream(self, "base64stream", blockNo, iBuilder->getInt32(0), iBuilder->getInt32(i));
     504        iBuilder->CreateBlockAlignedStore(iBuilder->bitCast(base64pack), base64stream_ptr);
    497505    }
    498506    Value * produced = getProducedItemCount(self, "base64stream");
     
    526534    Constant * packSize = iBuilder->getSize(iBuilder->getStride() / 8);
    527535    Value * blockNo = getScalarField(self, blockNoScalar);
    528     Value * radix64stream_ptr = getStreamSetBlockPtr(self, "radix64stream", blockNo);
    529     Value * base64stream_ptr = getStreamSetBlockPtr(self, "base64stream", blockNo);
    530    
     536
    531537    // Enter the loop only if there is at least one byte remaining to process.
    532538    iBuilder->CreateCondBr(iBuilder->CreateICmpEQ(remainingBytes, iBuilder->getSize(0)), fbExit, base64_loop);
     
    537543    idx->addIncoming(ConstantInt::getNullValue(iBuilder->getInt32Ty()), base64_fb_entry);
    538544    loopRemain->addIncoming(remainingBytes, base64_fb_entry);
    539     Value * bytepack = iBuilder->CreateBlockAlignedLoad(radix64stream_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(0), idx});
     545    Value * radix64streamPtr = getStream(self, "radix64stream", blockNo, iBuilder->getInt32(0), idx);
     546    Value * bytepack = iBuilder->CreateBlockAlignedLoad(radix64streamPtr);
    540547    Value * mask_gt_25 = iBuilder->simd_ugt(8, bytepack, iBuilder->simd_fill(8, iBuilder->getInt8(25)));
    541548    Value * mask_gt_51 = iBuilder->simd_ugt(8, bytepack, iBuilder->simd_fill(8, iBuilder->getInt8(51)));
     
    547554    Value * t0_62 = iBuilder->simd_sub(8, t0_61, iBuilder->simd_and(mask_eq_62, iBuilder->simd_fill(8, iBuilder->getInt8(15))));
    548555    Value * base64pack = iBuilder->simd_sub(8, t0_62, iBuilder->simd_and(mask_eq_63, iBuilder->simd_fill(8, iBuilder->getInt8(2))));
    549     iBuilder->CreateBlockAlignedStore(iBuilder->bitCast(base64pack), base64stream_ptr, {iBuilder->getInt32(0), iBuilder->getInt32(0), idx});
     556    Value * base64streamPtr = getStream(self, "base64stream", blockNo, iBuilder->getInt32(0), idx);
     557    iBuilder->CreateBlockAlignedStore(iBuilder->bitCast(base64pack), base64streamPtr);
    550558    idx->addIncoming(iBuilder->CreateAdd(idx, ConstantInt::get(iBuilder->getInt32Ty(), 1)), base64_loop);
    551559    Value* remainAfterLoop = iBuilder->CreateSub(loopRemain, packSize);
     
    557565    iBuilder->SetInsertPoint(loopExit);
    558566    iBuilder->CreateCondBr(iBuilder->CreateICmpEQ(padBytes, iBuilder->getSize(0)), fbExit, doPadding);
     567
    559568    iBuilder->SetInsertPoint(doPadding);
    560     Value * i8output_ptr = iBuilder->CreatePointerCast(base64stream_ptr, iBuilder->getInt8PtrTy());
     569    Value * i8output_ptr = getStreamView(iBuilder->getInt8PtrTy(), self, "base64stream", blockNo, iBuilder->getInt32(0));
    561570    iBuilder->CreateStore(ConstantInt::get(iBuilder->getInt8Ty(), '='), iBuilder->CreateGEP(i8output_ptr, remainingBytes));
    562571    iBuilder->CreateCondBr(iBuilder->CreateICmpEQ(remainMod4, iBuilder->getSize(3)), fbExit, doPadding2);
  • icGREP/icgrep-devel/icgrep/kernels/radix64.h

    r5246 r5260  
    1111
    1212namespace llvm { class Module; }
     13namespace llvm { class Value; }
    1314
    1415namespace IDISA { class IDISA_Builder; }
     
    4546
    4647private:
    47     virtual void generateDoBlockLogic(Value * self, Value * blockNo) const override;
     48    virtual void generateDoBlockLogic(llvm::Value * self, llvm::Value * blockNo) const override;
    4849    virtual void generateDoBlockMethod() const override;
    4950    virtual void generateFinalBlockMethod() const override;
     
    6061   
    6162private:
    62     virtual void generateDoBlockLogic(Value * self, Value * blockNo) const override;
     63    virtual void generateDoBlockLogic(llvm::Value * self, llvm::Value * blockNo) const override;
    6364    virtual void generateFinalBlockMethod() const override;
    6465    virtual void generateDoBlockMethod() const override;
  • icGREP/icgrep-devel/icgrep/kernels/s2p_kernel.cpp

    r5252 r5260  
    33 *  This software is licensed to the public under the Open Software License 3.0.
    44 */
     5
    56#include "s2p_kernel.h"
    6 #include <kernels/kernel.h>
    7 #include <IR_Gen/idisa_builder.h>
    8 #include <llvm/Support/raw_ostream.h>
     7#include <IR_Gen/idisa_builder.h>  // for IDISA_Builder
     8#include <llvm/IR/Constant.h>      // for Constant
     9#include <llvm/IR/Module.h>
     10namespace llvm { class BasicBlock; }
     11namespace llvm { class Function; }
     12namespace llvm { class Value; }
     13
     14using namespace llvm;
    915
    1016namespace kernel {
    11 using namespace llvm;
    1217
    1318const int PACK_LANES = 1;
     
    151156    iBuilder->SetInsertPoint(finalEmptyBlock);
    152157    Value * blockNo = getScalarField(self, blockNoScalar);
    153     Value * basisBitsBlock_ptr = getStreamSetBlockPtr(self, "basisBits", blockNo);
    154     iBuilder->CreateStore(Constant::getNullValue(basisBitsBlock_ptr->getType()->getPointerElementType()), basisBitsBlock_ptr);
     158    Value * basisBitsPtr = getStreamView(self, "basisBits", blockNo, iBuilder->getInt64(0));
     159    iBuilder->CreateStore(Constant::getNullValue(basisBitsPtr->getType()->getPointerElementType()), basisBitsPtr);
    155160    iBuilder->CreateBr(exitBlock);
    156161   
     
    159164    iBuilder->restoreIP(savePoint);
    160165}
    161 
    162166   
    163167void S2PKernel::generateDoBlockLogic(Value * self, Value * blockNo) const {
    164 
    165     Value * byteStream = getStreamSetBlockPtr(self, "byteStream", blockNo);
    166     Value * basisBits = getStreamSetBlockPtr(self, "basisBits", blockNo);
    167 
    168168    Value * bytepack[8];
    169169    for (unsigned i = 0; i < 8; i++) {
    170         Value * ptr = iBuilder->CreateGEP(byteStream, {iBuilder->getInt32(0), iBuilder->getInt32(0), iBuilder->getInt32(i)});
    171         bytepack[i] = iBuilder->CreateBlockAlignedLoad(ptr);
     170        Value * byteStream = getStream(self, "byteStream", blockNo, iBuilder->getInt32(0), iBuilder->getInt32(i));
     171        bytepack[i] = iBuilder->CreateBlockAlignedLoad(byteStream);
    172172    }
    173173    Value * basisbits[8];
    174174    s2p(iBuilder, bytepack, basisbits);
    175175    for (unsigned i = 0; i < 8; ++i) {
    176         iBuilder->CreateBlockAlignedStore(basisbits[i], basisBits, {iBuilder->getInt32(0), iBuilder->getInt32(i)});
     176        Value * basisBits = getStream(self, "basisBits", blockNo, iBuilder->getInt32(i));
     177        iBuilder->CreateBlockAlignedStore(basisbits[i], basisBits);
    177178    }
    178179    Value * produced = getProducedItemCount(self, "basisBits");
  • icGREP/icgrep-devel/icgrep/kernels/s2p_kernel.h

    r5246 r5260  
    66#define S2P_KERNEL_H
    77
    8 #include "streamset.h"
    9 #include "interface.h"
    10 #include "kernel.h"
    11 
    12 namespace llvm { class Module; }
    13 
    14 namespace IDISA { class IDISA_Builder; }
     8#include "kernel.h"  // for KernelBuilder
     9namespace IDISA { class IDISA_Builder; }  // lines 14-14
     10namespace llvm { class Value; }
    1511
    1612namespace kernel {
     
    2420       
    2521private:
    26     void generateDoBlockLogic(Value * self, Value * blockNo) const override;
     22    void generateDoBlockLogic(llvm::Value * self, llvm::Value * blockNo) const override;
    2723    void generateDoBlockMethod() const override;
    2824    void generateFinalBlockMethod() const override;
  • icGREP/icgrep-devel/icgrep/kernels/scanmatchgen.cpp

    r5246 r5260  
    88#include <llvm/IR/Intrinsics.h>
    99#include <IR_Gen/idisa_builder.h>
     10#include <llvm/IR/Module.h>
    1011#include <llvm/Support/raw_os_ostream.h>
    1112
     
    5859    Value * recordStart = getScalarField(kernelStuctParam, "LineStart");
    5960    Value * recordNum = getScalarField(kernelStuctParam, "LineNum");
    60     Value * matchResultsPtr = getStreamSetBlockPtr(kernelStuctParam, "matchResults", blockNo);   
    61     Value * matches = iBuilder->CreateBlockAlignedLoad(matchResultsPtr, {iBuilder->getInt32(0), iBuilder->getInt32(0)});
    62     Value * linebreaks = iBuilder->CreateBlockAlignedLoad(matchResultsPtr, {iBuilder->getInt32(0), iBuilder->getInt32(1)});
     61    Value * matches = iBuilder->CreateBlockAlignedLoad(getStream(kernelStuctParam, "matchResults", blockNo, iBuilder->getInt32(0)));
     62    Value * linebreaks = iBuilder->CreateBlockAlignedLoad(getStream(kernelStuctParam, "matchResults", blockNo, iBuilder->getInt32(1)));
    6363    Value * matchWordVector = iBuilder->CreateBitCast(matches, scanwordVectorType);
    6464    Value * breakWordVector = iBuilder->CreateBitCast(linebreaks, scanwordVectorType);
  • icGREP/icgrep-devel/icgrep/kernels/stdout_kernel.cpp

    r5256 r5260  
    33 *  This software is licensed to the public under the Open Software License 3.0.
    44 */
    5 #include <kernels/stdout_kernel.h>
    6 #include <kernels/kernel.h>
     5#include "stdout_kernel.h"
     6#include <llvm/IR/Module.h>
    77#include <IR_Gen/idisa_builder.h>
     8
     9using namespace llvm;
    810
    911namespace kernel {
     
    5153   
    5254    Value * blockNo = getScalarField(self, blockNoScalar);
    53     //iBuilder->CallPrintInt("blockNo", blockNo);
    54     Value * basePtr = getStreamSetBlockPtr(self, "codeUnitBuffer", blockNo);
    55     //iBuilder->CallPrintInt("basePtr", iBuilder->CreatePtrToInt(basePtr, iBuilder->getInt64Ty()));
    5655    Value * byteOffset = iBuilder->CreateMul(iBuilder->CreateURem(processed, blockItems), itemBytes);
    57     Value * bytePtr = iBuilder->CreateGEP(iBuilder->CreateBitCast(basePtr, i8PtrTy), byteOffset);
    58 
     56    Value * bytePtr = getStreamView(i8PtrTy, self, "codeUnitBuffer", blockNo, byteOffset);
    5957    iBuilder->CreateWriteCall(iBuilder->getInt32(1), bytePtr, iBuilder->CreateMul(itemsToDo, itemBytes));
    6058
     
    7270    Module * m = iBuilder->getModule();
    7371    Function * finalBlockFunction = m->getFunction(mKernelName + finalBlock_suffix);
    74     Type * i8PtrTy = iBuilder->getInt8PtrTy();
    75    
     72    Type * i8PtrTy = iBuilder->getInt8PtrTy();   
    7673    iBuilder->SetInsertPoint(BasicBlock::Create(iBuilder->getContext(), "fb_flush", finalBlockFunction, 0));
    7774    Constant * blockItems = iBuilder->getSize(iBuilder->getBitBlockWidth());
     
    8380    Value * itemsAvail = iBuilder->CreateSub(producerPos, processed);
    8481    Value * blockNo = getScalarField(self, blockNoScalar);
    85     Value * basePtr = getStreamSetBlockPtr(self, "codeUnitBuffer", blockNo);
    8682    Value * byteOffset = iBuilder->CreateMul(iBuilder->CreateURem(processed, blockItems), itemBytes);
    87     Value * bytePtr = iBuilder->CreateGEP(iBuilder->CreateBitCast(basePtr, i8PtrTy), byteOffset);
     83    Value * bytePtr = getStreamView(i8PtrTy, self, "codeUnitBuffer", blockNo, byteOffset);
    8884    iBuilder->CreateWriteCall(iBuilder->getInt32(1), bytePtr, iBuilder->CreateMul(itemsAvail, itemBytes));
    8985    setProcessedItemCount(self, "codeUnitBuffer", producerPos);
  • icGREP/icgrep-devel/icgrep/kernels/streamset.cpp

    r5246 r5260  
    44 */
    55
    6    
    7 #include <kernels/streamset.h>
    8 #include <vector>
    9 #include <IR_Gen/idisa_builder.h>
    10 #include <llvm/IR/Type.h>
     6#include "streamset.h"
     7#include <IR_Gen/idisa_builder.h>  // for IDISA_Builder
     8#include <IR_Gen/types/streamtype.h>
     9#include <assert.h>                // for assert
     10#include <llvm/IR/Type.h>          // for Type
     11#include <stdexcept>               // for runtime_error
     12#include <llvm/IR/BasicBlock.h>    // for BasicBlock
     13#include <llvm/IR/Constants.h>     // for ConstantInt
     14#include <llvm/IR/DataLayout.h>    // for DataLayout
     15#include <llvm/IR/DerivedTypes.h>  // for IntegerType (ptr only), PointerType
     16#include <llvm/IR/Module.h>        // for Module
     17#include <llvm/IR/Value.h>         // for Value
     18namespace llvm { class Constant; }
     19namespace llvm { class Function; }
    1120
    1221using namespace parabix;
     22using namespace llvm;
     23using namespace IDISA;
    1324
    1425enum SS_struct_index {iProducer_pos = 0, iConsumer_pos = 1, iEnd_of_input = 2, iBuffer_ptr = 3};
    1526
    16 llvm::Value * StreamSetBuffer::getProducerPosPtr(Value * bufferStructPtr) {
    17     return iBuilder->CreateGEP(bufferStructPtr, {iBuilder->getInt32(0), iBuilder->getInt32(iProducer_pos)});
    18 }
    19 
    20 void StreamSetBuffer::setProducerPos(Value * bufferStructPtr, llvm::Value * pos) {
    21     iBuilder->CreateStore(pos, getProducerPosPtr(bufferStructPtr));
    22 }
    23 
    24 llvm::Value * StreamSetBuffer::getConsumerPosPtr(Value * bufferStructPtr) {
    25     return iBuilder->CreateGEP(bufferStructPtr, {iBuilder->getInt32(0), iBuilder->getInt32(iConsumer_pos)});
    26 }
    27 
    28 void StreamSetBuffer::setConsumerPos(Value * bufferStructPtr, Value * pos) {
    29     iBuilder->CreateStore(pos, getConsumerPosPtr(bufferStructPtr));
    30 }
    31 
    32 llvm::Value * StreamSetBuffer::getEndOfInputPtr(Value * bufferStructPtr) {
    33     return iBuilder->CreateGEP(bufferStructPtr, {iBuilder->getInt32(0), iBuilder->getInt32(iEnd_of_input)});
    34 }
    35 
    36 void StreamSetBuffer::setEndOfInput(Value * bufferStructPtr) {
    37     iBuilder->CreateStore(ConstantInt::get(iBuilder->getInt1Ty(), 1), getEndOfInputPtr(bufferStructPtr));
     27Value * StreamSetBuffer::getProducerPosPtr(Value * self) const {
     28    return iBuilder->CreateGEP(self, {iBuilder->getInt32(0), iBuilder->getInt32(iProducer_pos)});
     29}
     30
     31void StreamSetBuffer::setProducerPos(Value * self, Value * pos) const {
     32    iBuilder->CreateStore(pos, getProducerPosPtr(self));
     33}
     34
     35Value * StreamSetBuffer::getConsumerPosPtr(Value * self) const {
     36    return iBuilder->CreateGEP(self, {iBuilder->getInt32(0), iBuilder->getInt32(iConsumer_pos)});
     37}
     38
     39void StreamSetBuffer::setConsumerPos(Value * self, Value * pos) const {
     40    iBuilder->CreateStore(pos, getConsumerPosPtr(self));
     41}
     42
     43Value * StreamSetBuffer::getEndOfInputPtr(Value * self) const {
     44    return iBuilder->CreateGEP(self, {iBuilder->getInt32(0), iBuilder->getInt32(iEnd_of_input)});
     45}
     46
     47void StreamSetBuffer::setEndOfInput(Value * self) const {
     48    iBuilder->CreateStore(ConstantInt::get(iBuilder->getInt1Ty(), 1), getEndOfInputPtr(self));
     49}
     50
     51Type * StreamSetBuffer::resolveStreamTypes(Type * type) {
     52    if (auto ty = dyn_cast<ArrayType>(type)) {
     53        unsigned numElems = ty->getNumElements();
     54        auto elemTy = ty->getElementType();
     55        if (isa<StreamType>(elemTy)) {
     56            return ArrayType::get(cast<StreamType>(elemTy)->resolveType(iBuilder), numElems);
     57        }
     58    }
     59    else if (auto ty = dyn_cast<StreamType>(type)) {
     60        return ty->resolveType(iBuilder);
     61    }
     62    return type;
    3863}
    3964
    4065void StreamSetBuffer::allocateBuffer() {
    41     Type * const size_ty = iBuilder->getSizeTy();
     66    Type * const sizeTy = iBuilder->getSizeTy();
    4267    Type * const int1ty = iBuilder->getInt1Ty();
    4368    mStreamSetBufferPtr = iBuilder->CreateCacheAlignedAlloca(mStreamSetType, iBuilder->getSize(mBufferBlocks));
    4469    mStreamSetStructPtr = iBuilder->CreateCacheAlignedAlloca(mStreamSetStructType);
    45     iBuilder->CreateStore(ConstantInt::get(size_ty, 0), iBuilder->CreateGEP(mStreamSetStructPtr, {iBuilder->getInt32(0), iBuilder->getInt32(iProducer_pos)}));
    46     iBuilder->CreateStore(ConstantInt::get(size_ty, 0), iBuilder->CreateGEP(mStreamSetStructPtr, {iBuilder->getInt32(0), iBuilder->getInt32(iConsumer_pos)}));
     70    iBuilder->CreateStore(ConstantInt::get(sizeTy, 0), iBuilder->CreateGEP(mStreamSetStructPtr, {iBuilder->getInt32(0), iBuilder->getInt32(iProducer_pos)}));
     71    iBuilder->CreateStore(ConstantInt::get(sizeTy, 0), iBuilder->CreateGEP(mStreamSetStructPtr, {iBuilder->getInt32(0), iBuilder->getInt32(iConsumer_pos)}));
    4772    iBuilder->CreateStore(ConstantInt::get(int1ty, 0), iBuilder->CreateGEP(mStreamSetStructPtr, {iBuilder->getInt32(0), iBuilder->getInt32(iEnd_of_input)}));
    4873    iBuilder->CreateStore(mStreamSetBufferPtr, iBuilder->CreateGEP(mStreamSetStructPtr, {iBuilder->getInt32(0), iBuilder->getInt32(iBuffer_ptr)}));
    4974}
    5075
     76
     77Value * StreamSetBuffer::getStream(Value * self, Value * blockNo, Value * index) const {
     78    return iBuilder->CreateGEP(getStreamSetPtr(self, blockNo), {iBuilder->getInt32(0), index});
     79}
     80
     81Value * StreamSetBuffer::getStream(Value * self, Value * blockNo, Value * index1, Value * index2) const {
     82    return iBuilder->CreateGEP(getStreamSetPtr(self, blockNo), {iBuilder->getInt32(0), index1, index2});
     83}
     84
     85Value * StreamSetBuffer::getStreamView(llvm::Value * self, Value * blockNo, llvm::Value * index) const {
     86    return iBuilder->CreateGEP(getStreamSetPtr(self, blockNo), index, "view");
     87}
     88
     89Value * StreamSetBuffer::getStreamView(llvm::Type * type, llvm::Value * self, Value * blockNo, llvm::Value * index) const {
     90    return iBuilder->CreateGEP(iBuilder->CreatePointerCast(getStreamSetPtr(self, blockNo), type), index, "view");
     91}
     92
    5193// Single Block Buffer
     94
    5295// For a single block buffer, the block pointer is always the buffer base pointer.
    53 llvm::Value * SingleBlockBuffer::getStreamSetBlockPointer(llvm::Value * bufferStructPtr, llvm::Value *) {
    54     Value * handle = iBuilder->CreateGEP(bufferStructPtr, {iBuilder->getInt32(0), iBuilder->getInt32(iBuffer_ptr)});
    55     return iBuilder->CreateLoad(handle);
    56 }
    57 
     96Value * SingleBlockBuffer::getStreamSetPtr(Value * self, Value *) const {
     97    return iBuilder->CreateLoad(iBuilder->CreateGEP(self, {iBuilder->getInt32(0), iBuilder->getInt32(iBuffer_ptr)}), "sb");
     98}
    5899
    59100// External File Buffer
    60 
    61 void ExternalFileBuffer::setStreamSetBuffer(llvm::Value * ptr, Value * fileSize) {
     101void ExternalFileBuffer::setStreamSetBuffer(Value * ptr, Value * fileSize) {
    62102   
    63103    Type * const size_ty = iBuilder->getSizeTy();
     
    74114}
    75115
    76 void ExternalFileBuffer::setEmptyBuffer(llvm::Value * ptr) {
     116void ExternalFileBuffer::setEmptyBuffer(Value * ptr) {
    77117   
    78118    Type * const size_ty = iBuilder->getSizeTy();
     
    93133}
    94134
    95 llvm::Value * ExternalFileBuffer::getStreamSetBlockPointer(llvm::Value * bufferStructPtr, llvm::Value * blockNo) {
    96     Value * handle = iBuilder->CreateGEP(bufferStructPtr, {iBuilder->getInt32(0), iBuilder->getInt32(iBuffer_ptr)});
     135Value * ExternalFileBuffer::getStreamSetPtr(Value * self, Value * blockNo) const {
     136    Value * handle = iBuilder->CreateGEP(self, {iBuilder->getInt32(0), iBuilder->getInt32(iBuffer_ptr)}, "ef");
    97137    Value * bufPtr = iBuilder->CreateLoad(handle);
    98138    return iBuilder->CreateGEP(bufPtr, blockNo);
    99139}
    100140
    101 // Circular Stack Allocated Buffer
    102 
    103 llvm::Value * CircularBuffer::getStreamSetBlockPointer(llvm::Value * bufferStructPtr, llvm::Value * blockNo) {
     141// Circular Buffer
     142
     143Value * CircularBuffer::getStreamSetPtr(Value * self, Value * blockNo) const {
    104144    assert (blockNo->getType()->isIntegerTy());
    105145
    106     Value * handle = iBuilder->CreateGEP(bufferStructPtr, {iBuilder->getInt32(0), iBuilder->getInt32(iBuffer_ptr)});
     146    Value * handle = iBuilder->CreateGEP(self, {iBuilder->getInt32(0), iBuilder->getInt32(iBuffer_ptr)}, "cb");
    107147    Value * bufPtr = iBuilder->CreateLoad(handle);
    108148    Value * offset = nullptr;
     
    117157}
    118158
    119 llvm::Value * LinearCopybackBuffer::getStreamSetBlockPointer(llvm::Value * bufferStructPtr, llvm::Value * blockNo) {
    120     Constant * blockWidth = iBuilder->getSize(iBuilder->getStride());
    121     Value * consumerPos_ptr = iBuilder->CreateGEP(bufferStructPtr, {iBuilder->getInt32(0), iBuilder->getInt32(iConsumer_pos)});
     159// Linear Copyback Buffer
     160
     161Value * LinearCopybackBuffer::getStreamSetPtr(Value * self, Value * blockNo) const {
     162    Value * consumerPos_ptr = iBuilder->CreateGEP(self, {iBuilder->getInt32(0), iBuilder->getInt32(iConsumer_pos)}, "lcb.c");
    122163    Value * consumerPos = iBuilder->CreateLoad(consumerPos_ptr);
    123     Value * consumerBlock = iBuilder->CreateUDiv(consumerPos, blockWidth);
    124     Value * handle = iBuilder->CreateGEP(bufferStructPtr, {iBuilder->getInt32(0), iBuilder->getInt32(iBuffer_ptr)});
     164    Value * consumerBlock = iBuilder->CreateUDiv(consumerPos, iBuilder->getSize(iBuilder->getStride()));
     165    consumerBlock = iBuilder->CreateZExtOrTrunc(consumerBlock, blockNo->getType());
     166    Value * handle = iBuilder->CreateGEP(self, {iBuilder->getInt32(0), iBuilder->getInt32(iBuffer_ptr)}, "lcb.p");
    125167    Value * bufPtr = iBuilder->CreateLoad(handle);
    126168    return iBuilder->CreateGEP(bufPtr, iBuilder->CreateSub(blockNo, consumerBlock));
    127169}
    128170
    129 void LinearCopybackBuffer::setConsumerPos(Value * bufferStructPtr, Value * newConsumerPos) {
     171void LinearCopybackBuffer::setConsumerPos(Value * self, Value * newConsumerPos) const {
    130172    Type * const i8 = iBuilder->getInt8Ty();
    131173    Type * const i8_ptr = i8->getPointerTo(mAddrSpace);
     
    141183    Constant * const one = ConstantInt::get(sizeTy, 1);
    142184
    143     Value * const consumerPosPtr = getConsumerPosPtr(bufferStructPtr);
     185    Value * const consumerPosPtr = getConsumerPosPtr(self);
    144186    Value * const consumerPos = iBuilder->CreateLoad(consumerPosPtr);
    145187
    146188    // Ensure that the new consumer position is no less than the current position.
    147189    newConsumerPos = iBuilder->CreateSelect(iBuilder->CreateICmpULT(newConsumerPos, consumerPos), consumerPos, newConsumerPos);
    148     Value * producerPos = iBuilder->CreateLoad(iBuilder->CreateGEP(bufferStructPtr, {iBuilder->getInt32(0), iBuilder->getInt32(iProducer_pos)}));
     190    Value * producerPos = iBuilder->CreateLoad(iBuilder->CreateGEP(self, {iBuilder->getInt32(0), iBuilder->getInt32(iProducer_pos)}));
    149191
    150192    // Ensure that the new consumer position is no greater than the current producer position.
     
    167209
    168210    // Must copy back one full block for each of the streams in the stream set.
    169     Value * handle = iBuilder->CreateGEP(bufferStructPtr, {iBuilder->getInt32(0), iBuilder->getInt32(iBuffer_ptr)});
     211    Value * handle = iBuilder->CreateGEP(self, {iBuilder->getInt32(0), iBuilder->getInt32(iBuffer_ptr)});
    170212    Value * bufferPtr = iBuilder->CreateLoad(handle);
    171213    Value * const consumerBlock = iBuilder->CreateUDiv(consumerPos, blockWidth);
     
    178220
    179221    iBuilder->CreateStore(newConsumerPos, consumerPosPtr);
    180 }   
     222}
     223
     224// Expandable Buffer
     225
     226Value * ExpandableBuffer::getStreamSetPtr(Value * self, Value * blockNo) const {
     227    return nullptr;
     228}
     229
     230llvm::Value * ExpandableBuffer::getStream(llvm::Value * self, llvm::Value * blockNo, llvm::Value * index) const {
     231    return nullptr;
     232}
     233
     234llvm::Value * ExpandableBuffer::getStream(llvm::Value * self, llvm::Value * blockNo, Value *index1, Value *index2) const {
     235    return nullptr;
     236}
     237
     238llvm::Value * ExpandableBuffer::getStreamView(llvm::Value * self, llvm::Value * blockNo, llvm::Value * index) const {
     239    return nullptr;
     240}
     241
     242llvm::Value * ExpandableBuffer::getStreamView(llvm::Type * type, llvm::Value * self, llvm::Value * blockNo, llvm::Value * index) const {
     243    return nullptr;
     244}
     245
     246
     247// Constructors
     248
     249SingleBlockBuffer::SingleBlockBuffer(IDISA::IDISA_Builder * b, llvm::Type * type)
     250: StreamSetBuffer(BufferKind::BlockBuffer, b, type, 1, 0) {
     251
     252}
     253
     254ExternalFileBuffer::ExternalFileBuffer(IDISA::IDISA_Builder * b, llvm::Type * type, unsigned AddressSpace)
     255: StreamSetBuffer(BufferKind::ExternalFileBuffer, b, type, 0, AddressSpace) {
     256
     257}
     258
     259CircularBuffer::CircularBuffer(IDISA::IDISA_Builder * b, llvm::Type * type, size_t bufferBlocks, unsigned AddressSpace)
     260: StreamSetBuffer(BufferKind::CircularBuffer, b, type, bufferBlocks, AddressSpace) {
     261
     262}
     263
     264LinearCopybackBuffer::LinearCopybackBuffer(IDISA::IDISA_Builder * b, llvm::Type * type, size_t bufferBlocks, unsigned AddressSpace)
     265: StreamSetBuffer(BufferKind::LinearCopybackBuffer, b, type, bufferBlocks, AddressSpace) {
     266
     267}
     268
     269ExpandableBuffer::ExpandableBuffer(IDISA::IDISA_Builder * b, llvm::Type * type, size_t bufferBlocks, unsigned AddressSpace)
     270: StreamSetBuffer(BufferKind::ExpandableBuffer, b, type, bufferBlocks, AddressSpace) {
     271
     272}
     273
     274StreamSetBuffer::StreamSetBuffer(BufferKind k, IDISA::IDISA_Builder * b, Type * type, unsigned blocks, unsigned AddressSpace)
     275: mBufferKind(k)
     276, iBuilder(b)
     277, mStreamSetType(resolveStreamTypes(type))
     278, mBufferBlocks(blocks)
     279, mAddrSpace(AddressSpace)
     280, mStreamSetBufferPtr(nullptr)
     281, mStreamSetStructPtr(nullptr)
     282, mStreamSetStructType(StructType::get(b->getContext(),
     283                        {{b->getSizeTy(),
     284                          b->getSizeTy(),
     285                          b->getInt1Ty(),
     286                          PointerType::get(mStreamSetType, AddressSpace)}})) {
     287
     288}
  • icGREP/icgrep-devel/icgrep/kernels/streamset.h

    r5246 r5260  
    77#define STREAMSET_H
    88
    9 #include <string>
    10 #include <vector>
    11 #include <IR_Gen/idisa_builder.h>
    12 #include <llvm/IR/Type.h>
     9#include <llvm/IR/Type.h>  // for Type
     10namespace IDISA { class IDISA_Builder; }
     11namespace llvm { class PointerType; }
     12namespace llvm { class Value; }
     13namespace kernel { class KernelBuilder; }
    1314
    1415namespace parabix {
    1516   
    16 enum FieldType {i1 = 1, i2 = 2, i4 = 4, i8 = 8, i16 = 16, i32 = 32, i64 = 64, i128 = 128, i256 = 256};
    17 
    1817// Stream Set Structs hold information about the current state of a stream set buffer.
    1918
    20 llvm::Value * getProducerPosPtr(IDISA::IDISA_Builder * b, Value * bufferStructPtr);
    21 llvm::Value * getConsumerPosPtr(IDISA::IDISA_Builder * b, Value * bufferStructPtr);
    22 llvm::Value * getEndOfInputPtr(IDISA::IDISA_Builder * b, Value * bufferStructPtr);
    23 llvm::Value * getStreamSetBufferPtr(IDISA::IDISA_Builder * b, Value * bufferStructPtr);
     19llvm::Value * getProducerPosPtr(IDISA::IDISA_Builder * b, llvm::Value * self);
     20llvm::Value * getConsumerPosPtr(IDISA::IDISA_Builder * b, llvm::Value * self);
     21llvm::Value * getEndOfInputPtr(IDISA::IDISA_Builder * b, llvm::Value * self);
     22llvm::Value * getStreamSetBufferPtr(IDISA::IDISA_Builder * b, llvm::Value * self);
    2423
    2524class StreamSetBuffer {
     25    friend class kernel::KernelBuilder;
     26
    2627public:
    2728
    28     enum class BufferKind : unsigned {BlockBuffer, ExternalFileBuffer, CircularBuffer, LinearCopybackBuffer};
     29    enum class BufferKind : unsigned {BlockBuffer, ExternalFileBuffer, CircularBuffer, LinearCopybackBuffer, ExpandableBuffer};
    2930
    3031    inline BufferKind getBufferKind() const {
     
    4546
    4647    size_t getBufferSize() const { return mBufferBlocks; }
    47        
     48
    4849    llvm::Value * getStreamSetBasePtr() const { return mStreamSetBufferPtr; }
    49    
     50
    5051    llvm::Value * getStreamSetStructPtr() const { return mStreamSetStructPtr; }
    5152
    5253    virtual void allocateBuffer();
    5354
    54     // Get the buffer pointer for a given block of the stream.
    55     virtual llvm::Value * getStreamSetBlockPointer(llvm::Value * bufferStructPtr, llvm::Value * blockNo) = 0;
     55    virtual llvm::Value * getStream(llvm::Value * self, llvm::Value * blockNo, llvm::Value * index) const;
     56
     57    virtual llvm::Value * getStream(llvm::Value * self, llvm::Value * blockNo, llvm::Value * index1, llvm::Value * index2) const;
    5658   
    57     llvm::Value * getProducerPosPtr(Value * bufferStructPtr);
     59    virtual llvm::Value * getStreamView(llvm::Value * self, llvm::Value * blockNo, llvm::Value * index) const;
    5860
    59     void setProducerPos(Value * bufferStructPtr, Value * pos);
     61    virtual llvm::Value * getStreamView(llvm::Type * type, llvm::Value * self, llvm::Value * blockNo, llvm::Value * index) const;
    6062
    61     llvm::Value * getConsumerPosPtr(Value * bufferStructPtr);
     63    llvm::Value * getProducerPosPtr(llvm::Value * self) const;
    6264
    63     virtual void setConsumerPos(Value * bufferStructPtr, Value * pos);
     65    void setProducerPos(llvm::Value * self, llvm::Value * pos) const;
    6466
    65     llvm::Value * getEndOfInputPtr(Value * bufferStructPtr);
     67    llvm::Value * getConsumerPosPtr(llvm::Value * self) const;
    6668
    67     void setEndOfInput(Value * bufferStructPtr);
     69    virtual void setConsumerPos(llvm::Value * self, llvm::Value * pos) const;
     70
     71    llvm::Value * getEndOfInputPtr(llvm::Value * self) const;
     72
     73    void setEndOfInput(llvm::Value * self) const;
    6874   
    69     llvm::Type * resolveStreamTypes(llvm::Type * type) {
    70         if (auto ty = dyn_cast<ArrayType>(type)) {
    71             unsigned numElems = ty->getNumElements();
    72             auto elemTy = ty->getElementType();
    73             if (isa<IDISA::StreamType>(elemTy)) {
    74                 return ArrayType::get(cast<IDISA::StreamType>(elemTy)->resolveType(iBuilder), numElems);
    75             }
    76         }
    77         else if (auto ty = dyn_cast<IDISA::StreamType>(type)) {
    78             return ty->resolveType(iBuilder);
    79         }
    80         return type;
    81     }
     75    llvm::Type * resolveStreamTypes(llvm::Type * type);
    8276   
    8377protected:
    84     StreamSetBuffer(BufferKind k, IDISA::IDISA_Builder * b, llvm::Type * type, unsigned blocks, unsigned AddressSpace = 0)
    85     : mBufferKind(k)
    86     , iBuilder(b)
    87     , mStreamSetType(resolveStreamTypes(type))
    88     , mBufferBlocks(blocks)
    89     , mAddrSpace(AddressSpace)
    90     , mStreamSetBufferPtr(nullptr)
    91     , mStreamSetStructPtr(nullptr)
    92     , mStreamSetStructType(StructType::get(b->getContext(),
    93                             {{b->getSizeTy(),
    94                               b->getSizeTy(),
    95                               b->getInt1Ty(),
    96                               PointerType::get(mStreamSetType, AddressSpace)}})) {
    9778
    98     }
     79    StreamSetBuffer(BufferKind k, IDISA::IDISA_Builder * b, llvm::Type * type, unsigned blocks, unsigned AddressSpace);
     80
     81    // Get the buffer pointer for a given block of the stream.
     82    virtual llvm::Value * getStreamSetPtr(llvm::Value * self, llvm::Value * blockNo) const = 0;
     83
    9984protected:
    100     const BufferKind        mBufferKind;
    101     IDISA::IDISA_Builder * iBuilder;
    102     llvm::Type * const      mStreamSetType;
    103     size_t                  mBufferBlocks;
    104     int                     mAddrSpace;
    105     llvm::Value *           mStreamSetBufferPtr;
    106     llvm::Value *           mStreamSetStructPtr;
    107     llvm::Type * const      mStreamSetStructType;
     85    const BufferKind                mBufferKind;
     86    IDISA::IDISA_Builder * const    iBuilder;
     87    llvm::Type * const              mStreamSetType;
     88    const size_t                    mBufferBlocks;
     89    const int                       mAddrSpace;
     90    llvm::Value *                   mStreamSetBufferPtr;
     91    llvm::Value *                   mStreamSetStructPtr;
     92    llvm::Type * const              mStreamSetStructType;
    10893};   
    10994
     
    11398        return b->getBufferKind() == BufferKind::BlockBuffer;
    11499    }   
    115     SingleBlockBuffer(IDISA::IDISA_Builder * b, llvm::Type * type)
    116     : StreamSetBuffer(BufferKind::BlockBuffer, b, type, 1, 0) {
    117100
    118     }
    119     llvm::Value * getStreamSetBlockPointer(llvm::Value * bufferBasePtr, llvm::Value * blockNo) override;
     101    SingleBlockBuffer(IDISA::IDISA_Builder * b, llvm::Type * type);
     102
     103protected:
     104    llvm::Value * getStreamSetPtr(llvm::Value * self, llvm::Value * blockNo) const override;
    120105};
    121106
     
    126111    }
    127112   
    128     ExternalFileBuffer(IDISA::IDISA_Builder * b, llvm::Type * type, unsigned AddressSpace = 0)
    129     : StreamSetBuffer(BufferKind::ExternalFileBuffer, b, type, 0, AddressSpace) {
    130 
    131     }
     113    ExternalFileBuffer(IDISA::IDISA_Builder * b, llvm::Type * type, unsigned AddressSpace = 0);
    132114
    133115    void setStreamSetBuffer(llvm::Value * ptr, llvm::Value * fileSize);
     116
    134117    void setEmptyBuffer(llvm::Value * buffer_ptr);
    135118
     
    137120    void allocateBuffer() override;
    138121
    139     llvm::Value * getStreamSetBlockPointer(llvm::Value * bufferStructPtr, llvm::Value * blockNo) override;
     122protected:
     123    llvm::Value * getStreamSetPtr(llvm::Value * self, llvm::Value * blockNo) const override;
    140124};
    141125   
     
    146130    }
    147131 
    148     CircularBuffer(IDISA::IDISA_Builder * b, llvm::Type * type, size_t bufferBlocks, unsigned AddressSpace = 0)
    149     : StreamSetBuffer(BufferKind::CircularBuffer, b, type, bufferBlocks, AddressSpace) {
     132    CircularBuffer(IDISA::IDISA_Builder * b, llvm::Type * type, size_t bufferBlocks, unsigned AddressSpace = 0);
    150133
    151     }
    152 
    153     llvm::Value * getStreamSetBlockPointer(llvm::Value * bufferBasePtr, llvm::Value * blockNo) override;
     134protected:
     135    llvm::Value * getStreamSetPtr(llvm::Value * bufferBasePtr, llvm::Value * blockNo) const override;
    154136};
    155137   
     
    161143    static inline bool classof(const StreamSetBuffer * b) {return b->getBufferKind() == BufferKind::LinearCopybackBuffer;}
    162144   
    163     LinearCopybackBuffer(IDISA::IDISA_Builder * b, llvm::Type * type, size_t bufferBlocks, unsigned AddressSpace = 0) :
    164         StreamSetBuffer(BufferKind::LinearCopybackBuffer, b, type, bufferBlocks, AddressSpace) {}
    165    
    166     llvm::Value * getStreamSetBlockPointer(llvm::Value * bufferStructPtr, llvm::Value * blockNo) override;
    167    
     145    LinearCopybackBuffer(IDISA::IDISA_Builder * b, llvm::Type * type, size_t bufferBlocks, unsigned AddressSpace = 0);
     146
    168147    // Reset the buffer to contain data starting at the base block of new_consumer_pos,
    169148    // copying back any data beyond that position.
    170     void setConsumerPos(Value * bufferStructPtr, Value * newConsumerPos) override;
     149    void setConsumerPos(llvm::Value * self, llvm::Value * newConsumerPos) const override;
     150
     151protected:
     152    llvm::Value * getStreamSetPtr(llvm::Value * self, llvm::Value * blockNo) const override;
     153};
     154
     155// ExpandableBuffers do not allow access to the base stream set but will automatically increase the number of streams
     156// within their set whenever the index exceeds its capacity
     157//
     158class ExpandableBuffer : public StreamSetBuffer {
     159public:
     160    static inline bool classof(const StreamSetBuffer * b) {return b->getBufferKind() == BufferKind::ExpandableBuffer;}
     161
     162    ExpandableBuffer(IDISA::IDISA_Builder * b, llvm::Type * type, size_t bufferBlocks, unsigned AddressSpace = 0);
     163
     164    llvm::Value * getStream(llvm::Value * self, llvm::Value * blockNo, llvm::Value * index) const override;
     165
     166    llvm::Value * getStream(llvm::Value * self, llvm::Value * blockNo, llvm::Value * index1, llvm::Value * index2) const override;
     167
     168    llvm::Value * getStreamView(llvm::Value * self, llvm::Value * blockNo, llvm::Value * index) const override;
     169
     170    llvm::Value * getStreamView(llvm::Type * type, llvm::Value * self, llvm::Value * blockNo, llvm::Value * index) const override;
     171
     172protected:
     173
     174    llvm::Value * getStreamSetPtr(llvm::Value * self, llvm::Value * blockNo) const override;
    171175};
    172176
  • icGREP/icgrep-devel/icgrep/object_cache.h

    r4964 r5260  
    22#define OBJECT_CACHE_H
    33
    4 #include <string>
    5 
     4#include <llvm/ADT/SmallString.h>
     5#include <llvm/ExecutionEngine/ObjectCache.h>
    66#include <llvm/IR/Module.h>
    77#include <llvm/Support/MemoryBuffer.h>
    8 #include <llvm/ExecutionEngine/ObjectCache.h>
    9 #include <llvm/ADT/SmallString.h>
     8#include <string>
    109
    1110class ICGrepObjectCache : public llvm::ObjectCache {
     
    1918
    2019    private:
    21         const static size_t mPathInitLength = 256;
    22         typedef llvm::SmallString<mPathInitLength> Path;
     20        using Path = llvm::SmallString<256>;
    2321        Path CacheDir;
    2422
    25         bool getCacheFilename(const std::string &ModID, Path &CacheName);
     23        bool getCacheFilename(const std::string & ModID, Path & CacheName);
    2624};
    2725
  • icGREP/icgrep-devel/icgrep/pablo/pablo_compiler.cpp

    r5246 r5260  
    66
    77#include <pablo/pablo_compiler.h>
     8#include <pablo/pablo_kernel.h>
    89#include <pablo/pablo_toolchain.h>
    910#include <pablo/codegenstate.h>
     
    1112#include <pablo/printer_pablos.h>
    1213#include <pablo/prototype.h>
    13 #include <re/re_name.h>
    1414#include <stdexcept>
    1515#include <sstream>
     
    5454            input = mKernel->getScalarFieldPtr(mSelf, name);
    5555        } else {
    56             input = mKernel->getStreamSetBlockPtr(mSelf, name, blockNo);
     56            input = mKernel->getStreamSetPtr(mSelf, name, blockNo);
    5757        }
    5858        mMarkerMap.emplace(var, input);
     
    6666            output = mKernel->getScalarFieldPtr(mSelf, name);
    6767        } else {
    68             output = mKernel->getStreamSetBlockPtr(mSelf, name, blockNo);
     68            output = mKernel->getStreamSetPtr(mSelf, name, blockNo);
    6969        }
    7070        mMarkerMap.emplace(var, output);
     
    452452            std::string inputName = var->getName()->to_string();;
    453453            Value * blockNo = mKernel->getScalarField(mSelf, blockNoScalar);
    454             Value * lookAhead_blockPtr  = mKernel->getStreamSetBlockPtr(mSelf, inputName, iBuilder->CreateAdd(blockNo, iBuilder->getSize(block_shift)));
     454            Value * lookAhead_blockPtr  = mKernel->getStreamSetPtr(mSelf, inputName, iBuilder->CreateAdd(blockNo, iBuilder->getSize(block_shift)));
    455455            Value * lookAhead_inputPtr = iBuilder->CreateGEP(lookAhead_blockPtr, {iBuilder->getInt32(0), iBuilder->getInt32(index)});
    456456            Value * lookAhead = iBuilder->CreateBlockAlignedLoad(lookAhead_inputPtr);
     
    458458                value = lookAhead;
    459459            } else { // Need to form shift result from two adjacent blocks.
    460                 Value * lookAhead_blockPtr1  = mKernel->getStreamSetBlockPtr(mSelf, inputName, iBuilder->CreateAdd(blockNo, iBuilder->getSize(block_shift + 1)));
     460                Value * lookAhead_blockPtr1  = mKernel->getStreamSetPtr(mSelf, inputName, iBuilder->CreateAdd(blockNo, iBuilder->getSize(block_shift + 1)));
    461461                Value * lookAhead_inputPtr1 = iBuilder->CreateGEP(lookAhead_blockPtr1, {iBuilder->getInt32(0), iBuilder->getInt32(index)});
    462462                Value * lookAhead1 = iBuilder->CreateBlockAlignedLoad(lookAhead_inputPtr1);
  • icGREP/icgrep-devel/icgrep/pablo/pablo_compiler.h

    r5238 r5260  
    88#define PABLO_COMPILER_H
    99
    10 //Pablo Expressions
    11 #include <string>
    12 #include <vector>
    1310#include <unordered_map>
    14 #include <pablo/pablo_kernel.h>
    15 #include <llvm/ADT/Twine.h>
    1611#include <llvm/IR/IRBuilder.h>
    1712#include <IR_Gen/idisa_builder.h>
     
    2015
    2116namespace llvm {
    22     class Value;
    23     class Module;
    24     class ExecutionEngine;
    25     class VectorType;
    26     class PointerType;
    27     class Constant;
    28     class FunctionType;
    29     class Function;
    30     class BasicBlock;
     17class Value;
     18class Function;
    3119}
    3220
     
    3523class PabloAST;
    3624class PabloBlock;
     25class PabloKernel;
    3726class String;
    38 class Var;
    3927class Statement;
    40 class StatementList;
    4128class If;
    4229class While;
    4330class CarryManager;
    44 class Extract;
    4531
    4632class PabloCompiler {
     
    4834
    4935    using IntSet = boost::container::flat_set<unsigned>;
    50     using MarkerMap = std::unordered_map<const PabloAST *, Value *>;
     36    using MarkerMap = std::unordered_map<const PabloAST *, llvm::Value *>;
    5137
    5238public:
     
    5440    ~PabloCompiler();
    5541    void initializeKernelData();
    56     void compile(Value * const self, Function * doBlockFunction);
     42    void compile(llvm::Value * const self, llvm::Function * doBlockFunction);
    5743
    5844private:
     
    6652    void compileWhile(const While * whileStmt);
    6753
    68     Value * compileExpression(const PabloAST * expr, const bool ensureLoaded = true) const;
     54    llvm::Value * compileExpression(const PabloAST * expr, const bool ensureLoaded = true) const;
    6955
    7056private:
     
    7359    CarryManager *          mCarryManager;
    7460    PabloKernel *           mKernel;
    75     Value *                 mSelf;
    76     Function *              mFunction;
     61    llvm::Value *           mSelf;
     62    llvm::Function *        mFunction;
    7763    MarkerMap               mMarkerMap;
    7864    IntSet                  mInputStreamOffset;
  • icGREP/icgrep-devel/icgrep/pablo/pablo_kernel.cpp

    r5250 r5260  
    44 */
    55
    6 #include <pablo/pablo_kernel.h>
     6#include "pablo_kernel.h"
    77#include <pablo/codegenstate.h>
    88#include <pablo/pablo_compiler.h>
    9 // #include <llvm/Support/Debug.h>
    109#include <pablo/pe_var.h>
     10#include <llvm/IR/Module.h>
    1111#include <llvm/IR/Verifier.h>
    1212#include <IR_Gen/idisa_builder.h>
  • icGREP/icgrep-devel/icgrep/u8u16.cpp

    r5255 r5260  
    313313    delK.generateKernel({&U8u16Bits, &DelMask}, {&U16Bits, &DeletionCounts});
    314314
    315     p2s_16Kernel_withCompressedOutput p2sk(iBuilder);
     315    P2S16KernelWithCompressedOutput p2sk(iBuilder);
    316316
    317317    StdOutKernel stdoutK(iBuilder, 16);
Note: See TracChangeset for help on using the changeset viewer.