Ignore:
Timestamp:
Sep 26, 2016, 12:05:51 PM (3 years ago)
Author:
lindanl
Message:

Add segment pipeline parallel strategy. Move ballot function to IDISA NVPTX.

File:
1 edited

Legend:

Unmodified
Added
Removed
  • icGREP/icgrep-devel/icgrep/grep_engine.cpp

    r5151 r5165  
    2424#include <llvm/IR/Verifier.h>
    2525#include <llvm/IR/TypeBuilder.h>
    26 #include <llvm/IR/InlineAsm.h>
    2726#include <UCD/UnicodeNameData.h>
    2827
     
    8786static cl::opt<bool> pipelineParallel("enable-pipeline-parallel", cl::desc("Enable multithreading with pipeline parallelism."), cl::cat(bGrepOutputOptions));
    8887
     88static cl::opt<bool> segmentPipelineParallel("enable-segment-pipeline-parallel", cl::desc("Enable multithreading with segment pipeline parallelism."), cl::cat(bGrepOutputOptions));
    8989
    9090bool isUTF_16 = false;
     
    147147using namespace parabix;
    148148
    149 void createBallotFunction(Module * m, IDISA::IDISA_Builder * iBuilder){
    150     Type * const int32ty = iBuilder->getInt32Ty();
    151     Type * const int1ty = iBuilder->getInt1Ty();
    152     Function * const ballotFn = cast<Function>(m->getOrInsertFunction("ballot_nvptx", int32ty, int1ty, nullptr));
    153     ballotFn->setCallingConv(CallingConv::C);
    154     Function::arg_iterator args = ballotFn->arg_begin();
    155 
    156     Value * const input = &*(args++);
    157     input->setName("input");
    158 
    159     iBuilder->SetInsertPoint(BasicBlock::Create(m->getContext(), "entry", ballotFn, 0));
    160 
    161     Value * conv = iBuilder->CreateZExt(input, int32ty);
    162 
    163     std::ostringstream AsmStream;
    164     AsmStream << "{.reg .pred %p1; ";
    165     AsmStream << "setp.ne.u32 %p1, $1, 0; ";
    166     AsmStream << "vote.ballot.b32  $0, %p1;}";
    167     FunctionType * AsmFnTy = FunctionType::get(int32ty, int32ty, false);
    168     llvm::InlineAsm *IA = llvm::InlineAsm::get(AsmFnTy, AsmStream.str(), "=r,r", true, false);
    169     llvm::CallInst * result = iBuilder->CreateCall(IA, conv);
    170     result->addAttribute(llvm::AttributeSet::FunctionIndex, llvm::Attribute::NoUnwind);
    171 
    172     iBuilder->CreateRet(result);
    173 }
    174 
    175149Function * generateGPUKernel(Module * m, IDISA::IDISA_Builder * iBuilder, bool CountOnly){
    176150    Type * const int64ty = iBuilder->getInt64Ty();
     
    272246
    273247    const unsigned segmentSize = codegen::SegmentSize;
     248    if (segmentPipelineParallel && codegen::BufferSegments < 2) codegen::BufferSegments = 2;
    274249    const unsigned bufferSegments = codegen::BufferSegments;
     250
    275251
    276252    unsigned encodingBits = UTF_16 ? 16 : 8;
     
    380356            generatePipelineParallel(iBuilder, {&s2pk, &icgrepK}, {s2pInstance, icgrepInstance});
    381357        }
     358        else if (segmentPipelineParallel){           
     359            generateSegmentParallelPipeline(iBuilder, {&s2pk, &icgrepK}, {s2pInstance, icgrepInstance}, fileSize);
     360        }
    382361        else{
    383362            generatePipelineLoop(iBuilder, {&s2pk, &icgrepK}, {s2pInstance, icgrepInstance}, fileSize);
     
    417396                generatePipelineParallel(iBuilder, {&s2pk, &icgrepK, &scanMatchK}, {s2pInstance, icgrepInstance, scanMatchInstance});
    418397            }
     398            else if (segmentPipelineParallel){
     399                generateSegmentParallelPipeline(iBuilder, {&s2pk, &icgrepK, &scanMatchK}, {s2pInstance, icgrepInstance, scanMatchInstance}, fileSize);
     400            }
    419401            else{
    420402                generatePipelineLoop(iBuilder, {&s2pk, &icgrepK, &scanMatchK}, {s2pInstance, icgrepInstance, scanMatchInstance}, fileSize);
Note: See TracChangeset for help on using the changeset viewer.