Skip to content

Commit df7cd25

Browse files
sletzsletz
committed
Add NewDSPInst FIR instruction.
1 parent e372a3b commit df7cd25

10 files changed

+200
-74
lines changed

‎compiler/generator/c/c_instructions.hh

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -263,6 +263,8 @@ class CInstVisitor : public TextInstVisitor {
263263
EndLine();
264264
}
265265

266+
virtual void visit(NewDSPInst* inst) { *fOut << "calloc(1, sizeof(" << inst->fName << "))"; }
267+
266268
virtual void visit(DeclareVarInst* inst)
267269
{
268270
if (inst->fAddress->isStaticStruct()) {
@@ -468,7 +470,6 @@ class CInstVisitor : public TextInstVisitor {
468470
// fTab--;
469471
inst->fGlobals->accept(this);
470472
tab(fTab, *fOut);
471-
472473
for (const auto& it : inst->fFunctions) {
473474
tab(fTab, *fOut);
474475
it->accept(this);

‎compiler/generator/cpp/cpp_code_container.hh

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -155,7 +155,7 @@ class CPPCodeContainer : public virtual CodeContainer {
155155

156156
printMathHeader();
157157

158-
fCodeProducer = new CPPInstVisitor(out);
158+
fCodeProducer = new CPPInstVisitor(out, "");
159159
}
160160

161161
virtual ~CPPCodeContainer()

‎compiler/generator/cpp/cpp_gpu_code_container.cpp

Lines changed: 12 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -165,15 +165,15 @@ void CPPOpenCLCodeContainer::produceClass()
165165
// Separate control and non-controls fields in 2 separeted structures
166166
tab1(n, *fGPUOut);
167167
*fGPUOut << "typedef struct {";
168-
DSPOpenCLInstVisitor dsp_visitor(fGPUOut, n + 1);
168+
DSPOpenCLInstVisitor dsp_visitor(fGPUOut, fKlassName, n + 1);
169169
fDeclarationInstructions->accept(&dsp_visitor);
170170
tab1(n, *fGPUOut);
171171
*fGPUOut << "} faustdsp;";
172172
tab1(n, *fGPUOut);
173173

174174
tab1(n, *fGPUOut);
175175
*fGPUOut << "typedef struct {";
176-
ControlOpenCLInstVisitor control_visitor(fGPUOut, n + 1);
176+
ControlOpenCLInstVisitor control_visitor(fGPUOut, fKlassName, n + 1);
177177
fDeclarationInstructions->accept(&control_visitor);
178178
tab1(n, *fGPUOut);
179179
*fGPUOut << "} faustcontrol;";
@@ -221,15 +221,15 @@ void CPPOpenCLCodeContainer::produceClass()
221221
// Separate control and non-controls fields in 2 structures
222222
tab(n + 1, *fOut);
223223
*fOut << "typedef struct {";
224-
DSPInstVisitor dsp_visitor1(fOut, n + 2);
224+
DSPInstVisitor dsp_visitor1(fOut, fKlassName, n + 2);
225225
fDeclarationInstructions->accept(&dsp_visitor1);
226226
tab(n + 1, *fOut);
227227
*fOut << "} faustdsp;";
228228
tab(n + 1, *fOut);
229229

230230
tab(n + 1, *fOut);
231231
*fOut << "typedef struct {";
232-
ControlInstVisitor control_visitor1(fOut, n + 2);
232+
ControlInstVisitor control_visitor1(fOut, fKlassName, n + 2);
233233
fDeclarationInstructions->accept(&control_visitor1);
234234
tab(n + 1, *fOut);
235235
*fOut << "} faustcontrol;";
@@ -943,7 +943,7 @@ void CPPOpenCLCodeContainer::produceClass()
943943
if (fUserInterfaceInstructions->fCode.size() > 0) {
944944
tab(n + 2, *fOut);
945945
fCodeProducer->Tab(n + 2);
946-
UIInstVisitor ui_visitor(fOut, n + 2);
946+
UIInstVisitor ui_visitor(fOut, fKlassName, n + 2);
947947
fUserInterfaceInstructions->accept(&ui_visitor);
948948
}
949949
tab(n + 1, *fOut);
@@ -1065,7 +1065,7 @@ void CPPOpenCLVectorCodeContainer::generateComputeKernel(int n)
10651065
tab1(n + 1, *fGPUOut);
10661066

10671067
// Generates local variables declaration and setup
1068-
BlockKernelInstVisitor block_visitor(fGPUOut, n + 1);
1068+
BlockKernelInstVisitor block_visitor(fGPUOut, fKlassName, n + 1);
10691069
fComputeBlockInstructions->accept(&block_visitor);
10701070

10711071
lclgraph dag;
@@ -1318,15 +1318,15 @@ void CPPCUDACodeContainer::produceClass()
13181318
tab(n, *fGPUOut);
13191319
tab(n, *fGPUOut);
13201320
*fGPUOut << "typedef struct {";
1321-
DSPInstVisitor dsp_visitor(fGPUOut, n + 1);
1321+
DSPInstVisitor dsp_visitor(fGPUOut, fKlassName, n + 1);
13221322
fDeclarationInstructions->accept(&dsp_visitor);
13231323
tab(n, *fGPUOut);
13241324
*fGPUOut << "} faustdsp;";
13251325
tab(n, *fGPUOut);
13261326

13271327
tab(n, *fGPUOut);
13281328
*fGPUOut << "typedef struct {";
1329-
ControlInstVisitor control_visitor(fGPUOut, n + 1);
1329+
ControlInstVisitor control_visitor(fGPUOut, fKlassName, n + 1);
13301330
fDeclarationInstructions->accept(&control_visitor);
13311331
tab(n, *fGPUOut);
13321332
*fGPUOut << "} faustcontrol;";
@@ -1358,15 +1358,15 @@ void CPPCUDACodeContainer::produceClass()
13581358
// Separate control and non-controls fields in 2 structures
13591359
tab(n, *fOut);
13601360
*fOut << "typedef struct {";
1361-
DSPInstVisitor dsp_visitor1(fOut, n + 1);
1361+
DSPInstVisitor dsp_visitor1(fOut, fKlassName, n + 1);
13621362
fDeclarationInstructions->accept(&dsp_visitor1);
13631363
tab(n, *fOut);
13641364
*fOut << "} faustdsp;";
13651365
tab(n, *fOut);
13661366

13671367
tab(n, *fOut);
13681368
*fOut << "typedef struct {";
1369-
ControlInstVisitor control_visitor1(fOut, n + 1);
1369+
ControlInstVisitor control_visitor1(fOut, fKlassName, n + 1);
13701370
fDeclarationInstructions->accept(&control_visitor1);
13711371
tab(n, *fOut);
13721372
*fOut << "} faustcontrol;";
@@ -1876,7 +1876,7 @@ void CPPCUDACodeContainer::produceClass()
18761876
if (fUserInterfaceInstructions->fCode.size() > 0) {
18771877
tab(n + 2, *fOut);
18781878
fCodeProducer->Tab(n + 2);
1879-
UIInstVisitor ui_visitor(fOut, n + 2);
1879+
UIInstVisitor ui_visitor(fOut, fKlassName, n + 2);
18801880
fUserInterfaceInstructions->accept(&ui_visitor);
18811881
}
18821882
tab(n + 1, *fOut);
@@ -2012,7 +2012,7 @@ void CPPCUDAVectorCodeContainer::generateComputeKernel(int n)
20122012
tab(n + 1, *fGPUOut);
20132013

20142014
// Generates local variables declaration and setup
2015-
BlockKernelInstVisitor block_visitor(fGPUOut, n + 1);
2015+
BlockKernelInstVisitor block_visitor(fGPUOut, fKlassName, n + 1);
20162016
fComputeBlockInstructions->accept(&block_visitor);
20172017

20182018
lclgraph dag;

‎compiler/generator/cpp/cpp_gpu_code_container.hh

Lines changed: 40 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -31,7 +31,10 @@ class CPPGPUCodeContainer : public CPPCodeContainer {
3131
protected:
3232
// To access control inside fControl field
3333
struct UIInstVisitor : public CPPInstVisitor {
34-
UIInstVisitor(std::ostream* out, int tab) : CPPInstVisitor(out, tab) {}
34+
UIInstVisitor(std::ostream* out, const std::string& struct_name, int tab)
35+
: CPPInstVisitor(out, struct_name, tab)
36+
{
37+
}
3538

3639
virtual void visit(AddMetaDeclareInst* inst)
3740
{
@@ -106,7 +109,10 @@ class CPPGPUCodeContainer : public CPPCodeContainer {
106109

107110
// Visitor that only generates non-control fields
108111
struct DSPInstVisitor : public CPPInstVisitor {
109-
DSPInstVisitor(std::ostream* out, int tab) : CPPInstVisitor(out, tab) {}
112+
DSPInstVisitor(std::ostream* out, const std::string& struct_name, int tab)
113+
: CPPInstVisitor(out, struct_name, tab)
114+
{
115+
}
110116

111117
virtual void visit(DeclareVarInst* inst)
112118
{
@@ -119,7 +125,10 @@ class CPPGPUCodeContainer : public CPPCodeContainer {
119125

120126
// Visitor that only generates control fields
121127
struct ControlInstVisitor : public CPPInstVisitor {
122-
ControlInstVisitor(std::ostream* out, int tab) : CPPInstVisitor(out, tab) {}
128+
ControlInstVisitor(std::ostream* out, const std::string& struct_name, int tab)
129+
: CPPInstVisitor(out, struct_name, tab)
130+
{
131+
}
123132

124133
virtual void visit(DeclareVarInst* inst)
125134
{
@@ -136,7 +145,10 @@ class CPPGPUCodeContainer : public CPPCodeContainer {
136145
using CPPInstVisitor::visit;
137146

138147
std::map<std::string, std::string> fFunctionTable;
139-
KernelInstVisitor(std::ostream* out, int tab) : CPPInstVisitor(out, tab) {}
148+
KernelInstVisitor(std::ostream* out, const std::string& struct_name, int tab)
149+
: CPPInstVisitor(out, struct_name, tab)
150+
{
151+
}
140152

141153
virtual void visit(LoadVarInst* inst)
142154
{
@@ -300,7 +312,8 @@ class CPPOpenCLCodeContainer : public CPPGPUCodeContainer {
300312
}
301313
}
302314

303-
OpenCLKernelInstVisitor(std::ostream* out, int tab) : KernelInstVisitor(out, tab)
315+
OpenCLKernelInstVisitor(std::ostream* out, const std::string& struct_name, int tab)
316+
: KernelInstVisitor(out, struct_name, tab)
304317
{
305318
fFunctionTable["sin"] = "native_sin";
306319
fFunctionTable["sinf"] = "native_sin";
@@ -337,7 +350,10 @@ class CPPOpenCLCodeContainer : public CPPGPUCodeContainer {
337350
}
338351
}
339352

340-
ControlOpenCLInstVisitor(std::ostream* out, int tab) : ControlInstVisitor(out, tab) {}
353+
ControlOpenCLInstVisitor(std::ostream* out, const std::string& struct_name, int tab)
354+
: ControlInstVisitor(out, struct_name, tab)
355+
{
356+
}
341357
};
342358

343359
// To be used when generating GPU kernel string
@@ -352,7 +368,10 @@ class CPPOpenCLCodeContainer : public CPPGPUCodeContainer {
352368
}
353369
}
354370

355-
DSPOpenCLInstVisitor(std::ostream* out, int tab) : DSPInstVisitor(out, tab) {}
371+
DSPOpenCLInstVisitor(std::ostream* out, const std::string& struct_name, int tab)
372+
: DSPInstVisitor(out, struct_name, tab)
373+
{
374+
}
356375
};
357376

358377
// Add __local keyword for stack variables
@@ -369,7 +388,10 @@ class CPPOpenCLCodeContainer : public CPPGPUCodeContainer {
369388
}
370389
}
371390

372-
BlockKernelInstVisitor(std::ostream* out, int tab) : KernelInstVisitor(out, tab) {}
391+
BlockKernelInstVisitor(std::ostream* out, const std::string& struct_name, int tab)
392+
: KernelInstVisitor(out, struct_name, tab)
393+
{
394+
}
373395

374396
virtual void visit(DeclareVarInst* inst)
375397
{
@@ -413,7 +435,7 @@ class CPPOpenCLCodeContainer : public CPPGPUCodeContainer {
413435
: CPPGPUCodeContainer(name, super, numInputs, numOutputs, out)
414436
{
415437
fGPUOut = new std::ostringstream();
416-
fKernelCodeProducer = new OpenCLKernelInstVisitor(fGPUOut, 0);
438+
fKernelCodeProducer = new OpenCLKernelInstVisitor(fGPUOut, name, 0);
417439
}
418440
virtual ~CPPOpenCLCodeContainer() { delete fGPUOut; }
419441

@@ -438,14 +460,20 @@ class CPPOpenCLVectorCodeContainer : public CPPOpenCLCodeContainer {
438460
class CPPCUDACodeContainer : public CPPGPUCodeContainer {
439461
protected:
440462
struct CUDAKernelInstVisitor : public KernelInstVisitor {
441-
CUDAKernelInstVisitor(std::ostream* out, int tab) : KernelInstVisitor(out, tab) {}
463+
CUDAKernelInstVisitor(std::ostream* out, const std::string& struct_name, int tab)
464+
: KernelInstVisitor(out, struct_name, tab)
465+
{
466+
}
442467
};
443468

444469
// Add __shared__ keyword for stack variables
445470
struct BlockKernelInstVisitor : public KernelInstVisitor {
446471
using KernelInstVisitor::visit;
447472

448-
BlockKernelInstVisitor(std::ostream* out, int tab) : KernelInstVisitor(out, tab) {}
473+
BlockKernelInstVisitor(std::ostream* out, const std::string& struct_name, int tab)
474+
: KernelInstVisitor(out, struct_name, tab)
475+
{
476+
}
449477

450478
virtual void visit(DeclareVarInst* inst)
451479
{
@@ -477,7 +505,7 @@ class CPPCUDACodeContainer : public CPPGPUCodeContainer {
477505
{
478506
std::string filename = gGlobal->gOutputFile + ".cu";
479507
fGPUOut = new std::ofstream(filename.c_str());
480-
fKernelCodeProducer = new CUDAKernelInstVisitor(fGPUOut, 0);
508+
fKernelCodeProducer = new CUDAKernelInstVisitor(fGPUOut, name, 0);
481509
fNumInputs = numInputs;
482510
fNumOutputs = numOutputs;
483511
}

‎compiler/generator/cpp/cpp_instructions.hh

Lines changed: 8 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -46,8 +46,8 @@ class CPPInstVisitor : public TextInstVisitor {
4646
public:
4747
using TextInstVisitor::visit;
4848

49-
CPPInstVisitor(std::ostream* out, int tab = 0)
50-
: TextInstVisitor(out, "->", new CStringTypeManager(xfloat(), "*"), tab)
49+
CPPInstVisitor(std::ostream* out, const std::string& struct_name, int tab = 0)
50+
: TextInstVisitor(out, "->", new CStringTypeManager(xfloat(), "*", struct_name), tab)
5151
{
5252
// Mark all math.h functions as generated...
5353
gFunctionSymbolTable["abs"] = true;
@@ -363,6 +363,8 @@ class CPPInstVisitor : public TextInstVisitor {
363363
EndLine();
364364
}
365365

366+
virtual void visit(NewDSPInst* inst) { *fOut << "new " << inst->fName << "()"; }
367+
366368
virtual void visit(DeclareVarInst* inst)
367369
{
368370
if (inst->getAccess() & Address::kConst) {
@@ -539,7 +541,10 @@ class CPPInstVisitor : public TextInstVisitor {
539541

540542
class CPPVecInstVisitor : public CPPInstVisitor {
541543
public:
542-
CPPVecInstVisitor(std::ostream* out, int tab = 0) : CPPInstVisitor(out, tab) {}
544+
CPPVecInstVisitor(std::ostream* out, const std::string& struct_name, int tab = 0)
545+
: CPPInstVisitor(out, struct_name, tab)
546+
{
547+
}
543548
};
544549

545550
#endif

0 commit comments

Comments
 (0)