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/IDISA/idisa_nvptx_builder.cpp

    r5128 r5165  
    1010#include <llvm/IR/Intrinsics.h>
    1111#include <llvm/IR/Function.h>
     12#include <llvm/IR/InlineAsm.h>
     13#include <sstream>
    1214
    1315namespace IDISA {
     
    240242}
    241243
    242 }
     244void IDISA_NVPTX20_Builder::CreateBallotFunc(){
     245    Type * const int32ty = getInt32Ty();
     246    Type * const int1ty = getInt1Ty();
     247    Function * const ballotFn = cast<Function>(mMod->getOrInsertFunction("ballot_nvptx", int32ty, int1ty, nullptr));
     248    ballotFn->setCallingConv(CallingConv::C);
     249    Function::arg_iterator args = ballotFn->arg_begin();
     250
     251    Value * const input = &*(args++);
     252    input->setName("input");
     253
     254    SetInsertPoint(BasicBlock::Create(mMod->getContext(), "entry", ballotFn, 0));
     255
     256    Value * conv = CreateZExt(input, int32ty);
     257
     258    std::ostringstream AsmStream;
     259    AsmStream << "{.reg .pred %p1; ";
     260    AsmStream << "setp.ne.u32 %p1, $1, 0; ";
     261    AsmStream << "vote.ballot.b32  $0, %p1;}";
     262    FunctionType * AsmFnTy = FunctionType::get(int32ty, int32ty, false);
     263    llvm::InlineAsm *IA = llvm::InlineAsm::get(AsmFnTy, AsmStream.str(), "=r,r", true, false);
     264    llvm::CallInst * result = CreateCall(IA, conv);
     265    result->addAttribute(llvm::AttributeSet::FunctionIndex, llvm::Attribute::NoUnwind);
     266
     267    CreateRet(result);
     268}
     269
     270}
Note: See TracChangeset for help on using the changeset viewer.