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

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

Symbol table work

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