Pooled IDs (#2116)

Use pooled identifiers for magic UDAs and dcompute. Fixes #2011

ddmd/idgen.d

@@ -397,6 +397,25 @@ Msgtable[] msgtable =
     // IN_LLVM: LDC-specific traits.
     { "targetCPU" },
     { "targetHasFeature" },
+    
+    // IN_LLVM: LDC-specific attributes
+    { "ldc" },
+    { "attributes" },
+    { "udaAllocSize", "allocSize" },
+    // fastmath is an AliasSeq of llvmAttr and llvmFastMathFlag
+    { "udaOptStrategy", "optStrategy" },
+    { "udaLLVMAttr", "llvmAttr" },
+    { "udaLLVMFastMathFlag", "llvmFastMathFlag" },
+    { "udaSection", "section" },
+    { "udaTarget", "target" },
+    { "udaWeak", "_weak" },
+    { "udaCompute", "compute" },
+    { "udaKernel", "_kernel" },
+    
+    // IN_LLVM: DCompute specific types and functionss
+    { "dcompute" },
+    { "dcPointer", "Pointer" },
+    { "dcReflect", "__dcompute_reflect" },
 ];
 
 

gen/uda.cpp

@@ -8,42 +8,33 @@
 #include "expression.h"
 #include "ir/irfunction.h"
 #include "module.h"
