/*
* This file is part of the X10 project (http://x10-lang.org).
*
* This file is licensed to You under the Eclipse Public License (EPL);
* You may not use this file except in compliance with the License.
* You may obtain a copy of the License at
* http://www.opensource.org/licenses/eclipse-1.0.php
*
* (C) Copyright IBM Corporation 2006-2010.
*/
package x10cuda.types;
import java.util.ArrayList;
import java.util.Iterator;
import polyglot.ast.Expr;
import polyglot.ast.LocalDecl;
import polyglot.types.Name;
import polyglot.types.Type;
import polyglot.visit.Translator;
import polyglot.types.TypeSystem;
import x10.util.ClassifiedStream;
import x10.util.StreamWrapper;
public class ConstMem {
ArrayList<Decl> decls = new ArrayList<Decl>();
private abstract static class Decl {
public final LocalDecl ast;
public Decl (LocalDecl ast) { this.ast = ast; }
abstract public void generateDef(StreamWrapper out, String offset, Translator tr);
abstract public void generateInit(StreamWrapper out, String offset, Translator tr);
abstract public void generateSize(StreamWrapper inc, Translator tr);
}
private static class Rail extends Decl {
public final Expr numElements;
public final Expr init;
public Rail (LocalDecl ast, Expr numElements, Expr init) {
super(ast);
this.numElements = numElements;
this.init = init;
}
public void generateDef(StreamWrapper out, String offset, Translator tr) {
String name = ast.name().id().toString();
// FIXME: x10_float is baked in here
out.write("x10_float *"+name+" = (x10_float*) &__shm["+offset+"];"); out.newline();
}
public void generateInit(StreamWrapper out, String offset, Translator tr) {
out.write("{"); out.newline(4); out.begin(0);
out.write("x10_int __len = ");
tr.print(null, numElements, out);
out.write(";"); out.newline();
out.write("for (int i=0 ; i<__len ; ++i) {"); out.newline(4); out.begin(0);
// TODO: assumes rail initialised with another rail -- closure version also possible
out.write(ast.name().id()+"[i] = ");
tr.print(null, init, out);
out.write("[i];"); out.newline();
out.end(); out.newline();
out.write("}");
out.end(); out.newline();
out.write("}");
}
public void generateSize(StreamWrapper inc, Translator tr) {
tr.print(null, numElements, inc);
// FIXME: x10_float is baked in here
inc.write("*sizeof(x10_float)");
}
}
private static class Var extends Decl {
public Var (LocalDecl ast) { super(ast); }
public void generateDef(StreamWrapper out, String offset, Translator tr) {
// TODO: not implemented
assert false: "not implemented";
}
public void generateInit(StreamWrapper out, String offset, Translator tr) {
// TODO: not implemented
assert false: "not implemented";
}
public void generateSize(StreamWrapper inc, Translator tr) {
// TODO: not implemented
assert false: "not implemented";
}
}
public void addRail(LocalDecl ast, Expr numElements, Expr init) {
decls.add(new Rail(ast,numElements,init));
}
public void addVar(LocalDecl ast) {
decls.add(new Var(ast));
}
public boolean has(Name n) {
for (Decl d : decls) {
if (d.ast.name().id() == n) {
return true;
}
}
return false;
}
public void generateCode(StreamWrapper out, Translator tr) {
if (decls.size()==0) return;
out.write("// shm");
out.newline();
for (ConstMem.Decl d : decls) {
// FIXME: offset is broken when more than one shm definition
String offset = "0";
d.generateDef(out, offset, tr);
out.write("if (threadIdx.x == 0) {"); out.newline(4); out.begin(0);
d.generateInit(out, offset, tr);
out.end(); out.newline();
out.write("}"); out.newline();
}
out.write("__syncthreads();"); out.newline();
out.forceNewline();
}
public void generateSize(StreamWrapper inc, Translator tr) {
// TODO Auto-generated method stub
String prefix = "";
for (ConstMem.Decl d : decls) {
inc.write(prefix);
d.generateSize(inc, tr);
prefix = " + ";
}
if (prefix.equals("")) inc.write("0");
}
}