# HG changeset patch
# User Doug Simon <doug.simon@oracle.com>
# Date 1381138609 -7200
# Node ID 43bf803203c0f659f1b83fb1ee515e6a93d7f693
# Parent  234aea2460bbca8e0517f6a72c7e539275a48801# Parent  cf4dd10ced32f18370af0627063ce733f946e3e2
Merge.

diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java
--- a/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java	Mon Oct 07 11:36:49 2013 +0200
@@ -22,6 +22,7 @@
  */
 package com.oracle.graal.asm.ptx;
 
+import static com.oracle.graal.asm.ptx.PTXStateSpace.*;
 import static com.oracle.graal.api.code.ValueUtil.*;
 
 import com.oracle.graal.asm.Label;
@@ -521,8 +522,8 @@
         emitString("@%p" + pred + " " + "bra" + " " + tgt + ";");
     }
 
-    public final void bra(String target) {
-        emitString("bra " + target + ";");
+    public final void bra(String src) {
+        emitString("bra " + src + ";");
     }
 
     public final void bra_uni(String tgt) {
@@ -536,7 +537,13 @@
         }
 
         public void emit(PTXAssembler asm) {
-            asm.emitString("cvt." + super.emit());
+            if (dest.getKind() == Kind.Float ||
+                dest.getKind() == Kind.Double) {
+                // round-to-zero - might not be right
+                asm.emitString("cvt.rz." + super.emit());
+            } else {
+                asm.emitString("cvt." + super.emit());
+            }
         }
     }
     
@@ -628,7 +635,7 @@
             this.targets = targets;
         }
 