+#include "id.h"
 
 #include "llvm/ADT/StringExtras.h"
 
 namespace {
 
-/// Names of the attribute structs we recognize.
+/// Checks whether `moduleDecl` is in the ldc package and it's identifier is `id`.
-namespace attr {
+bool isMagicModule(const ModuleDeclaration *moduleDecl, const Identifier* id) {
-const std::string allocSize = "allocSize";
-const std::string llvmAttr = "llvmAttr";
-const std::string llvmFastMathFlag = "llvmFastMathFlag";
-const std::string optStrategy = "optStrategy";
-const std::string section = "section";
-const std::string target = "target";
-const std::string weak = "_weak";
-}
-
-/// Checks whether `moduleDecl` is the ldc.attributes module.
-bool isLdcAttibutes(const ModuleDeclaration *moduleDecl) {
   if (!moduleDecl)
     return false;
 
-  if (strcmp("attributes", moduleDecl->id->toChars())) {
+  if (moduleDecl->id != id) {
     return false;
   }
 
   if (moduleDecl->packages->dim != 1 ||
-      strcmp("ldc", (*moduleDecl->packages)[0]->toChars())) {
+      (*moduleDecl->packages)[0] != Id::ldc) {
     return false;
   }
   return true;
 }
 
-/// Checks whether the type of `e` is a struct from the ldc.attributes module.
+/// Checks whether the type of `e` is a struct from an ldc recognised module,
-bool isFromLdcAttibutes(const StructLiteralExp *e) {
+/// i.e. ldc.attributes or ldc.dcompute.
+bool isFromMagicModule(const StructLiteralExp *e, const Identifier* id) {
   auto moduleDecl = e->sd->getModule()->md;
-  return isLdcAttibutes(moduleDecl);
+  return isMagicModule(moduleDecl,id);
 }
 
 StructLiteralExp *getLdcAttributesStruct(Expression *attr) {
@@ -61,7 +52,7 @@ StructLiteralExp *getLdcAttributesStruct(Expression *attr) {
   }
 
   auto sle = static_cast<StructLiteralExp *>(e);
-  if (isFromLdcAttibutes(sle)) {
+  if (isFromMagicModule(sle,Id::attributes)) {
     return sle;
   }
 
@@ -71,25 +62,29 @@ StructLiteralExp *getLdcAttributesStruct(Expression *attr) {
 void checkStructElems(StructLiteralExp *sle, ArrayParam<Type *> elemTypes) {
   if (sle->elements->dim != elemTypes.size()) {
     sle->error(
-        "unexpected field count in 'ldc.attributes.%s'; does druntime not "
+        "unexpected field count in 'ldc.%s.%s'; does druntime not "
         "match compiler version?",
+        sle->sd->getModule()->md->id->toChars(),
         sle->sd->ident->toChars());
     fatal();
   }
 
   for (size_t i = 0; i < sle->elements->dim; ++i) {
     if ((*sle->elements)[i]->type->toBasetype() != elemTypes[i]) {
-      sle->error("invalid field type in 'ldc.attributes.%s'; does druntime not "
+      sle->error("invalid field type in 'ldc.%s.%s'; does druntime not "
                  "match compiler version?",
+                 sle->sd->getModule()->md->id->toChars(),
                  sle->sd->ident->toChars());
       fatal();
     }
   }
 }
 
-/// Returns the StructLiteralExp magic attribute with name `name` if it is
+/// Returns the StructLiteralExp magic attribute with identifier `id` from
-/// applied to `sym`, otherwise returns nullptr.
+/// the ldc magic module with identifier `from` (attributes or dcompute)
-StructLiteralExp *getMagicAttribute(Dsymbol *sym, std::string name) {
+/// if it is applied to `sym`, otherwise returns nullptr.
+StructLiteralExp *getMagicAttribute(Dsymbol *sym, const Identifier* id,
+                                    const Identifier* from) {
   if (!sym->userAttribDecl)
     return nullptr;
 
@@ -97,11 +92,13 @@ StructLiteralExp *getMagicAttribute(Dsymbol *sym, std::string name) {
   Expressions *attrs = sym->userAttribDecl->getAttributes();
   expandTuples(attrs);
   for (auto &attr : *attrs) {
-    auto sle = getLdcAttributesStruct(attr);
+    if (attr->op != TOKstructliteral)
-    if (!sle)
+      continue;
+    auto sle = static_cast<StructLiteralExp *>(attr);
+    if (!isFromMagicModule(sle,from))
       continue;
 
-    if (name == sle->sd->ident->toChars()) {
+    if (id == sle->sd->ident) {
       return sle;
     }
   }
@@ -348,18 +345,18 @@ void applyVarDeclUDAs(VarDeclaration *decl, llvm::GlobalVariable *gvar) {
     if (!sle)
       continue;
 
-    auto name = sle->sd->ident->toChars();
+    auto ident = sle->sd->ident;
-    if (name == attr::section) {
+    if (ident == Id::udaSection) {
       applyAttrSection(sle, gvar);
-    } else if (name == attr::optStrategy || name == attr::target) {
+    } else if (ident == Id::udaOptStrategy || ident == Id::udaTarget) {
       sle->error(
           "Special attribute 'ldc.attributes.%s' is only valid for functions",
-          name);
+          ident->toChars());
-    } else if (name == attr::weak) {
+    } else if (ident == Id::udaWeak) {
       // @weak is applied elsewhere
     } else {
       sle->warning(
-          "Ignoring unrecognized special attribute 'ldc.attributes.%s'", name);
+          "Ignoring unrecognized special attribute 'ldc.attributes.%s'", ident->toChars());
     }
   }
 }
@@ -378,31 +375,31 @@ void applyFuncDeclUDAs(FuncDeclaration *decl, IrFunction *irFunc) {
     if (!sle)
       continue;
 
-    auto name = sle->sd->ident->toChars();
+    auto ident = sle->sd->ident;
-    if (name == attr::allocSize) {
+    if (ident == Id::udaAllocSize) {
       applyAttrAllocSize(sle, irFunc);
-    } else if (name == attr::llvmAttr) {
+    } else if (ident == Id::udaLLVMAttr) {
       applyAttrLLVMAttr(sle, func);
-    } else if (name == attr::llvmFastMathFlag) {
+    } else if (ident == Id::udaLLVMFastMathFlag) {
       applyAttrLLVMFastMathFlag(sle, irFunc);
-    } else if (name == attr::optStrategy) {
+    } else if (ident == Id::udaOptStrategy) {
       applyAttrOptStrategy(sle, irFunc);
-    } else if (name == attr::section) {
+    } else if (ident == Id::udaSection) {
       applyAttrSection(sle, func);
-    } else if (name == attr::target) {
+    } else if (ident == Id::udaTarget) {
       applyAttrTarget(sle, func);
-    } else if (name == attr::weak) {
+    } else if (ident == Id::udaWeak || ident == Id::udaKernel) {
-      // @weak is applied elsewhere
+      // @weak and @kernel are applied elsewhere
     } else {
       sle->warning(
-          "Ignoring unrecognized special attribute 'ldc.attributes.%s'", name);
+          "Ignoring unrecognized special attribute 'ldc.attributes.%s'", ident->toChars());
     }
   }
 }
 
 /// Checks whether 'sym' has the @ldc.attributes._weak() UDA applied.
 bool hasWeakUDA(Dsymbol *sym) {
-  auto sle = getMagicAttribute(sym, attr::weak);
+  auto sle = getMagicAttribute(sym, Id::udaWeak, Id::attributes);
   if (!sle)
     return false;
 
@@ -413,3 +410,32 @@ bool hasWeakUDA(Dsymbol *sym) {
                "global variables");
   return true;
 }
+
+/// Returns 0 if 'sym' does not have the @ldc.dcompute.compute() UDA applied.
+/// Returns 1 + n if 'sym' does and is @compute(n).
+int hasComputeAttr(Dsymbol *sym) {
+
+  auto sle = getMagicAttribute(sym, Id::udaCompute, Id::dcompute);
+  if (!sle)
+    return 0;
+
+  checkStructElems(sle, {Type::tint32});
+
+  return 1 + (*sle->elements)[0]->toInteger();
+}
+
+/// Checks whether 'sym' has the @ldc.dcompute._kernel() UDA applied.
+bool hasKernelAttr(Dsymbol *sym) {
+  auto sle = getMagicAttribute(sym, Id::udaKernel, Id::dcompute);
+  if (!sle)
+    return false;
+
+  checkStructElems(sle, {});
+
+  if (!sym->isFuncDeclaration() && !hasComputeAttr(sym->getModule()))
+    sym->error("@ldc.dcompute.kernel can only be applied to functions"
+               " in modules marked @ldc.dcompute.compute");
+
+  return true;
+}
+

gen/uda.h

@@ -27,5 +27,7 @@ void applyFuncDeclUDAs(FuncDeclaration *decl, IrFunction *irFunc);
 void applyVarDeclUDAs(VarDeclaration *decl, llvm::GlobalVariable *gvar);
 
 bool hasWeakUDA(Dsymbol *sym);
+int hasComputeAttr(Dsymbol *sym);
+bool hasKernelAttr(Dsymbol *sym);
 
 #endif