source: icGREP/icgrep-devel/icgrep/kernels/symboltablepipeline.cpp @ 5000

Last change on this file since 5000 was 5000, checked in by nmedfort, 3 years ago

Redesigned buffer system to allow the pipeline to control selection of the current input and output streams; DoBlock? functions containing lookahead now take multiple input stream arguments. Selection and passing occurs automatically. Some work on Symbol Table.

File size: 41.4 KB
Line 
1#include "symboltablepipeline.h"
2
3/*
4 *  Copyright (c) 2016 International Characters.
5 *  This software is licensed to the public under the Open Software License 3.0.
6 */
7
8#include "pipeline.h"
9#include "toolchain.h"
10#include "utf_encoding.h"
11
12#include <kernels/s2p_kernel.h>
13#include <kernels/instance.h>
14#include <kernels/stdout_kernel.h>
15
16#include <pablo/function.h>
17#include <pablo/pablo_compiler.h>
18#include <pablo/analysis/pabloverifier.hpp>
19
20#include <re/re_cc.h>
21#include <re/re_rep.h>
22#include <re/re_name.h>
23#include <re/re_compiler.h>
24#include <re/printer_re.h>
25
26#include <cc/cc_compiler.h>
27
28#include <pablo/printer_pablos.h>
29#include <iostream>
30
31#include <llvm/IR/Intrinsics.h>
32
33using namespace re;
34using namespace pablo;
35
36namespace kernel {
37
38SymbolTableBuilder::SymbolTableBuilder(Module * m, IDISA::IDISA_Builder * b)
39: mMod(m)
40, iBuilder(b)
41, mLongestLookahead(0)
42, mBitBlockType(b->getBitBlockType())
43, mBlockSize(b->getBitBlockWidth()) {
44
45}
46
47/** ------------------------------------------------------------------------------------------------------------- *
48 * @brief generateLeadingFunction
49 ** ------------------------------------------------------------------------------------------------------------- */
50PabloFunction * SymbolTableBuilder::generateLeadingFunction(const std::vector<unsigned> & endpoints) {
51    PabloFunction * const function = PabloFunction::Create("leading", 8, endpoints.size() + 2);
52    Encoding enc(Encoding::Type::ASCII, 8);
53    cc::CC_Compiler ccCompiler(*function, enc);
54    re::RE_Compiler reCompiler(*function, ccCompiler);
55    RE * cc = makeName(makeCC(makeCC(makeCC('a', 'z'), makeCC('A', 'Z')), makeCC('0', '9')));
56    reCompiler.compileUnicodeNames(cc);
57    PabloAST * const matches = reCompiler.compile(cc).stream;
58    PabloBlock * const entry = function->getEntryBlock();
59    PabloAST * const adv = entry->createAdvance(matches, 1);
60    PabloAST * const starts = entry->createAnd(matches, entry->createNot(adv));
61    PabloAST * const ends = entry->createAnd(adv, entry->createNot(matches));
62
63    function->setResult(0, entry->createAssign("l.S", starts));
64    function->setResult(1, entry->createAssign("l.E", ends));
65
66    PabloAST * M = ends;
67    unsigned step = 1;
68    unsigned i = 0;
69    for (unsigned endpoint : endpoints) {
70        assert (endpoint >= step);
71        unsigned span = endpoint - step;
72        while (span > step) {
73            M = entry->createOr(entry->createAdvance(M, step), M);
74            span = span - step;
75            step *= 2;
76        }
77        M = entry->createOr(entry->createAdvance(M, span), M);
78        function->setResult(i + 2, entry->createAssign("l.M" + std::to_string(i), M));
79        ++i;
80        step += span;
81    }
82
83    return function;
84}
85
86/** ------------------------------------------------------------------------------------------------------------- *
87 * @brief generateSortingFunction
88 ** ------------------------------------------------------------------------------------------------------------- */
89PabloFunction * SymbolTableBuilder::generateSortingFunction(const PabloFunction * const leading, const std::vector<unsigned> & endpoints) {
90    PabloFunction * const function = PabloFunction::Create("sorting", leading->getNumOfResults(), (leading->getNumOfResults() - 1) * 2);
91    PabloBlock * entry = function->getEntryBlock();
92    function->setParameter(0, entry->createVar("l.S"));
93    function->setParameter(1, entry->createVar("l.E"));
94    for (unsigned i = 2; i < leading->getNumOfResults(); ++i) {
95        function->setParameter(i, entry->createVar("l.M" + std::to_string(i - 2)));
96    }
97    PabloAST * R = function->getParameter(0);
98    PabloAST * const E = entry->createNot(function->getParameter(1));
99    unsigned i = 0;
100    unsigned lowerbound = 0;
101    for (unsigned endpoint : endpoints) {
102        PabloAST * const M = function->getParameter(i + 2);
103        PabloAST * const L = entry->createLookahead(M, endpoint, "lookahead" + std::to_string(endpoint));
104        PabloAST * S = entry->createAnd(L, R);
105        Assign * Si = entry->createAssign("s.S_" + std::to_string(i + 1), S);
106        PabloAST * F = entry->createScanThru(S, E);
107        Assign * Ei = entry->createAssign("s.E_" + std::to_string(i + 1), F);
108        function->setResult(i * 2, Si);
109        function->setResult(i * 2 + 1, Ei);
110        R = entry->createXor(R, S);
111        ++i;
112        lowerbound = endpoint;
113    }
114    Assign * Si = entry->createAssign("s.S_n", R);
115    PabloAST * F = entry->createScanThru(R, E);
116    Assign * Ei = entry->createAssign("s.E_n", F);
117    function->setResult(i * 2, Si);
118    function->setResult(i * 2 + 1, Ei);
119    mLongestLookahead = lowerbound;
120
121    return function;
122}
123
124/** ------------------------------------------------------------------------------------------------------------- *
125 * @brief generateCountForwardZeroes
126 ** ------------------------------------------------------------------------------------------------------------- */
127inline Value * generateCountForwardZeroes(IDISA::IDISA_Builder * iBuilder, Value * bits) {
128    Value * cttzFunc = Intrinsic::getDeclaration(iBuilder->getModule(), Intrinsic::cttz, bits->getType());
129    return iBuilder->CreateCall(cttzFunc, std::vector<Value *>({bits, ConstantInt::get(iBuilder->getInt1Ty(), 0)}));
130}
131
132/** ------------------------------------------------------------------------------------------------------------- *
133 * @brief generateMaskedGather
134 ** ------------------------------------------------------------------------------------------------------------- */
135inline Value * SymbolTableBuilder::generateMaskedGather(Value * const base, Value * const vindex, Value * const mask) {
136
137    /*
138        From Intel:
139
140        extern __m256i _mm256_mask_i32gather_epi32(__m256i def_vals, int const * base, __m256i vindex, __m256i vmask, const int scale);
141
142        From Clang avx2intrin.h:
143
144        #define _mm256_mask_i32gather_epi32(a, m, i, mask, s) __extension__ ({ \
145           (__m256i)__builtin_ia32_gatherd_d256((__v8si)(__m256i)(a), \
146                                                (int const *)(m), \
147                                                (__v8si)(__m256i)(i), \
148                                                (__v8si)(__m256i)(mask), (s)); })
149        From llvm IntrinsicsX86.td:
150
151        def llvm_ptr_ty        : LLVMPointerType<llvm_i8_ty>;             // i8*
152
153        def int_x86_avx2_gather_d_d_256 : GCCBuiltin<"__builtin_ia32_gatherd_d256">,
154           Intrinsic<[llvm_v8i32_ty],
155           [llvm_v8i32_ty, llvm_ptr_ty, llvm_v8i32_ty, llvm_v8i32_ty, llvm_i8_ty],
156           [IntrReadArgMem]>;
157
158     */
159
160    VectorType * const vecType = VectorType::get(iBuilder->getInt32Ty(), 8);
161    Function * const vgather = Intrinsic::getDeclaration(iBuilder->getModule(), Intrinsic::x86_avx2_gather_d_d_256);
162    return iBuilder->CreateCall(vgather, {Constant::getNullValue(vecType), base, iBuilder->CreateBitCast(vindex, vecType), iBuilder->CreateBitCast(mask, vecType), iBuilder->getInt8(1)});
163}
164
165/** ------------------------------------------------------------------------------------------------------------- *
166 * @brief generateResetLowestBit
167 ** ------------------------------------------------------------------------------------------------------------- */
168inline Value * generateResetLowestBit(IDISA::IDISA_Builder * iBuilder, Value * bits) {
169    Value * bits_minus1 = iBuilder->CreateSub(bits, ConstantInt::get(bits->getType(), 1));
170    return iBuilder->CreateAnd(bits_minus1, bits);
171}
172
173/** ------------------------------------------------------------------------------------------------------------- *
174 * @brief generateGatherKernel
175 ** ------------------------------------------------------------------------------------------------------------- */
176void SymbolTableBuilder::generateGatherKernel(KernelBuilder * kBuilder, const std::vector<unsigned> & endpoints, const unsigned scanWordBitWidth) {
177
178    Type * const intScanWordTy = iBuilder->getIntNTy(scanWordBitWidth);
179    const unsigned fieldCount = iBuilder->getBitBlockWidth() / scanWordBitWidth;
180    Type * const scanWordVectorType = VectorType::get(intScanWordTy, fieldCount);
181    const unsigned vectorWidth = iBuilder->getBitBlockWidth() / 32;
182    const unsigned gatherCount = vectorWidth * 4;
183    Type * const transposedVectorType = VectorType::get(iBuilder->getInt8Ty(), iBuilder->getBitBlockWidth() / 8);
184
185    Type * startArrayType = ArrayType::get(iBuilder->getInt32Ty(), iBuilder->getBitBlockWidth() + gatherCount);
186    Type * endArrayType = ArrayType::get(iBuilder->getInt32Ty(), gatherCount);
187    Type * groupType = StructType::get(iBuilder->getInt32Ty(), startArrayType, iBuilder->getInt32Ty(), endArrayType, nullptr);
188    const unsigned baseIdx = kBuilder->addInternalState(iBuilder->getInt8PtrTy(), "Base");
189    const unsigned gatherPositionArrayIdx = kBuilder->addInternalState(ArrayType::get(groupType, endpoints.size()), "Positions");
190
191    for (unsigned maxKeyLength : endpoints) {
192        kBuilder->addInputStream(1, "startStream" + std::to_string(maxKeyLength));
193        kBuilder->addInputStream(1, "endStream" + std::to_string(maxKeyLength));
194        kBuilder->addOutputStream(4); // ((maxKeyLength + 3) / 4) * 4
195    }
196    kBuilder->addInputStream(1, "startStreamN");
197    kBuilder->addInputStream(1, "endStreamN");
198
199    Function * const function = kBuilder->prepareFunction();
200
201    BasicBlock * const entry = iBuilder->GetInsertBlock();
202
203    BasicBlock * groupCond = BasicBlock::Create(mMod->getContext(), "groupCond", function, 0);
204    BasicBlock * groupBody = BasicBlock::Create(mMod->getContext(), "groupBody", function, 0);
205
206    BasicBlock * startOuterCond = BasicBlock::Create(mMod->getContext(), "startOuterCond", function, 0);
207    BasicBlock * startOuterBody = BasicBlock::Create(mMod->getContext(), "startOuterBody", function, 0);
208    BasicBlock * startInnerCond = BasicBlock::Create(mMod->getContext(), "startInnerCond", function, 0);
209    BasicBlock * startInnerBody = BasicBlock::Create(mMod->getContext(), "startInnerBody", function, 0);
210
211    BasicBlock * endOuterCond = BasicBlock::Create(mMod->getContext(), "endOuterCond", function, 0);
212    BasicBlock * endOuterBody = BasicBlock::Create(mMod->getContext(), "endOuterBody", function, 0);
213    BasicBlock * endInnerCond = BasicBlock::Create(mMod->getContext(), "endInnerCond", function, 0);
214    BasicBlock * endInnerBody = BasicBlock::Create(mMod->getContext(), "endInnerBody", function, 0);
215
216    BasicBlock * gather = BasicBlock::Create(mMod->getContext(), "gather", function, 0);
217
218    BasicBlock * nextGroup = BasicBlock::Create(mMod->getContext(), "nextGroup", function, 0);
219
220    BasicBlock * exit = BasicBlock::Create(mMod->getContext(), "exit", function, 0);
221
222
223    // ENTRY BLOCK
224    iBuilder->SetInsertPoint(entry);
225    Type * const int32PtrTy = PointerType::get(iBuilder->getInt32Ty(), 0);
226    FunctionType * const gatherFunctionType = FunctionType::get(iBuilder->getVoidTy(), {iBuilder->getInt8PtrTy(), int32PtrTy, int32PtrTy, iBuilder->getInt32Ty(), iBuilder->getInt8PtrTy()}, false);
227    Value * const gatherFunctionPtrArray = iBuilder->CreateAlloca(PointerType::get(gatherFunctionType, 0), iBuilder->getInt32(endpoints.size()));
228
229    unsigned i = 0;
230    unsigned minKeyLength = 0;
231    for (unsigned maxKeyLength : endpoints) {
232        Function * f = generateGatherFunction(minKeyLength, maxKeyLength, transposedVectorType);
233        mGatherFunction.push_back(f);
234        iBuilder->CreateStore(f, iBuilder->CreateGEP(gatherFunctionPtrArray, iBuilder->getInt32(i++)));
235        minKeyLength = maxKeyLength;
236    }
237
238    //TODO: this won't work on files > 2^32 bytes yet; needs an intermediate flush then a recalculation of the base pointer.
239    Value * const base = iBuilder->CreateLoad(kBuilder->getInternalState(baseIdx), "base");
240    Value * const positionArray = kBuilder->getInternalState(gatherPositionArrayIdx);
241
242    Value * blockPos = iBuilder->CreateLoad(kBuilder->getBlockNo());
243    blockPos = iBuilder->CreateMul(blockPos, iBuilder->getInt64(iBuilder->getBitBlockWidth()));
244
245    iBuilder->CreateBr(groupCond);
246
247    // GROUP COND
248    iBuilder->SetInsertPoint(groupCond);
249    PHINode * groupIV = iBuilder->CreatePHI(iBuilder->getInt32Ty(), 2);
250    groupIV->addIncoming(iBuilder->getInt32(0), entry);
251    Value * groupTest = iBuilder->CreateICmpNE(groupIV, iBuilder->getInt32(endpoints.size()));
252    iBuilder->CreateCondBr(groupTest, groupBody, exit);
253
254    // GROUP BODY
255    iBuilder->SetInsertPoint(groupBody);
256    // if two positions cannot be in the same vector element, we could possibly do some work in parallel here.
257
258    Value * index = iBuilder->CreateMul(groupIV, iBuilder->getInt32(2));
259    Value * startStreamPtr = kBuilder->getInputStream(index);
260    Value * startStream = iBuilder->CreateBlockAlignedLoad(startStreamPtr);
261    startStream = iBuilder->CreateBitCast(startStream, scanWordVectorType, "startStream");
262
263    index = iBuilder->CreateAdd(index, iBuilder->getInt32(1));
264    Value * endStreamPtr = kBuilder->getInputStream(index);
265    Value * endStream = iBuilder->CreateBlockAlignedLoad(endStreamPtr);
266    endStream = iBuilder->CreateBitCast(endStream, scanWordVectorType, "endStream");
267
268    Value * startIndexPtr = iBuilder->CreateGEP(positionArray, {iBuilder->getInt32(0), groupIV, iBuilder->getInt32(0)}, "startIndexPtr");
269    Value * startIndex = iBuilder->CreateLoad(startIndexPtr, "startIndex");
270    Value * startArray = iBuilder->CreateGEP(positionArray, {iBuilder->getInt32(0), groupIV, iBuilder->getInt32(1)}, "startArray");
271    Value * endIndexPtr = iBuilder->CreateGEP(positionArray, {iBuilder->getInt32(0), groupIV, iBuilder->getInt32(2)}, "endIndexPtr");
272    Value * endIndex = iBuilder->CreateLoad(endIndexPtr, "endIndex");
273    Value * endArray = iBuilder->CreateGEP(positionArray, {iBuilder->getInt32(0), groupIV, iBuilder->getInt32(3)}, "endArray");
274
275    iBuilder->CreateBr(startOuterCond);
276
277    // START OUTER COND
278    iBuilder->SetInsertPoint(startOuterCond);
279    PHINode * startBlockOffset = iBuilder->CreatePHI(iBuilder->getInt64Ty(), 2);
280    startBlockOffset->addIncoming(blockPos, groupBody);
281    PHINode * startIndexPhi1 = iBuilder->CreatePHI(startIndex->getType(), 2);
282    startIndexPhi1->addIncoming(startIndex, groupBody);
283    PHINode * startIV = iBuilder->CreatePHI(iBuilder->getInt64Ty(), 2);
284    startIV->addIncoming(iBuilder->getInt64(0), groupBody);
285    Value * startOuterTest = iBuilder->CreateICmpNE(startIV, iBuilder->getInt64(fieldCount));
286    iBuilder->CreateCondBr(startOuterTest, startOuterBody, endOuterCond);
287
288    // START OUTER BODY
289    iBuilder->SetInsertPoint(startOuterBody);
290    Value * startField = iBuilder->CreateExtractElement(startStream, startIV);
291    startIV->addIncoming(iBuilder->CreateAdd(startIV, iBuilder->getInt64(1)), startInnerCond);
292    startBlockOffset->addIncoming(iBuilder->CreateAdd(startBlockOffset, iBuilder->getInt64(scanWordBitWidth)), startInnerCond);
293    iBuilder->CreateBr(startInnerCond);
294
295    // START INNER COND
296    iBuilder->SetInsertPoint(startInnerCond);
297    PHINode * startIndexPhi3 = iBuilder->CreatePHI(startIndex->getType(), 2);
298    startIndexPhi3->addIncoming(startIndexPhi1, startOuterBody);
299    startIndexPhi1->addIncoming(startIndexPhi3, startInnerCond);
300    PHINode * startFieldPhi = iBuilder->CreatePHI(intScanWordTy, 2);
301    startFieldPhi->addIncoming(startField, startOuterBody);
302    Value * test = iBuilder->CreateICmpNE(startFieldPhi, ConstantInt::getNullValue(intScanWordTy));
303    iBuilder->CreateCondBr(test, startInnerBody, startOuterCond);
304
305    // START INNER BODY
306    iBuilder->SetInsertPoint(startInnerBody);
307    Value * startPos = generateCountForwardZeroes(iBuilder, startFieldPhi);
308    startFieldPhi->addIncoming(generateResetLowestBit(iBuilder, startFieldPhi), startInnerBody);
309    startPos = iBuilder->CreateTruncOrBitCast(iBuilder->CreateOr(startPos, startBlockOffset), iBuilder->getInt32Ty());
310    iBuilder->CreateStore(startPos, iBuilder->CreateGEP(startArray, {iBuilder->getInt32(0), startIndexPhi3}));
311    startIndexPhi3->addIncoming(iBuilder->CreateAdd(startIndexPhi3, ConstantInt::get(startIndexPhi3->getType(), 1)), startInnerBody);
312    iBuilder->CreateBr(startInnerCond);
313
314    // END POINT OUTER COND
315    iBuilder->SetInsertPoint(endOuterCond);
316    PHINode * endBlockOffset = iBuilder->CreatePHI(iBuilder->getInt64Ty(), 2);
317    endBlockOffset->addIncoming(blockPos, startOuterCond);
318    PHINode * endIndexPhi1 = iBuilder->CreatePHI(endIndex->getType(), 2);
319    endIndexPhi1->addIncoming(endIndex, startOuterCond);
320    PHINode * startIndexPhi2 = iBuilder->CreatePHI(startIndex->getType(), 2);
321    startIndexPhi2->addIncoming(startIndexPhi1, startOuterCond);
322    PHINode * endIV = iBuilder->CreatePHI(iBuilder->getInt64Ty(), 2);
323    endIV->addIncoming(iBuilder->getInt64(0), startOuterCond);
324    Value * endOuterTest = iBuilder->CreateICmpNE(endIV, iBuilder->getInt64(fieldCount));
325    iBuilder->CreateCondBr(endOuterTest, endOuterBody, nextGroup);
326
327    // END POINT OUTER BODY
328    iBuilder->SetInsertPoint(endOuterBody);
329    Value * endField = iBuilder->CreateExtractElement(endStream, endIV);
330    endIV->addIncoming(iBuilder->CreateAdd(endIV, iBuilder->getInt64(1)), endInnerCond);
331    endBlockOffset->addIncoming(iBuilder->CreateAdd(endBlockOffset, iBuilder->getInt64(scanWordBitWidth)), endInnerCond);
332    iBuilder->CreateBr(endInnerCond);
333
334    // END POINT INNER COND
335    iBuilder->SetInsertPoint(endInnerCond);
336    startIndexPhi3 = iBuilder->CreatePHI(startIndexPhi2->getType(), 3);
337    startIndexPhi3->addIncoming(startIndexPhi2, endOuterBody);
338    startIndexPhi3->addIncoming(startIndexPhi3, endInnerBody);
339    startIndexPhi2->addIncoming(startIndexPhi3, endInnerCond);
340    PHINode * endIndexPhi2 = iBuilder->CreatePHI(endIndex->getType(), 3);
341    endIndexPhi2->addIncoming(endIndexPhi1, endOuterBody);
342    endIndexPhi1->addIncoming(endIndexPhi2, endInnerCond);
343    endIndexPhi2->addIncoming(ConstantInt::getNullValue(endIndex->getType()), gather);
344    PHINode * endFieldPhi = iBuilder->CreatePHI(intScanWordTy, 3);
345    endFieldPhi->addIncoming(endField, endOuterBody);
346    Value * endInnerTest = iBuilder->CreateICmpNE(endFieldPhi, ConstantInt::getNullValue(intScanWordTy));
347    iBuilder->CreateCondBr(endInnerTest, endInnerBody, endOuterCond);
348
349    // END POINT INNER BODY
350    iBuilder->SetInsertPoint(endInnerBody);
351    Value * endPos = generateCountForwardZeroes(iBuilder, endFieldPhi);
352    Value * updatedEndFieldPhi = generateResetLowestBit(iBuilder, endFieldPhi);
353    endFieldPhi->addIncoming(updatedEndFieldPhi, endInnerBody);
354    endFieldPhi->addIncoming(updatedEndFieldPhi, gather);
355    endPos = iBuilder->CreateTruncOrBitCast(iBuilder->CreateOr(endPos, endBlockOffset), iBuilder->getInt32Ty());
356    iBuilder->CreateStore(endPos, iBuilder->CreateGEP(endArray, {iBuilder->getInt32(0), endIndexPhi2}));
357    Value * updatedEndIndexPhi = iBuilder->CreateAdd(endIndexPhi2, ConstantInt::get(endIndexPhi2->getType(), 1));
358    endIndexPhi2->addIncoming(updatedEndIndexPhi, endInnerBody);
359    Value * filledEndPosBufferTest = iBuilder->CreateICmpEQ(updatedEndIndexPhi, ConstantInt::get(updatedEndIndexPhi->getType(), gatherCount));
360    iBuilder->CreateCondBr(filledEndPosBufferTest, gather, endInnerCond);
361
362    // GATHER
363    iBuilder->SetInsertPoint(gather);
364
365    Value * startArrayPtr = iBuilder->CreatePointerCast(startArray, PointerType::get(iBuilder->getInt32Ty(), 0));
366    Value * endArrayPtr = iBuilder->CreatePointerCast(endArray, PointerType::get(iBuilder->getInt32Ty(), 0));
367    Value * gatherFunctionPtr = iBuilder->CreateLoad(iBuilder->CreateGEP(gatherFunctionPtrArray, groupIV));
368    Value * outputBuffer = iBuilder->CreatePointerCast(kBuilder->getOutputStream(groupIV), iBuilder->getInt8PtrTy());
369    iBuilder->CreateCall5(gatherFunctionPtr, base, startArrayPtr, endArrayPtr, iBuilder->getInt32(32), outputBuffer);
370
371    Value * remainingArrayPtr = iBuilder->CreateGEP(startArrayPtr, iBuilder->getInt32(gatherCount));
372    Value * remainingCount = iBuilder->CreateSub(startIndexPhi3, iBuilder->getInt32(gatherCount));
373    iBuilder->CreateMemCpy(startArrayPtr, remainingArrayPtr, remainingCount, 4);
374    startIndexPhi3->addIncoming(remainingCount, gather);
375    iBuilder->CreateBr(endInnerCond);
376
377    // NEXT GROUP
378    iBuilder->SetInsertPoint(nextGroup);
379    iBuilder->CreateStore(startIndexPhi2, startIndexPtr);
380    iBuilder->CreateStore(endIndexPhi1, endIndexPtr);
381    groupIV->addIncoming(iBuilder->CreateAdd(groupIV, ConstantInt::get(groupIV->getType(), 1)), nextGroup);
382    iBuilder->CreateBr(groupCond);
383
384    iBuilder->SetInsertPoint(exit);
385    kBuilder->finalize();
386}
387
388/** ------------------------------------------------------------------------------------------------------------- *
389 * @brief generateGatherFunction
390 ** ------------------------------------------------------------------------------------------------------------- */
391Function * SymbolTableBuilder::generateGatherFunction(const unsigned minKeyLength, const unsigned maxKeyLength, Type * const resultType) {
392
393    assert (minKeyLength < maxKeyLength);
394
395    const std::string functionName = "gather_" + std::to_string(minKeyLength) + "_to_" + std::to_string(maxKeyLength);
396    Function * function = mMod->getFunction(functionName);
397    if (function == nullptr) {
398
399        const auto ip = iBuilder->saveIP();
400
401        const unsigned minCount = (minKeyLength / 4);
402        const unsigned maxCount = ((maxKeyLength + 3) / 4);
403
404        const unsigned vectorWidth = iBuilder->getBitBlockWidth() / 32;
405        Type * const gatherVectorType =  VectorType::get(iBuilder->getInt32Ty(), vectorWidth);
406        Type * const gatherVectorArrayType = ArrayType::get(gatherVectorType, maxCount);
407
408        Type * const int32PtrTy = PointerType::get(iBuilder->getInt32Ty(), 0);
409        FunctionType * const functionType = FunctionType::get(iBuilder->getVoidTy(), {iBuilder->getInt8PtrTy(), int32PtrTy, int32PtrTy, iBuilder->getInt32Ty(), iBuilder->getInt8PtrTy()}, false);
410        function = Function::Create(functionType, GlobalValue::ExternalLinkage, functionName, mMod);
411        function->setCallingConv(CallingConv::C);
412        function->setDoesNotCapture(1);
413        function->setDoesNotCapture(2);
414        function->setDoesNotCapture(3);
415        function->setDoesNotThrow();
416
417        Function::arg_iterator args = function->arg_begin();
418        Value * const base = args++;
419        base->setName("base");
420        Value * startArray = args++;
421        startArray->setName("startArray");
422        Value * endArray = args++;
423        endArray->setName("endArray");
424        Value * const numOfKeys = args++;
425        numOfKeys->setName("numOfKeys");
426        Value * result = args++;
427        result->setName("result");
428
429        BasicBlock * entry = BasicBlock::Create(mMod->getContext(), "entry", function, 0);
430        BasicBlock * gatherCond = BasicBlock::Create(mMod->getContext(), "gatherCond", function, 0);
431        BasicBlock * partialGatherCond = BasicBlock::Create(mMod->getContext(), "partialGatherCond", function, 0);
432        BasicBlock * partialGatherBody = BasicBlock::Create(mMod->getContext(), "partialGatherBody", function, 0);
433        BasicBlock * gatherBody = BasicBlock::Create(mMod->getContext(), "gatherBody", function, 0);
434        BasicBlock * transposeCond = BasicBlock::Create(mMod->getContext(), "transposeCond", function, 0);
435        BasicBlock * transposeBody = BasicBlock::Create(mMod->getContext(), "transposeBody", function, 0);
436        BasicBlock * exit = BasicBlock::Create(mMod->getContext(), "exit", function, 0);
437
438        Value * const four = iBuilder->CreateVectorSplat(vectorWidth, iBuilder->getInt32(4));
439
440        // ENTRY
441        iBuilder->SetInsertPoint(entry);
442        AllocaInst * const buffer = iBuilder->CreateAlloca(resultType, iBuilder->getInt32(maxCount * 4), "buffer");
443        iBuilder->CreateStore(Constant::getNullValue(buffer->getAllocatedType()), buffer);
444        AllocaInst * const untransposedBuffer = iBuilder->CreateAlloca(gatherVectorArrayType, iBuilder->getInt32(4), "tmp");
445        iBuilder->CreateStore(Constant::getNullValue(untransposedBuffer->getAllocatedType()), untransposedBuffer);
446        iBuilder->CreateBr(gatherCond);
447
448        // FULL GATHER COND
449        iBuilder->SetInsertPoint(gatherCond);
450        PHINode * remainingLanes = iBuilder->CreatePHI(iBuilder->getInt32Ty(), 2);
451        remainingLanes->addIncoming(numOfKeys, entry);
452        PHINode * gatherIV = iBuilder->CreatePHI(iBuilder->getInt32Ty(), 2);
453        gatherIV->addIncoming(iBuilder->getInt32(0), entry);
454        Value * gatherLoopTest = iBuilder->CreateICmpNE(gatherIV, iBuilder->getInt32(4));
455        iBuilder->CreateCondBr(gatherLoopTest, partialGatherCond, transposeCond);
456
457        // PARTIAL GATHER COND
458        iBuilder->SetInsertPoint(partialGatherCond);
459        Value * partialGatherLoopTest = iBuilder->CreateICmpUGE(remainingLanes, iBuilder->getInt32(vectorWidth));
460        iBuilder->CreateCondBr(partialGatherLoopTest, gatherBody, partialGatherBody);
461
462        // PARTIAL GATHER BODY
463        iBuilder->SetInsertPoint(partialGatherBody);
464        iBuilder->CallPrintInt(functionName + ".remainingLanes", remainingLanes);
465        Type * registerType = iBuilder->getIntNTy(iBuilder->getBitBlockWidth());
466        Value * maskedLanes = iBuilder->CreateSub(iBuilder->getInt32(vectorWidth), remainingLanes);
467        maskedLanes = iBuilder->CreateMul(maskedLanes, iBuilder->getInt32(32));
468        maskedLanes = iBuilder->CreateZExt(maskedLanes, registerType);
469        maskedLanes = iBuilder->CreateLShr(Constant::getAllOnesValue(registerType), maskedLanes);
470        maskedLanes = iBuilder->CreateBitCast(maskedLanes, gatherVectorType);
471
472        iBuilder->CreateBr(gatherBody);
473
474        // FULL GATHER BODY
475        iBuilder->SetInsertPoint(gatherBody);
476        PHINode * activeLanes = iBuilder->CreatePHI(gatherVectorType, 2, "activeLanes");
477        activeLanes->addIncoming(Constant::getAllOnesValue(gatherVectorType), partialGatherCond);
478        activeLanes->addIncoming(maskedLanes, partialGatherBody);
479
480        iBuilder->CallPrintRegister(functionName + ".activeLanes", activeLanes);
481
482        startArray = iBuilder->CreateBitCast(startArray, PointerType::get(gatherVectorType, 0));
483        Value * startPos = iBuilder->CreateAlignedLoad(iBuilder->CreateGEP(startArray, gatherIV), 4);
484        for (unsigned blockCount = 0; blockCount < minCount; ++blockCount) {
485            Value * tokenData = generateMaskedGather(base, startPos, activeLanes);
486            startPos = iBuilder->CreateAdd(startPos, four);
487            iBuilder->CreateAlignedStore(tokenData, iBuilder->CreateGEP(untransposedBuffer, {iBuilder->getInt32(blockCount), gatherIV}), 4);
488        }
489
490        endArray = iBuilder->CreateBitCast(endArray, PointerType::get(gatherVectorType, 0));
491        Value * const endPos = iBuilder->CreateAlignedLoad(iBuilder->CreateGEP(endArray, gatherIV), 4);
492        for (unsigned blockCount = minCount; blockCount < maxCount; ++blockCount) {
493
494            // if we have not fully gathered the data for this key
495            Value * atLeastOneByte = iBuilder->CreateSExt(iBuilder->CreateICmpULT(startPos, endPos), startPos->getType());
496            atLeastOneByte = iBuilder->CreateAnd(atLeastOneByte, activeLanes);
497
498            // gather it ...
499            Value * tokenData = generateMaskedGather(base, startPos, atLeastOneByte);
500            // and compute how much data is remaining.
501            Value * remaining = iBuilder->CreateSub(endPos, startPos);
502
503            // if this token only has 1 to 3 bytes remaining ...
504            Value * atLeastFourBytes = iBuilder->CreateSExt(iBuilder->CreateICmpUGE(remaining, four), remaining->getType());
505
506            // determine how many bits do *not* belong to the token
507            remaining = iBuilder->CreateSub(four, remaining);
508            remaining = iBuilder->CreateShl(remaining, ConstantInt::get(remaining->getType(), 3));
509
510            // then mask them out prior to storing the value
511            Value * partialTokenMask = iBuilder->CreateLShr(ConstantInt::getAllOnesValue(remaining->getType()), remaining);
512            partialTokenMask = iBuilder->CreateOr(partialTokenMask, atLeastFourBytes);
513
514            tokenData = iBuilder->CreateAnd(partialTokenMask, tokenData);
515            Value * untransposedBufferPtr = iBuilder->CreateGEP(untransposedBuffer, {iBuilder->getInt32(blockCount), gatherIV});
516            iBuilder->CreateAlignedStore(tokenData, untransposedBufferPtr, 4);
517            if (blockCount < (maxCount - 1)) {
518                startPos = iBuilder->CreateAdd(startPos, four);
519            }
520        }
521        gatherIV->addIncoming(iBuilder->CreateAdd(gatherIV, iBuilder->getInt32(1)), gatherBody);
522        remainingLanes->addIncoming(iBuilder->CreateSub(remainingLanes, iBuilder->getInt32(vectorWidth)), gatherBody);
523        iBuilder->CreateBr(gatherCond);
524
525        // TRANSPOSE COND
526        iBuilder->SetInsertPoint(transposeCond);
527        PHINode * transposeIV = iBuilder->CreatePHI(iBuilder->getInt32Ty(), 2);
528        transposeIV->addIncoming(iBuilder->getInt32(0), gatherCond);
529        Value * transposeLoopTest = iBuilder->CreateICmpNE(transposeIV, iBuilder->getInt32(maxCount));
530        iBuilder->CreateCondBr(transposeLoopTest, transposeBody, exit);
531
532        // TRANSPOSE BODY
533        iBuilder->SetInsertPoint(transposeBody);
534
535        Value * value[4];
536        Value * temporary[4];
537        for (unsigned i = 0; i < 4; ++i) {
538            Value * const ptr = iBuilder->CreateGEP(untransposedBuffer, {transposeIV, iBuilder->getInt32(i)});
539            value[i] = iBuilder->CreateAlignedLoad(ptr, 4);
540        }
541        for (unsigned fieldWidth = 16; fieldWidth != 4; fieldWidth /= 2) {
542            const unsigned fieldCount = iBuilder->getBitBlockWidth() / fieldWidth;
543            VectorType * const vecType = VectorType::get(IntegerType::get(mMod->getContext(), fieldWidth), fieldCount);
544            std::vector<Constant *> lowFields(fieldCount);
545            std::vector<Constant *> highFields(fieldCount);
546            for (unsigned j = 0; j < fieldCount; ++j) {
547                lowFields[j] = iBuilder->getInt32(j * 2);
548                highFields[j] = iBuilder->getInt32(j * 2 + 1);
549            }
550            Constant * const lowVector = ConstantVector::get(lowFields);
551            Constant * const highVector = ConstantVector::get(highFields);
552            for (unsigned i = 0; i < 4; i += 2) {
553                value[i] = iBuilder->CreateBitCast(value[i], vecType);
554                value[i + 1] = iBuilder->CreateBitCast(value[i + 1], vecType);
555                temporary[i / 2] = iBuilder->CreateShuffleVector(value[i], value[i + 1], lowVector);
556                temporary[(i / 2) + 2] = iBuilder->CreateShuffleVector(value[i], value[i + 1], highVector);
557            }
558            std::swap(value, temporary);
559        }
560        Value * offset = iBuilder->CreateShl(transposeIV, ConstantInt::get(transposeIV->getType(), 2));
561        transposeIV->addIncoming(iBuilder->CreateAdd(transposeIV, iBuilder->getInt32(1)), transposeBody);
562
563        for (unsigned i = 0; i < 4; ++i) {
564            Value * index = offset;
565            if (i) {
566                index = iBuilder->CreateAdd(offset, iBuilder->getInt32(i));
567            }
568            iBuilder->CallPrintRegister(functionName, value[i]);
569            iBuilder->CreateAlignedStore(value[i], iBuilder->CreateGEP(buffer, index), 4);
570        }
571
572        Value * emptyGatherTest = iBuilder->CreateICmpUGT(remainingLanes, iBuilder->getInt32(0));
573        iBuilder->CreateCondBr(emptyGatherTest, transposeCond, exit);
574
575        // EXIT
576        iBuilder->SetInsertPoint(exit);
577
578        // ... call hashing function ...
579
580
581        iBuilder->CreateRetVoid();
582
583        iBuilder->restoreIP(ip);
584    }
585
586    return function;
587}
588
589
590/** ------------------------------------------------------------------------------------------------------------- *
591 * @brief createKernels
592 ** ------------------------------------------------------------------------------------------------------------- */
593void SymbolTableBuilder::createKernels() {
594
595    std::vector<unsigned> endpoints;
596    endpoints.push_back(1);
597    endpoints.push_back(2);
598    endpoints.push_back(4);
599    endpoints.push_back(8);
600    endpoints.push_back(16);
601
602    PabloCompiler pablo_compiler(mMod, iBuilder);
603    PabloFunction * const leading = generateLeadingFunction(endpoints);
604    PabloFunction * const sorting = generateSortingFunction(leading, endpoints);
605
606    const auto bufferSize = ((mLongestLookahead + iBuilder->getBitBlockWidth() - 1) / iBuilder->getBitBlockWidth()) + 1;
607
608    mS2PKernel = new KernelBuilder(iBuilder, "s2p", 1);
609    mLeadingKernel = new KernelBuilder(iBuilder, "leading", bufferSize);
610    mSortingKernel = new KernelBuilder(iBuilder, "sorting", bufferSize);
611    mGatherKernel = new KernelBuilder(iBuilder, "gathering", 1);
612    mStdOutKernel = new KernelBuilder(iBuilder, "stddout", 1);
613
614    generateS2PKernel(mMod, iBuilder, mS2PKernel);
615
616    pablo_compiler.setKernel(mLeadingKernel);
617    pablo_compiler.compile(leading);
618    pablo_compiler.setKernel(mSortingKernel);
619    pablo_compiler.compile(sorting);
620
621    delete leading;
622    delete sorting;
623
624    releaseSlabAllocatorMemory();
625
626    generateGatherKernel(mGatherKernel, endpoints, 64);
627    generateStdOutKernel(mMod, iBuilder, mStdOutKernel);
628}
629
630Function * SymbolTableBuilder::ExecuteKernels(){
631
632    Type * intType = iBuilder->getInt64Ty();
633
634    Type * inputType = PointerType::get(ArrayType::get(StructType::get(mMod->getContext(), std::vector<Type *>({ArrayType::get(mBitBlockType, 8)})), 1), 0);
635    Function * const main = cast<Function>(mMod->getOrInsertFunction("Main", Type::getVoidTy(mMod->getContext()), inputType, intType, nullptr));
636    main->setCallingConv(CallingConv::C);
637    Function::arg_iterator args = main->arg_begin();
638
639    Value * const inputStream = args++;
640    inputStream->setName("inputStream");
641
642    Value * const bufferSize = args++;
643    bufferSize->setName("bufferSize");
644
645    iBuilder->SetInsertPoint(BasicBlock::Create(mMod->getContext(), "entry", main,0));
646
647    BasicBlock * entryBlock = iBuilder->GetInsertBlock();
648
649    BasicBlock * leadingTestBlock = BasicBlock::Create(mMod->getContext(), "leadingCond", main, 0);
650    BasicBlock * safetyCheckBlock = BasicBlock::Create(mMod->getContext(), "safetyCheck", main, 0);
651    BasicBlock * leadingBodyBlock = BasicBlock::Create(mMod->getContext(), "leadingBody", main, 0);
652
653    BasicBlock * regularTestBlock = BasicBlock::Create(mMod->getContext(), "fullCond", main, 0);
654    BasicBlock * regularBodyBlock = BasicBlock::Create(mMod->getContext(), "fullBody", main, 0);
655    BasicBlock * regularExitBlock = BasicBlock::Create(mMod->getContext(), "fullExit", main, 0);
656
657    BasicBlock * partialBlock = BasicBlock::Create(mMod->getContext(),  "partialBlock", main, 0);
658
659    BasicBlock * finalTestBlock = BasicBlock::Create(mMod->getContext(),  "finalCond", main, 0);
660    BasicBlock * finalBodyBlock = BasicBlock::Create(mMod->getContext(),  "finalBody", main, 0);
661
662    BasicBlock * remainingBlock = BasicBlock::Create(mMod->getContext(), "remaining", main, 0);
663
664    Instance * s2pInstance = mS2PKernel->instantiate(inputStream);
665    Instance * leadingInstance = mLeadingKernel->instantiate(s2pInstance->getResultSet());
666    Instance * sortingInstance = mSortingKernel->instantiate(leadingInstance->getResultSet());
667    Instance * gatheringInstance = mGatherKernel->instantiate(sortingInstance->getResultSet());
668    Instance * stdOutInstance = mStdOutKernel->instantiate(gatheringInstance->getResultSet());
669
670    gatheringInstance->setInternalState("Base", iBuilder->CreateBitCast(inputStream, iBuilder->getInt8PtrTy()));
671
672    stdOutInstance->setInternalState("RemainingBytes", bufferSize);  // The total number of bytes to be sent to stdout.
673
674    const unsigned leadingBlocks = (mLongestLookahead + iBuilder->getBitBlockWidth() - 1) / iBuilder->getBitBlockWidth();
675
676    Value * const requiredBytes = iBuilder->getInt64(mBlockSize * leadingBlocks);
677    Value * const blockSize = iBuilder->getInt64(mBlockSize);
678
679    // If the buffer size is smaller than our largest length group, only check up to the buffer size.
680    Value * safetyCheck = iBuilder->CreateICmpUGE(bufferSize, blockSize);
681    if (blockSize == requiredBytes) {
682        iBuilder->CreateCondBr(safetyCheck, leadingTestBlock, remainingBlock); // fix this to be a special case
683    } else {
684        throw std::runtime_error("Not supported yet!");
685    }
686
687    // First compute any necessary leading blocks to allow the sorting kernel access to the "future" data produced by
688    // the leading kernel ...
689    iBuilder->SetInsertPoint(leadingTestBlock);
690    PHINode * blockNo = iBuilder->CreatePHI(intType, 2);
691    blockNo->addIncoming(iBuilder->getInt64(0), entryBlock);
692    PHINode * remainingBytes = iBuilder->CreatePHI(intType, 2);
693    remainingBytes->addIncoming(bufferSize, entryBlock);
694    Value * leadingBlocksCond = iBuilder->CreateICmpULT(blockNo, iBuilder->getInt64(leadingBlocks));
695    iBuilder->CreateCondBr(leadingBlocksCond, safetyCheckBlock, regularTestBlock);
696
697    iBuilder->SetInsertPoint(safetyCheckBlock);
698    Value * safetyCheckCond = iBuilder->CreateICmpULT(remainingBytes, blockSize);
699    iBuilder->CreateCondBr(safetyCheckCond, regularExitBlock, leadingBodyBlock);
700
701    iBuilder->SetInsertPoint(leadingBodyBlock);
702
703    s2pInstance->CreateDoBlockCall();
704    leadingInstance->CreateDoBlockCall();
705    blockNo->addIncoming(iBuilder->CreateAdd(blockNo, iBuilder->getInt64(1)), leadingBodyBlock);
706    remainingBytes->addIncoming(iBuilder->CreateSub(remainingBytes, blockSize), leadingBodyBlock);
707    iBuilder->CreateBr(leadingTestBlock);
708
709    // Now all the data for which we can produce and consume a full leading block...
710    iBuilder->SetInsertPoint(regularTestBlock);
711    PHINode * remainingBytes2 = iBuilder->CreatePHI(intType, 2);
712    remainingBytes2->addIncoming(remainingBytes, leadingTestBlock);
713    Value * remainingBytesCond = iBuilder->CreateICmpULT(remainingBytes2, requiredBytes);
714    iBuilder->CreateCondBr(remainingBytesCond, regularExitBlock, regularBodyBlock);
715
716    iBuilder->SetInsertPoint(regularBodyBlock);
717
718    s2pInstance->CreateDoBlockCall();
719    leadingInstance->CreateDoBlockCall();
720    sortingInstance->CreateDoBlockCall();
721    gatheringInstance->CreateDoBlockCall();
722//    stdOutInstance->CreateDoBlockCall();
723    remainingBytes2->addIncoming(iBuilder->CreateSub(remainingBytes2, blockSize), regularBodyBlock);
724    iBuilder->CreateBr(regularTestBlock);
725
726    // Check if we have a partial blocks worth of leading data remaining
727    iBuilder->SetInsertPoint(regularExitBlock);
728    PHINode * remainingBytes3 = iBuilder->CreatePHI(intType, 2);
729    remainingBytes3->addIncoming(remainingBytes, safetyCheckBlock);
730    remainingBytes3->addIncoming(remainingBytes2, regularTestBlock);
731    Value * partialBlockCond = iBuilder->CreateICmpNE(remainingBytes3, ConstantInt::getNullValue(intType));
732    iBuilder->CreateCondBr(partialBlockCond, finalTestBlock, partialBlock);
733
734    // If we do, process it and mask out the data
735    iBuilder->SetInsertPoint(partialBlock);
736    s2pInstance->CreateDoBlockCall();
737    leadingInstance->CreateDoBlockCall();
738    leadingInstance->clearOutputStreamSet();
739    sortingInstance->CreateDoBlockCall();
740    gatheringInstance->CreateDoBlockCall();
741//    stdOutInstance->CreateDoBlockCall();
742    iBuilder->CreateBr(finalTestBlock);
743
744    // Now clear the leading data and test the final blocks
745    iBuilder->SetInsertPoint(finalTestBlock);
746    PHINode * remainingFullBlocks = iBuilder->CreatePHI(iBuilder->getInt64Ty(), 3);
747    remainingFullBlocks->addIncoming(iBuilder->getInt64(leadingBlocks), regularExitBlock);
748    remainingFullBlocks->addIncoming(iBuilder->getInt64(leadingBlocks), partialBlock);
749    Value * remainingFullBlocksCond = iBuilder->CreateICmpUGT(remainingFullBlocks, ConstantInt::getNullValue(intType));
750    iBuilder->CreateCondBr(remainingFullBlocksCond, finalBodyBlock, remainingBlock);
751
752    iBuilder->SetInsertPoint(finalBodyBlock);
753
754    leadingInstance->clearOutputStreamSet();
755    sortingInstance->CreateDoBlockCall();
756    gatheringInstance->CreateDoBlockCall();
757//    stdOutInstance->CreateDoBlockCall();
758    remainingFullBlocks->addIncoming(iBuilder->CreateSub(remainingFullBlocks, iBuilder->getInt64(1)), finalBodyBlock);
759
760
761    iBuilder->CreateBr(finalTestBlock);
762
763
764    // perform a final partial gather on all length groups ...
765    iBuilder->SetInsertPoint(remainingBlock);
766
767    Value * const base = iBuilder->CreateLoad(gatheringInstance->getInternalState("Base"));
768    Value * positionArray = gatheringInstance->getInternalState("Positions");
769
770    for (unsigned i = 0; i < mGatherFunction.size(); ++i) {
771        BasicBlock * nonEmptyGroup = BasicBlock::Create(mMod->getContext(), "", main, 0);
772
773        BasicBlock * nextNonEmptyGroup = BasicBlock::Create(mMod->getContext(), "", main, 0);
774
775        ConstantInt * groupIV = iBuilder->getInt32(i);
776        Value * startIndexPtr = iBuilder->CreateGEP(positionArray, {iBuilder->getInt32(0), groupIV, iBuilder->getInt32(0)}, "startIndexPtr");
777        Value * startIndex = iBuilder->CreateLoad(startIndexPtr, "remaining");
778        Value * cond = iBuilder->CreateICmpNE(startIndex, ConstantInt::getNullValue(startIndex->getType()));
779        iBuilder->CreateCondBr(cond, nonEmptyGroup, nextNonEmptyGroup);
780
781        iBuilder->SetInsertPoint(nonEmptyGroup);
782        Value * startArray = iBuilder->CreateGEP(positionArray, {iBuilder->getInt32(0), groupIV, iBuilder->getInt32(1)}, "startArray");
783        Value * startArrayPtr = iBuilder->CreatePointerCast(startArray, PointerType::get(iBuilder->getInt32Ty(), 0));
784        Value * endArray = iBuilder->CreateGEP(positionArray, {iBuilder->getInt32(0), groupIV, iBuilder->getInt32(3)}, "endArray");
785        Value * endArrayPtr = iBuilder->CreatePointerCast(endArray, PointerType::get(iBuilder->getInt32Ty(), 0));
786        Value * outputBuffer = iBuilder->CreatePointerCast(gatheringInstance->getOutputStream(groupIV), iBuilder->getInt8PtrTy());
787        iBuilder->CreateCall5(mGatherFunction.at(i), base, startArrayPtr, endArrayPtr, startIndex, outputBuffer);
788        iBuilder->CreateBr(nextNonEmptyGroup);
789
790        iBuilder->SetInsertPoint(nextNonEmptyGroup);
791    }
792    iBuilder->CreateRetVoid();
793
794    return main;
795}
796
797SymbolTableBuilder::~SymbolTableBuilder() {
798    delete mS2PKernel;
799    delete mLeadingKernel;
800    delete mSortingKernel;
801    delete mGatherKernel;
802    delete mStdOutKernel;
803}
804
805
806}
Note: See TracBrowser for help on using the repository browser.