Mercurial > hg > truffle

changeset 11901:61767ccd4600

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
PTX boolean return value, emitIntegerTestMove, warnings
author Morris Meyer <morris.meyer@oracle.com>
date Sun, 06 Oct 2013 18:15:56 -0400
parents 1f82cda83ced
children 67a1e27a8dbb df1d665ca846 b18d4089535d
files graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ControlPTXTest.java graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXTestOp.java src/gpu/ptx/vm/gpu_ptx.cpp src/gpu/ptx/vm/ptxKernelArguments.cpp src/gpu/ptx/vm/ptxKernelArguments.hpp
diffstat 9 files changed, 218 insertions(+), 23 deletions(-) [+]
line wrap: on
line diff
--- a/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java	Sun Oct 06 13:55:09 2013 -0400
+++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java	Sun Oct 06 18:15:56 2013 -0400
@@ -521,8 +521,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) {
@@ -628,7 +628,7 @@
             this.targets = targets;
         }
 
-        private String valueForKind(Kind k) {
+        private static String valueForKind(Kind k) {
             switch (k.getTypeChar()) {
                 case 'i':
                     return "s32";
@@ -639,7 +639,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 +722,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 +757,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 +784,10 @@
             operator = co;
         }
 
+        public void setBooleanOperator(BooleanOperator bo) {
+            booleanOperator = bo;
+        }
+
         private ConditionOperator operatorForConditon(Condition condition) {
             char typeChar = kind.getTypeChar();
 
@@ -873,8 +908,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
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ControlPTXTest.java	Sun Oct 06 13:55:09 2013 -0400
+++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ControlPTXTest.java	Sun Oct 06 18:15:56 2013 -0400
@@ -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
         }
     }
 
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java	Sun Oct 06 13:55:09 2013 -0400
+++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java	Sun Oct 06 18:15:56 2013 -0400
@@ -231,6 +231,7 @@
         if (maxPredRegNum > 0) {
             codeBuffer.emitString(".reg .pred %p<" + maxPredRegNum + ">;");
         }
+        codeBuffer.emitString(".reg .pred %r;");  // used for setp bool
     }
 
     @Override
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java	Sun Oct 06 13:55:09 2013 -0400
+++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java	Sun Oct 06 18:15:56 2013 -0400
@@ -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
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java	Sun Oct 06 13:55:09 2013 -0400
+++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java	Sun Oct 06 18:15:56 2013 -0400
@@ -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)) :
--- /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	Sun Oct 06 18:15:56 2013 -0400
@@ -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);
+    }
+}
--- a/src/gpu/ptx/vm/gpu_ptx.cpp	Sun Oct 06 13:55:09 2013 -0400
+++ b/src/gpu/ptx/vm/gpu_ptx.cpp	Sun Oct 06 18:15:56 2013 -0400
@@ -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;
--- a/src/gpu/ptx/vm/ptxKernelArguments.cpp	Sun Oct 06 13:55:09 2013 -0400
+++ b/src/gpu/ptx/vm/ptxKernelArguments.cpp	Sun Oct 06 18:15:56 2013 -0400
@@ -168,6 +168,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;
--- a/src/gpu/ptx/vm/ptxKernelArguments.hpp	Sun Oct 06 13:55:09 2013 -0400
+++ b/src/gpu/ptx/vm/ptxKernelArguments.hpp	Sun Oct 06 18:15:56 2013 -0400
@@ -29,6 +29,7 @@
 #include "runtime/signature.hpp"
 
 #define T_BYTE_SIZE       1
+#define T_BOOLEAN_SIZE    4
 #define T_INT_BYTE_SIZE   4
 #define T_FLOAT_BYTE_SIZE 4
 #define T_LONG_BYTE_SIZE  8
@@ -99,16 +100,13 @@
 
 
   void do_byte();
+  void do_bool();
   void do_int();
   void do_float();
   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");