[pkg-d-commits] [ldc] 64/74: [dcomptue] semantic analysis & lit tests (#2143)
Matthias Klumpp
mak at moszumanska.debian.org
Thu Jul 13 20:54:19 UTC 2017
This is an automated email from the git hooks/post-receive script.
mak pushed a commit to annotated tag v1.3.0-beta2
in repository ldc.
commit 69ad69e872f53c14c101e2c029c4757c4073f487
Author: Nicholas Wilson <thewilsonator at users.noreply.github.com>
Date: Mon May 29 16:02:17 2017 +0800
[dcomptue] semantic analysis & lit tests (#2143)
* CI & semantic analysis
* semantic and codegen tests
---
.travis.yml | 25 +++-
ddmd/mars.d | 8 ++
ddmd/statementsem.d | 8 ++
gen/dcompute/abi-rewrites.h | 4 +-
gen/dcompute/druntime.h | 6 +-
gen/dcompute/targetOCL.cpp | 4 +-
gen/functions.cpp | 15 +-
gen/pgo.cpp | 1 +
gen/recursivevisitor.h | 12 ++
gen/semantic-dcompute.cpp | 245 +++++++++++++++++++++++++++++++++
gen/semantic.d | 27 ++++
gen/uda.cpp | 2 +-
gen/uda.h | 6 +-
runtime/druntime | 2 +-
tests/codegen/dcompute_cl_addrspaces.d | 57 ++++++++
tests/codegen/dcompute_cu_addrspaces.d | 37 +++++
tests/semantic/dcompute.d | 93 +++++++++++++
tests/semantic/inputs/notatcompute.d | 3 +
18 files changed, 530 insertions(+), 25 deletions(-)
diff --git a/.travis.yml b/.travis.yml
index 151b404..b578ead 100644
--- a/.travis.yml
+++ b/.travis.yml
@@ -23,7 +23,7 @@ matrix:
env: LLVM_VERSION=3.5.2 OPTS="-DBUILD_SHARED_LIBS=OFF -DTEST_COVERAGE=ON"
- os: osx
d: ldc-beta
- env: LLVM_VERSION=4.0.0 OPTS="-DBUILD_SHARED_LIBS=OFF"
+ env: LLVM_VERSION=5.0.0 OPTS="-DBUILD_SHARED_LIBS=OFF" LLVM_SPIRV_AVAILABLE=ON
- os: osx
d: ldc
env: LLVM_VERSION=4.0.0 OPTS="-DBUILD_SHARED_LIBS=ON"
@@ -32,6 +32,7 @@ matrix:
cache:
directories:
+ - llvm-spirv-5.0.0
- llvm-4.0.0
- llvm-3.9.1
- llvm-3.9.0
@@ -70,12 +71,24 @@ before_install:
export LLVM_ARCH="x86_64-apple-darwin";
fi;
if [ -z "$(ls -A llvm-$LLVM_VERSION)" ]; then
- wget -O llvm-$LLVM_VERSION.tar.xz http://llvm.org/releases/$LLVM_VERSION/clang+llvm-$LLVM_VERSION-${LLVM_ARCH}.tar.xz;
- mkdir llvm-$LLVM_VERSION;
- tar -xvf llvm-$LLVM_VERSION.tar.xz --strip 1 -C llvm-$LLVM_VERSION;
+ if [ -n "${LLVM_SPIRV_AVAILABLE}" ]; then
+ wget -O llvm-spirv-$LLVM_VERSION.tar.bz2 https://github.com/thewilsonator/llvm/releases/download/pre-intrinsics/LLVM-5.0.0svn-Darwin.tar.bz2;
+ mkdir llvm-spirv-$LLVM_VERSION;
+ tar -xjf llvm-spirv-$LLVM_VERSION.tar.bz2 --strip 1 -C llvm-spirv-$LLVM_VERSION;
+ else
+ wget -O llvm-$LLVM_VERSION.tar.xz http://llvm.org/releases/$LLVM_VERSION/clang+llvm-$LLVM_VERSION-${LLVM_ARCH}.tar.xz;
+ mkdir llvm-$LLVM_VERSION;
+ tar -xvf llvm-$LLVM_VERSION.tar.xz --strip 1 -C llvm-$LLVM_VERSION;
+ fi;
+ fi;
+
+ if [ -n "${LLVM_SPIRV_AVAILABLE}" ]; then
+ llvm-spirv-$LLVM_VERSION/bin/llvm-config --version;
+ export LLVM_CONFIG="llvm-spirv-$LLVM_VERSION/bin/llvm-config";
+ else
+ llvm-$LLVM_VERSION/bin/llvm-config --version;
+ export LLVM_CONFIG="llvm-$LLVM_VERSION/bin/llvm-config";
fi;
- llvm-$LLVM_VERSION/bin/llvm-config --version;
- export LLVM_CONFIG="llvm-$LLVM_VERSION/bin/llvm-config";
install:
- if [ "${TRAVIS_OS_NAME}" = "linux" ]; then export CC="gcc-4.9"; export CXX="g++-4.9"; fi
- if [ "${TRAVIS_OS_NAME}" = "osx" ]; then brew update; brew install ninja; fi;
diff --git a/ddmd/mars.d b/ddmd/mars.d
index 8bd8a9a..c32e019 100644
--- a/ddmd/mars.d
+++ b/ddmd/mars.d
@@ -58,6 +58,7 @@ import ddmd.utils;
version(IN_LLVM)
{
+ import gen.semantic : extraLDCSpecificSemanticAnalysis;
extern (C++):
void genCmain(Scope* sc);
@@ -1602,6 +1603,13 @@ extern (C++) int mars_mainBody(ref Strings files, ref Strings libmodules)
if (global.errors)
fatal();
+ version (IN_LLVM)
+ {
+ extraLDCSpecificSemanticAnalysis(modules);
+ if (global.errors)
+ fatal();
+ }
+
version (IN_LLVM) {} else
{
// Scan for functions to inline
diff --git a/ddmd/statementsem.d b/ddmd/statementsem.d
index a68552f..483e45e 100644
--- a/ddmd/statementsem.d
+++ b/ddmd/statementsem.d
@@ -3100,6 +3100,10 @@ version(IN_LLVM)
Expression e = new DotIdExp(ss.loc, new VarExp(ss.loc, tmp), Id.ptr);
e = e.semantic(sc);
e = new CallExp(ss.loc, new VarExp(ss.loc, fdenter, false), e);
+
+ //IN_LLVM - dmd#6840
+ (cast(CallExp)e).f = fdenter;
+
e.type = Type.tvoid; // do not run semantic on e
cs.push(new ExpStatement(ss.loc, e));
@@ -3107,6 +3111,10 @@ version(IN_LLVM)
e = new DotIdExp(ss.loc, new VarExp(ss.loc, tmp), Id.ptr);
e = e.semantic(sc);
e = new CallExp(ss.loc, new VarExp(ss.loc, fdexit, false), e);
+
+ //IN_LLVM - dmd#6840
+ (cast(CallExp)e).f = fdexit;
+
e.type = Type.tvoid; // do not run semantic on e
Statement s = new ExpStatement(ss.loc, e);
s = new TryFinallyStatement(ss.loc, ss._body, s);
diff --git a/gen/dcompute/abi-rewrites.h b/gen/dcompute/abi-rewrites.h
index 2ecd4e0..4a53316 100644
--- a/gen/dcompute/abi-rewrites.h
+++ b/gen/dcompute/abi-rewrites.h
@@ -15,12 +15,12 @@
#define LDC_GEN_DCOMPUTE_ABI_REWRITES_H
#include "gen/abi.h"
+#include "gen/dcompute/druntime.h"
#include "gen/irstate.h"
#include "gen/llvmhelpers.h"
#include "gen/logger.h"
#include "gen/structs.h"
#include "gen/tollvm.h"
-#include "gen/dcompute/druntime.h"
struct DComputePointerRewrite : ABIRewrite {
LLType *type(Type *t) override {
@@ -29,7 +29,7 @@ struct DComputePointerRewrite : ABIRewrite {
}
LLValue *getLVal(Type *dty, LLValue *v) override {
// TODO: Is this correct?
- return DtoBitCast(v, this->type(dty));
+ return DtoAllocaDump(v, this->type(dty));
}
LLValue *put(DValue *dv) override {
LLValue *address = getAddressOf(dv);
diff --git a/gen/dcompute/druntime.h b/gen/dcompute/druntime.h
index b4f9559..6ef029f 100644
--- a/gen/dcompute/druntime.h
+++ b/gen/dcompute/druntime.h
@@ -14,11 +14,11 @@
#include "ddmd/aggregate.h"
#include "ddmd/mtype.h"
-#include "llvm/ADT/Optional.h"
#include "gen/dcompute/target.h"
#include "gen/irstate.h"
#include "gen/llvm.h"
#include "gen/tollvm.h"
+#include "llvm/ADT/Optional.h"
class Dsymbol;
class Type;
@@ -30,11 +30,11 @@ struct DcomputePointer {
Type *type;
DcomputePointer(int as, Type *ty) : addrspace(as), type(ty) {}
LLType *toLLVMType(bool translate) {
- auto llType = DtoMemType(type);
+ auto llType = DtoType(type);
int as = addrspace;
if (translate)
as = gIR->dcomputetarget->mapping[as];
- return llType->getPointerElementType()->getPointerTo(as);
+ return llType->getPointerTo(as);
}
};
llvm::Optional<DcomputePointer> toDcomputePointer(StructDeclaration *sd);
diff --git a/gen/dcompute/targetOCL.cpp b/gen/dcompute/targetOCL.cpp
index f1f2903..821ecd4 100644
--- a/gen/dcompute/targetOCL.cpp
+++ b/gen/dcompute/targetOCL.cpp
@@ -108,14 +108,14 @@ public:
// MDNode for the kernel argument address space qualifiers.
std::array<llvm::SmallVector<llvm::Metadata *, 8>,count_KernArgMD> paramArgs;
- std::array<const char*,count_KernArgMD> args = {
+ std::array<const char*,count_KernArgMD> args = {{
"kernel_arg_addr_space",
"kernel_arg_access_qual",
"kernel_arg_type",
"kernel_arg_type_qual",
"kernel_arg_base_type",
"kernel_arg_name"
- };
+ }};
for (auto md : args) {
paramArgs[i].push_back(llvm::MDString::get(ctx, md));
diff --git a/gen/functions.cpp b/gen/functions.cpp
index 35bc8b0..3043762 100644
--- a/gen/functions.cpp
+++ b/gen/functions.cpp
@@ -433,13 +433,14 @@ void applyTargetMachineAttributes(llvm::Function &func,
const llvm::TargetOptions &TO = target.Options;
// TODO: implement commandline switches to change the default values.
-
- // Target CPU capabilities
- func.addFnAttr("target-cpu", target.getTargetCPU());
- auto featStr = target.getTargetFeatureString();
- if (!featStr.empty())
- func.addFnAttr("target-features", featStr);
-
+ // TODO: (correctly) apply these for NVPTX (but not for SPIRV).
+ if (!gIR->dcomputetarget) {
+ // Target CPU capabilities
+ func.addFnAttr("target-cpu", target.getTargetCPU());
+ auto featStr = target.getTargetFeatureString();
+ if (!featStr.empty())
+ func.addFnAttr("target-features", featStr);
+ }
// Floating point settings
func.addFnAttr("unsafe-fp-math", TO.UnsafeFPMath ? "true" : "false");
const bool lessPreciseFPMADOption =
diff --git a/gen/pgo.cpp b/gen/pgo.cpp
index 6333eb7..1dde5f8 100644
--- a/gen/pgo.cpp
+++ b/gen/pgo.cpp
@@ -184,6 +184,7 @@ struct MapRegionCounters : public StoppableVisitor {
void visit(Expression *exp) override {}
void visit(Declaration *decl) override {}
void visit(Initializer *init) override {}
+ void visit(Dsymbol *dsym) override {}
void visit(FuncDeclaration *fd) override {
if (NextCounter) {
diff --git a/gen/recursivevisitor.h b/gen/recursivevisitor.h
index 9019d16..28d01b1 100644
--- a/gen/recursivevisitor.h
+++ b/gen/recursivevisitor.h
@@ -29,6 +29,7 @@
#ifndef LDC_GEN_RECURSIVEVISITOR_H
#define LDC_GEN_RECURSIVEVISITOR_H
+#include "ddmd/attrib.h"
#include "ddmd/declaration.h"
#include "ddmd/init.h"
#include "ddmd/statement.h"
@@ -285,6 +286,10 @@ public:
using Visitor::visit;
+ void visit(AttribDeclaration* ad) override {
+ call_visitor(ad) || recurse(ad->decl);
+ }
+
void visit(FuncDeclaration *fd) override {
call_visitor(fd) || recurse(fd->fbody);
}
@@ -383,6 +388,9 @@ public:
recurse(stmt->_body);
}
+ void visit(PragmaStatement* stmt) override {
+ call_visitor(stmt) || recurse(stmt->_body);
+ }
void visit(DebugStatement *stmt) override {
call_visitor(stmt) || recurse(stmt->statement);
}
@@ -457,6 +465,10 @@ public:
call_visitor(init) || recurse(init->exp);
}
+ void visit(ScopeDsymbol *s) override {
+ call_visitor(s) || recurse(s->members);
+ }
+
// Override the default assert(0) behavior of Visitor:
void visit(Statement *stmt) override { call_visitor(stmt); }
void visit(Expression *exp) override { call_visitor(exp); }
diff --git a/gen/semantic-dcompute.cpp b/gen/semantic-dcompute.cpp
new file mode 100644
index 0000000..00ce5db
--- /dev/null
+++ b/gen/semantic-dcompute.cpp
@@ -0,0 +1,245 @@
+//===-- 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 "ddmd/declaration.h"
+#include "ddmd/identifier.h"
+#include "ddmd/module.h"
+#include "ddmd/template.h"
+#include "gen/dcompute/target.h"
+#include "gen/logger.h"
+#include "gen/recursivevisitor.h"
+#include "gen/uda.h"
+#include "id.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) {
+ if (currentFunction == nullptr)
+ return false;
+ TemplateInstance *inst = currentFunction->isInstantiated();
+ if (!inst)
+ return false;
+
+ FuncDeclaration *f = ce->f;
+ Objects *tiargs = inst->tiargs;
+ size_t i = 0, len = tiargs->dim;
+ 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 != TOKfunction)
+ continue;
+ if (f->equals((((FuncExp *)e)->fd))) {
+ IF_LOG Logger::println("match");
+ return true;
+ }
+ }
+ return false;
+ }
+
+ 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() && strncmp(decl->toChars(), "__critsec", 9)) {
+ decl->error("global variables not allowed in @compute code");
+ stop = true;
+ return;
+ }
+
+ if (decl->type->ty == Taarray) {
+ decl->error("associative arrays not allowed in @compute code");
+ stop = true;
+ }
+ // includes interfaces
+ else if (decl->type->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 != Tarray || !e->elements || !e->elements->dim)
+ 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 == TOKarraylength) {
+ 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 @compue 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 (!e->condition->type->isintegral()) {
+ e->error("cannot switch on strings in @compute code");
+ stop = true;
+ }
+ }
+
+ void visit(IfStatement *stmt) override {
+ // 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 (stmt->condition->op == TOKcall) {
+ auto ce = (CallExp *)stmt->condition;
+ 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 msut not require 'this'");
+ stop = true;
+ return;
+ }
+
+ IF_LOG Logger::println("current function = %s", fd->toChars());
+ currentFunction = fd;
+ }
+ // 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->dim; 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);
+ }
+}
diff --git a/gen/semantic.d b/gen/semantic.d
new file mode 100644
index 0000000..75163e4
--- /dev/null
+++ b/gen/semantic.d
@@ -0,0 +1,27 @@
+//===-- gen/semantic.d - Additional LDC semantic analysis ---------*- D -*-===//
+//
+// LDC – the LLVM D compiler
+//
+// This file is distributed under the BSD-style LDC license. See the LICENSE
+// file for details.
+//
+//===----------------------------------------------------------------------===//
+
+module gen.semantic;
+
+import ddmd.arraytypes;
+import ddmd.dsymbol;
+import ddmd.dmodule;
+
+extern(C++) void dcomputeSemanticAnalysis(Module m);
+extern(C) int hasComputeAttr(Dsymbol m);
+
+extern(C++) void extraLDCSpecificSemanticAnalysis(ref Modules modules)
+{
+ foreach(m; modules[])
+ {
+ if (hasComputeAttr(m))
+ dcomputeSemanticAnalysis(m);
+ }
+
+}
diff --git a/gen/uda.cpp b/gen/uda.cpp
index edf3e7a..806bd4f 100644
--- a/gen/uda.cpp
+++ b/gen/uda.cpp
@@ -418,7 +418,7 @@ bool hasWeakUDA(Dsymbol *sym) {
/// Returns 0 if 'sym' does not have the @ldc.dcompute.compute() UDA applied.
/// Returns 1 + n if 'sym' does and is @compute(n).
-DComputeCompileFor hasComputeAttr(Dsymbol *sym) {
+extern "C" DComputeCompileFor hasComputeAttr(Dsymbol *sym) {
auto sle = getMagicAttribute(sym, Id::udaCompute, Id::dcompute);
if (!sle)
diff --git a/gen/uda.h b/gen/uda.h
index b47f24f..74a3dc1 100644
--- a/gen/uda.h
+++ b/gen/uda.h
@@ -28,12 +28,12 @@ void applyVarDeclUDAs(VarDeclaration *decl, llvm::GlobalVariable *gvar);
bool hasWeakUDA(Dsymbol *sym);
bool hasKernelAttr(Dsymbol *sym);
-/// Must match ldc.attributes.Compilefor + 1 == DComputeCompileFor
-enum class DComputeCompileFor
+/// Must match ldc.dcompute.Compilefor + 1 == DComputeCompileFor
+enum class DComputeCompileFor : int
{
hostOnly = 0,
deviceOnly = 1,
hostAndDevice = 2
};
-DComputeCompileFor hasComputeAttr(Dsymbol *sym);
+extern "C" DComputeCompileFor hasComputeAttr(Dsymbol *sym);
#endif
diff --git a/runtime/druntime b/runtime/druntime
index bc714d6..5818ecb 160000
--- a/runtime/druntime
+++ b/runtime/druntime
@@ -1 +1 @@
-Subproject commit bc714d688f35189fe0a20691d2d8f067fe1b0bb0
+Subproject commit 5818ecb7ddeaf726d0a2a90861623bc713615852
diff --git a/tests/codegen/dcompute_cl_addrspaces.d b/tests/codegen/dcompute_cl_addrspaces.d
new file mode 100644
index 0000000..ca19845
--- /dev/null
+++ b/tests/codegen/dcompute_cl_addrspaces.d
@@ -0,0 +1,57 @@
+// REQUIRES: target_SPIRV
+// RUN: %ldc -c -mdcompute-targets=ocl-220 -m64 -output-ll -output-o %s && FileCheck %s --check-prefix=LL < kernels_ocl220_64.ll && llvm-spirv -to-text kernels_ocl220_64.spv && FileCheck %s --check-prefix=SPT < kernels_ocl220_64.spt
+ at compute(CompileFor.deviceOnly) module dcompute_cl_addrspaces;
+import ldc.dcompute;
+
+// LL: %"ldc.dcompute.Pointer!(cast(AddrSpace)0u, float).Pointer" = type { float* }
+// LL: %"ldc.dcompute.Pointer!(cast(AddrSpace)1u, float).Pointer" = type { float addrspace(1)* }
+// LL: %"ldc.dcompute.Pointer!(cast(AddrSpace)2u, float).Pointer" = type { float addrspace(2)* }
+// LL: %"ldc.dcompute.Pointer!(cast(AddrSpace)3u, immutable(float)).Pointer" = type { float addrspace(3)* }
+// LL: %"ldc.dcompute.Pointer!(cast(AddrSpace)4u, float).Pointer" = type { float addrspace(4)* }
+
+// SPT-DAG: 2 TypeVoid [[VOID_ID:[0-9]+]]
+// SPT-DAG: 3 TypeFloat [[FLOAT_ID:[0-9]+]] 32
+
+//See section 3.7 of the SPIR-V Specification for the numbers in the 4th column.
+// SPT-DAG: 4 TypePointer [[SHARED_FLOAT_POINTER_ID:[0-9]+]] 4 [[FLOAT_ID]]
+// SPT-DAG: 4 TypePointer [[CONSTANT_FLOAT_POINTER_ID:[0-9]+]] 0 [[FLOAT_ID]]
+// SPT-DAG: 4 TypePointer [[GLOBAL_FLOAT_POINTER_ID:[0-9]+]] 5 [[FLOAT_ID]]
+// SPT-DAG: 4 TypePointer [[GENERIC_FLOAT_POINTER_ID:[0-9]+]] 8 [[FLOAT_ID]]
+// SPT-DAG: 4 TypePointer [[PRIVATE_FLOAT_POINTER_ID:[0-9]+]] 7 [[FLOAT_ID]]
+
+//void function({ T addrspace(n)* })
+// SPT-DAG: 4 TypeFunction [[FOO_PRIVATE:[0-9]+]] [[VOID_ID]] [[PRIVATE_FLOAT_POINTER_ID]]
+// SPT-DAG: 4 TypeFunction [[FOO_GLOBAL:[0-9]+]] [[VOID_ID]] [[GLOBAL_FLOAT_POINTER_ID]]
+// SPT-DAG: 4 TypeFunction [[FOO_SHARED:[0-9]+]] [[VOID_ID]] [[SHARED_FLOAT_POINTER_ID]]
+// SPT-DAG: 4 TypeFunction [[FOO_CONSTANT:[0-9]+]] [[VOID_ID]] [[CONSTANT_FLOAT_POINTER_ID]]
+// SPT-DAG: 4 TypeFunction [[FOO_GENERIC:[0-9]+]] [[VOID_ID]] [[GENERIC_FLOAT_POINTER_ID]]
+
+void foo(PrivatePointer!float f) {
+ // LL: load float, float*
+ // SPT-DAG: 5 Function [[VOID_ID]] {{[0-9]+}} 0 [[FOO_PRIVATE]]
+ float g = *f;
+}
+
+void foo(GlobalPointer!float f) {
+ // LL: load float, float addrspace(1)*
+ // SPT-DAG: 5 Function [[VOID_ID]] {{[0-9]+}} 0 [[FOO_GLOBAL]]
+ float g = *f;
+}
+
+void foo(SharedPointer!float f) {
+ // LL: load float, float addrspace(2)*
+ // SPT-DAG: 5 Function [[VOID_ID]] {{[0-9]+}} 0 [[FOO_SHARED]]
+ float g = *f;
+}
+
+void foo(ConstantPointer!float f) {
+ // LL: load float, float addrspace(3)*
+ // SPT-DAG: 5 Function [[VOID_ID]] {{[0-9]+}} 0 [[FOO_CONSTANT]]
+ float g = *f;
+}
+
+void foo(GenericPointer!float f) {
+ // LL: load float, float addrspace(4)*
+ // SPT-DAG: 5 Function [[VOID_ID]] {{[0-9]+}} 0 [[FOO_GENERIC]]
+ float g = *f;
+}
diff --git a/tests/codegen/dcompute_cu_addrspaces.d b/tests/codegen/dcompute_cu_addrspaces.d
new file mode 100644
index 0000000..77fa475
--- /dev/null
+++ b/tests/codegen/dcompute_cu_addrspaces.d
@@ -0,0 +1,37 @@
+// REQUIRES: atleast_llvm309
+// REQUIRES: target_NVPTX
+// RUN: %ldc -c -mdcompute-targets=cuda-350 -m64 -output-ll -output-o %s && FileCheck %s --check-prefix=LL < kernels_cuda350_64.ll && FileCheck %s --check-prefix=PTX < kernels_cuda350_64.ptx
+ at compute(CompileFor.deviceOnly) module dcompute_cu_addrspaces;
+import ldc.dcompute;
+
+// LL: %"ldc.dcompute.Pointer!(cast(AddrSpace)0u, float).Pointer" = type { float addrspace(5)* }
+// LL: %"ldc.dcompute.Pointer!(cast(AddrSpace)1u, float).Pointer" = type { float addrspace(1)* }
+// LL: %"ldc.dcompute.Pointer!(cast(AddrSpace)2u, float).Pointer" = type { float addrspace(3)* }
+// LL: %"ldc.dcompute.Pointer!(cast(AddrSpace)3u, immutable(float)).Pointer" = type { float addrspace(4)* }
+// LL: %"ldc.dcompute.Pointer!(cast(AddrSpace)4u, float).Pointer" = type { float* }
+
+void foo(PrivatePointer!float f) {
+ // LL: load float, float addrspace(5)*
+ // PTX: ld.local.f32
+ float g = *f;
+}
+void foo(GlobalPointer!float f) {
+ // LL: load float, float addrspace(1)*
+ // PTX: ld.global.f32
+ float g = *f;
+}
+void foo(SharedPointer!float f) {
+ // LL: load float, float addrspace(3)*
+ // PTX: ld.shared.f32
+ float g = *f;
+}
+void foo(ConstantPointer!float f) {
+ // LL: load float, float addrspace(4)*
+ // PTX: ld.const.f32
+ float g = *f;
+}
+void foo(GenericPointer!float f) {
+ // LL: load float, float*
+ // PTX: ld.f32
+ float g = *f;
+}
diff --git a/tests/semantic/dcompute.d b/tests/semantic/dcompute.d
new file mode 100644
index 0000000..11aa980
--- /dev/null
+++ b/tests/semantic/dcompute.d
@@ -0,0 +1,93 @@
+// RUN: not %ldc -o- -verrors=0 -I%S %s 2>&1 | FileCheck %s
+
+
+ at compute(CompileFor.deviceOnly) module tests.semaintic.dcompute;
+import ldc.dcompute;
+import inputs.notatcompute : somefunc;
+
+extern(C) bool perhaps();
+//CHECK: dcompute.d([[@LINE+1]]): Error: {{.*}} interfaces and classes not allowed in @compute code
+interface I {}
+
+//CHECK: dcompute.d([[@LINE+1]]): Error: {{.*}} interfaces and classes not allowed in @compute code
+class C : Throwable { this() { super(""); } }
+
+//CHECK: dcompute.d([[@LINE+1]]): Error: {{.*}} global variables not allowed in @compute code
+C c;
+
+void func()
+{
+ //CHECK: dcompute.d([[@LINE+1]]): Error: {{.*}} associative arrays not allowed in @compute code
+ int[int] foo;
+ //CHECK: dcompute.d([[@LINE+1]]): Error: array literal in @compute code not allowed
+ auto bar = [0, 1, 2];
+ //CHECK: dcompute.d([[@LINE+1]]): Error: cannot use 'new' in @compute code
+ auto baz = new int;
+ //CHECK: dcompute.d([[@LINE+1]]): Error: cannot use 'delete' in @compute code
+ delete baz;
+
+ //CHECK: dcompute.d([[@LINE+1]]): Error: {{.*}} interfaces and classes not allowed in @compute code
+ I i;
+ //CHECK: dcompute.d([[@LINE+1]]): Error: {{.*}} interfaces and classes not allowed in @compute code
+ C cc;
+ int[] quux;
+ //CHECK: dcompute.d([[@LINE+1]]): Error: setting 'length' in @compute code not allowed
+ quux.length = 1;
+ //CHECK: dcompute.d([[@LINE+1]]): Error: cannot use operator ~= in @compute code
+ quux ~= 42;
+ //CHECK: dcompute.d([[@LINE+1]]): Error: cannot use operator ~ in @compute code
+ cast(void) (quux ~ 1);
+ //CHECK: dcompute.d([[@LINE+1]]): Error: typeinfo not available in @compute code
+ cast(void) typeid(int);
+ //CHECK: dcompute.d([[@LINE+1]]): Error: cannot use 'synchronized' in @compute code
+ synchronized {}
+ //CHECK: dcompute.d([[@LINE+1]]): Error: string literals not allowed in @compue code
+ auto s = "geaxsese";
+ //CHECK: dcompute.d([[@LINE+1]]): Error: cannot switch on strings in @compute code
+ switch(s)
+ {
+ default:
+ break;
+ }
+
+ //CHECK: dcompute.d([[@LINE+1]]): Error: can only call functions from other @compute modules in @compute code
+ somefunc();
+ if (__dcompute_reflect(ReflectTarget.Host,0))
+ //CHECK-NOT: Error:
+ somefunc();
+
+ //CHECK: dcompute.d([[@LINE+1]]): Error: no exceptions in @compute code
+ try
+ {
+ func1();
+ }
+ catch(C c)
+ {
+ }
+
+ if (perhaps())
+ //CHECK: dcompute.d([[@LINE+1]]): Error: no exceptions in @compute code
+ throw c;
+
+ //CHECK-NOT: Error:
+ try
+ {
+ func1();
+ }
+ finally
+ {
+ func2();
+ }
+ //CHECK-NOT: Error:
+ scope(exit)
+ func2();
+
+ //CHECK: dcompute.d([[@LINE+1]]): Error: asm not allowed in @compute code
+ asm {ret;}
+}
+
+void func1() {}
+void func2() {}
+
+//CHECK: dcompute.d([[@LINE+1]]): Error: pragma lib linking additional libraries not supported in @compute code
+pragma(lib, "bar");
diff --git a/tests/semantic/inputs/notatcompute.d b/tests/semantic/inputs/notatcompute.d
new file mode 100644
index 0000000..b204bb8
--- /dev/null
+++ b/tests/semantic/inputs/notatcompute.d
@@ -0,0 +1,3 @@
+module tests.semantic.inputs.notatcompute;
+
+void somefunc() {}
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/pkg-d/ldc.git
More information about the pkg-d-commits
mailing list