-        private String valueForKind(Kind k) {
+        private static String valueForKind(Kind k) {
             switch (k.getTypeChar()) {
                 case 'i':
                     return "s32";
@@ -639,7 +646,7 @@
             }
         }
 
-        private String emitTargets(PTXAssembler asm, LabelRef[] refs) {
+        private static String emitTargets(PTXAssembler asm, LabelRef[] refs) {
             StringBuffer sb = new StringBuffer();
 
             for (int i = 0; i < refs.length; i++) {
@@ -722,14 +729,34 @@
         emitString("ret.uni;" + " " + "");
     }
 
+    public enum BooleanOperator {
+        AND("and"),
+        OR("or"),
+        XOR("xor");
+
+        private final String output;
+
+        private BooleanOperator(String out) {
+            this.output = out;
+        }
+
+        public String getOperator() {
+            return output + ".";
+        }
+    }
+
     public static class Setp  {
 
+
+        private BooleanOperator    booleanOperator;
         private ConditionOperator  operator;
         private Value first, second;
         private Kind kind;
         private int predicate;
 
-        public Setp(Condition condition, Value first, Value second, int predicateRegisterNumber) {
+        public Setp(Condition condition,
+                    Value first, Value second,
+                    int predicateRegisterNumber) {
             setFirst(first);
             setSecond(second);
             setPredicate(predicateRegisterNumber);
@@ -737,6 +764,17 @@
             setConditionOperator(operatorForConditon(condition));
         }
 
+        public Setp(Condition condition, BooleanOperator operator,
+                    Value first, Value second,
+                    int predicateRegisterNumber) {
+            setFirst(first);
+            setSecond(second);
+            setPredicate(predicateRegisterNumber);
+            setKind();
+            setConditionOperator(operatorForConditon(condition));
+            setBooleanOperator(operator);
+        }
+
         public void setFirst(Value v) {
             first = v;
         }
@@ -753,6 +791,10 @@
             operator = co;
         }
 
+        public void setBooleanOperator(BooleanOperator bo) {
+            booleanOperator = bo;
+        }
+
         private ConditionOperator operatorForConditon(Condition condition) {
             char typeChar = kind.getTypeChar();
 
@@ -873,8 +915,21 @@
         }
 
         public void emit(PTXAssembler asm) {
-            asm.emitString("setp." + operator.getOperator() + "." + typeForKind(kind) +
-                           " %p" + predicate + emitValue(first) + emitValue(second) + ";");
+
+            if (booleanOperator != null) {
+                asm.emitString("setp." +
+                               operator.getOperator() + "." +
+                               booleanOperator.getOperator() +
+                               typeForKind(kind) + " %p" + predicate +
+                               emitValue(first) + emitValue(second) +
+                               ", %r;"); // Predicates need to be objects
+
+            } else {
+                asm.emitString("setp." +
+                               operator.getOperator() + "." +
+                               typeForKind(kind) + " %p" + predicate +
+                               emitValue(first) + emitValue(second) + ";");
+            }
         }
     }
     @Override
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ControlPTXTest.java
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ControlPTXTest.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ControlPTXTest.java	Mon Oct 07 11:36:49 2013 +0200
@@ -60,11 +60,24 @@
         } else {
             printReport("testIfElse2I: no VALUE");
         }
+        Boolean bret = (Boolean) invoke(compile("testIntegerTestBranch2I"),
+                                        0xff00, 0x00ff);
+        if (bret != null) {
+            printReport("testIntegerTestBranch2I: " + bret);
+            printReport("testIntegerTestBranch2I: actual: " +
+                                testIntegerTestBranch2I(0xff00, 0x00ff));
+        } else {
+            printReport("testIntegerTestBranch2I: no VALUE");
+        }
         compile("testStatic");
         compile("testCall");
         compile("testLookupSwitch1I");
     }
 
+    public static boolean testIntegerTestBranch2I(int x, int y) {
+        return (x & y) == 0;
+    }
+
     public static int testLoop(int n) {
         int sum = 0;
 
@@ -84,11 +97,11 @@
 
     public static int testIfElse2I(int c, int y) {
         if  (c > 19) {
-            return (int) 'M';    // millenial
+            return 'M';    // millenial
         } else if (y > 84) {
-            return (int) 'Y';    // young
+            return 'Y';    // young
         } else {
-            return (int) 'O';    // old
+            return 'O';    // old
         }
     }
 
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/FloatPTXTest.java
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/FloatPTXTest.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/FloatPTXTest.java	Mon Oct 07 11:36:49 2013 +0200
@@ -31,41 +31,44 @@
 /* PTX ISA 3.1 - 8.7.3 Floating-Point Instructions */
 public class FloatPTXTest extends PTXTestBase {
 
-    @Ignore
     @Test
     public void testAdd() {
-        CompilationResult r = compile("testAdd2F");
-        if (r.getTargetCode() == null) {
-            printReport("Compilation of testAdd2F FAILED");
+        Float ret = (Float) invoke(compile("testAdd2I"), 42, 43);
+        if (ret != null) {
+            printReport("testAdd2I: " + ret);
+        } else {
+            printReport("testAdd2I: no VALUE");
         }
 
-        /*
-        r = compile("testAdd2D");
-        if (r.getTargetCode() == null) {
-            printReport("Compilation of testAdd2D FAILED");
+        ret = (Float) invoke(compile("testAdd2F"), 42.1F, 43.5F);
+        if (ret != null) {
+            printReport("testAdd2F: " + ret);
+        } else {
+            printReport("testAdd2F: no VALUE");
         }
 
-        r = compile("testAddFConst");
-        if (r.getTargetCode() == null) {
-            printReport("Compilation of testAddFConst FAILED");
-        }
-        r = compile("testAddConstF");
-        if (r.getTargetCode() == null) {
-            printReport("Compilation of testConstF FAILED");
+        ret = (Float) invoke(compile("testAddFConst"), 42.1F);
+        if (ret != null) {
+            printReport("testAddFConst: " + ret);
+        } else {
+            printReport("testAddFConst: no VALUE");
         }
-        r = compile("testAddDConst");
-        if (r.getTargetCode() == null) {
-            printReport("Compilation of testAddDConst FAILED");
+
+        Double dret = (Double) invoke(compile("testAdd2D"), 42.1, 43.5);
+        if (dret != null) {
+            printReport("testAdd2D: " + dret);
+        } else {
+            printReport("testAdd2D: no VALUE");
         }
-        r = compile("testAddConstD");
-        if (r.getTargetCode() == null) {
-            printReport("Compilation of testConstD FAILED");
-        }
-        */
+
+    }
+
+    public static float testAdd2I(int a, int b) {
+        return (float) (a + b);
     }
 
     public static float testAdd2F(float a, float b) {
-        return a + b;
+        return (a + b);
     }
 
     public static double testAdd2D(double a, double b) {
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java	Mon Oct 07 11:36:49 2013 +0200
@@ -85,7 +85,8 @@
     }
 
     @Override
-    public TargetMethodAssembler newAssembler(LIRGenerator lirGen, CompilationResult compilationResult) {
+    public TargetMethodAssembler newAssembler(LIRGenerator lirGen,
+                                              CompilationResult compilationResult) {
         // Omit the frame of the method:
         // - has no spill slots or other slots allocated during register allocation
         // - has no callee-saved registers
@@ -99,7 +100,9 @@
         return tasm;
     }
 
-    private static void emitKernelEntry(TargetMethodAssembler tasm, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) {
+    private static void emitKernelEntry(TargetMethodAssembler tasm,
+                                        LIRGenerator lirGen,
+                                        ResolvedJavaMethod codeCacheOwner) {
         // Emit PTX kernel entry text based on PTXParameterOp
         // instructions in the start block. Remove the instructions
         // once kernel entry text and directives are emitted to
@@ -109,8 +112,8 @@
         Buffer codeBuffer = tasm.asm.codeBuffer;
 
         // Emit initial boiler-plate directives.
-        codeBuffer.emitString(".version 2.1");
-        codeBuffer.emitString(".target sm_20");
+        codeBuffer.emitString(".version 3.0");
+        codeBuffer.emitString(".target sm_30");
         codeBuffer.emitString0(".entry " + name + " (");
         codeBuffer.emitString("");
 
@@ -140,9 +143,13 @@
     }
 
     // Emit .reg space declarations
-    private static void emitRegisterDecl(TargetMethodAssembler tasm, LIRGenerator lirGen,
+    private static void emitRegisterDecl(TargetMethodAssembler tasm,
+                                         LIRGenerator lirGen,
                                          ResolvedJavaMethod codeCacheOwner) {
-        assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
+
+        assert codeCacheOwner != null :
+               lirGen.getGraph() + " is not associated with a method";
+
         Buffer codeBuffer = tasm.asm.codeBuffer;
 
         final SortedSet<Integer> signed32 = new TreeSet<>();
@@ -231,6 +238,7 @@
         if (maxPredRegNum > 0) {
             codeBuffer.emitString(".reg .pred %p<" + maxPredRegNum + ">;");
         }
+        codeBuffer.emitString(".reg .pred %r;");  // used for setp bool
     }
 
     @Override
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java	Mon Oct 07 11:36:49 2013 +0200
@@ -333,7 +333,9 @@
 
     @Override
     public void emitIntegerTestBranch(Value left, Value right, boolean negated, LabelRef label) {
-        throw GraalInternalError.unimplemented("PTXLIRGenerator.emitIntegerTestBranch()");
+        /// emitIntegerTest(left, right);
+       //  append(new BranchOp(negated ? Condition.NE : Condition.EQ, label));
+        throw GraalInternalError.unimplemented("emitIntegerTestBranch()");
     }
 
     @Override
@@ -343,8 +345,7 @@
 
         Condition finalCondition = LIRValueUtil.isVariable(right) ? cond.mirror() : cond;
 
-        boolean mirrored;
-        mirrored = emitCompare(finalCondition, left, right);
+        emitCompare(finalCondition, left, right);
 
         Variable result = newVariable(trueValue.getKind());
         switch (left.getKind().getStackKind()) {
@@ -414,8 +415,30 @@
 
 
     @Override
-    public Variable emitIntegerTestMove(Value left, Value right, Value trueValue, Value falseValue) {
-        throw new InternalError("NYI");
+    public Variable emitIntegerTestMove(Value left, Value right,
+                                        Value trueValue, Value falseValue) {
+
+        emitIntegerTest(left, right);
+        Variable result = newVariable(trueValue.getKind());
+        append(new CondMoveOp(result, Condition.EQ,
+                              load(trueValue), loadNonConst(falseValue),
+                              nextPredRegNum));
+        nextPredRegNum++;
+
+        return result;
+    }
+
+
+    private void emitIntegerTest(Value a, Value b) {
+
+        assert a.getKind().getStackKind() == Kind.Int ||
+               a.getKind() == Kind.Long;
+
+        if (LIRValueUtil.isVariable(b)) {
+            append(new PTXTestOp(load(b), loadNonConst(a), nextPredRegNum));
+        } else {
+            append(new PTXTestOp(load(a), loadNonConst(b), nextPredRegNum));
+        }
     }
 
     @Override
@@ -889,24 +912,46 @@
         throw GraalInternalError.unimplemented("PTXLIRGenerator.visitInfopointNode()");
     }
 
-    public Variable emitLoadParam(Kind kind, Value address, DeoptimizingNode deopting) {
+    public Variable emitLoadParam(Kind kind, Value address,
+                                  DeoptimizingNode deopting) {
+
         PTXAddressValue loadAddress = asAddress(address);
         Variable result = newVariable(kind);
-        append(new LoadParamOp(kind, result, loadAddress, deopting != null ? state(deopting) : null));
+        append(new LoadParamOp(kind, result, loadAddress,
+                               deopting != null ? state(deopting) : null));
+
         return result;
     }
 
-    public Variable emitLoadReturnAddress(Kind kind, Value address, DeoptimizingNode deopting) {
+    public Variable emitLoadReturnAddress(Kind kind, Value address,
+                                          DeoptimizingNode deopting) {
+
         PTXAddressValue loadAddress = asAddress(address);
-        Variable result = newVariable(kind);
-        append(new LoadReturnAddrOp(kind, result, loadAddress, deopting != null ? state(deopting) : null));
+        Variable result;
+        switch (kind) {
+            case Float:
+                result = newVariable(Kind.Int);
+                break;
+            case Double:
+                result = newVariable(Kind.Long);
+                break;
+            default:
+                result = newVariable(kind);
+
+        }
+        append(new LoadReturnAddrOp(kind, result, loadAddress,
+                                    deopting != null ? state(deopting) : null));
+
         return result;
     }
 
-    public void emitStoreReturnValue(Kind kind, Value address, Value inputVal, DeoptimizingNode deopting) {
+    public void emitStoreReturnValue(Kind kind, Value address, Value inputVal,
+                                     DeoptimizingNode deopting) {
+
         PTXAddressValue storeAddress = asAddress(address);
         Variable input = load(inputVal);
-        append(new StoreReturnValOp(kind, storeAddress, input, deopting != null ? state(deopting) : null));
+        append(new StoreReturnValOp(kind, storeAddress, input,
+                                    deopting != null ? state(deopting) : null));
     }
 
     @Override
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXTargetMethodAssembler.java
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXTargetMethodAssembler.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXTargetMethodAssembler.java	Mon Oct 07 11:36:49 2013 +0200
@@ -33,10 +33,13 @@
 
 public class PTXTargetMethodAssembler extends TargetMethodAssembler {
 
-    private static CompilerToGPU toGPU = HotSpotGraalRuntime.graalRuntime().getCompilerToGPU();
+    private static CompilerToGPU toGPU =
+                    HotSpotGraalRuntime.graalRuntime().getCompilerToGPU();
+
     private static boolean validDevice = toGPU.deviceInit();
 
-    private static final int totalProcessors = (validDevice ? toGPU.availableProcessors() : 0);
+    private static final int totalProcessors =
+                            (validDevice ? toGPU.availableProcessors() : 0);
 
     public static int getAvailableProcessors() {
         return totalProcessors;
@@ -44,8 +47,12 @@
 
     // detach ??
 
-    public PTXTargetMethodAssembler(TargetDescription target, CodeCacheProvider runtime, FrameMap frameMap,
-                                    AbstractAssembler asm, FrameContext frameContext, CompilationResult compilationResult) {
+    public PTXTargetMethodAssembler(TargetDescription target,
+                                    CodeCacheProvider runtime,
+                                    FrameMap frameMap,
+                                    AbstractAssembler asm,
+                                    FrameContext frameContext,
+                                    CompilationResult compilationResult) {
         super(target, runtime, frameMap, asm, frameContext, compilationResult);
     }
 
@@ -53,11 +60,14 @@
     public CompilationResult finishTargetMethod(StructuredGraph graph) {
         ResolvedJavaMethod method = graph.method();
         assert method != null : graph + " is not associated wth a method";
-        ExternalCompilationResult graalCompile = (ExternalCompilationResult) super.finishTargetMethod(graph);
+
+        ExternalCompilationResult graalCompile =
+            (ExternalCompilationResult) super.finishTargetMethod(graph);
 
         try {
             if ((validDevice) && (graalCompile.getTargetCode() != null)) {
-                long kernel = toGPU.generateKernel(graalCompile.getTargetCode(), method.getName());
+                long kernel = toGPU.generateKernel(graalCompile.getTargetCode(),
+                                                   method.getName());
                 graalCompile.setEntryPoint(kernel);
             }
         } catch (Throwable th) {
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotRegisterConfig.java
--- a/graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotRegisterConfig.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotRegisterConfig.java	Mon Oct 07 11:36:49 2013 +0200
@@ -104,7 +104,14 @@
         int currentStackOffset = 0;
 
         Kind returnKind = returnType == null ? Kind.Void : returnType.getKind();
-        AllocatableValue returnLocation = returnKind == Kind.Void ? Value.ILLEGAL : new Variable(returnKind, currentGeneral++);
+
+        AllocatableValue returnLocation;
+        if (returnKind == Kind.Void) {
+            returnLocation = Value.ILLEGAL;
+        } else {
+            returnLocation = new Variable(returnKind, currentGeneral++);
+        }
+
         AllocatableValue[] locations = new AllocatableValue[parameterTypes.length];
 
         for (int i = 0; i < parameterTypes.length; i++) {
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java	Mon Oct 07 11:36:49 2013 +0200
@@ -22,7 +22,6 @@
  */
 package com.oracle.graal.lir.ptx;
 
-import static com.oracle.graal.api.code.ValueUtil.*;
 import static com.oracle.graal.asm.ptx.PTXAssembler.*;
 import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*;
 import static com.oracle.graal.lir.LIRValueUtil.*;
@@ -161,7 +160,7 @@
         assert !result.equals(trueValue);
 
         PTXMove.move(tasm, asm, result, falseValue);
-        cmove(tasm, asm, result, condition, trueValue, predicateRegister);
+        cmove(asm, result, trueValue, predicateRegister);
 
         if (isFloat) {
             if (unorderedIsTrue && !trueOnUnordered(condition)) {
@@ -187,8 +186,8 @@
         }
     }
 
-    private static void cmove(TargetMethodAssembler tasm, PTXAssembler asm,
-                              Value result, Condition cond, Value other,
+    private static void cmove(PTXAssembler asm,
+                              Value result, Value other,
                               int predicateRegister) {
         if (isVariable(other)) {
             assert !asVariable(other).equals(asVariable(result)) :
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMemOp.java
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMemOp.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMemOp.java	Mon Oct 07 11:36:49 2013 +0200
@@ -43,7 +43,8 @@
         @Use({COMPOSITE}) protected PTXAddressValue address;
         @State protected LIRFrameState state;
 
-        public LoadOp(Kind kind, Variable result, PTXAddressValue address, LIRFrameState state) {
+        public LoadOp(Kind kind, Variable result, PTXAddressValue address,
+                      LIRFrameState state) {
             this.kind = kind;
             this.result = result;
             this.address = address;
@@ -62,7 +63,8 @@
                 case Float:
                 case Double:
                 case Object:
-                    new Ld(Global, result, addr.getBase(), Constant.forLong(addr.getDisplacement())).emit(masm);
+                    new Ld(Global, result, addr.getBase(),
+                           Constant.forLong(addr.getDisplacement())).emit(masm);
                     break;
                 default:
                     throw GraalInternalError.shouldNotReachHere();
@@ -97,7 +99,8 @@
                 case Float:
                 case Double:
                 case Object:
-                    new St(Global, input, addr.getBase(), Constant.forLong(addr.getDisplacement())).emit(masm);
+                    new St(Global, input, addr.getBase(),
+                           Constant.forLong(addr.getDisplacement())).emit(masm);
                     break;
                 default:
                     throw GraalInternalError.shouldNotReachHere("missing: " + address.getKind());
@@ -114,7 +117,8 @@
         @Use({COMPOSITE}) protected PTXAddressValue address;
         @State protected LIRFrameState state;
 
-        public LoadParamOp(Kind kind, Variable result, PTXAddressValue address, LIRFrameState state) {
+        public LoadParamOp(Kind kind, Variable result, PTXAddressValue address,
+                           LIRFrameState state) {
             this.kind = kind;
             this.result = result;
             this.address = address;
@@ -133,7 +137,8 @@
                 case Float:
                 case Double:
                 case Object:
-                    new Ld(Parameter, result, addr.getBase(), Constant.forLong(addr.getDisplacement())).emit(masm);
+                    new Ld(Parameter, result, addr.getBase(),
+                           Constant.forLong(addr.getDisplacement())).emit(masm);
                     break;
                 default:
                     throw GraalInternalError.shouldNotReachHere();
@@ -151,7 +156,8 @@
         @Use({COMPOSITE}) protected PTXAddressValue address;
         @State protected LIRFrameState state;
 
-        public LoadReturnAddrOp(Kind kind, Variable result, PTXAddressValue address, LIRFrameState state) {
+        public LoadReturnAddrOp(Kind kind, Variable result,
+                                PTXAddressValue address, LIRFrameState state) {
             this.kind = kind;
             this.result = result;
             this.address = address;
@@ -166,7 +172,8 @@
                 case Long:
                 case Float:
                 case Double:
-                    new Ld(Parameter, result, addr.getBase(), Constant.forLong(addr.getDisplacement())).emit(masm);
+                    new Ld(Parameter, result, addr.getBase(),
+                           Constant.forLong(addr.getDisplacement())).emit(masm);
                     break;
                 default:
                     throw GraalInternalError.shouldNotReachHere();
@@ -183,7 +190,8 @@
         @Use({REG}) protected Variable input;
         @State protected LIRFrameState state;
 
-        public StoreReturnValOp(Kind kind, PTXAddressValue address, Variable input, LIRFrameState state) {
+        public StoreReturnValOp(Kind kind, PTXAddressValue address,
+                                Variable input, LIRFrameState state) {
             this.kind = kind;
             this.address = address;
             this.input = input;
@@ -202,7 +210,8 @@
                 case Float:
                 case Double:
                 case Object:
-                    new St(Global, input, addr.getBase(), Constant.forLong(addr.getDisplacement())).emit(masm);
+                    new St(Global, input, addr.getBase(),
+                           Constant.forLong(addr.getDisplacement())).emit(masm);
                     break;
                 default:
                     throw GraalInternalError.shouldNotReachHere("missing: " + address.getKind());
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXTestOp.java
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXTestOp.java	Mon Oct 07 11:36:49 2013 +0200
@@ -0,0 +1,70 @@
+/*
+ * Copyright (c) 2011, 2012, Oracle and/or its affiliates. All rights reserved.
+ * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
+ *
+ * This code is free software; you can redistribute it and/or modify it
+ * under the terms of the GNU General Public License version 2 only, as
+ * published by the Free Software Foundation.
+ *
+ * This code is distributed in the hope that it will be useful, but WITHOUT
+ * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+ * FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+ * version 2 for more details (a copy is included in the LICENSE file that
+ * accompanied this code).
+ *
+ * You should have received a copy of the GNU General Public License version
+ * 2 along with this work; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
+ *
+ * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
+ * or visit www.oracle.com if you need additional information or have any
+ * questions.
+ */
+package com.oracle.graal.lir.ptx;
+
+import static com.oracle.graal.asm.ptx.PTXAssembler.BooleanOperator.*;
+import static com.oracle.graal.asm.ptx.PTXAssembler.*;
+import static com.oracle.graal.api.code.ValueUtil.*;
+import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*;
+
+import com.oracle.graal.api.meta.*;
+import com.oracle.graal.asm.ptx.*;
+import com.oracle.graal.graph.*;
+import com.oracle.graal.lir.asm.*;
+import com.oracle.graal.nodes.calc.Condition;
+
+public class PTXTestOp extends PTXLIRInstruction {
+
+    @Use({REG}) protected Value x;
+    @Use({REG, STACK, CONST}) protected Value y;
+    int predicate;
+
+    public PTXTestOp(Value x, Value y, int predicate) {
+        this.x = x;
+        this.y = y;
+        this.predicate = predicate;
+    }
+
+    @Override
+    public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) {
+        emit(tasm, masm, x, y, predicate);
+    }
+
+    @Override
+    protected void verify() {
+        super.verify();
+        assert (x.getKind() == Kind.Int &&
+                y.getKind().getStackKind() == Kind.Int) ||
+                (x.getKind() == Kind.Long && y.getKind() == Kind.Long) :
+                x + " " + y;
+    }
+
+    public static void emit(TargetMethodAssembler tasm, PTXAssembler masm,
+                            Value x, Value y, int predicate) {
+        /*
+         * This is not yet quite right - as the result for the equivalent in
+         * ControlPTXText.testIntegerTestBranch2I is wrong.
+         */
+        new Setp(Condition.EQ, AND, x, y, predicate).emit(masm);
+    }
+}
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.graal.truffle.test/src/com/oracle/graal/truffle/test/nodes/AbstractTestNode.java
--- a/graal/com.oracle.graal.truffle.test/src/com/oracle/graal/truffle/test/nodes/AbstractTestNode.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.graal.truffle.test/src/com/oracle/graal/truffle/test/nodes/AbstractTestNode.java	Mon Oct 07 11:36:49 2013 +0200
@@ -27,5 +27,9 @@
 
 public abstract class AbstractTestNode extends Node {
 
+    protected AbstractTestNode() {
+        super(null);
+    }
+
     public abstract int execute(VirtualFrame frame);
 }
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.graal.truffle.test/src/com/oracle/graal/truffle/test/nodes/RootTestNode.java
--- a/graal/com.oracle.graal.truffle.test/src/com/oracle/graal/truffle/test/nodes/RootTestNode.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.graal.truffle.test/src/com/oracle/graal/truffle/test/nodes/RootTestNode.java	Mon Oct 07 11:36:49 2013 +0200
@@ -31,6 +31,7 @@
     @Child AbstractTestNode node;
 
     public RootTestNode(String name, AbstractTestNode node) {
+        super(null);
         this.name = name;
         this.node = node;
     }
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.truffle.api.dsl.test/src/com/oracle/truffle/api/dsl/test/TypeSystemTest.java
--- a/graal/com.oracle.truffle.api.dsl.test/src/com/oracle/truffle/api/dsl/test/TypeSystemTest.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.truffle.api.dsl.test/src/com/oracle/truffle/api/dsl/test/TypeSystemTest.java	Mon Oct 07 11:36:49 2013 +0200
@@ -58,6 +58,10 @@
     @TypeSystemReference(SimpleTypes.class)
     public abstract static class ValueNode extends Node {
 
+        public ValueNode() {
+            super(null);
+        }
+
         public int executeInt(VirtualFrame frame) throws UnexpectedResultException {
             return SimpleTypesGen.SIMPLETYPES.expectInteger(execute(frame));
         }
@@ -105,6 +109,7 @@
         @Child private E node;
 
         public TestRootNode(E node) {
+            super(null);
             this.node = adoptChild(node);
         }
 
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/ArgumentsTest.java
--- a/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/ArgumentsTest.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/ArgumentsTest.java	Mon Oct 07 11:36:49 2013 +0200
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2012, Oracle and/or its affiliates. All rights reserved.
+ * Copyright (c) 2012, 2013, Oracle and/or its affiliates. All rights reserved.
  * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
  *
  * This code is free software; you can redistribute it and/or modify it
@@ -75,6 +75,7 @@
         @Children private TestArgumentNode[] children;
 
         TestRootNode(TestArgumentNode[] children) {
+            super(null);
             this.children = adoptChildren(children);
         }
 
@@ -93,6 +94,7 @@
         private final int index;
 
         TestArgumentNode(int index) {
+            super(null);
             this.index = index;
         }
 
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/CallTest.java
--- a/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/CallTest.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/CallTest.java	Mon Oct 07 11:36:49 2013 +0200
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2012, Oracle and/or its affiliates. All rights reserved.
+ * Copyright (c) 2012, 2013, Oracle and/or its affiliates. All rights reserved.
  * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
  *
  * This code is free software; you can redistribute it and/or modify it
@@ -60,6 +60,7 @@
         private final CallTarget secondTarget;
 
         DualCallNode(CallTarget firstTarget, CallTarget secondTarget) {
+            super(null);
             this.firstTarget = firstTarget;
             this.secondTarget = secondTarget;
         }
@@ -75,6 +76,7 @@
         private final int value;
 
         public ConstantRootNode(int value) {
+            super(null);
             this.value = value;
         }
 
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/ChildNodeTest.java
--- a/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/ChildNodeTest.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/ChildNodeTest.java	Mon Oct 07 11:36:49 2013 +0200
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2012, Oracle and/or its affiliates. All rights reserved.
+ * Copyright (c) 2012, 2013, Oracle and/or its affiliates. All rights reserved.
  * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
  *
  * This code is free software; you can redistribute it and/or modify it
@@ -78,6 +78,7 @@
         @Child private TestChildNode right;
 
         public TestRootNode(TestChildNode left, TestChildNode right) {
+            super(null);
             this.left = adoptChild(left);
             this.right = adoptChild(right);
         }
@@ -90,6 +91,10 @@
 
     class TestChildNode extends Node {
 
+        public TestChildNode() {
+            super(null);
+        }
+
         public int execute() {
             return 21;
         }
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/ChildrenNodesTest.java
--- a/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/ChildrenNodesTest.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/ChildrenNodesTest.java	Mon Oct 07 11:36:49 2013 +0200
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2012, Oracle and/or its affiliates. All rights reserved.
+ * Copyright (c) 2012, 2013, Oracle and/or its affiliates. All rights reserved.
  * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
  *
  * This code is free software; you can redistribute it and/or modify it
@@ -71,6 +71,7 @@
         @Children private final TestChildNode[] children;
 
         public TestRootNode(TestChildNode[] children) {
+            super(null);
             this.children = adoptChildren(children);
         }
 
@@ -86,6 +87,10 @@
 
     class TestChildNode extends Node {
 
+        public TestChildNode() {
+            super(null);
+        }
+
         public int execute() {
             return 21;
         }
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/FinalFieldTest.java
--- a/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/FinalFieldTest.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/FinalFieldTest.java	Mon Oct 07 11:36:49 2013 +0200
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2012, Oracle and/or its affiliates. All rights reserved.
+ * Copyright (c) 2012, 2013, Oracle and/or its affiliates. All rights reserved.
  * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
  *
  * This code is free software; you can redistribute it and/or modify it
@@ -65,6 +65,7 @@
         @Children TestChildNode[] children;
 
         public TestRootNode(TestChildNode[] children) {
+            super(null);
             this.children = adoptChildren(children);
         }
 
@@ -83,6 +84,7 @@
         private final int value;
 
         public TestChildNode(int value) {
+            super(null);
             this.value = value;
         }
 
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/FrameSlotTypeSpecializationTest.java
--- a/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/FrameSlotTypeSpecializationTest.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/FrameSlotTypeSpecializationTest.java	Mon Oct 07 11:36:49 2013 +0200
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2012, Oracle and/or its affiliates. All rights reserved.
+ * Copyright (c) 2012, 2013, Oracle and/or its affiliates. All rights reserved.
  * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
  *
  * This code is free software; you can redistribute it and/or modify it
@@ -63,6 +63,7 @@
         @Child TestChildNode right;
 
         public TestRootNode(TestChildNode left, TestChildNode right) {
+            super(null);
             this.left = adoptChild(left);
             this.right = adoptChild(right);
         }
@@ -76,6 +77,10 @@
 
     abstract class TestChildNode extends Node {
 
+        protected TestChildNode() {
+            super(null);
+        }
+
         abstract Object execute(VirtualFrame frame);
     }
 
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/FrameTest.java
--- a/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/FrameTest.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/FrameTest.java	Mon Oct 07 11:36:49 2013 +0200
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2012, Oracle and/or its affiliates. All rights reserved.
+ * Copyright (c) 2012, 2013, Oracle and/or its affiliates. All rights reserved.
  * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
  *
  * This code is free software; you can redistribute it and/or modify it
@@ -77,6 +77,7 @@
         @Child TestChildNode right;
 
         public TestRootNode(TestChildNode left, TestChildNode right) {
+            super(null);
             this.left = adoptChild(left);
             this.right = adoptChild(right);
         }
@@ -89,6 +90,10 @@
 
     abstract class TestChildNode extends Node {
 
+        public TestChildNode() {
+            super(null);
+        }
+
         abstract int execute(VirtualFrame frame);
     }
 
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/ReplaceTest.java
--- a/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/ReplaceTest.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/ReplaceTest.java	Mon Oct 07 11:36:49 2013 +0200
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2012, Oracle and/or its affiliates. All rights reserved.
+ * Copyright (c) 2012, 2013, Oracle and/or its affiliates. All rights reserved.
  * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
  *
  * This code is free software; you can redistribute it and/or modify it
@@ -84,6 +84,7 @@
         @Children private final ValueNode[] children;
 
         public TestRootNode(ValueNode[] children) {
+            super(null);
             this.children = adoptChildren(children);
         }
 
@@ -99,6 +100,10 @@
 
     abstract class ValueNode extends Node {
 
+        public ValueNode() {
+            super(null);
+        }
+
         abstract int execute();
     }
 
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/ReturnTypeSpecializationTest.java
--- a/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/ReturnTypeSpecializationTest.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/ReturnTypeSpecializationTest.java	Mon Oct 07 11:36:49 2013 +0200
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2012, Oracle and/or its affiliates. All rights reserved.
+ * Copyright (c) 2012, 2013, Oracle and/or its affiliates. All rights reserved.
  * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
  *
  * This code is free software; you can redistribute it and/or modify it
@@ -62,6 +62,7 @@
         @Child TestChildNode right;
 
         public TestRootNode(TestChildNode left, TestChildNode right) {
+            super(null);
             this.left = adoptChild(left);
             this.right = adoptChild(right);
         }
@@ -75,6 +76,10 @@
 
     abstract class TestChildNode extends Node {
 
+        public TestChildNode() {
+            super(null);
+        }
+
         abstract Object execute(VirtualFrame frame);
 
         int executeInt(VirtualFrame frame) throws UnexpectedResultException {
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/RootNodeTest.java
--- a/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/RootNodeTest.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.truffle.api.test/src/com/oracle/truffle/api/test/RootNodeTest.java	Mon Oct 07 11:36:49 2013 +0200
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2012, Oracle and/or its affiliates. All rights reserved.
+ * Copyright (c) 2012, 2013, Oracle and/or its affiliates. All rights reserved.
  * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
  *
  * This code is free software; you can redistribute it and/or modify it
@@ -59,6 +59,10 @@
 
     class TestRootNode extends RootNode {
 
+        public TestRootNode() {
+            super(null);
+        }
+
         @Override
         public Object execute(VirtualFrame frame) {
             return 42;
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.truffle.api/src/com/oracle/truffle/api/nodes/NodeUtil.java
--- a/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/nodes/NodeUtil.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/nodes/NodeUtil.java	Mon Oct 07 11:36:49 2013 +0200
@@ -657,7 +657,7 @@
             final SourceSection sourceSection = node.getSourceSection();
             if (sourceSection != null) {
                 final String txt = sourceSection.getSource().getCode();
-                p.println("Full source len=(" + txt.length() + ")  txt=___" + txt + "___");
+                p.println("Full source len=(" + txt.length() + ")  ___" + txt + "___");
                 p.println("AST source attribution:");
             }
         }
@@ -793,7 +793,7 @@
             final StringBuilder sb = new StringBuilder();
             sb.append("source:  len=" + srcText.length());
             sb.append(" (" + section.getCharIndex() + "," + (section.getCharEndIndex() - 1) + ")");
-            sb.append(" txt=___" + srcText + "___");
+            sb.append(" ___" + srcText + "___");
             return sb.toString();
         }
         return "";
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.truffle.sl/src/com/oracle/truffle/sl/SLNode.java
--- a/graal/com.oracle.truffle.sl/src/com/oracle/truffle/sl/SLNode.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.truffle.sl/src/com/oracle/truffle/sl/SLNode.java	Mon Oct 07 11:36:49 2013 +0200
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2012, 2012, Oracle and/or its affiliates. All rights reserved.
+ * Copyright (c) 2012, 2013, Oracle and/or its affiliates. All rights reserved.
  * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
  *
  * This code is free software; you can redistribute it and/or modify it
@@ -28,6 +28,11 @@
 @TypeSystemReference(SLTypes.class)
 public class SLNode extends Node {
 
+    public SLNode() {
+        // No source attribution
+        super(null);
+    }
+
     @Override
     public String toString() {
         return NodeUtil.printTreeToString(this);
diff -r 234aea2460bb -r 43bf803203c0 graal/com.oracle.truffle.sl/src/com/oracle/truffle/sl/nodes/FunctionDefinitionNode.java
--- a/graal/com.oracle.truffle.sl/src/com/oracle/truffle/sl/nodes/FunctionDefinitionNode.java	Sun Oct 06 23:45:02 2013 +0200
+++ b/graal/com.oracle.truffle.sl/src/com/oracle/truffle/sl/nodes/FunctionDefinitionNode.java	Mon Oct 07 11:36:49 2013 +0200
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2012, 2012, Oracle and/or its affiliates. All rights reserved.
+ * Copyright (c) 2012, 2013, Oracle and/or its affiliates. All rights reserved.
  * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
  *
  * This code is free software; you can redistribute it and/or modify it
@@ -35,6 +35,7 @@
     private final String name;
 
     public FunctionDefinitionNode(StatementNode body, FrameDescriptor frameDescriptor, String name, TypedNode returnValue) {
+        super(null);
         this.body = adoptChild(body);
         this.frameDescriptor = frameDescriptor;
         this.name = name;
diff -r 234aea2460bb -r 43bf803203c0 src/gpu/ptx/vm/gpu_ptx.cpp
--- a/src/gpu/ptx/vm/gpu_ptx.cpp	Sun Oct 06 23:45:02 2013 +0200
+++ b/src/gpu/ptx/vm/gpu_ptx.cpp	Mon Oct 07 11:36:49 2013 +0200
@@ -393,6 +393,17 @@
          ret.set_jint(return_val);
        }
        break;
+     case T_BOOLEAN:
+       {
+         int return_val;
+         status = gpu::Ptx::_cuda_cu_memcpy_dtoh(&return_val, ptxka._return_value_ptr, T_INT_BYTE_SIZE);
+         if (status != GRAAL_CUDA_SUCCESS) {
+           tty->print_cr("[CUDA] *** Error (%d) Failed to copy value to device argument", status);
+           return false;
+         }
+         ret.set_jint(return_val);
+       }
+       break;
      case T_FLOAT:
        {
          float return_val;
@@ -404,6 +415,17 @@
          ret.set_jfloat(return_val);
        }
        break;
+     case T_DOUBLE:
+       {
+         double return_val;
+         status = gpu::Ptx::_cuda_cu_memcpy_dtoh(&return_val, ptxka._return_value_ptr, T_DOUBLE_BYTE_SIZE);
+         if (status != GRAAL_CUDA_SUCCESS) {
+           tty->print_cr("[CUDA] *** Error (%d) Failed to copy value to device argument", status);
+           return false;
+         }
+         ret.set_jdouble(return_val);
+       }
+       break;
      case T_LONG:
        {
          long return_val;
diff -r 234aea2460bb -r 43bf803203c0 src/gpu/ptx/vm/ptxKernelArguments.cpp
--- a/src/gpu/ptx/vm/ptxKernelArguments.cpp	Sun Oct 06 23:45:02 2013 +0200
+++ b/src/gpu/ptx/vm/ptxKernelArguments.cpp	Mon Oct 07 11:36:49 2013 +0200
@@ -104,6 +104,39 @@
     return;
 }
 
+void PTXKernelArguments::do_double() {
+    if (is_after_invocation()) {
+        return;
+    }
+    // If the parameter is a return value,
+    jvalue doubleval;
+    if (is_return_type()) {
+        // Allocate device memory for T_INT return value pointer on device. Size in bytes
+        int status = gpu::Ptx::_cuda_cu_memalloc(&_return_value_ptr, T_DOUBLE_BYTE_SIZE);
+        if (status != GRAAL_CUDA_SUCCESS) {
+            tty->print_cr("[CUDA] *** Error (%d) Failed to allocate memory for return value pointer on device", status);
+            _success = false;
+            return;
+        }
+        // Push _return_value_ptr to _kernelBuffer
+        *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _return_value_ptr;
+        // _bufferOffset += sizeof(_return_value_ptr);
+        _bufferOffset += sizeof(doubleval.d);
+    } else {
+        // Get the next java argument and its value which should be a T_INT
+        oop arg = next_arg(T_FLOAT);
+        // Copy the java argument value to kernelArgBuffer
+        if (java_lang_boxing_object::get_value(arg, &doubleval) != T_DOUBLE) {
+            tty->print_cr("[CUDA] *** Error: Unexpected argument type; expecting T_INT");
+            _success = false;
+            return;
+        }
+        *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = doubleval.d;
+        _bufferOffset += sizeof(doubleval.d);
+    }
+    return;
+}
+
 void PTXKernelArguments::do_long() {
   if (is_after_invocation()) {
     return;
@@ -168,6 +201,38 @@
     return;
 }
 
+void PTXKernelArguments::do_bool() {
+    if (is_after_invocation()) {
+        return;
+    }
+    // If the parameter is a return value,
+    if (is_return_type()) {
+        // Allocate device memory for T_BYTE return value pointer on device. Size in bytes
+        int status = gpu::Ptx::_cuda_cu_memalloc(&_return_value_ptr, T_BOOLEAN_SIZE);
+        if (status != GRAAL_CUDA_SUCCESS) {
+            tty->print_cr("[CUDA] *** Error (%d) Failed to allocate memory for return value pointer on device", status);
+            _success = false;
+            return;
+        }
+        // Push _return_value_ptr to _kernelBuffer
+        *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _return_value_ptr;
+        _bufferOffset += sizeof(_return_value_ptr);
+    } else {
+        // Get the next java argument and its value which should be a T_BYTE
+        oop arg = next_arg(T_BYTE);
+        // Copy the java argument value to kernelArgBuffer
+        jvalue val;
+        if (java_lang_boxing_object::get_value(arg, &val) != T_BOOLEAN) {
+            tty->print_cr("[CUDA] *** Error: Unexpected argument type; expecting T_BYTE");
+            _success = false;
+            return;
+        }
+        *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = val.z;
+        _bufferOffset += sizeof(val.z);
+    }
+    return;
+}
+
 void PTXKernelArguments::do_array(int begin, int end) {
     gpu::Ptx::CUdeviceptr _array_ptr;
     int status;
diff -r 234aea2460bb -r 43bf803203c0 src/gpu/ptx/vm/ptxKernelArguments.hpp
--- a/src/gpu/ptx/vm/ptxKernelArguments.hpp	Sun Oct 06 23:45:02 2013 +0200
+++ b/src/gpu/ptx/vm/ptxKernelArguments.hpp	Mon Oct 07 11:36:49 2013 +0200
@@ -28,11 +28,13 @@
 #include "runtime/gpu.hpp"
 #include "runtime/signature.hpp"
 
-#define T_BYTE_SIZE       1
-#define T_INT_BYTE_SIZE   4
-#define T_FLOAT_BYTE_SIZE 4
-#define T_LONG_BYTE_SIZE  8
-#define T_ARRAY_BYTE_SIZE 8
+#define T_BYTE_SIZE        1
+#define T_BOOLEAN_SIZE     4
+#define T_INT_BYTE_SIZE    4
+#define T_FLOAT_BYTE_SIZE  4
+#define T_DOUBLE_BYTE_SIZE 8
+#define T_LONG_BYTE_SIZE   8
+#define T_ARRAY_BYTE_SIZE  8
 
 class PTXKernelArguments : public SignatureIterator {
 public:
@@ -99,16 +101,14 @@
 
 
   void do_byte();
+  void do_bool();
   void do_int();
   void do_float();
+  void do_double();
   void do_long();
   void do_array(int begin, int end);
   void do_void();
 
-  inline void do_bool()   {
-    /* TODO : To be implemented */
-    guarantee(false, "do_bool:NYI");
-  }
   inline void do_char()   {
     /* TODO : To be implemented */
     guarantee(false, "do_char:NYI");
@@ -117,11 +117,6 @@
     /* TODO : To be implemented */
     guarantee(false, "do_short:NYI");
   }
-  inline void do_double() {
-    /* TODO : To be implemented */
-    guarantee(false, "do_double:NYI");
-  }
-
   inline void do_object() {
     /* TODO : To be implemented */
     guarantee(false, "do_object:NYI");
diff -r 234aea2460bb -r 43bf803203c0 src/share/vm/graal/graalCompilerToGPU.cpp
--- a/src/share/vm/graal/graalCompilerToGPU.cpp	Sun Oct 06 23:45:02 2013 +0200
+++ b/src/share/vm/graal/graalCompilerToGPU.cpp	Mon Oct 07 11:36:49 2013 +0200
@@ -96,7 +96,20 @@
   } else {
     oop o = java_lang_boxing_object::create(ptxka.get_ret_type(), (jvalue *) result.get_value_addr(), CHECK_NULL);
     if (TraceGPUInteraction) {
-      tty->print_cr("GPU execution returned %d", result.get_jint());
+      switch (ptxka.get_ret_type()) {
+        case T_INT:
+          tty->print_cr("GPU execution returned %d", result.get_jint());
+          break;
+        case T_FLOAT:
+          tty->print_cr("GPU execution returned %f", result.get_jfloat());
+          break;
+        case T_DOUBLE:
+          tty->print_cr("GPU execution returned %f", result.get_jdouble());
+          break;
+        default:
+          tty->print_cr("GPU returned unhandled");
+          break;
+        }
     }
     return JNIHandles::make_local(o);
   }
@@ -135,7 +148,20 @@
   } else {
     oop o = java_lang_boxing_object::create(ptxka.get_ret_type(), (jvalue *) result.get_value_addr(), CHECK_NULL);
     if (TraceGPUInteraction) {
-      tty->print_cr("GPU execution returned %d", result.get_jint());
+      switch (ptxka.get_ret_type()) {
+        case T_INT:
+          tty->print_cr("GPU execution returned %d", result.get_jint());
+          break;
+        case T_FLOAT:
+          tty->print_cr("GPU execution returned %f", result.get_jfloat());
+          break;
+        case T_DOUBLE:
+          tty->print_cr("GPU execution returned %g", result.get_jdouble());
+          break;
+        default:
+          tty->print_cr("GPU returned unhandled");
+          break;
+      }
     }
     return JNIHandles::make_local(o);
   }