File: semantic-dcompute.cpp

package info (click to toggle)
ldc 1%3A1.30.0-1
  • links: PTS, VCS
  • area: main
  • in suites: bookworm
  • size: 59,248 kB
  • sloc: cpp: 61,598; ansic: 14,545; sh: 1,014; makefile: 972; asm: 510; objc: 135; exp: 48; python: 12
file content (284 lines) | stat: -rw-r--r-- 9,166 bytes parent folder | download
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
//===-- semantic-dcompute.cpp ---------------------------------------------===//
//
//                         LDC – the LLVM D compiler
//
// This file is distributed under the BSD-style LDC license. See the LICENSE
// file for details.
//
//===----------------------------------------------------------------------===//
//
// Validation for @compute code:
//   enforce: @nogc, nothrow, all function calls are to modules that are also
//   @compute. The enforcemnt of nothrow is simpler because all functions
//   are assumed to not throw. We only need to check for ThrowStatement.
//   We dont use dmd's nothrow detection because it still allows errors.
//
//   ban: classes, interfaces, asm, typeid, global variables, synhronized,
//        associative arrays, pragma(lib,...)
//
//===----------------------------------------------------------------------===//

#include "dmd/declaration.h"
#include "dmd/expression.h"
#include "dmd/id.h"
#include "dmd/identifier.h"
#include "dmd/module.h"
#include "dmd/template.h"
#include "gen/dcompute/target.h"
#include "gen/logger.h"
#include "gen/recursivevisitor.h"
#include "gen/uda.h"

struct DComputeSemanticAnalyser : public StoppableVisitor {
  FuncDeclaration *currentFunction;
  // In @compute code only calls to other functions in `@compute` code are
  // allowed.
  // However, a @kernel function taking a template alias function parameter is
  // allowed, but while the alias appears in the symbol table of the module of
  // the
  // template declaration, it's module of origin is the module at the point of
  // instansiation so we need to check for that.
  bool isNonComputeCallExpVaild(CallExp *ce) {
    FuncDeclaration *f = ce->f;
    if (f->ident == Id::dcReflect)
      return true;
    if (currentFunction == nullptr)
      return false;
    TemplateInstance *inst = currentFunction->isInstantiated();
    if (!inst)
      return false;

    Objects *tiargs = inst->tiargs;
    size_t i = 0, len = tiargs->length;
    IF_LOG Logger::println("checking against: %s (%p) (dyncast=%d)",
                           f->toPrettyChars(), (void *)f, f->dyncast());
    LOG_SCOPE
    for (; i < len; i++) {
      RootObject *o = (*tiargs)[i];
      if (o->dyncast() != DYNCAST_EXPRESSION)
        continue;
      Expression *e = (Expression *)o;
      if (e->op != EXP::function_)
        continue;
      if (f->equals((((FuncExp *)e)->fd))) {
        IF_LOG Logger::println("match");
        return true;
      }
    }
    return false;
  }

  using StoppableVisitor::visit;

  void visit(InterfaceDeclaration *decl) override {
    decl->error("interfaces and classes not allowed in `@compute` code");
    stop = true;
  }

  void visit(ClassDeclaration *decl) override {
    decl->error("interfaces and classes not allowed in `@compute` code");
    stop = true;
  }

  void visit(VarDeclaration *decl) override {
    // Don't print multiple errors for 'synchronized'. see visit(CallExp*)
    if (decl->isDataseg()) {
      if (strncmp(decl->toChars(), "__critsec", 9) &&
        strncmp(decl->toChars(), "typeid", 6)) {
        decl->error("global variables not allowed in `@compute` code");
      }
      // Ignore typeid: it is ignored by codegen.
      stop = true;
      return;
    }

    if (decl->type->ty == TY::Taarray) {
      decl->error("associative arrays not allowed in `@compute` code");
      stop = true;
    }
    // includes interfaces
    else if (decl->type->ty == TY::Tclass) {
      decl->error("interfaces and classes not allowed in `@compute` code");
    }
  }
  void visit(PragmaDeclaration *decl) override {
    if (decl->ident == Id::lib) {
      decl->error(
          "linking additional libraries not supported in `@compute` code");
      stop = true;
    }
  }

  // Nogc enforcement.
  // No need to check AssocArrayLiteral because AA's are banned anyway
  void visit(ArrayLiteralExp *e) override {
    if (e->type->ty != TY::Tarray || !e->elements || !e->elements->length)
      return;
    e->error("array literal in `@compute` code not allowed");
    stop = true;
  }
  void visit(NewExp *e) override {
    e->error("cannot use `new` in `@compute` code");
    stop = true;
  }

  void visit(DeleteExp *e) override {
    e->error("cannot use `delete` in `@compute` code");
    stop = true;
  }
  // No need to check IndexExp because AA's are banned anyway
  void visit(AssignExp *e) override {
    if (e->e1->op == EXP::arrayLength) {
      e->error("setting `length` in `@compute` code not allowed");
      stop = true;
    }
  }

