# HG changeset patch # User Morris Meyer # Date 1381097756 14400 # Node ID 61767ccd4600af886a36723048e29dc1f37e02e4 # Parent 1f82cda83ced7536795c46f9fe3f3f06b6a5426b PTX boolean return value, emitIntegerTestMove, warnings diff -r 1f82cda83ced -r 61767ccd4600 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 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 diff -r 1f82cda83ced -r 61767ccd4600 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 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 } } diff -r 1f82cda83ced -r 61767ccd4600 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 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 diff -r 1f82cda83ced -r 61767ccd4600 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 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 diff -r 1f82cda83ced -r 61767ccd4600 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 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)) : diff -r 1f82cda83ced -r 61767ccd4600 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 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); + } +} diff -r 1f82cda83ced -r 61767ccd4600 src/gpu/ptx/vm/gpu_ptx.cpp --- 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; diff -r 1f82cda83ced -r 61767ccd4600 src/gpu/ptx/vm/ptxKernelArguments.cpp --- 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; diff -r 1f82cda83ced -r 61767ccd4600 src/gpu/ptx/vm/ptxKernelArguments.hpp --- 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");