# HG changeset patch # User Thomas Wuerthinger # Date 1381108280 -7200 # Node ID b18d4089535db9b5f2ec3a6726ba8ee73c08ac8c # Parent 10755fbf8df8040562064ed60f7880d38aa1225c# Parent 61767ccd4600af886a36723048e29dc1f37e02e4 Merge. diff -r 10755fbf8df8 -r b18d4089535d 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 Mon Oct 07 03:08:50 2013 +0200 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Mon Oct 07 03:11:20 2013 +0200 @@ -517,8 +517,8 @@ emitString("@%p" + pred + " " + "bra" + " " + tgt + ";"); } - public final void bra(String targetIdentifier) { - emitString("bra " + targetIdentifier + ";"); + public final void bra(String src) { + emitString("bra " + src + ";"); } public final void bra_uni(String tgt) { @@ -715,14 +715,34 @@ emitString("ret.uni;" + " " + ""); } - public static class Setp { + public enum BooleanOperator { + AND("and"), + OR("or"), + XOR("xor"); + + private final String output; + + private BooleanOperator(String out) { + this.output = out; + } - private ConditionOperator operator; + 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); @@ -730,6 +750,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; } @@ -746,6 +777,10 @@ operator = co; } + public void setBooleanOperator(BooleanOperator bo) { + booleanOperator = bo; + } + private ConditionOperator operatorForConditon(Condition condition) { char typeChar = kind.getTypeChar(); @@ -883,7 +918,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) + ";"); + } } } diff -r 10755fbf8df8 -r b18d4089535d 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 Mon Oct 07 03:08:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ControlPTXTest.java Mon Oct 07 03:11:20 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; diff -r 10755fbf8df8 -r b18d4089535d 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 Mon Oct 07 03:08:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java Mon Oct 07 03:11:20 2013 +0200 @@ -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 10755fbf8df8 -r b18d4089535d 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 Mon Oct 07 03:08:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java Mon Oct 07 03:11:20 2013 +0200 @@ -328,7 +328,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 @@ -401,8 +403,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 10755fbf8df8 -r b18d4089535d graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java diff -r 10755fbf8df8 -r b18d4089535d 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 03:11:20 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 10755fbf8df8 -r b18d4089535d src/gpu/ptx/vm/gpu_ptx.cpp --- a/src/gpu/ptx/vm/gpu_ptx.cpp Mon Oct 07 03:08:50 2013 +0200 +++ b/src/gpu/ptx/vm/gpu_ptx.cpp Mon Oct 07 03:11:20 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; diff -r 10755fbf8df8 -r b18d4089535d src/gpu/ptx/vm/ptxKernelArguments.cpp --- a/src/gpu/ptx/vm/ptxKernelArguments.cpp Mon Oct 07 03:08:50 2013 +0200 +++ b/src/gpu/ptx/vm/ptxKernelArguments.cpp Mon Oct 07 03:11:20 2013 +0200 @@ -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 10755fbf8df8 -r b18d4089535d src/gpu/ptx/vm/ptxKernelArguments.hpp --- a/src/gpu/ptx/vm/ptxKernelArguments.hpp Mon Oct 07 03:08:50 2013 +0200 +++ b/src/gpu/ptx/vm/ptxKernelArguments.hpp Mon Oct 07 03:11:20 2013 +0200 @@ -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");