  void visit(CatAssignExp *e) override {
    e->error("cannot use operator `~=` in `@compute` code");
    stop = true;
  }
  void visit(CatExp *e) override {
    e->error("cannot use operator `~` in `@compute` code");
    stop = true;
  }
  // Ban typeid(T)
  void visit(TypeidExp *e) override {
    e->error("typeinfo not available in `@compute` code");
    stop = true;
  }

  void visit(StringExp *e) override {
    e->error("string literals not allowed in `@compute` code");
    stop = true;
  }
  void visit(CompoundAsmStatement *e) override {
    e->error("asm not allowed in `@compute` code");
    stop = true;
  }
  void visit(AsmStatement *e) override {
    e->error("asm not allowed in `@compute` code");
    stop = true;
  }

  // Enforce nothrow. Disallow 'catch' as it is dead code.
  // try...finally is allowed to facilitate scope(exit)
  void visit(TryCatchStatement *e) override {
    e->error("no exceptions in `@compute` code");
    stop = true;
  }
  void visit(ThrowStatement *e) override {
    e->error("no exceptions in `@compute` code");
    stop = true;
  }
  void visit(SwitchStatement *e) override {
    if (auto ce = e->condition->isCallExp()) {
      if (ce->f->ident == Id::__switch) {
        e->error("cannot `switch` on strings in `@compute` code");
        stop = true;
      }
    }
  }

  void visit(IfStatement *stmt) override {
    // Don't descend into ctfe only code
    if (auto ve = stmt->condition->isVarExp()) {
      if (ve->var->ident == Id::ctfe) {
        if (stmt->elsebody)
          visit(stmt->elsebody);
        stop = true;
      }
    } else if (auto ne = stmt->condition->isNotExp()) {
      if (auto ve = ne->e1->isVarExp()) {
        if (ve->var->ident == Id::ctfe) {
          visit(stmt->ifbody);
          stop = true;
        }
      }
    }
    // Code inside an if(__dcompute_reflect(0,0)) { ...} is explicitly
    // for the host and is therefore allowed to call non @compute functions.
    // Thus, the if-statement body's code should not be checked for
    // @compute semantics and the recursive visitor should stop here.
    if (auto ce = stmt->condition->isCallExp()) {
      if (ce->f && ce->f->ident == Id::dcReflect) {
        auto arg1 = (DComputeTarget::ID)(*ce->arguments)[0]->toInteger();
        if (arg1 == DComputeTarget::Host)
          stop = true;
      }
    }
  }
  void visit(CallExp *e) override {
    // SynchronizedStatement is lowered to
    //    Critsec __critsec105; // 105 == line number
    //    _d_criticalenter(& __critsec105); <--
    //    ...                                 |
    //    _d_criticalexit( & __critsec105);   |
    // So we intercept it with the CallExp ----

    if (e->f->ident == Id::criticalenter) {
      e->error("cannot use `synchronized` in `@compute` code");
      stop = true;
      return;
    }

    if (e->f->ident == Id::criticalexit) {
      stop = true;
      return;
    }
      
    Module *m = e->f->getModule();
    if ((m == nullptr || (hasComputeAttr(m) == DComputeCompileFor::hostOnly)) &&
        !isNonComputeCallExpVaild(e)) {
      e->error("can only call functions from other `@compute` modules in "
               "`@compute` code");
      stop = true;
    }
  }

  void visit(FuncDeclaration *fd) override {
    if (hasKernelAttr(fd) && fd->vthis) {
      fd->error("`@kernel` functions must not require `this`");
      stop = true;
      return;
    }

    IF_LOG Logger::println("current function = %s", fd->toChars());
    currentFunction = fd;
  }

  void visit(TemplateDeclaration*) override {
    // Don't try to analyse uninstansiated templates.
    stop = true;
  }

  void visit(TemplateInstance *ti) override {
    // object.RTInfo(Impl) template instantiations are skipped during codegen,
    // as they contain unsupported global variables.
    if (ti->tempdecl == Type::rtinfo || ti->tempdecl == Type::rtinfoImpl) {
      stop = true;
    }
  }

  // Override the default assert(0) behavior of Visitor:
  void visit(Statement *) override {}   // do nothing
  void visit(Expression *) override {}  // do nothing
  void visit(Declaration *) override {} // do nothing
  void visit(Initializer *) override {} // do nothing
  void visit(Dsymbol *) override {}     // do nothing
};

void dcomputeSemanticAnalysis(Module *m) {
  DComputeSemanticAnalyser v;
  RecursiveWalker r(&v);
  for (unsigned k = 0; k < m->members->length; k++) {
    Dsymbol *dsym = (*m->members)[k];
    assert(dsym);
    IF_LOG Logger::println("dcomputeSema: %s: %s", m->toPrettyChars(),
                           dsym->toPrettyChars());
    LOG_SCOPE
    v.currentFunction = nullptr;

    dsym->accept(&r);
  }
}