source: icGREP/icgrep-devel/icgrep/grep_engine.cpp @ 5151

Last change on this file since 5151 was 5151, checked in by lindanl, 3 years ago

Add NVPTX Arch to the framework. Fix directory bug in make check.

File size: 23.9 KB
Line 
1/*
2 *  Copyright (c) 2016 International Characters.
3 *  This software is licensed to the public under the Open Software License 3.0.
4 *  icgrep is a trademark of International Characters.
5 */
6
7#include <grep_engine.h>
8#include <IDISA/idisa_builder.h>
9#include <IDISA/idisa_target.h>
10#include <llvm/Support/CommandLine.h>
11#include <re/re_toolchain.h>
12#include <re/re_cc.h>
13
14#include <pablo/pablo_toolchain.h>
15#include <toolchain.h>
16#include <pablo/pablo_compiler.h>
17#include <kernels/pipeline.h>
18#include <llvm/IR/Function.h>
19#include <llvm/IR/Type.h>
20#include <llvm/IR/Module.h>
21#include <llvm/ExecutionEngine/MCJIT.h>
22#include <llvm/IRReader/IRReader.h>
23#include <llvm/Support/Debug.h>
24#include <llvm/IR/Verifier.h>
25#include <llvm/IR/TypeBuilder.h>
26#include <llvm/IR/InlineAsm.h>
27#include <UCD/UnicodeNameData.h>
28
29
30#include <kernels/streamset.h>
31#include <kernels/scanmatchgen.h>
32#include <kernels/s2p_kernel.h>
33#include <kernels/cc_kernel.h>
34#include <kernels/pipeline.h>
35
36#include <pablo/function.h>
37#include <pablo/pablo_kernel.h>
38#include <pablo/pablo_toolchain.h>
39
40#include <llvm/IR/Intrinsics.h>
41#include "llvm/Support/SourceMgr.h"
42#include "llvm/IRReader/IRReader.h"
43#include "llvm/Linker/Linker.h"
44
45
46#include <fstream>
47#include <sstream>
48#include <iostream>
49#include <string>
50#include <stdint.h>
51
52#include <stdio.h>
53#include <stdlib.h>
54#include <unistd.h>
55#include <errno.h>
56#include <sys/types.h>
57#include <sys/stat.h>
58#include <stdexcept>
59#include <cctype>
60
61
62#include <llvm/Support/raw_os_ostream.h>
63
64// mmap system
65#include <boost/filesystem.hpp>
66#include <boost/iostreams/device/mapped_file.hpp>
67
68#include <fcntl.h>
69
70#include <kernels/kernel.h>
71
72#ifdef CUDA_ENABLED
73#include <IDISA/CudaDriver.h>
74#endif
75
76static cl::OptionCategory bGrepOutputOptions("Output Options",
77                                             "These options control the output.");
78
79static cl::opt<bool> NormalizeLineBreaks("normalize-line-breaks", cl::desc("Normalize line breaks to std::endl."), cl::init(false),  cl::cat(bGrepOutputOptions));
80
81static cl::opt<bool> ShowFileNames("H", cl::desc("Show the file name with each matching line."), cl::cat(bGrepOutputOptions));
82static cl::alias ShowFileNamesLong("with-filename", cl::desc("Alias for -H"), cl::aliasopt(ShowFileNames));
83
84static cl::opt<bool> ShowLineNumbers("n", cl::desc("Show the line number with each matching line."), cl::cat(bGrepOutputOptions));
85static cl::alias ShowLineNumbersLong("line-number", cl::desc("Alias for -n"), cl::aliasopt(ShowLineNumbers));
86
87static cl::opt<bool> pipelineParallel("enable-pipeline-parallel", cl::desc("Enable multithreading with pipeline parallelism."), cl::cat(bGrepOutputOptions));
88
89
90bool isUTF_16 = false;
91std::string IRFilename = "icgrep.ll";
92std::string PTXFilename = "icgrep.ptx";
93
94void GrepEngine::doGrep(const std::string & fileName, const int fileIdx, bool CountOnly, std::vector<size_t> & total_CountOnly, bool UTF_16) {
95    boost::filesystem::path file(fileName);
96    if (exists(file)) {
97        if (is_directory(file)) {
98            return;
99        }
100    } else {
101        std::cerr << "Error: cannot open " << fileName << " for processing. Skipped.\n";
102        return;
103    }
104
105    const auto fileSize = file_size(file);
106    if (fileSize > 0) {
107        try {
108            boost::iostreams::mapped_file_source source(fileName, fileSize, 0);
109            char * fileBuffer = const_cast<char *>(source.data());
110#ifdef CUDA_ENABLED 
111            if(codegen::NVPTX){
112                ulong * rslt = RunPTX(PTXFilename, fileBuffer, fileSize, CountOnly);
113                if (CountOnly){
114                    exit(0);
115                }
116                else{
117                    mGrepFunction_CPU((char *)rslt, fileBuffer, fileSize, fileIdx);
118                    return;
119                }
120               
121            } 
122#endif
123            if (CountOnly) {
124                total_CountOnly[fileIdx] = mGrepFunction_CountOnly(fileBuffer, fileSize, fileIdx);
125            } else {
126                mGrepFunction(fileBuffer, fileSize, fileIdx);
127            }
128            source.close();
129        } catch (std::exception & e) {
130            throw std::runtime_error("Boost mmap error: " + fileName + ": " + e.what());
131        }
132    } else {
133#ifdef CUDA_ENABLED
134        if (codegen::NVPTX){
135            std::cout << 0 << std::endl;
136            exit(0);
137        }
138#endif
139        if (CountOnly) {
140            total_CountOnly[fileIdx] = mGrepFunction_CountOnly(nullptr, 0, fileIdx);
141        } else {
142            mGrepFunction(nullptr, 0, fileIdx);
143        }
144    }
145}
146
147using namespace parabix;
148
149void 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
175Function * generateGPUKernel(Module * m, IDISA::IDISA_Builder * iBuilder, bool CountOnly){
176    Type * const int64ty = iBuilder->getInt64Ty();
177    Type * const inputType = PointerType::get(int64ty, 1);
178    Type * const resultTy = iBuilder->getVoidTy();
179    Function * kernelFunc = cast<Function>(m->getOrInsertFunction("GPU_Main", resultTy, inputType, inputType, inputType, nullptr));
180    kernelFunc->setCallingConv(CallingConv::C);
181    Function::arg_iterator args = kernelFunc->arg_begin();
182
183    Value * const inputPtr = &*(args++);
184    inputPtr->setName("inputPtr");
185    Value * const bufferSizePtr = &*(args++);
186    bufferSizePtr->setName("bufferSizePtr");
187    Value * const outputPtr = &*(args++);
188    outputPtr->setName("resultPtr");
189
190    BasicBlock * entryBlock = BasicBlock::Create(m->getContext(), "entry", kernelFunc, 0);
191    iBuilder->SetInsertPoint(entryBlock);
192
193    Function * tidFunc = m->getFunction("llvm.nvvm.read.ptx.sreg.tid.x");
194    Value * id = iBuilder->CreateCall(tidFunc);
195
196    Function * mainFunc = m->getFunction("Main");
197    Value * inputThreadPtr = iBuilder->CreateGEP(inputPtr, iBuilder->CreateMul(id, iBuilder->getInt32(8)));
198    Type * const inputStreamType = PointerType::get(ArrayType::get(ArrayType::get(iBuilder->getBitBlockType(), 8), 1), 1);
199    Value * inputStream = iBuilder->CreateBitCast(inputThreadPtr, inputStreamType);   
200    Value * bufferSize = iBuilder->CreateLoad(bufferSizePtr);
201    if (CountOnly){
202        Value * outputThreadPtr = iBuilder->CreateGEP(outputPtr, id);
203        Value * result = iBuilder->CreateCall(mainFunc, {inputStream, bufferSize});
204        iBuilder->CreateStore(result, outputThreadPtr);
205    }
206    else {
207        Value * outputThreadPtr = iBuilder->CreateGEP(outputPtr, iBuilder->CreateMul(id, iBuilder->getInt32(2)));
208        Type * const outputStremType = PointerType::get(ArrayType::get(iBuilder->getBitBlockType(), 2), 1);
209        Value * outputStream = iBuilder->CreateBitCast(outputThreadPtr, outputStremType);
210        iBuilder->CreateCall(mainFunc, {inputStream, bufferSize, outputStream});
211    }   
212
213    iBuilder->CreateRetVoid();
214    return kernelFunc;
215}
216
217Function * generateCPUKernel(Module * m, IDISA::IDISA_Builder * iBuilder, bool isNameExpression){
218    Type * const size_ty = iBuilder->getSizeTy();
219    Type * const int8PtrTy = iBuilder->getInt8PtrTy();
220    Type * const rsltType = PointerType::get(ArrayType::get(iBuilder->getBitBlockType(), 2), 0);
221    Function * const mainCPUFn = cast<Function>(m->getOrInsertFunction("CPU_Main", iBuilder->getVoidTy(), rsltType, int8PtrTy, size_ty, size_ty, nullptr));
222    mainCPUFn->setCallingConv(CallingConv::C);
223    iBuilder->SetInsertPoint(BasicBlock::Create(m->getContext(), "entry", mainCPUFn, 0));
224    Function::arg_iterator args = mainCPUFn->arg_begin();
225   
226    Value * const rsltStream = &*(args++);
227    rsltStream->setName("rslt");
228    Value * const inputStream = &*(args++);
229    inputStream->setName("input");
230    Value * const fileSize = &*(args++);
231    fileSize->setName("fileSize");
232    Value * const fileIdx = &*(args++);
233    fileIdx->setName("fileIdx");
234
235    ExternalFileBuffer MatchResults(iBuilder, StreamSetType(2, i1));
236    MatchResults.setStreamSetBuffer(rsltStream, fileSize);
237
238    kernel::scanMatchKernel scanMatchK(iBuilder, isNameExpression);
239    scanMatchK.generateKernel({&MatchResults}, {});
240           
241    Value * scanMatchInstance = scanMatchK.createInstance({inputStream, fileSize, fileIdx});
242   
243    generatePipelineLoop(iBuilder, {&scanMatchK}, {scanMatchInstance}, fileSize);
244    iBuilder->CreateRetVoid();
245
246    return mainCPUFn;
247}
248
249void GrepEngine::grepCodeGen(std::string moduleName, re::RE * re_ast, bool CountOnly, bool UTF_16, bool isNameExpression) {
250    isUTF_16 = UTF_16;
251    int addrSpace = 0;
252    bool CPU_Only = true;
253
254    Module * cpuM = new Module(moduleName+":cpu", getGlobalContext());
255    IDISA::IDISA_Builder * CPUBuilder = IDISA::GetIDISA_Builder(cpuM); 
256    Module * M = cpuM; 
257    IDISA::IDISA_Builder * iBuilder = CPUBuilder; 
258
259#ifdef CUDA_ENABLED
260    setNVPTXOption(); 
261    if(codegen::NVPTX){     
262        Module * gpuM = new Module(moduleName+":gpu", getGlobalContext());
263        IDISA::IDISA_Builder * GPUBuilder = IDISA::GetIDISA_GPU_Builder(gpuM);
264        M = gpuM;
265        iBuilder = GPUBuilder;
266        M->setDataLayout("e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64");
267        M->setTargetTriple("nvptx64-nvidia-cuda");
268        addrSpace = 1;
269        CPU_Only = false;
270    }   
271#endif
272
273    const unsigned segmentSize = codegen::SegmentSize;
274    const unsigned bufferSegments = codegen::BufferSegments;
275
276    unsigned encodingBits = UTF_16 ? 16 : 8;
277
278    mIsNameExpression = isNameExpression;
279
280    Type * const int32ty = iBuilder->getInt32Ty();
281    Type * const size_ty = iBuilder->getSizeTy();
282    Type * const int8PtrTy = iBuilder->getInt8PtrTy();
283    Type * const voidTy = Type::getVoidTy(M->getContext());   
284    Type * const voidPtrTy = TypeBuilder<void *, false>::get(M->getContext());
285    Type * const inputType = PointerType::get(ArrayType::get(ArrayType::get(iBuilder->getBitBlockType(), (UTF_16 ? 16 : 8)), 1), addrSpace);
286    Type * const resultTy = CountOnly ? size_ty : iBuilder->getVoidTy();
287
288    Function * mainFn = nullptr;
289    Value * inputStream = nullptr;
290    Value * fileSize = nullptr;
291    Value * fileIdx = nullptr;
292
293#ifdef CUDA_ENABLED   
294    Value * outputStream = nullptr;
295    Type * const outputType = PointerType::get(ArrayType::get(iBuilder->getBitBlockType(), 2), addrSpace);
296    if (codegen::NVPTX){
297        if (CountOnly){
298            mainFn = cast<Function>(M->getOrInsertFunction("Main", resultTy, inputType, size_ty, nullptr));
299            mainFn->setCallingConv(CallingConv::C);
300            iBuilder->SetInsertPoint(BasicBlock::Create(M->getContext(), "entry", mainFn, 0));
301            Function::arg_iterator args = mainFn->arg_begin();
302           
303            inputStream = &*(args++);
304            inputStream->setName("input");
305            fileSize = &*(args++);
306            fileSize->setName("fileSize");
307        }
308        else{
309            mainFn = cast<Function>(M->getOrInsertFunction("Main", resultTy, inputType, size_ty, outputType, nullptr));
310            mainFn->setCallingConv(CallingConv::C);
311            iBuilder->SetInsertPoint(BasicBlock::Create(M->getContext(), "entry", mainFn, 0));
312            Function::arg_iterator args = mainFn->arg_begin();
313           
314            inputStream = &*(args++);
315            inputStream->setName("input");
316            fileSize = &*(args++);
317            fileSize->setName("fileSize");
318            outputStream = &*(args++);
319            outputStream->setName("output");
320        }
321    } 
322#endif
323    if (CPU_Only){
324        mainFn = cast<Function>(M->getOrInsertFunction("Main", resultTy, inputType, size_ty, size_ty, nullptr));   
325        mainFn->setCallingConv(CallingConv::C);
326        iBuilder->SetInsertPoint(BasicBlock::Create(M->getContext(), "entry", mainFn, 0));
327        Function::arg_iterator args = mainFn->arg_begin();
328       
329        inputStream = &*(args++);
330        inputStream->setName("input");
331        fileSize = &*(args++);
332        fileSize->setName("fileSize");
333        fileIdx = &*(args++);
334        fileIdx->setName("fileIdx");
335    }
336       
337    ExternalFileBuffer ByteStream(iBuilder, StreamSetType(1, i8));
338    CircularBuffer BasisBits(iBuilder, StreamSetType(8, i1), segmentSize * bufferSegments);
339
340    kernel::s2pKernel  s2pk(iBuilder);
341    s2pk.generateKernel({&ByteStream}, {&BasisBits});
342   
343    re_ast = re::regular_expression_passes(re_ast);
344    pablo::PabloFunction * function = re::re2pablo_compiler(encodingBits, re_ast, CountOnly);
345    pablo_function_passes(function);
346
347    ByteStream.setStreamSetBuffer(inputStream, fileSize);
348    BasisBits.allocateBuffer();
349
350    Value * s2pInstance = s2pk.createInstance({});
351 
352    Type * pthreadTy = size_ty;
353    FunctionType * funVoidPtrVoidTy = FunctionType::get(voidTy, int8PtrTy, false);   
354   
355    Function * pthreadCreateFunc = cast<Function>(M->getOrInsertFunction("pthread_create",
356                                        int32ty, 
357                                        pthreadTy->getPointerTo(), 
358                                        voidPtrTy, 
359                                        static_cast<Type *>(funVoidPtrVoidTy)->getPointerTo(),
360                                        voidPtrTy, nullptr));
361    pthreadCreateFunc->setCallingConv(llvm::CallingConv::C);
362    Function * pthreadJoinFunc = cast<Function>(M->getOrInsertFunction("pthread_join", 
363                                        int32ty, 
364                                        pthreadTy, 
365                                        PointerType::get(int8PtrTy, 0), nullptr));
366    pthreadJoinFunc->setCallingConv(llvm::CallingConv::C);
367
368    Function * pthreadExitFunc = cast<Function>(M->getOrInsertFunction("pthread_exit", 
369                                        voidTy, 
370                                        voidPtrTy, nullptr));
371    pthreadExitFunc->addFnAttr(llvm::Attribute::NoReturn);
372    pthreadExitFunc->setCallingConv(llvm::CallingConv::C);
373
374    if (CountOnly) {
375        pablo::PabloKernel  icgrepK(iBuilder, "icgrep", function, {"matchedLineCount"});
376        icgrepK.generateKernel({&BasisBits}, {});       
377        Value * icgrepInstance = icgrepK.createInstance({});
378
379        if (pipelineParallel){
380            generatePipelineParallel(iBuilder, {&s2pk, &icgrepK}, {s2pInstance, icgrepInstance});
381        }
382        else{
383            generatePipelineLoop(iBuilder, {&s2pk, &icgrepK}, {s2pInstance, icgrepInstance}, fileSize);
384        }
385       
386        Value * matchCount = icgrepK.createGetAccumulatorCall(icgrepInstance, "matchedLineCount");
387        iBuilder->CreateRet(matchCount);
388
389    }
390    else {
391#ifdef CUDA_ENABLED
392        if (codegen::NVPTX){
393            ExternalFileBuffer MatchResults(iBuilder, StreamSetType(2, i1), addrSpace);
394            MatchResults.setStreamSetBuffer(outputStream, fileSize);
395
396            pablo::PabloKernel  icgrepK(iBuilder, "icgrep", function, {});
397            icgrepK.generateKernel({&BasisBits},  {&MatchResults});
398            Value * icgrepInstance = icgrepK.createInstance({});
399
400            generatePipelineLoop(iBuilder, {&s2pk, &icgrepK}, {s2pInstance, icgrepInstance}, fileSize);
401
402        }
403#endif
404        if (CPU_Only){
405            CircularBuffer MatchResults(iBuilder, StreamSetType(2, i1), segmentSize * bufferSegments);
406            MatchResults.allocateBuffer();
407
408            pablo::PabloKernel  icgrepK(iBuilder, "icgrep", function, {});
409            icgrepK.generateKernel({&BasisBits},  {&MatchResults});
410            Value * icgrepInstance = icgrepK.createInstance({});
411
412            kernel::scanMatchKernel scanMatchK(iBuilder, mIsNameExpression);
413            scanMatchK.generateKernel({&MatchResults}, {});               
414            Value * scanMatchInstance = scanMatchK.createInstance({iBuilder->CreateBitCast(inputStream, int8PtrTy), fileSize, fileIdx});
415
416            if (pipelineParallel){
417                generatePipelineParallel(iBuilder, {&s2pk, &icgrepK, &scanMatchK}, {s2pInstance, icgrepInstance, scanMatchInstance});
418            }
419            else{
420                generatePipelineLoop(iBuilder, {&s2pk, &icgrepK, &scanMatchK}, {s2pInstance, icgrepInstance, scanMatchInstance}, fileSize);
421            }
422        }
423
424        iBuilder->CreateRetVoid();
425
426    }
427
428#ifdef CUDA_ENABLED
429    Function * mainCPUFn = nullptr;
430    if(codegen::NVPTX){
431        Function * kernelFunction = generateGPUKernel(M, iBuilder, CountOnly);
432        MDNode * Node = MDNode::get(M->getContext(),
433                                    {llvm::ValueAsMetadata::get(kernelFunction),
434                                     MDString::get(M->getContext(), "kernel"), 
435                                     ConstantAsMetadata::get(ConstantInt::get(iBuilder->getInt32Ty(), 1))});
436        NamedMDNode *NMD = M->getOrInsertNamedMetadata("nvvm.annotations");
437        NMD->addOperand(Node);
438   
439        Compile2PTX(M, IRFilename, PTXFilename);
440        mainCPUFn = generateCPUKernel(cpuM, CPUBuilder, mIsNameExpression);
441        if (CountOnly) return;
442    }
443#endif
444   
445     
446    mEngine = JIT_to_ExecutionEngine(cpuM);
447    ApplyObjectCache(mEngine);
448    icgrep_Linking(cpuM, mEngine);
449
450#ifndef NDEBUG
451    verifyModule(*M, &dbgs());
452#endif
453
454    mEngine->finalizeObject();
455    delete iBuilder;
456   
457    if (CountOnly) {
458        mGrepFunction_CountOnly = reinterpret_cast<GrepFunctionType_CountOnly>(mEngine->getPointerToFunction(mainFn));
459    } else {
460#ifdef CUDA_ENABLED
461        if(codegen::NVPTX){
462            mGrepFunction_CPU = reinterpret_cast<GrepFunctionType_CPU>(mEngine->getPointerToFunction(mainCPUFn));
463        }
464#endif
465        if (CPU_Only) {
466            mGrepFunction = reinterpret_cast<GrepFunctionType>(mEngine->getPointerToFunction(mainFn));
467        }
468    }
469
470}
471
472re::CC *  GrepEngine::grepCodepoints() {
473
474    setParsedCodePointSet();
475    char * mFileBuffer = getUnicodeNameDataPtr();
476    size_t mFileSize = getUnicodeNameDataSize();
477
478    mGrepFunction(mFileBuffer, mFileSize, 0);
479
480    return getParsedCodePointSet();
481}
482
483GrepEngine::~GrepEngine() {
484    delete mEngine;
485}
486
487
488static int * total_count;
489static std::stringstream * resultStrs = nullptr;
490static std::vector<std::string> inputFiles;
491
492void initResult(std::vector<std::string> filenames){
493    const int n = filenames.size();
494    if (n > 1) {
495        ShowFileNames = true;
496    }
497    inputFiles = filenames;
498    resultStrs = new std::stringstream[n];
499    total_count = new int[n];
500    for (unsigned i = 0; i < inputFiles.size(); ++i){
501        total_count[i] = 0;
502    }
503   
504}
505
506extern "C" {
507    void wrapped_report_match(size_t lineNum, size_t line_start, size_t line_end, const char * buffer, size_t filesize, int fileIdx) {
508        int index = isUTF_16 ? 2 : 1;
509        int idx = fileIdx;
510         
511        if (ShowFileNames) {
512            resultStrs[idx] << inputFiles[idx] << ':';
513        }
514        if (ShowLineNumbers) {
515            resultStrs[idx] << lineNum << ":";
516        }
517       
518        if ((!isUTF_16 && buffer[line_start] == 0xA) && (line_start != line_end)) {
519            // The line "starts" on the LF of a CRLF.  Really the end of the last line.
520            line_start++;
521        }
522        if (((isUTF_16 && buffer[line_start] == 0x0) && buffer[line_start + 1] == 0xA) && (line_start != line_end)) {
523            // The line "starts" on the LF of a CRLF.  Really the end of the last line.
524            line_start += 2;
525        }
526        if (line_end == filesize) {
527            // The match position is at end-of-file.   We have a final unterminated line.
528            resultStrs[idx].write(&buffer[line_start * index], (line_end - line_start) * index);
529            if (NormalizeLineBreaks) {
530                resultStrs[idx] << '\n';  // terminate it
531            }
532            return;
533        }
534        unsigned char end_byte = (unsigned char)buffer[line_end]; 
535        unsigned char penult_byte = (unsigned char)(buffer[line_end - 1]);
536        if (NormalizeLineBreaks) {
537            if (end_byte == 0x85) {
538                // Line terminated with NEL, on the second byte.  Back up 1.
539                line_end--;
540            } else if (end_byte > 0xD) {
541                // Line terminated with PS or LS, on the third byte.  Back up 2.
542                isUTF_16 ? line_end-- : line_end -= 2;
543            }
544            resultStrs[idx].write(&buffer[line_start * index], (line_end - line_start) * index);
545            resultStrs[idx] << '\n';
546        }
547        else {   
548            if ((!isUTF_16 && end_byte == 0x0D) || (isUTF_16 && (end_byte == 0x0D && penult_byte == 0x0))) {
549                // Check for line_end on first byte of CRLF;  note that we don't
550                // want to access past the end of buffer.
551                if (line_end + 1 < filesize) {
552                    if (!isUTF_16 && buffer[line_end + 1] == 0x0A) {
553                        // Found CRLF; preserve both bytes.
554                        line_end++;
555                    }
556                    if (isUTF_16 && buffer[line_end + 1] == 0x0 && buffer[line_end + 2] == 0x0A) {
557                        // Found CRLF; preserve both bytes.
558                        line_end += 2;
559                    }
560                }
561            }
562            resultStrs[idx].write(&buffer[line_start * index], (line_end - line_start + 1) * index);
563        }
564    }
565}
566
567void PrintResult(bool CountOnly, std::vector<size_t> & total_CountOnly){
568    if(CountOnly){
569        if (!ShowFileNames) {
570            for (unsigned i = 0; i < inputFiles.size(); ++i){
571                std::cout << total_CountOnly[i] << std::endl;
572            }
573        }
574        else {
575            for (unsigned i = 0; i < inputFiles.size(); ++i){
576                std::cout << inputFiles[i] << ':' << total_CountOnly[i] << std::endl;
577            };
578        }
579        return;
580    }
581   
582    for (unsigned i = 0; i < inputFiles.size(); ++i){
583        std::cout << resultStrs[i].str();
584    }
585}
586
587re::CC * parsedCodePointSet;
588
589extern "C" {
590    void insert_codepoints(size_t lineNum, size_t line_start, size_t line_end, const char * buffer) {
591        re::codepoint_t c = 0;
592        ssize_t line_pos = line_start;
593        while (isxdigit(buffer[line_pos])) {
594            if (isdigit(buffer[line_pos])) {
595                c = (c << 4) | (buffer[line_pos] - '0');
596            }
597            else {
598                c = (c << 4) | (tolower(buffer[line_pos]) - 'a' + 10);
599            }
600            line_pos++;
601        }
602        assert(((line_pos - line_start) >= 4) && ((line_pos - line_start) <= 6)); // UCD format 4 to 6 hex digits.       
603        parsedCodePointSet->insert(c);
604    }
605}
606
607void setParsedCodePointSet(){
608    parsedCodePointSet = re::makeCC();
609}
610
611re::CC * getParsedCodePointSet(){
612    return parsedCodePointSet;
613}
614
615
616void icgrep_Linking(Module * m, ExecutionEngine * e) {
617    Module::FunctionListType & fns = m->getFunctionList();
618    for (Module::FunctionListType::iterator it = fns.begin(), it_end = fns.end(); it != it_end; ++it) {
619        std::string fnName = it->getName().str();
620        if (fnName == "s2p_block") continue;
621        if (fnName == "process_block") continue;
622        if (fnName == "process_block_initialize_carries") continue;
623       
624        if (fnName == "wrapped_report_match") {
625            e->addGlobalMapping(cast<GlobalValue>(it), (void *)&wrapped_report_match);
626        }
627        if (fnName == "insert_codepoints") {
628            e->addGlobalMapping(cast<GlobalValue>(it), (void *)&insert_codepoints);
629        }
630#ifndef DISABLE_PREGENERATED_UCD_FUNCTIONS
631        else {
632            const UCD::ExternalProperty & ep = UCD::resolveExternalProperty(fnName);
633            e->addGlobalMapping(cast<GlobalValue>(it), std::get<0>(ep));
634        }
635#endif
636    }
637}
638
Note: See TracBrowser for help on using the repository browser.