Menu
▾
▴
[X10-commits] SF.net SVN: x10:[11033] trunk
|
From: <spa...@us...> - 2009-08-14 23:09:38
|
Revision: 11033
http://x10.svn.sourceforge.net/x10/?rev=11033&view=rev
Author: sparksparkspark
Date: 2009-08-14 23:09:29 +0000 (Fri, 14 Aug 2009)
Log Message:
-----------
Add function to get size of a primitive from type system.
Remove dead ArrayType class
Add "cuda" native annotation, used only when compiling a kernel, falls back to "c++" if "cuda" not found
Add x10cuda to build.xml
Modified Paths:
--------------
trunk/x10.compiler/build.xml
trunk/x10.compiler/src/x10/types/X10TypeSystem.java
trunk/x10.compiler/src/x10/types/X10TypeSystem_c.java
trunk/x10.compiler/src/x10cpp/visit/ASTQuery.java
trunk/x10.compiler/src/x10cpp/visit/MessagePassingCodeGenerator.java
trunk/x10.compiler/src/x10cpp/visit/SharedVarsMethods.java
trunk/x10.compiler/src/x10cuda/types/SharedMem.java
trunk/x10.compiler/src/x10cuda/visit/CUDACodeGenerator.java
trunk/x10.runtime/src-x10/x10/lang/Float.x10
trunk/x10.runtime/src-x10/x10/lang/Rail.x10
trunk/x10.runtime/src-x10/x10/lang/ValRail.x10
Removed Paths:
-------------
trunk/x10.compiler/src/x10/types/X10ArrayType.java
trunk/x10.compiler/src/x10/types/X10ArrayType_c.java
Modified: trunk/x10.compiler/build.xml
===================================================================
--- trunk/x10.compiler/build.xml 2009-08-14 21:43:40 UTC (rev 11032)
+++ trunk/x10.compiler/build.xml 2009-08-14 23:09:29 UTC (rev 11033)
@@ -66,7 +66,7 @@
</target>
<target name="jar" depends="build">
<jar jarfile="${build}/${jar}">
- <fileset dir="${build}" includes="org/**,x10/**,x10cpp/**,x10c/**,data/**" excludes="${jar}"/>
+ <fileset dir="${build}" includes="org/**,x10/**,x10cpp/**,x10c/**,x10cuda/**,data/**" excludes="${jar}"/>
<fileset dir="${x10.constraints.location}/classes" includes="x10/constraint/**" excludes="x10/constraint/test/**"/>
<fileset dir="${x10.common.location}/classes" includes="x10/**"/>
</jar>
@@ -83,7 +83,7 @@
</target>
<target name="build" depends="init,prereq-jars">
<echo message="${ant.project.name}: ${ant.file}"/>
- <javac destdir="${build}" source="1.5" target="1.5" debug="on" includes="x10/**,x10c/**,x10cpp/**," excludes="polyglot/ext/x10/dom/**,polyglot/ext/x10/plugin/**,polyglot/ext/x10/visit/Propagate*AnnotationsVisitor.java">
+ <javac destdir="${build}" source="1.5" target="1.5" debug="on" includes="x10/**,x10c/**,x10cpp/**,x10cuda/**" excludes="polyglot/ext/x10/dom/**,polyglot/ext/x10/plugin/**,polyglot/ext/x10/visit/Propagate*AnnotationsVisitor.java">
<src path="${src}"/>
<classpath refid="project.classpath"/>
</javac>
Deleted: trunk/x10.compiler/src/x10/types/X10ArrayType.java
===================================================================
--- trunk/x10.compiler/src/x10/types/X10ArrayType.java 2009-08-14 21:43:40 UTC (rev 11032)
+++ trunk/x10.compiler/src/x10/types/X10ArrayType.java 2009-08-14 23:09:29 UTC (rev 11033)
@@ -1,15 +0,0 @@
-/*
- *
- * (C) Copyright IBM Corporation 2006-2008.
- *
- * This file is part of X10 Language.
- *
- */
-
-package x10.types;
-
-import polyglot.types.ArrayType;
-
-public interface X10ArrayType extends X10NamedType, ArrayType {
-
-}
Deleted: trunk/x10.compiler/src/x10/types/X10ArrayType_c.java
===================================================================
--- trunk/x10.compiler/src/x10/types/X10ArrayType_c.java 2009-08-14 21:43:40 UTC (rev 11032)
+++ trunk/x10.compiler/src/x10/types/X10ArrayType_c.java 2009-08-14 23:09:29 UTC (rev 11033)
@@ -1,70 +0,0 @@
-/*
- *
- * (C) Copyright IBM Corporation 2006-2008.
- *
- * This file is part of X10 Language.
- *
- */
-
-package x10.types;
-
-import polyglot.types.ArrayType_c;
-import polyglot.types.Flags;
-import polyglot.types.Named;
-import polyglot.types.QName;
-import polyglot.types.Ref;
-import polyglot.types.Name;
-import polyglot.types.Type;
-import polyglot.types.TypeSystem;
-import polyglot.util.Position;
-
-public class X10ArrayType_c extends ArrayType_c implements X10ArrayType {
- public X10ArrayType_c(TypeSystem ts, Position pos, Ref<? extends Type> base) {
- super(ts, pos, base);
- }
-
- public boolean isGloballyAccessible() {
- return false;
- }
-
- public Name name() { return ((Named) base).name();}
- public QName fullName() { return ((Named) base).fullName();}
-
- public boolean safe() {
- return true;
- }
-
- public String toString() {
- StringBuffer sb = new StringBuffer();
- sb.append(super.toString());
- return sb.toString();
- }
- // TODO: vj check if this can have any flags in reality.
- // begin Flagged mixin
- Flags flags;
- public Flags flags() { return flags;}
-
- public X10Type setFlags(Flags f) {
- X10Flags xf = (X10Flags) f;
- X10ArrayType_c c = (X10ArrayType_c) this.copy();
- if (c.flags == null)
- c.flags = X10Flags.toX10Flags(Flags.NONE);
- c.flags = xf.isRooted()
- ? (xf.isStruct() ? ((X10Flags) c.flags).Rooted().Struct() : ((X10Flags) c.flags).Rooted())
- : ((xf.isStruct()) ? ((X10Flags) c.flags).Struct() : c.flags);
- return c;
- }
- public X10Type clearFlags(Flags f) {
- X10ArrayType_c c = (X10ArrayType_c) this.copy();
- if (c.flags == null)
- c.flags = X10Flags.toX10Flags(Flags.NONE);
- c.flags = c.flags.clear(f);
- return c;
- }
- public boolean isRooted() { return flags == null ? false : ((X10Flags) flags).isRooted(); }
- public boolean isX10Struct() { return flags == null ? false : ((X10Flags) flags).isStruct(); }
-
- public boolean equalsNoFlag(X10Type t2) {
- return this == t2;
- }
-}
Modified: trunk/x10.compiler/src/x10/types/X10TypeSystem.java
===================================================================
--- trunk/x10.compiler/src/x10/types/X10TypeSystem.java 2009-08-14 21:43:40 UTC (rev 11032)
+++ trunk/x10.compiler/src/x10/types/X10TypeSystem.java 2009-08-14 23:09:29 UTC (rev 11033)
@@ -345,5 +345,7 @@
boolean isSigned(Type l);
boolean numericConversionValid(Type toType, Type fromType, Object constantValue, Context context);
+
+ public Long size(Type t);
}
Modified: trunk/x10.compiler/src/x10/types/X10TypeSystem_c.java
===================================================================
--- trunk/x10.compiler/src/x10/types/X10TypeSystem_c.java 2009-08-14 21:43:40 UTC (rev 11032)
+++ trunk/x10.compiler/src/x10/types/X10TypeSystem_c.java 2009-08-14 23:09:29 UTC (rev 11033)
@@ -2478,5 +2478,18 @@
X10Context xc = (X10Context) context;
return env(context).isSubtypeWithValueInterfaces(t1, t2);
}
+
+ // Returns the number of bytes required to represent the type, or null if unknown (e.g. involves an address somehow)
+ public Long size(Type t) {
+ if (t.isFloat()) return 4l;
+ if (t.isDouble()) return 8l;
+ if (t.isChar()) return 2l;
+ if (t.isByte()) return 1l;
+ if (t.isShort()) return 2l;
+ if (t.isInt()) return 4l;
+ if (t.isLong()) return 8l;
+ // TODO: rails & valrails
+ return null;
+ }
}
Modified: trunk/x10.compiler/src/x10cpp/visit/ASTQuery.java
===================================================================
--- trunk/x10.compiler/src/x10cpp/visit/ASTQuery.java 2009-08-14 21:43:40 UTC (rev 11032)
+++ trunk/x10.compiler/src/x10cpp/visit/ASTQuery.java 2009-08-14 23:09:29 UTC (rev 11033)
@@ -1,6 +1,6 @@
package x10cpp.visit;
-import static x10cpp.visit.SharedVarsMethods.NATIVE_STRING;
+import static x10cpp.visit.SharedVarsMethods.CPP_NATIVE_STRING;
import java.util.ArrayList;
import java.util.Arrays;
@@ -296,7 +296,7 @@
for (Type at : as) {
assertNumberOfInitializers(at, 4);
String lang = getPropertyInit(at, 0);
- if (lang != null && lang.equals(NATIVE_STRING)) {
+ if (lang != null && lang.equals(CPP_NATIVE_STRING)) {
return getPropertyInit(at, i);
}
}
Modified: trunk/x10.compiler/src/x10cpp/visit/MessagePassingCodeGenerator.java
===================================================================
--- trunk/x10.compiler/src/x10cpp/visit/MessagePassingCodeGenerator.java 2009-08-14 21:43:40 UTC (rev 11032)
+++ trunk/x10.compiler/src/x10cpp/visit/MessagePassingCodeGenerator.java 2009-08-14 23:09:29 UTC (rev 11033)
@@ -26,7 +26,7 @@
import static x10cpp.visit.SharedVarsMethods.DESERIALIZE_METHOD;
import static x10cpp.visit.SharedVarsMethods.INSTANCE_INIT;
import static x10cpp.visit.SharedVarsMethods.MAKE;
-import static x10cpp.visit.SharedVarsMethods.NATIVE_STRING;
+import static x10cpp.visit.SharedVarsMethods.CPP_NATIVE_STRING;
import static x10cpp.visit.SharedVarsMethods.SAVED_THIS;
import static x10cpp.visit.SharedVarsMethods.SERIALIZATION_BUFFER;
import static x10cpp.visit.SharedVarsMethods.SERIALIZATION_ID_FIELD;
@@ -3491,20 +3491,26 @@
assert (false) : ("Function assign should have been desugared earlier");
}
+ // allow overriding in subclasses
+ // [DC] FIXME: ASTQuery.getCppRepParam still uses CPP_NATIVE_STRING directly
+ protected String[] getCurrentNativeStrings() { return new String[] {CPP_NATIVE_STRING}; }
- private static String getCppImplForDef(X10Def o) {
+ private String getCppImplForDef(X10Def o) {
X10TypeSystem xts = (X10TypeSystem) o.typeSystem();
try {
- Type java = (Type) xts.systemResolver().find(QName.make("x10.compiler.Native"));
- List<Type> as = o.annotationsMatching(java);
- for (Type at : as) {
- assertNumberOfInitializers(at, 2);
- String lang = getPropertyInit(at, 0);
- if (lang != null && lang.equals(NATIVE_STRING)) {
- String lit = getPropertyInit(at, 1);
- return lit;
- }
- }
+ Type annotation = (Type) xts.systemResolver().find(QName.make("x10.compiler.Native"));
+ String[] our_langs = getCurrentNativeStrings();
+ for (String our_lang : our_langs) {
+ List<Type> as = o.annotationsMatching(annotation);
+ for (Type at : as) {
+ assertNumberOfInitializers(at, 2);
+ String lang = getPropertyInit(at, 0);
+ if (lang != null && lang.equals(our_lang)) {
+ String lit = getPropertyInit(at, 1);
+ return lit;
+ }
+ }
+ }
}
catch (SemanticException e) {}
return null;
Modified: trunk/x10.compiler/src/x10cpp/visit/SharedVarsMethods.java
===================================================================
--- trunk/x10.compiler/src/x10cpp/visit/SharedVarsMethods.java 2009-08-14 21:43:40 UTC (rev 11032)
+++ trunk/x10.compiler/src/x10cpp/visit/SharedVarsMethods.java 2009-08-14 23:09:29 UTC (rev 11033)
@@ -63,8 +63,8 @@
static final boolean refsAsPointers = false;
static final String VOID = "void";
static final String VOID_PTR = "void*";
- static final String SAVED_THIS = "saved_this";
- static final String THIS = "this";
+ public static final String SAVED_THIS = "saved_this";
+ public static final String THIS = "this";
static final String INSTANCE_INIT = "_instance_init"; // instance field initialisers
static final String CONSTRUCTOR = "_constructor";
static final String MAKE = "_make";
@@ -84,7 +84,8 @@
static final String VIM_MODELINE = "vim:tabstop=4:shiftwidth=4:expandtab";
- static final String NATIVE_STRING = "c++";
+ public static final String CPP_NATIVE_STRING = "c++";
+ public static final String CUDA_NATIVE_STRING = "cuda";
public static String chevrons(String type) {
return "<" + type + (type.endsWith(">")?" ":"")+">";
Modified: trunk/x10.compiler/src/x10cuda/types/SharedMem.java
===================================================================
--- trunk/x10.compiler/src/x10cuda/types/SharedMem.java 2009-08-14 21:43:40 UTC (rev 11032)
+++ trunk/x10.compiler/src/x10cuda/types/SharedMem.java 2009-08-14 23:09:29 UTC (rev 11033)
@@ -5,6 +5,8 @@
import polyglot.ast.Expr;
import polyglot.ast.LocalDecl;
import polyglot.types.Name;
+import polyglot.types.Type;
+import x10.types.X10TypeSystem;
public class SharedMem {
@@ -12,6 +14,13 @@
private abstract static class Decl {
public final LocalDecl ast;
+ public long bytes() {
+ Type t = ast.type().type();
+ X10TypeSystem xts = (X10TypeSystem) t.typeSystem();
+ Long bytes = xts.size(t);
+ assert bytes != null : t;
+ return bytes.longValue();
+ }
public Decl (LocalDecl ast) { this.ast = ast; }
}
Modified: trunk/x10.compiler/src/x10cuda/visit/CUDACodeGenerator.java
===================================================================
--- trunk/x10.compiler/src/x10cuda/visit/CUDACodeGenerator.java 2009-08-14 21:43:40 UTC (rev 11032)
+++ trunk/x10.compiler/src/x10cuda/visit/CUDACodeGenerator.java 2009-08-14 23:09:29 UTC (rev 11033)
@@ -13,7 +13,10 @@
package x10cuda.visit;
-
+import static x10cpp.visit.SharedVarsMethods.CUDA_NATIVE_STRING;
+import static x10cpp.visit.SharedVarsMethods.CPP_NATIVE_STRING;
+import static x10cpp.visit.SharedVarsMethods.THIS;
+import static x10cpp.visit.SharedVarsMethods.SAVED_THIS;
import polyglot.ast.ArrayInit_c;
import polyglot.ast.Assert_c;
import polyglot.ast.Assign_c;
@@ -123,6 +126,10 @@
super(sw,tr);
}
+ protected String[] getCurrentNativeStrings() {
+ if (!generatingKernel()) return new String[] { CPP_NATIVE_STRING };
+ return new String[] { CUDA_NATIVE_STRING, CPP_NATIVE_STRING };
+ }
private X10CUDAContext_c context() {
return (X10CUDAContext_c) tr.context();
@@ -184,52 +191,58 @@
sw.write("/* block split-compiled to cuda as "+kernel_name+" */ ");
ClassifiedStream out = cudaStream();
- // disable name-mangling which seems to be inconsistent across cuda versions
+
+ // environment (passed into kernel via pointer)
out.write("struct "+kernel_name+"_env {"); out.newline(4); out.begin(0);
- emitter.printDeclarationList(out, context(), context().kernelParams());
- /*
+ //emitter.printDeclarationList(out, context(), context().kernelParams());
for (VarInstance var : context().kernelParams()) {
Type t = var.type();
- String type = emitter.translateType(t, true);
+ String type = Emitter.translateType(t, true);
if (isIntRail(t)) {
- type = "x10_int **";
+ type = "x10_int *";
} else if (isFloatRail(t)) {
- type = "x10_float **";
+ type = "x10_float *";
+ } else {
+ type = type + " ";
}
String name = var.name().toString();
if (name.equals(THIS)) {
name = SAVED_THIS;
+ } else {
+ name = Emitter.mangled_non_method_name(name);
}
-
- else {
- name = emitter.mangled_non_method_name(name);
- }
out.write(type + name + ";");
out.newline();
}
- */
out.end(); out.newline();
- out.write("};"); out.newline(); out.forceNewline();
- out.write("extern \"C\" __global__ void "+kernel_name+"("+kernel_name+"_env *"+env+")"); out.newline();
+ out.write("};"); out.newline();
- // decode buffer
- //out.write(" var = *(*T)buf; buf += sizeof(T);")
+ out.forceNewline();
+ // kernel (extern "C" to disable name-mangling which seems to be inconsistent across cuda versions)
+ out.write("extern \"C\" __global__ void "+kernel_name+"("+kernel_name+"_env *"+env+") {"); out.newline(4); out.begin(0);
+
+ // shm
+ out.write("// shm"); out.newline();
+ /*
+ for (context().shm()) {
+ out.write();
+ }
+ */
+/*
+ float *new_clusterv = (float*) dyn_shm; // [DIM*clusterc_odd]
+ int *new_counterv = (int*)&new_clusterv[DIM*clusterc_odd]; // [clusterc]
+*/
+
// body
- //X10CUDAContext_c save_ctxt = context();
-
- //TypeSystem ts = tr.typeSystem();
- //context(ts.emptyContext());
- //context((Context)save_ctxt.copy());
sw.pushCurrentStream(out);
super.visit(b);
sw.popCurrentStream();
- out.write(" // "+kernel_name);
- //context(save_ctxt);
-
// end
+ out.end(); out.newline();
+ out.write("} // "+kernel_name);
out.forceNewline();
}
Modified: trunk/x10.runtime/src-x10/x10/lang/Float.x10
===================================================================
--- trunk/x10.runtime/src-x10/x10/lang/Float.x10 2009-08-14 21:43:40 UTC (rev 11032)
+++ trunk/x10.runtime/src-x10/x10/lang/Float.x10 2009-08-14 23:09:29 UTC (rev 11033)
@@ -93,10 +93,12 @@
@Native("java", "java.lang.Float.MAX_VALUE")
@Native("c++", "x10aux::float_utils::fromIntBits(0x7f7fffff)")
+ @Native("cuda", "FLT_MAX")
public const MAX_VALUE: Float = Float.fromIntBits(0x7f7fffff);
@Native("java", "java.lang.Float.MIN_VALUE")
@Native("c++", "x10aux::float_utils::fromIntBits(0x00800000)")
+ @Native("cuda", "FLT_MIN")
public const MIN_VALUE: Float = Float.fromIntBits(0x00800000);
@Native("java", "java.lang.Float.toHexString(#0)")
Modified: trunk/x10.runtime/src-x10/x10/lang/Rail.x10
===================================================================
--- trunk/x10.runtime/src-x10/x10/lang/Rail.x10 2009-08-14 21:43:40 UTC (rev 11032)
+++ trunk/x10.runtime/src-x10/x10/lang/Rail.x10 2009-08-14 23:09:29 UTC (rev 11033)
@@ -50,10 +50,12 @@
@Native("java", "#0.apply(#1)")
@Native("c++", "(*#0)[#1]")
+ @Native("cuda", "(#0)[#1]")
public native safe def apply(i: nat): T;
@Native("java", "#0.set(#1, #2)")
@Native("c++", "(*#0)[#2] = #1")
+ @Native("cuda", "(#0)[#2] = #1")
public native safe def set(v: T, i: nat): T;
@Native("java", "#0.iterator()")
Modified: trunk/x10.runtime/src-x10/x10/lang/ValRail.x10
===================================================================
--- trunk/x10.runtime/src-x10/x10/lang/ValRail.x10 2009-08-14 21:43:40 UTC (rev 11032)
+++ trunk/x10.runtime/src-x10/x10/lang/ValRail.x10 2009-08-14 23:09:29 UTC (rev 11033)
@@ -37,6 +37,7 @@
@Native("java", "#0.apply(#1)")
@Native("c++", "(*#0)[#1]")
+ @Native("cuda", "(#0)[#1]")
public native def apply(i: nat): T;
@Native("java", "#0.iterator()")
This was sent by the SourceForge.net collaborative development platform, the world's largest Open Source development site.
|
