Mercurial > hg > truffle
changeset 7805:0e58445d54df
Integration fixes.
line wrap: on
line diff
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/Register.java Mon Feb 18 14:47:54 2013 -0800 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/Register.java Mon Feb 18 16:04:25 2013 -0800 @@ -67,7 +67,7 @@ /** * The assembler calls this method to get the register's encoding. */ - public final int encoding() { + public int encoding() { return encoding; }
--- a/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/AbstractPTXAssembler.java Mon Feb 18 14:47:54 2013 -0800 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/AbstractPTXAssembler.java Mon Feb 18 16:04:25 2013 -0800 @@ -24,7 +24,6 @@ import com.oracle.graal.api.code.*; import com.oracle.graal.asm.*; -import com.oracle.graal.lir.ptx.*; /** * The platform-dependent base class for the PTX assembler. @@ -35,10 +34,12 @@ super(target); } + public static final String UNBOUND_TARGET = "L" + Integer.MAX_VALUE; + @Override public final void bind(Label l) { super.bind(l); - emitString0("L"+l.toString() + ":\n"); + emitString0("L" + l.toString() + ":\n"); } @Override @@ -53,10 +54,10 @@ @Override protected void patchJumpTarget(int branch, int jumpTarget) { - final int spaces = PTXControlFlow.BranchOp.UNBOUND_TARGET.length(); - String target = String.format("L%-" + spaces + "s", jumpTarget+";"); + final int spaces = UNBOUND_TARGET.length(); + String targetString = String.format("L%-" + spaces + "s", jumpTarget + ";"); int offset = "\tbra ".length(); // XXX we need a better way to figure this out - codeBuffer.emitString(target, branch + offset); + codeBuffer.emitString(targetString, branch + offset); } @Override
--- a/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Mon Feb 18 14:47:54 2013 -0800 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Mon Feb 18 16:04:25 2013 -0800 @@ -1,550 +1,740 @@ +/* + * Copyright (c) 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 + * 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.asm.ptx; -import static com.oracle.graal.ptx.PTX.*; -import static com.oracle.graal.api.code.MemoryBarriers.*; -import static com.oracle.graal.api.code.ValueUtil.*; -import static com.oracle.graal.asm.NumUtil.*; -import static com.oracle.graal.asm.ptx.PTXAsmOptions.*; - -import com.oracle.graal.ptx.*; import com.oracle.graal.api.code.*; -import com.oracle.graal.api.meta.*; -import com.oracle.graal.asm.*; public class PTXAssembler extends AbstractPTXAssembler { + + @SuppressWarnings("unused") public PTXAssembler(TargetDescription target, RegisterConfig registerConfig) { super(target); - //this.frameRegister = registerConfig == null ? null : registerConfig.getFrameRegister(); } public final void at() { emitString("@%p" + " " + ""); } + public final void add_s16(Register d, Register a, Register b) { emitString("add.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void add_s32(Register d, Register a, Register b) { emitString("add.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void add_s64(Register d, Register a, Register b) { emitString("add.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void add_s16(Register d, Register a, short s16) { emitString("add.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + ""); } + public final void add_s32(Register d, Register a, int s32) { emitString("add.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + ""); } + public final void add_s64(Register d, Register a, long s64) { emitString("add.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + ""); } + public final void add_u16(Register d, Register a, Register b) { emitString("add.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void add_u32(Register d, Register a, Register b) { emitString("add.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void add_u64(Register d, Register a, Register b) { emitString("add.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void add_u16(Register d, Register a, short u16) { emitString("add.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + ""); } + public final void add_u32(Register d, Register a, int u32) { emitString("add.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } + public final void add_u64(Register d, Register a, long u64) { emitString("add.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + ""); } + public final void add_sat_s32(Register d, Register a, Register b) { emitString("add.sat.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void add_sat_s32(Register d, Register a, int s32) { emitString("add.sat.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + ""); } + public final void and_b16(Register d, Register a, Register b) { emitString("and.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void and_b32(Register d, Register a, Register b) { emitString("and.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void and_b64(Register d, Register a, Register b) { emitString("and.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void and_b16(Register d, Register a, short b16) { emitString("and.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b16 + ";" + ""); } + public final void and_b32(Register d, Register a, int b32) { emitString("and.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b32 + ";" + ""); } + public final void and_b64(Register d, Register a, long b64) { emitString("and.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b64 + ";" + ""); } + public final void bra(String tgt) { emitString("bra" + " " + tgt + ";" + ""); } + public final void bra_uni(String tgt) { emitString("bra.uni" + " " + tgt + ";" + ""); } + public final void div_s16(Register d, Register a, Register b) { emitString("div.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void div_s32(Register d, Register a, Register b) { emitString("div.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void div_s64(Register d, Register a, Register b) { emitString("div.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void div_s16(Register d, Register a, short s16) { emitString("div.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + ""); } + public final void div_s32(Register d, Register a, int s32) { emitString("div.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + ""); } + public final void div_s64(Register d, Register a, long s64) { emitString("div.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + ""); } + public final void div_u16(Register d, Register a, Register b) { emitString("div.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void div_u32(Register d, Register a, Register b) { emitString("div.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void div_u64(Register d, Register a, Register b) { emitString("div.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void div_u16(Register d, Register a, short u16) { emitString("div.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + ""); } + public final void div_u32(Register d, Register a, int u32) { emitString("div.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } + public final void div_u64(Register d, Register a, long u64) { emitString("div.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + ""); } + public final void exit() { emitString("exit;" + " " + ""); } + public final void ld_global_b8(Register d, Register a, int immOff) { emitString("ld.global.b8" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } + public final void ld_global_b16(Register d, Register a, int immOff) { emitString("ld.global.b16" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } + public final void ld_global_b32(Register d, Register a, int immOff) { emitString("ld.global.b32" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } + public final void ld_global_b64(Register d, Register a, int immOff) { emitString("ld.global.b64" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } + public final void ld_global_u8(Register d, Register a, int immOff) { emitString("ld.global.u8" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } + public final void ld_global_u16(Register d, Register a, int immOff) { emitString("ld.global.u16" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } + public final void ld_global_u32(Register d, Register a, int immOff) { emitString("ld.global.u32" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } + public final void ld_global_u64(Register d, Register a, int immOff) { emitString("ld.global.u64" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } + public final void ld_global_s8(Register d, Register a, int immOff) { emitString("ld.global.s8" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } + public final void ld_global_s16(Register d, Register a, int immOff) { emitString("ld.global.s16" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } + public final void ld_global_s32(Register d, Register a, int immOff) { emitString("ld.global.s32" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } + public final void ld_global_s64(Register d, Register a, int immOff) { emitString("ld.global.s64" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } + public final void ld_global_f32(Register d, Register a, int immOff) { emitString("ld.global.f32" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } + public final void ld_global_f64(Register d, Register a, int immOff) { emitString("ld.global.f64" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } + public final void mov_b16(Register d, Register a) { emitString("mov.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void mov_b32(Register d, Register a) { emitString("mov.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void mov_b64(Register d, Register a) { emitString("mov.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void mov_u16(Register d, Register a) { emitString("mov.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void mov_u32(Register d, Register a) { emitString("mov.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void mov_u64(Register d, Register a) { emitString("mov.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void mov_s16(Register d, Register a) { emitString("mov.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void mov_s32(Register d, Register a) { emitString("mov.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void mov_s64(Register d, Register a) { emitString("mov.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void mov_f32(Register d, Register a) { emitString("mov.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void mov_f64(Register d, Register a) { emitString("mov.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void mov_b16(Register d, short b16) { emitString("mov.b16" + " " + "%r" + d.encoding() + ", " + b16 + ";" + ""); } + public final void mov_b32(Register d, int b32) { emitString("mov.b32" + " " + "%r" + d.encoding() + ", " + b32 + ";" + ""); } + public final void mov_b64(Register d, long b64) { emitString("mov.b64" + " " + "%r" + d.encoding() + ", " + b64 + ";" + ""); } + public final void mov_u16(Register d, short u16) { emitString("mov.u16" + " " + "%r" + d.encoding() + ", " + u16 + ";" + ""); } + public final void mov_u32(Register d, int u32) { emitString("mov.u32" + " " + "%r" + d.encoding() + ", " + u32 + ";" + ""); } + public final void mov_u64(Register d, long u64) { emitString("mov.u64" + " " + "%r" + d.encoding() + ", " + u64 + ";" + ""); } + public final void mov_s16(Register d, short s16) { emitString("mov.s16" + " " + "%r" + d.encoding() + ", " + s16 + ";" + ""); } + public final void mov_s32(Register d, int s32) { emitString("mov.s32" + " " + "%r" + d.encoding() + ", " + s32 + ";" + ""); } + public final void mov_s64(Register d, long s64) { emitString("mov.s64" + " " + "%r" + d.encoding() + ", " + s64 + ";" + ""); } + public final void mov_f32(Register d, float f32) { emitString("mov.f32" + " " + "%r" + d.encoding() + ", " + f32 + ";" + ""); } + public final void mov_f64(Register d, double f64) { emitString("mov.f64" + " " + "%r" + d.encoding() + ", " + f64 + ";" + ""); } + public final void mul_s16(Register d, Register a, Register b) { emitString("mul.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void mul_s32(Register d, Register a, Register b) { emitString("mul.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void mul_s64(Register d, Register a, Register b) { emitString("mul.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void mul_s16(Register d, Register a, short s16) { emitString("mul.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + ""); } + public final void mul_s32(Register d, Register a, int s32) { emitString("mul.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + ""); } + public final void mul_s64(Register d, Register a, long s64) { emitString("mul.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + ""); } + public final void mul_u16(Register d, Register a, Register b) { emitString("mul.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void mul_u32(Register d, Register a, Register b) { emitString("mul.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void mul_u64(Register d, Register a, Register b) { emitString("mul.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void mul_u16(Register d, Register a, short u16) { emitString("mul.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + ""); } + public final void mul_u32(Register d, Register a, int u32) { emitString("mul.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } + public final void mul_u64(Register d, Register a, long u64) { emitString("mul.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + ""); } + public final void neg_s16(Register d, Register a) { emitString("neg.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void neg_s32(Register d, Register a) { emitString("neg.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void neg_s64(Register d, Register a) { emitString("neg.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void popc_b32(Register d, Register a) { emitString("popc.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void popc_b64(Register d, Register a) { emitString("popc.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void rem_s16(Register d, Register a, Register b) { emitString("rem.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void rem_s32(Register d, Register a, Register b) { emitString("rem.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void rem_s64(Register d, Register a, Register b) { emitString("rem.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void rem_s16(Register d, Register a, short s16) { emitString("rem.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + ""); } + public final void rem_s32(Register d, Register a, int s32) { emitString("rem.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + ""); } + public final void rem_s64(Register d, Register a, long s64) { emitString("rem.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + ""); } + public final void rem_u16(Register d, Register a, Register b) { emitString("rem.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void rem_u32(Register d, Register a, Register b) { emitString("rem.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void rem_u64(Register d, Register a, Register b) { emitString("rem.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void rem_u16(Register d, Register a, short u16) { emitString("rem.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + ""); } + public final void rem_u32(Register d, Register a, int u32) { emitString("rem.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } + public final void rem_u64(Register d, Register a, long u64) { emitString("rem.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + ""); } + public final void ret() { emitString("ret;" + " " + ""); } + public final void ret_uni() { emitString("ret.uni;" + " " + ""); } + public final void setp_eq_s32(Register a, Register b) { emitString("setp.eq.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void setp_ne_s32(Register a, Register b) { emitString("setp.ne.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void setp_lt_s32(Register a, Register b) { emitString("setp.lt.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void setp_le_s32(Register a, Register b) { emitString("setp.le.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void setp_gt_s32(Register a, Register b) { emitString("setp.gt.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void setp_ge_s32(Register a, Register b) { emitString("setp.ge.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void setp_eq_s32(Register a, int s32) { emitString("setp.eq.s32" + " " + "%p" + ", %r" + a.encoding() + ", " + s32 + ";" + ""); } + public final void setp_ne_s32(Register a, int s32) { emitString("setp.ne.s32" + " " + "%p" + ", %r" + a.encoding() + ", " + s32 + ";" + ""); } + public final void setp_lt_s32(Register a, int s32) { emitString("setp.lt.s32" + " " + "%p" + ", %r" + a.encoding() + ", " + s32 + ";" + ""); } + public final void setp_le_s32(Register a, int s32) { emitString("setp.le.s32" + " " + "%p" + ", %r" + a.encoding() + ", " + s32 + ";" + ""); } + public final void setp_gt_s32(Register a, int s32) { emitString("setp.gt.s32" + " " + "%p" + ", %r" + a.encoding() + ", " + s32 + ";" + ""); } + public final void setp_ge_s32(Register a, int s32) { emitString("setp.ge.s32" + " " + "%p" + ", %r" + a.encoding() + ", " + s32 + ";" + ""); } + public final void setp_eq_s32(int s32, Register b) { emitString("setp.eq.s32" + " " + "%p" + ", " + s32 + ", %r" + b.encoding() + ";" + ""); } + public final void setp_ne_s32(int s32, Register b) { emitString("setp.ne.s32" + " " + "%p" + ", " + s32 + ", %r" + b.encoding() + ";" + ""); } + public final void setp_lt_s32(int s32, Register b) { emitString("setp.lt.s32" + " " + "%p" + ", " + s32 + ", %r" + b.encoding() + ";" + ""); } + public final void setp_le_s32(int s32, Register b) { emitString("setp.le.s32" + " " + "%p" + ", " + s32 + ", %r" + b.encoding() + ";" + ""); } + public final void setp_gt_s32(int s32, Register b) { emitString("setp.gt.s32" + " " + "%p" + ", " + s32 + ", %r" + b.encoding() + ";" + ""); } + public final void setp_ge_s32(int s32, Register b) { emitString("setp.ge.s32" + " " + "%p" + ", " + s32 + ", %r" + b.encoding() + ";" + ""); } + public final void setp_eq_u32(Register a, Register b) { emitString("setp.eq.u32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void setp_ne_u32(Register a, Register b) { emitString("setp.ne.u32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void setp_lt_u32(Register a, Register b) { emitString("setp.lt.u32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void setp_le_u32(Register a, Register b) { emitString("setp.le.u32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void setp_gt_u32(Register a, Register b) { emitString("setp.gt.u32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void setp_ge_u32(Register a, Register b) { emitString("setp.ge.u32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void setp_eq_u32(Register a, int u32) { emitString("setp.eq.u32" + " " + "%p" + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } + public final void setp_ne_u32(Register a, int u32) { emitString("setp.ne.u32" + " " + "%p" + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } + public final void setp_lt_u32(Register a, int u32) { emitString("setp.lt.u32" + " " + "%p" + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } + public final void setp_le_u32(Register a, int u32) { emitString("setp.le.u32" + " " + "%p" + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } + public final void setp_gt_u32(Register a, int u32) { emitString("setp.gt.u32" + " " + "%p" + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } + public final void setp_ge_u32(Register a, int u32) { emitString("setp.ge.u32" + " " + "%p" + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } + public final void setp_eq_u32(int u32, Register b) { emitString("setp.eq.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + ""); } + public final void setp_ne_u32(int u32, Register b) { emitString("setp.ne.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + ""); } + public final void setp_lt_u32(int u32, Register b) { emitString("setp.lt.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + ""); } + public final void setp_le_u32(int u32, Register b) { emitString("setp.le.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + ""); } + public final void setp_gt_u32(int u32, Register b) { emitString("setp.gt.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + ""); } + public final void setp_ge_u32(int u32, Register b) { emitString("setp.ge.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + ""); } + public final void shr_s16(Register d, Register a, Register b) { emitString("shr.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void shr_s32(Register d, Register a, Register b) { emitString("shr.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void shr_s64(Register d, Register a, Register b) { emitString("shr.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void shr_s16(Register d, Register a, int u32) { emitString("shr.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } + public final void shr_s32(Register d, Register a, int u32) { emitString("shr.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } + public final void shr_s64(Register d, Register a, int u32) { emitString("shr.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } + public final void shr_u16(Register d, Register a, Register b) { emitString("shr.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void shr_u32(Register d, Register a, Register b) { emitString("shr.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void shr_u64(Register d, Register a, Register b) { emitString("shr.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void shr_u16(Register d, Register a, int u32) { emitString("shr.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } + public final void shr_u32(Register d, Register a, int u32) { emitString("shr.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } + public final void shr_u64(Register d, Register a, int u32) { emitString("shr.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } + public final void st_global_b8(Register a, int immOff, Register b) { emitString("st.global.b8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } + public final void st_global_b16(Register a, int immOff, Register b) { emitString("st.global.b16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } + public final void st_global_b32(Register a, int immOff, Register b) { emitString("st.global.b32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } + public final void st_global_b64(Register a, int immOff, Register b) { emitString("st.global.b64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } + public final void st_global_u8(Register a, int immOff, Register b) { emitString("st.global.u8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } + public final void st_global_u16(Register a, int immOff, Register b) { emitString("st.global.u16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } + public final void st_global_u32(Register a, int immOff, Register b) { emitString("st.global.u32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } + public final void st_global_u64(Register a, int immOff, Register b) { emitString("st.global.u64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } + public final void st_global_s8(Register a, int immOff, Register b) { emitString("st.global.s8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } + public final void st_global_s16(Register a, int immOff, Register b) { emitString("st.global.s16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } + public final void st_global_s32(Register a, int immOff, Register b) { emitString("st.global.s32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } + public final void st_global_s64(Register a, int immOff, Register b) { emitString("st.global.s64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } + public final void st_global_f32(Register a, int immOff, Register b) { emitString("st.global.f32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } + public final void st_global_f64(Register a, int immOff, Register b) { emitString("st.global.f64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } + public final void sub_s16(Register d, Register a, Register b) { emitString("sub.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void sub_s32(Register d, Register a, Register b) { emitString("sub.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void sub_s64(Register d, Register a, Register b) { emitString("sub.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void sub_s16(Register d, Register a, short s16) { emitString("sub.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + ""); } + public final void sub_s32(Register d, Register a, int s32) { emitString("sub.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + ""); } + public final void sub_s64(Register d, Register a, long s64) { emitString("sub.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + ""); } + public final void sub_s16(Register d, short s16, Register b) { emitString("sub.s16" + " " + "%r" + d.encoding() + ", " + s16 + ", %r" + b.encoding() + ";" + ""); } + public final void sub_s32(Register d, int s32, Register b) { emitString("sub.s32" + " " + "%r" + d.encoding() + ", " + s32 + ", %r" + b.encoding() + ";" + ""); } + public final void sub_s64(Register d, long s64, Register b) { emitString("sub.s64" + " " + "%r" + d.encoding() + ", " + s64 + ", %r" + b.encoding() + ";" + ""); } + public final void sub_sat_s32(Register d, Register a, Register b) { emitString("sub.sat.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); } + public final void sub_sat_s32(Register d, Register a, int s32) { emitString("sub.sat.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + ""); } + public final void sub_sat_s32(Register d, int s32, Register b) { emitString("sub.sat.s32" + " " + "%r" + d.encoding() + ", " + s32 + ", %r" + b.encoding() + ";" + ""); }
--- a/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXMacroAssembler.java Mon Feb 18 14:47:54 2013 -0800 +++ /dev/null Thu Jan 01 00:00:00 1970 +0000 @@ -1,36 +0,0 @@ -/* - * Copyright (c) 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 - * 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.asm.ptx; - -import com.oracle.graal.api.code.*; - -/** - * This class implements commonly used PTX code patterns. - */ -public class PTXMacroAssembler extends PTXAssembler { - - public PTXMacroAssembler(TargetDescription target, RegisterConfig registerConfig) { - super(target, registerConfig); - } - -}
--- a/graal/com.oracle.graal.asm/src/com/oracle/graal/asm/AbstractAssembler.java Mon Feb 18 14:47:54 2013 -0800 +++ b/graal/com.oracle.graal.asm/src/com/oracle/graal/asm/AbstractAssembler.java Mon Feb 18 16:04:25 2013 -0800 @@ -86,7 +86,7 @@ protected final void emitString(String x) { codeBuffer.emitString(x); } - + // XXX for pretty-printing protected final void emitString0(String x) { codeBuffer.emitString0(x);
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java Mon Feb 18 14:47:54 2013 -0800 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java Mon Feb 18 16:04:25 2013 -0800 @@ -22,24 +22,16 @@ */ package com.oracle.graal.compiler.ptx; -import static com.oracle.graal.api.code.CallingConvention.Type.JavaCallee; - import com.oracle.graal.api.code.*; -import com.oracle.graal.api.code.CallingConvention.Type; import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.*; import com.oracle.graal.asm.ptx.*; -import com.oracle.graal.compiler.gen.LIRGenerator; -import com.oracle.graal.compiler.target.Backend; -import com.oracle.graal.hotspot.HotSpotLIRGenerator; -import com.oracle.graal.hotspot.meta.HotSpotRuntime; -import com.oracle.graal.hotspot.nodes.DirectCompareAndSwapNode; -import com.oracle.graal.hotspot.stubs.Stub; +import com.oracle.graal.compiler.gen.*; +import com.oracle.graal.compiler.target.*; import com.oracle.graal.lir.*; import com.oracle.graal.lir.asm.*; import com.oracle.graal.nodes.*; -import com.oracle.graal.nodes.java.*; -import com.oracle.graal.phases.GraalOptions; +import com.oracle.graal.phases.*; /** * PTX specific backend. @@ -50,174 +42,23 @@ super(runtime, target); } + @Override public LIRGenerator newLIRGenerator(StructuredGraph graph, FrameMap frameMap, ResolvedJavaMethod method, LIR lir) { - return new HotSpotPTXLIRGenerator(graph, runtime(), target, frameMap, method, lir); - } - - static final class HotSpotPTXLIRGenerator extends PTXLIRGenerator implements HotSpotLIRGenerator { - - private HotSpotRuntime runtime() { - return (HotSpotRuntime) runtime; - } - - private HotSpotPTXLIRGenerator(StructuredGraph graph, CodeCacheProvider runtime, TargetDescription target, FrameMap frameMap, ResolvedJavaMethod method, LIR lir) { - super(graph, runtime, target, frameMap, method, lir); - } - - @Override - protected boolean needOnlyOopMaps() { - // Stubs only need oop maps - return runtime().asStub(method) != null; - } - - @Override - protected CallingConvention createCallingConvention() { - Stub stub = runtime().asStub(method); - if (stub != null) { - return stub.getLinkage().getCallingConvention(); - } - - if (graph.getEntryBCI() == StructuredGraph.INVOCATION_ENTRY_BCI) { - return super.createCallingConvention(); - } else { - return frameMap.registerConfig.getCallingConvention(JavaCallee, method.getSignature().getReturnType(null), new JavaType[]{runtime.lookupJavaType(long.class)}, target, false); - } - } - - @Override - public void visitSafepointNode(SafepointNode i) { - // We don't need safepoints in GPU code. - } - - @Override - public void visitExceptionObject(ExceptionObjectNode x) { -// HotSpotVMConfig config = runtime().config; -// RegisterValue thread = runtime().threadRegister().asValue(); -// Address exceptionAddress = new Address(Kind.Object, thread, config.threadExceptionOopOffset); -// Address pcAddress = new Address(Kind.Long, thread, config.threadExceptionPcOffset); -// Value exception = emitLoad(exceptionAddress, false); -// emitStore(exceptionAddress, Constant.NULL_OBJECT, false); -// emitStore(pcAddress, Constant.LONG_0, false); -// setResult(x, exception); - throw new InternalError("NYI"); - } - - @SuppressWarnings("hiding") - @Override - public void visitDirectCompareAndSwap(DirectCompareAndSwapNode x) { -// Kind kind = x.newValue().kind(); -// assert kind == x.expectedValue().kind(); -// -// Value expected = loadNonConst(operand(x.expectedValue())); -// Variable newVal = load(operand(x.newValue())); -// -// int disp = 0; -// Address address; -// Value index = operand(x.offset()); -// if (ValueUtil.isConstant(index) && NumUtil.isInt(ValueUtil.asConstant(index).asLong() + disp)) { -// assert !runtime.needsDataPatch(asConstant(index)); -// disp += (int) ValueUtil.asConstant(index).asLong(); -// address = new Address(kind, load(operand(x.object())), disp); -// } else { -// address = new Address(kind, load(operand(x.object())), load(index), Address.Scale.Times1, disp); -// } -// -// RegisterValue rax = AMD64.rax.asValue(kind); -// emitMove(expected, rax); -// append(new CompareAndSwapOp(rax, address, rax, newVal)); -// -// Variable result = newVariable(x.kind()); -// emitMove(rax, result); -// setResult(x, result); - throw new InternalError("NYI"); - } - - @Override - public void emitTailcall(Value[] args, Value address) { - //append(new AMD64TailcallOp(args, address)); - throw new InternalError("NYI"); - } - - @Override - protected void emitDirectCall(DirectCallTargetNode callTarget, Value result, Value[] parameters, Value[] temps, LIRFrameState callState) { -// append(new AMD64DirectCallOp(callTarget.target(), result, parameters, temps, callState, ((HotSpotDirectCallTargetNode) callTarget).invokeKind(), lir)); - throw new InternalError("NYI"); - } - - @Override - protected void emitIndirectCall(IndirectCallTargetNode callTarget, Value result, Value[] parameters, Value[] temps, LIRFrameState callState) { -// Value metaspaceMethod = AMD64.rbx.asValue(); -// emitMove(operand(((HotSpotIndirectCallTargetNode) callTarget).metaspaceMethod()), metaspaceMethod); -// Value targetAddress = AMD64.rax.asValue(); -// emitMove(operand(callTarget.computedAddress()), targetAddress); -// append(new AMD64IndirectCallOp(callTarget.target(), result, parameters, temps, metaspaceMethod, targetAddress, callState)); - throw new InternalError("NYI"); - } + return new PTXLIRGenerator(graph, runtime(), target, frameMap, method, lir); } class HotSpotFrameContext implements FrameContext { @Override public void enter(TargetMethodAssembler tasm) { - FrameMap frameMap = tasm.frameMap; - int frameSize = frameMap.frameSize(); - - PTXMacroAssembler asm = (PTXMacroAssembler) tasm.asm; -// emitStackOverflowCheck(tasm, false); -// asm.push(rbp); -// asm.movq(rbp, rsp); -// asm.decrementq(rsp, frameSize - 8); // account for the push of RBP above -// if (GraalOptions.ZapStackOnMethodEntry) { -// final int intSize = 4; -// for (int i = 0; i < frameSize / intSize; ++i) { -// asm.movl(new Address(Kind.Int, rsp.asValue(), i * intSize), 0xC1C1C1C1); -// } -// } -// CalleeSaveLayout csl = frameMap.registerConfig.getCalleeSaveLayout(); -// if (csl != null && csl.size != 0) { -// int frameToCSA = frameMap.offsetToCalleeSaveArea(); -// assert frameToCSA >= 0; -// asm.save(csl, frameToCSA); -// } + Buffer codeBuffer = tasm.asm.codeBuffer; + codeBuffer.emitString(".version 1.4"); + codeBuffer.emitString(".target sm_10"); + // codeBuffer.emitString(".address_size 32"); // PTX ISA version 2.3 } @Override public void leave(TargetMethodAssembler tasm) { - int frameSize = tasm.frameMap.frameSize(); - PTXMacroAssembler asm = (PTXMacroAssembler) tasm.asm; - CalleeSaveLayout csl = tasm.frameMap.registerConfig.getCalleeSaveLayout(); - RegisterConfig regConfig = tasm.frameMap.registerConfig; - -// if (csl != null && csl.size != 0) { -// tasm.targetMethod.setRegisterRestoreEpilogueOffset(asm.codeBuffer.position()); -// // saved all registers, restore all registers -// int frameToCSA = tasm.frameMap.offsetToCalleeSaveArea(); -// asm.restore(csl, frameToCSA); -// } -// -// asm.incrementq(rsp, frameSize - 8); // account for the pop of RBP below -// asm.pop(rbp); -// -// if (GraalOptions.GenSafepoints) { -// HotSpotVMConfig config = runtime().config; -// -// // If at the return point, then the frame has already been popped -// // so deoptimization cannot be performed here. The HotSpot runtime -// // detects this case - see the definition of frame::should_be_deoptimized() -// -// Register scratch = regConfig.getScratchRegister(); -// int offset = SafepointPollOffset % target.pageSize; -// if (config.isPollingPageFar) { -// asm.movq(scratch, config.safepointPollingAddress + offset); -// tasm.recordMark(Marks.MARK_POLL_RETURN_FAR); -// asm.movq(scratch, new Address(tasm.target.wordKind, scratch.asValue())); -// } else { -// tasm.recordMark(Marks.MARK_POLL_RETURN_NEAR); -// // The C++ code transforms the polling page offset into an RIP displacement -// // to the real address at that offset in the polling page. -// asm.movq(scratch, new Address(tasm.target.wordKind, rip.asValue(), offset)); -// } -// } } } @@ -231,7 +72,7 @@ boolean omitFrame = GraalOptions.CanOmitFrame && frameMap.frameSize() == frameMap.initialFrameSize && frameMap.registerConfig.getCalleeSaveLayout().registers.length == 0 && !lir.hasArgInCallerFrame() && !lir.hasDebugInfo(); - AbstractAssembler masm = new PTXMacroAssembler(target, frameMap.registerConfig); + AbstractAssembler masm = new PTXAssembler(target, frameMap.registerConfig); HotSpotFrameContext frameContext = omitFrame ? null : new HotSpotFrameContext(); TargetMethodAssembler tasm = new TargetMethodAssembler(target, runtime(), frameMap, masm, frameContext, lir.stubs); tasm.setFrameSize(frameMap.frameSize()); @@ -244,34 +85,30 @@ // Emit the prologue final String name = method.getName(); Buffer codeBuffer = tasm.asm.codeBuffer; - - codeBuffer.emitString(".version 1.4"); - codeBuffer.emitString(".target sm_10"); - //codeBuffer.emitString(".address_size 32"); // PTX ISA version 2.3 codeBuffer.emitString0(".entry " + name + " ("); codeBuffer.emitString(""); - + Signature signature = method.getSignature(); for (int i = 0; i < signature.getParameterCount(false); i++) { - System.err.println(i+": "+signature.getParameterKind(i)); - String param = ".param .u32 param" + i; + System.err.println(i + ": " + signature.getParameterKind(i)); + String param = ".param .u32 param" + i; codeBuffer.emitString(param); } - + codeBuffer.emitString0(") {"); codeBuffer.emitString(""); - + // XXX For now declare one predicate and all registers codeBuffer.emitString(" .reg .pred %p;"); codeBuffer.emitString(" .reg .u32 %r<16>;"); - + // Emit code for the LIR lir.emitCode(tasm); // Emit the epilogue codeBuffer.emitString0("}"); codeBuffer.emitString(""); - + byte[] data = codeBuffer.copyData(0, codeBuffer.position()); System.err.println(new String(data)); }
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXDeoptimizationStub.java Mon Feb 18 14:47:54 2013 -0800 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXDeoptimizationStub.java Mon Feb 18 16:04:25 2013 -0800 @@ -22,8 +22,6 @@ */ package com.oracle.graal.compiler.ptx; -import java.util.*; - import com.oracle.graal.api.code.*; import com.oracle.graal.api.code.RuntimeCallTarget.*; import com.oracle.graal.api.meta.*; @@ -32,7 +30,6 @@ import com.oracle.graal.lir.*; import com.oracle.graal.lir.ptx.*; import com.oracle.graal.lir.asm.*; -import com.oracle.graal.phases.*; public class PTXDeoptimizationStub extends PTXCode { @@ -52,24 +49,8 @@ this.deoptInfo = deoptInfo; } - private static ArrayList<Object> keepAlive = new ArrayList<>(); - @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { - Register scratch = tasm.frameMap.registerConfig.getScratchRegister(); - - masm.bind(label); - if (GraalOptions.CreateDeoptInfo && deoptInfo != null) { -// masm.nop(); - keepAlive.add(deoptInfo.toString()); -// AMD64Move.move(tasm, masm, scratch.asValue(), Constant.forObject(deoptInfo)); -// AMD64Call.directCall(tasm, masm, tasm.runtime.lookupRuntimeCall(SET_DEOPT_INFO), info); - masm.exit(); - } - -// masm.movl(scratch, tasm.runtime.encodeDeoptActionAndReason(action, reason)); -// AMD64Call.directCall(tasm, masm, tasm.runtime.lookupRuntimeCall(DEOPTIMIZE), info); -// AMD64Call.shouldNotReachHere(tasm, masm); + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { masm.exit(); }
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java Mon Feb 18 14:47:54 2013 -0800 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java Mon Feb 18 16:04:25 2013 -0800 @@ -27,11 +27,8 @@ import static com.oracle.graal.lir.ptx.PTXArithmetic.*; import static com.oracle.graal.lir.ptx.PTXCompare.*; import static com.oracle.graal.lir.ptx.PTXBitManipulationOp.IntrinsicOpcode.*; -//import static com.oracle.graal.lir.ptx.PTXMathIntrinsicOp.IntrinsicOpcode.*; -import com.oracle.graal.ptx.*; import com.oracle.graal.api.code.*; -import com.oracle.graal.api.code.RuntimeCallTarget.Descriptor; import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.*; import com.oracle.graal.compiler.gen.*; @@ -46,30 +43,23 @@ import com.oracle.graal.lir.ptx.PTXMove.LoadOp; import com.oracle.graal.lir.ptx.PTXMove.MoveFromRegOp; import com.oracle.graal.lir.ptx.PTXMove.MoveToRegOp; -import com.oracle.graal.lir.ptx.PTXMove.NullCheckOp; import com.oracle.graal.lir.ptx.PTXMove.StoreOp; -import com.oracle.graal.nodes.BreakpointNode; -import com.oracle.graal.nodes.DirectCallTargetNode; -import com.oracle.graal.nodes.IndirectCallTargetNode; -import com.oracle.graal.nodes.StructuredGraph; -import com.oracle.graal.nodes.ValueNode; +import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.calc.Condition; import com.oracle.graal.nodes.calc.ConvertNode; import com.oracle.graal.nodes.extended.IndexedLocationNode; import com.oracle.graal.nodes.extended.LocationNode; -import com.oracle.graal.nodes.java.CompareAndSwapNode; -import com.oracle.graal.phases.util.Util; +import com.oracle.graal.nodes.java.*; /** * This class implements the PTX specific portion of the LIR generator. */ -public abstract class PTXLIRGenerator extends LIRGenerator { +public class PTXLIRGenerator extends LIRGenerator { public static class PTXSpillMoveFactory implements LIR.SpillMoveFactory { @Override public LIRInstruction createMove(Value result, Value input) { -// return new SpillMoveOp(result, input); throw new InternalError("NYI"); } } @@ -188,9 +178,6 @@ @Override public Variable emitLea(Value address) { -// Variable result = newVariable(target().wordKind); -// append(new LeaOp(result, address)); -// return result; throw new InternalError("NYI"); } @@ -201,22 +188,15 @@ @Override public void emitCompareBranch(Value left, Value right, Condition cond, boolean unorderedIsTrue, LabelRef label, LIRFrameState info) { -// boolean mirrored = emitCompare(left, right); -// Condition finalCondition = mirrored ? cond.mirror() : cond; switch (left.getKind().getStackKind()) { case Int: append(new CompareOp(ICMP, cond, left, right)); append(new BranchOp(cond, label, info)); break; -// case Long: case Object: append(new CompareOp(ACMP, cond, left, right)); append(new BranchOp(cond, label, info)); break; -// case Float: -// case Double: -// append(new FloatBranchOp(finalCondition, unorderedIsTrue, label, info)); -// break; default: throw GraalInternalError.shouldNotReachHere("" + left.getKind()); } @@ -224,97 +204,19 @@ @Override public void emitIntegerTestBranch(Value left, Value right, boolean negated, LabelRef label, LIRFrameState info) { -// emitIntegerTest(left, right); -// append(new BranchOp(negated ? Condition.NE : Condition.EQ, label, info)); throw new InternalError("NYI"); } @Override public Variable emitConditionalMove(Value left, Value right, Condition cond, boolean unorderedIsTrue, Value trueValue, Value falseValue) { -// boolean mirrored = emitCompare(left, right); -// Condition finalCondition = mirrored ? cond.mirror() : cond; -// -// Variable result = newVariable(trueValue.getKind()); -// switch (left.getKind().getStackKind()) { -// case Int: -// case Long: -// case Object: -// append(new CondMoveOp(result, finalCondition, load(trueValue), loadNonConst(falseValue))); -// break; -// case Float: -// case Double: -// append(new FloatCondMoveOp(result, finalCondition, unorderedIsTrue, load(trueValue), load(falseValue))); -// break; -// default: -// throw GraalInternalError.shouldNotReachHere("" + left.getKind()); -// } -// return result; throw new InternalError("NYI"); } @Override 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))); -// return result; - throw new InternalError("NYI"); - } - - private void emitIntegerTest(Value a, Value b) { -// assert a.getKind().getStackKind() == Kind.Int || a.getKind() == Kind.Long; -// if (LIRValueUtil.isVariable(b)) { -// append(new AMD64TestOp(load(b), loadNonConst(a))); -// } else { -// append(new AMD64TestOp(load(a), loadNonConst(b))); -// } throw new InternalError("NYI"); } -// /** -// * This method emits the compare instruction, and may reorder the operands. It returns true if -// * it did so. -// * -// * @param a the left operand of the comparison -// * @param b the right operand of the comparison -// * @return true if the left and right operands were switched, false otherwise -// */ -// private boolean emitCompare(Value a, Value b) { -// Variable left; -// Value right; -// boolean mirrored; -// if (LIRValueUtil.isVariable(b)) { -// left = load(b); -// right = loadNonConst(a); -// mirrored = true; -// } else { -// left = load(a); -// right = loadNonConst(b); -// mirrored = false; -// } -// switch (left.getKind().getStackKind()) { -//// case Jsr: -// case Int: -// append(new CompareOp(ICMP, left, right)); -// break; -//// case Long: -//// append(new CompareOp(LCMP, left, right)); -//// break; -//// case Object: -//// append(new CompareOp(ACMP, left, right)); -//// break; -//// case Float: -//// append(new CompareOp(FCMP, left, right)); -//// break; -//// case Double: -//// append(new CompareOp(DCMP, left, right)); -//// break; -// default: -// throw GraalInternalError.shouldNotReachHere(); -// } -// return mirrored; -// } - @Override public Variable emitNegate(Value input) { Variable result = newVariable(input.getKind()); @@ -322,15 +224,6 @@ case Int: append(new Op1Stack(INEG, result, input)); break; -// case Long: -// append(new Op1Stack(LNEG, result, input)); -// break; -// case Float: -// append(new Op2Reg(FXOR, result, input, Constant.forFloat(Float.intBitsToFloat(0x80000000)))); -// break; -// case Double: -// append(new Op2Reg(DXOR, result, input, Constant.forDouble(Double.longBitsToDouble(0x8000000000000000L)))); -// break; default: throw GraalInternalError.shouldNotReachHere(); } @@ -344,15 +237,6 @@ case Int: append(new Op2Stack(IADD, result, a, loadNonConst(b))); break; -// case Long: -// append(new Op2Stack(LADD, result, a, loadNonConst(b))); -// break; -// case Float: -// append(new Op2Stack(FADD, result, a, loadNonConst(b))); -// break; -// case Double: -// append(new Op2Stack(DADD, result, a, loadNonConst(b))); -// break; default: throw GraalInternalError.shouldNotReachHere(); } @@ -366,15 +250,6 @@ case Int: append(new Op2Stack(ISUB, result, a, loadNonConst(b))); break; -// case Long: -// append(new Op2Stack(LSUB, result, a, loadNonConst(b))); -// break; -// case Float: -// append(new Op2Stack(FSUB, result, a, loadNonConst(b))); -// break; -// case Double: -// append(new Op2Stack(DSUB, result, a, loadNonConst(b))); -// break; default: throw GraalInternalError.shouldNotReachHere(); } @@ -388,15 +263,6 @@ case Int: append(new Op2Reg(IMUL, result, a, loadNonConst(b))); break; -// case Long: -// append(new Op2Reg(LMUL, result, a, loadNonConst(b))); -// break; -// case Float: -// append(new Op2Stack(FMUL, result, a, loadNonConst(b))); -// break; -// case Double: -// append(new Op2Stack(DMUL, result, a, loadNonConst(b))); -// break; default: throw GraalInternalError.shouldNotReachHere(); } @@ -409,105 +275,23 @@ return false; } - public Value[] emitIntegerDivRem(Value a, Value b) { -// switch (a.getKind()) { -// case Int: -// emitMove(a, RAX_I); -// append(new DivRemOp(IDIVREM, RAX_I, load(b), state())); -// return new Value[]{emitMove(RAX_I), emitMove(RDX_I)}; -// case Long: -// emitMove(a, RAX_L); -// append(new DivRemOp(LDIVREM, RAX_L, load(b), state())); -// return new Value[]{emitMove(RAX_L), emitMove(RDX_L)}; -// default: -// throw GraalInternalError.shouldNotReachHere(); -// } - throw new InternalError("NYI"); - } - @Override public Value emitDiv(Value a, Value b) { -// switch (a.getKind()) { -// case Int: -// emitMove(a, RAX_I); -// append(new DivOp(IDIV, RAX_I, RAX_I, load(b), state())); -// return emitMove(RAX_I); -// case Long: -// emitMove(a, RAX_L); -// append(new DivOp(LDIV, RAX_L, RAX_L, load(b), state())); -// return emitMove(RAX_L); -// case Float: { -// Variable result = newVariable(a.getKind()); -// append(new Op2Stack(FDIV, result, a, loadNonConst(b))); -// return result; -// } -// case Double: { -// Variable result = newVariable(a.getKind()); -// append(new Op2Stack(DDIV, result, a, loadNonConst(b))); -// return result; -// } -// default: -// throw GraalInternalError.shouldNotReachHere(); -// } throw new InternalError("NYI"); } @Override public Value emitRem(Value a, Value b) { -// switch (a.getKind()) { -// case Int: -// emitMove(a, RAX_I); -// append(new DivOp(IREM, RDX_I, RAX_I, load(b), state())); -// return emitMove(RDX_I); -// case Long: -// emitMove(a, RAX_L); -// append(new DivOp(LREM, RDX_L, RAX_L, load(b), state())); -// return emitMove(RDX_L); -// case Float: { -// RuntimeCallTarget stub = runtime.lookupRuntimeCall(ARITHMETIC_FREM); -// return emitCall(stub, stub.getCallingConvention(), false, a, b); -// } -// case Double: { -// RuntimeCallTarget stub = runtime.lookupRuntimeCall(ARITHMETIC_DREM); -// return emitCall(stub, stub.getCallingConvention(), false, a, b); -// } -// default: -// throw GraalInternalError.shouldNotReachHere(); -// } throw new InternalError("NYI"); } @Override public Variable emitUDiv(Value a, Value b) { -// switch (a.getKind()) { -// case Int: -// emitMove(a, RAX_I); -// append(new DivOp(IUDIV, RAX_I, RAX_I, load(b), state())); -// return emitMove(RAX_I); -// case Long: -// emitMove(a, RAX_L); -// append(new DivOp(LUDIV, RAX_L, RAX_L, load(b), state())); -// return emitMove(RAX_L); -// default: -// throw GraalInternalError.shouldNotReachHere(); -// } throw new InternalError("NYI"); } @Override public Variable emitURem(Value a, Value b) { -// switch (a.getKind()) { -// case Int: -// emitMove(a, RAX_I); -// append(new DivOp(IUREM, RDX_I, RAX_I, load(b), state())); -// return emitMove(RDX_I); -// case Long: -// emitMove(a, RAX_L); -// append(new DivOp(LUREM, RDX_L, RAX_L, load(b), state())); -// return emitMove(RDX_L); -// default: -// throw GraalInternalError.shouldNotReachHere(); -// } throw new InternalError("NYI"); } @@ -518,9 +302,6 @@ case Int: append(new Op2Stack(IAND, result, a, loadNonConst(b))); break; -// case Long: -// append(new Op2Stack(LAND, result, a, loadNonConst(b))); -// break; default: throw GraalInternalError.shouldNotReachHere(); } @@ -529,69 +310,21 @@ @Override public Variable emitOr(Value a, Value b) { -// Variable result = newVariable(a.getKind()); -// switch (a.getKind()) { -// case Int: -// append(new Op2Stack(IOR, result, a, loadNonConst(b))); -// break; -// case Long: -// append(new Op2Stack(LOR, result, a, loadNonConst(b))); -// break; -// default: -// throw GraalInternalError.shouldNotReachHere(); -// } -// return result; throw new InternalError("NYI"); } @Override public Variable emitXor(Value a, Value b) { -// Variable result = newVariable(a.getKind()); -// switch (a.getKind()) { -// case Int: -// append(new Op2Stack(IXOR, result, a, loadNonConst(b))); -// break; -// case Long: -// append(new Op2Stack(LXOR, result, a, loadNonConst(b))); -// break; -// default: -// throw GraalInternalError.shouldNotReachHere(); -// } -// return result; throw new InternalError("NYI"); } @Override public Variable emitShl(Value a, Value b) { -// Variable result = newVariable(a.getKind()); -// switch (a.getKind()) { -// case Int: -// append(new ShiftOp(ISHL, result, a, b)); -// break; -// case Long: -// append(new ShiftOp(LSHL, result, a, b)); -// break; -// default: -// GraalInternalError.shouldNotReachHere(); -// } -// return result; throw new InternalError("NYI"); } @Override public Variable emitShr(Value a, Value b) { -// Variable result = newVariable(a.getKind()); -// switch (a.getKind()) { -// case Int: -// append(new ShiftOp(ISHR, result, a, b)); -// break; -// case Long: -// append(new ShiftOp(LSHR, result, a, b)); -// break; -// default: -// GraalInternalError.shouldNotReachHere(); -// } -// return result; throw new InternalError("NYI"); } @@ -602,9 +335,6 @@ case Int: append(new ShiftOp(IUSHR, result, a, b)); break; -// case Long: -// append(new ShiftOp(LUSHR, result, a, b)); -// break; default: GraalInternalError.shouldNotReachHere(); } @@ -613,127 +343,36 @@ @Override public Variable emitConvert(ConvertNode.Op opcode, Value inputVal) { -// Variable input = load(inputVal); -// Variable result = newVariable(opcode.to); -// switch (opcode) { -// case I2L: -// append(new Op1Reg(I2L, result, input)); -// break; -// case L2I: -// append(new Op1Stack(L2I, result, input)); -// break; -// case I2B: -// append(new Op1Stack(I2B, result, input)); -// break; -// case I2C: -// append(new Op1Stack(I2C, result, input)); -// break; -// case I2S: -// append(new Op1Stack(I2S, result, input)); -// break; -// case F2D: -// append(new Op1Reg(F2D, result, input)); -// break; -// case D2F: -// append(new Op1Reg(D2F, result, input)); -// break; -// case I2F: -// append(new Op1Reg(I2F, result, input)); -// break; -// case I2D: -// append(new Op1Reg(I2D, result, input)); -// break; -// case F2I: -// append(new Op1Reg(F2I, result, input)); -// break; -// case D2I: -// append(new Op1Reg(D2I, result, input)); -// break; -// case L2F: -// append(new Op1Reg(L2F, result, input)); -// break; -// case L2D: -// append(new Op1Reg(L2D, result, input)); -// break; -// case F2L: -// append(new Op1Reg(F2L, result, input)); -// break; -// case D2L: -// append(new Op1Reg(D2L, result, input)); -// break; -// case MOV_I2F: -// append(new Op1Reg(MOV_I2F, result, input)); -// break; -// case MOV_L2D: -// append(new Op1Reg(MOV_L2D, result, input)); -// break; -// case MOV_F2I: -// append(new Op1Reg(MOV_F2I, result, input)); -// break; -// case MOV_D2L: -// append(new Op1Reg(MOV_D2L, result, input)); -// break; -// case UNSIGNED_I2L: -// // Instructions that move or generate 32-bit register values also set the upper 32 -// // bits of the register to zero. -// // Consequently, there is no need for a special zero-extension move. -// emitMove(input, result); -// break; -// default: -// throw GraalInternalError.shouldNotReachHere(); -// } -// return result; throw new InternalError("NYI"); } @Override public void emitDeoptimizeOnOverflow(DeoptimizationAction action, DeoptimizationReason reason, Object deoptInfo) { -// LIRFrameState info = state(); -// LabelRef stubEntry = createDeoptStub(action, reason, info, deoptInfo); -// append(new BranchOp(ConditionFlag.overflow, stubEntry, info)); throw new InternalError("NYI"); } @Override public void emitDeoptimize(DeoptimizationAction action, DeoptimizationReason reason, Object deoptInfo) { -// LIRFrameState info = state(); -// LabelRef stubEntry = createDeoptStub(action, reason, info, deoptInfo); -// append(new JumpOp(stubEntry, info)); - System.err.println("deopt"); + throw new InternalError("NYI"); } @Override public void emitMembar(int barriers) { -// int necessaryBarriers = target.arch.requiredBarriers(barriers); -// if (target.isMP && necessaryBarriers != 0) { -// append(new MembarOp(necessaryBarriers)); -// } throw new InternalError("NYI"); } @Override protected void emitDirectCall(DirectCallTargetNode callTarget, Value result, Value[] parameters, Value[] temps, LIRFrameState callState) { -// append(new DirectCallOp(callTarget.target(), result, parameters, temps, callState)); throw new InternalError("NYI"); } @Override protected void emitIndirectCall(IndirectCallTargetNode callTarget, Value result, Value[] parameters, Value[] temps, LIRFrameState callState) { -// // The current register allocator cannot handle variables at call sites, need a fixed -// // register. -// Value targetAddress = AMD64.rax.asValue(); -// emitMove(operand(callTarget.computedAddress()), targetAddress); -// append(new IndirectCallOp(callTarget.target(), result, parameters, temps, targetAddress, callState)); throw new InternalError("NYI"); } @Override protected void emitCall(RuntimeCallTarget callTarget, Value result, Value[] arguments, Value[] temps, Value targetAddress, LIRFrameState info) { -// if (isConstant(targetAddress)) { -// append(new DirectCallOp(callTarget, result, arguments, temps, info)); -// } else { -// append(new IndirectCallOp(callTarget, result, arguments, temps, targetAddress, info)); -// } throw new InternalError("NYI"); } @@ -748,59 +387,46 @@ @Override public void emitBitScanForward(Variable result, Value value) { -// append(new AMD64BitManipulationOp(BSF, result, value)); throw new InternalError("NYI"); } @Override public void emitBitScanReverse(Variable result, Value value) { -// if (value.getKind().getStackKind() == Kind.Int) { -// append(new AMD64BitManipulationOp(IBSR, result, value)); -// } else { -// append(new AMD64BitManipulationOp(LBSR, result, value)); -// } throw new InternalError("NYI"); } @Override public void emitMathAbs(Variable result, Variable input) { -// append(new Op2Reg(DAND, result, input, Constant.forDouble(Double.longBitsToDouble(0x7FFFFFFFFFFFFFFFL)))); throw new InternalError("NYI"); } @Override public void emitMathSqrt(Variable result, Variable input) { -// append(new AMD64MathIntrinsicOp(SQRT, result, input)); throw new InternalError("NYI"); } @Override public void emitMathLog(Variable result, Variable input, boolean base10) { -// append(new AMD64MathIntrinsicOp(base10 ? LOG10 : LOG, result, input)); throw new InternalError("NYI"); } @Override public void emitMathCos(Variable result, Variable input) { -// append(new AMD64MathIntrinsicOp(COS, result, input)); throw new InternalError("NYI"); } @Override public void emitMathSin(Variable result, Variable input) { -// append(new AMD64MathIntrinsicOp(SIN, result, input)); throw new InternalError("NYI"); } @Override public void emitMathTan(Variable result, Variable input) { -// append(new AMD64MathIntrinsicOp(TAN, result, input)); throw new InternalError("NYI"); } @Override public void emitByteSwap(Variable result, Value input) { -// append(new AMD64ByteSwapOp(result, input)); throw new InternalError("NYI"); } @@ -811,29 +437,16 @@ @Override protected void emitSequentialSwitch(Constant[] keyConstants, LabelRef[] keyTargets, LabelRef defaultTarget, Value key) { -// // Making a copy of the switch value is necessary because jump table destroys the input -// // value -// if (key.getKind() == Kind.Int || key.getKind() == Kind.Long) { -// append(new SequentialSwitchOp(keyConstants, keyTargets, defaultTarget, key, Value.ILLEGAL)); -// } else { -// assert key.getKind() == Kind.Object : key.getKind(); -// append(new SequentialSwitchOp(keyConstants, keyTargets, defaultTarget, key, newVariable(Kind.Object))); -// } throw new InternalError("NYI"); } @Override protected void emitSwitchRanges(int[] lowKeys, int[] highKeys, LabelRef[] targets, LabelRef defaultTarget, Value key) { -// append(new SwitchRangesOp(lowKeys, highKeys, targets, defaultTarget, key)); throw new InternalError("NYI"); } @Override protected void emitTableSwitch(int lowKey, LabelRef defaultTarget, LabelRef[] targets, Value key) { -// // Making a copy of the switch value is necessary because jump table destroys the input -// // value -// Variable tmp = emitMove(key); -// append(new TableSwitchOp(lowKey, defaultTarget, targets, tmp, newVariable(target.wordKind))); throw new InternalError("NYI"); } @@ -847,50 +460,26 @@ @Override protected void emitNullCheckGuard(ValueNode object) { - Variable value = load(operand(object)); - LIRFrameState info = state(); - append(new NullCheckOp(value, info)); + throw new InternalError("NYI"); } @Override public void visitCompareAndSwap(CompareAndSwapNode node) { -// Kind kind = node.newValue().kind(); -// assert kind == node.expected().kind(); -// -// Value expected = loadNonConst(operand(node.expected())); -// Variable newValue = load(operand(node.newValue())); -// -// Address address; -// int displacement = node.displacement(); -// Value index = operand(node.offset()); -// if (isConstant(index) && NumUtil.isInt(asConstant(index).asLong() + displacement)) { -// assert !runtime.needsDataPatch(asConstant(index)); -// displacement += (int) asConstant(index).asLong(); -// address = new Address(kind, load(operand(node.object())), displacement); -// } else { -// address = new Address(kind, load(operand(node.object())), load(index), Address.Scale.Times1, displacement); -// } -// -// RegisterValue rax = AMD64.rax.asValue(kind); -// emitMove(expected, rax); -// append(new CompareAndSwapOp(rax, address, rax, newValue)); -// -// Variable result = newVariable(node.kind()); -// append(new CondMoveOp(result, Condition.EQ, load(Constant.TRUE), Constant.FALSE)); -// setResult(node, result); throw new InternalError("NYI"); } @Override public void visitBreakpointNode(BreakpointNode node) { -// JavaType[] sig = new JavaType[node.arguments.size()]; -// for (int i = 0; i < sig.length; i++) { -// sig[i] = node.arguments.get(i).stamp().javaType(runtime); -// } -// -// CallingConvention cc = frameMap.registerConfig.getCallingConvention(CallingConvention.Type.JavaCall, null, sig, target(), false); -// Value[] parameters = visitInvokeArguments(cc, node.arguments); -// append(new AMD64BreakpointOp(parameters)); + throw new InternalError("NYI"); + } + + @Override + public void visitExceptionObject(ExceptionObjectNode i) { + throw new InternalError("NYI"); + } + + @Override + public void visitSafepointNode(SafepointNode i) { throw new InternalError("NYI"); } }
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXArithmetic.java Mon Feb 18 14:47:54 2013 -0800 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXArithmetic.java Mon Feb 18 16:04:25 2013 -0800 @@ -25,12 +25,9 @@ import static com.oracle.graal.api.code.ValueUtil.*; import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; -import com.oracle.graal.ptx.*; import com.oracle.graal.api.code.*; import com.oracle.graal.api.meta.*; -import com.oracle.graal.asm.*; import com.oracle.graal.asm.ptx.*; -import com.oracle.graal.asm.ptx.PTXAssembler.*; import com.oracle.graal.graph.*; import com.oracle.graal.lir.*; import com.oracle.graal.lir.asm.*; @@ -61,7 +58,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { emit(tasm, masm, opcode, result, x, null); } } @@ -78,7 +75,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { emit(tasm, masm, opcode, result, x, null); } } @@ -97,7 +94,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { emit(tasm, masm, opcode, result, x, y, null); } @@ -122,7 +119,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { emit(tasm, masm, opcode, result, x, y, null); } @@ -147,7 +144,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { if (sameRegister(result, y)) { emit(tasm, masm, opcode, result, x, null); } else { @@ -177,7 +174,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { emit(tasm, masm, opcode, result, x, y, null); } @@ -189,45 +186,11 @@ } } - public static class DivRemOp extends PTXLIRInstruction { - @Opcode private final PTXArithmetic opcode; - @Def protected Value divResult; - @Def protected Value remResult; - @Use protected Value x; - @Alive protected Value y; - @State protected LIRFrameState state; - - public DivRemOp(PTXArithmetic opcode, Value x, Value y, LIRFrameState state) { - this.opcode = opcode; -// this.divResult = PTX.rax.asValue(x.getKind()); -// this.remResult = PTX.rdx.asValue(x.getKind()); - this.x = x; - this.y = y; - this.state = state; - } - - @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { - emit(tasm, masm, opcode, null, y, state); - } - - @Override - protected void verify() { - super.verify(); - // left input in rax, right input in any register but rax and rdx, result quotient in rax, result remainder in rdx -// assert asRegister(x) == PTX.rax; -// assert differentRegisters(y, PTX.rax.asValue(), PTX.rdx.asValue()); - verifyKind(opcode, divResult, x, y); - verifyKind(opcode, remResult, x, y); - } - } - public static class DivOp extends PTXLIRInstruction { @Opcode private final PTXArithmetic opcode; @Def protected Value result; @Use protected Value x; @Alive protected Value y; - @Temp protected Value temp; @State protected LIRFrameState state; public DivOp(PTXArithmetic opcode, Value result, Value x, Value y, LIRFrameState state) { @@ -235,200 +198,47 @@ this.result = result; this.x = x; this.y = y; -// this.temp = asRegister(result) == PTX.rax ? PTX.rdx.asValue(result.getKind()) : PTX.rax.asValue(result.getKind()); this.state = state; } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { emit(tasm, masm, opcode, result, y, state); } @Override protected void verify() { super.verify(); -// // left input in rax, right input in any register but rax and rdx, result quotient in rax, result remainder in rdx -// assert asRegister(x) == PTX.rax; -// assert differentRegisters(y, PTX.rax.asValue(), PTX.rdx.asValue()); -// assert (name().endsWith("DIV") && asRegister(result) == PTX.rax) || (name().endsWith("REM") && asRegister(result) == PTX.rdx); verifyKind(opcode, result, x, y); } } @SuppressWarnings("unused") - protected static void emit(TargetMethodAssembler tasm, PTXMacroAssembler masm, PTXArithmetic opcode, Value result) { + protected static void emit(TargetMethodAssembler tasm, PTXAssembler masm, PTXArithmetic opcode, Value result) { switch (opcode) { -// case LNEG: masm.negq(asLongReg(result)); break; -// case L2I: masm.andl(asIntReg(result), 0xFFFFFFFF); break; -// case I2B: masm.signExtendByte(asIntReg(result)); break; -// case I2C: masm.andl(asIntReg(result), 0xFFFF); break; -// case I2S: masm.signExtendShort(asIntReg(result)); break; default: throw GraalInternalError.shouldNotReachHere(); } } - public static void emit(TargetMethodAssembler tasm, PTXMacroAssembler masm, PTXArithmetic opcode, Value dst, Value src, LIRFrameState info) { + public static void emit(TargetMethodAssembler tasm, PTXAssembler masm, PTXArithmetic opcode, Value dst, Value src, LIRFrameState info) { int exceptionOffset = -1; if (isRegister(src)) { Register a = asIntReg(src); Register d = asIntReg(dst); switch (opcode) { case INEG: masm.neg_s32(d, a); break; -// case IOR: masm.orl(asIntReg(dst), asIntReg(src)); break; -// case IXOR: masm.xorl(asIntReg(dst), asIntReg(src)); break; -// case ISHL: masm.shll(asIntReg(dst)); break; -// case ISHR: masm.sarl(asIntReg(dst)); break; -// -// case LADD: masm.addq(asLongReg(dst), asLongReg(src)); break; -// case LSUB: masm.subq(asLongReg(dst), asLongReg(src)); break; -// case LMUL: masm.imulq(asLongReg(dst), asLongReg(src)); break; -// case LAND: masm.andq(asLongReg(dst), asLongReg(src)); break; -// case LOR: masm.orq(asLongReg(dst), asLongReg(src)); break; -// case LXOR: masm.xorq(asLongReg(dst), asLongReg(src)); break; -// case LSHL: masm.shlq(asLongReg(dst)); break; -// case LSHR: masm.sarq(asLongReg(dst)); break; -// case LUSHR:masm.shrq(asLongReg(dst)); break; -// -// case FADD: masm.addss(asFloatReg(dst), asFloatReg(src)); break; -// case FSUB: masm.subss(asFloatReg(dst), asFloatReg(src)); break; -// case FMUL: masm.mulss(asFloatReg(dst), asFloatReg(src)); break; -// case FDIV: masm.divss(asFloatReg(dst), asFloatReg(src)); break; -// case FAND: masm.andps(asFloatReg(dst), asFloatReg(src)); break; -// case FOR: masm.orps(asFloatReg(dst), asFloatReg(src)); break; -// case FXOR: masm.xorps(asFloatReg(dst), asFloatReg(src)); break; -// -// case DADD: masm.addsd(asDoubleReg(dst), asDoubleReg(src)); break; -// case DSUB: masm.subsd(asDoubleReg(dst), asDoubleReg(src)); break; -// case DMUL: masm.mulsd(asDoubleReg(dst), asDoubleReg(src)); break; -// case DDIV: masm.divsd(asDoubleReg(dst), asDoubleReg(src)); break; -// case DAND: masm.andpd(asDoubleReg(dst), asDoubleReg(src)); break; -// case DOR: masm.orpd(asDoubleReg(dst), asDoubleReg(src)); break; -// case DXOR: masm.xorpd(asDoubleReg(dst), asDoubleReg(src)); break; -// -// case I2L: masm.movslq(asLongReg(dst), asIntReg(src)); break; -// case F2D: masm.cvtss2sd(asDoubleReg(dst), asFloatReg(src)); break; -// case D2F: masm.cvtsd2ss(asFloatReg(dst), asDoubleReg(src)); break; -// case I2F: masm.cvtsi2ssl(asFloatReg(dst), asIntReg(src)); break; -// case I2D: masm.cvtsi2sdl(asDoubleReg(dst), asIntReg(src)); break; -// case L2F: masm.cvtsi2ssq(asFloatReg(dst), asLongReg(src)); break; -// case L2D: masm.cvtsi2sdq(asDoubleReg(dst), asLongReg(src)); break; -// case F2I: -// masm.cvttss2sil(asIntReg(dst), asFloatReg(src)); -// emitConvertFixup(tasm, masm, dst, src); -// break; -// case D2I: -// masm.cvttsd2sil(asIntReg(dst), asDoubleReg(src)); -// emitConvertFixup(tasm, masm, dst, src); -// break; -// case F2L: -// masm.cvttss2siq(asLongReg(dst), asFloatReg(src)); -// emitConvertFixup(tasm, masm, dst, src); -// break; -// case D2L: -// masm.cvttsd2siq(asLongReg(dst), asDoubleReg(src)); -// emitConvertFixup(tasm, masm, dst, src); -// break; -// case MOV_I2F: masm.movdl(asFloatReg(dst), asIntReg(src)); break; -// case MOV_L2D: masm.movdq(asDoubleReg(dst), asLongReg(src)); break; -// case MOV_F2I: masm.movdl(asIntReg(dst), asFloatReg(src)); break; -// case MOV_D2L: masm.movdq(asLongReg(dst), asDoubleReg(src)); break; -// -// case IDIVREM: -// case IDIV: -// case IREM: -// masm.cdql(); -// exceptionOffset = masm.codeBuffer.position(); -// masm.idivl(asRegister(src)); -// break; -// -// case LDIVREM: -// case LDIV: -// case LREM: -// masm.cdqq(); -// exceptionOffset = masm.codeBuffer.position(); -// masm.idivq(asRegister(src)); -// break; -// -// case IUDIV: -// case IUREM: -// // Must zero the high 64-bit word (in RDX) of the dividend -// masm.xorq(PTX.rdx, PTX.rdx); -// exceptionOffset = masm.codeBuffer.position(); -// masm.divl(asRegister(src)); -// break; -// -// case LUDIV: -// case LUREM: -// // Must zero the high 64-bit word (in RDX) of the dividend -// masm.xorq(PTX.rdx, PTX.rdx); -// exceptionOffset = masm.codeBuffer.position(); -// masm.divq(asRegister(src)); -// break; default: throw GraalInternalError.shouldNotReachHere(); } } else if (isConstant(src)) { switch (opcode) { -// case IADD: masm.incrementl(asIntReg(dst), tasm.asIntConst(src)); break; case ISUB: masm.sub_s32(asIntReg(dst), asIntReg(dst), tasm.asIntConst(src)); break; -// case IMUL: masm.imull(asIntReg(dst), asIntReg(dst), tasm.asIntConst(src)); break; case IAND: masm.and_b32(asIntReg(dst), asIntReg(dst), tasm.asIntConst(src)); break; -// case IOR: masm.orl(asIntReg(dst), tasm.asIntConst(src)); break; -// case IXOR: masm.xorl(asIntReg(dst), tasm.asIntConst(src)); break; -// case ISHL: masm.shll(asIntReg(dst), tasm.asIntConst(src) & 31); break; -// case ISHR: masm.sarl(asIntReg(dst), tasm.asIntConst(src) & 31); break; - -// case LADD: masm.addq(asLongReg(dst), tasm.asIntConst(src)); break; -// case LSUB: masm.subq(asLongReg(dst), tasm.asIntConst(src)); break; -// case LMUL: masm.imulq(asLongReg(dst), asLongReg(dst), tasm.asIntConst(src)); break; -// case LAND: masm.andq(asLongReg(dst), tasm.asIntConst(src)); break; -// case LOR: masm.orq(asLongReg(dst), tasm.asIntConst(src)); break; -// case LXOR: masm.xorq(asLongReg(dst), tasm.asIntConst(src)); break; -// case LSHL: masm.shlq(asLongReg(dst), tasm.asIntConst(src) & 63); break; -// case LSHR: masm.sarq(asLongReg(dst), tasm.asIntConst(src) & 63); break; -// case LUSHR:masm.shrq(asLongReg(dst), tasm.asIntConst(src) & 63); break; -// -// case FADD: masm.addss(asFloatReg(dst), tasm.asFloatConstRef(src)); break; -// case FSUB: masm.subss(asFloatReg(dst), tasm.asFloatConstRef(src)); break; -// case FMUL: masm.mulss(asFloatReg(dst), tasm.asFloatConstRef(src)); break; -// case FAND: masm.andps(asFloatReg(dst), tasm.asFloatConstRef(src, 16)); break; -// case FOR: masm.orps(asFloatReg(dst), tasm.asFloatConstRef(src, 16)); break; -// case FXOR: masm.xorps(asFloatReg(dst), tasm.asFloatConstRef(src, 16)); break; -// case FDIV: masm.divss(asFloatReg(dst), tasm.asFloatConstRef(src)); break; -// -// case DADD: masm.addsd(asDoubleReg(dst), tasm.asDoubleConstRef(src)); break; -// case DSUB: masm.subsd(asDoubleReg(dst), tasm.asDoubleConstRef(src)); break; -// case DMUL: masm.mulsd(asDoubleReg(dst), tasm.asDoubleConstRef(src)); break; -// case DDIV: masm.divsd(asDoubleReg(dst), tasm.asDoubleConstRef(src)); break; -// case DAND: masm.andpd(asDoubleReg(dst), tasm.asDoubleConstRef(src, 16)); break; -// case DOR: masm.orpd(asDoubleReg(dst), tasm.asDoubleConstRef(src, 16)); break; -// case DXOR: masm.xorpd(asDoubleReg(dst), tasm.asDoubleConstRef(src, 16)); break; default: throw GraalInternalError.shouldNotReachHere(); } } else { switch (opcode) { -// case IADD: masm.addl(asIntReg(dst), tasm.asIntAddr(src)); break; -// case ISUB: masm.subl(asIntReg(dst), tasm.asIntAddr(src)); break; -// case IAND: masm.andl(asIntReg(dst), tasm.asIntAddr(src)); break; -// case IOR: masm.orl(asIntReg(dst), tasm.asIntAddr(src)); break; -// case IXOR: masm.xorl(asIntReg(dst), tasm.asIntAddr(src)); break; -// -// case LADD: masm.addq(asLongReg(dst), tasm.asLongAddr(src)); break; -// case LSUB: masm.subq(asLongReg(dst), tasm.asLongAddr(src)); break; -// case LAND: masm.andq(asLongReg(dst), tasm.asLongAddr(src)); break; -// case LOR: masm.orq(asLongReg(dst), tasm.asLongAddr(src)); break; -// case LXOR: masm.xorq(asLongReg(dst), tasm.asLongAddr(src)); break; -// -// case FADD: masm.addss(asFloatReg(dst), tasm.asFloatAddr(src)); break; -// case FSUB: masm.subss(asFloatReg(dst), tasm.asFloatAddr(src)); break; -// case FMUL: masm.mulss(asFloatReg(dst), tasm.asFloatAddr(src)); break; -// case FDIV: masm.divss(asFloatReg(dst), tasm.asFloatAddr(src)); break; -// -// case DADD: masm.addsd(asDoubleReg(dst), tasm.asDoubleAddr(src)); break; -// case DSUB: masm.subsd(asDoubleReg(dst), tasm.asDoubleAddr(src)); break; -// case DMUL: masm.mulsd(asDoubleReg(dst), tasm.asDoubleAddr(src)); break; -// case DDIV: masm.divsd(asDoubleReg(dst), tasm.asDoubleAddr(src)); break; default: throw GraalInternalError.shouldNotReachHere(); } } @@ -439,7 +249,7 @@ } } - public static void emit(TargetMethodAssembler tasm, PTXMacroAssembler masm, PTXArithmetic opcode, Value dst, Value src1, Value src2, LIRFrameState info) { + public static void emit(TargetMethodAssembler tasm, PTXAssembler masm, PTXArithmetic opcode, Value dst, Value src1, Value src2, LIRFrameState info) { int exceptionOffset = -1; if (isConstant(src1)) { int a = tasm.asIntConst(src1); @@ -447,7 +257,7 @@ Register d = asIntReg(dst); switch (opcode) { case ISUB: masm.sub_s32(d, a, b); break; - case IAND: masm.and_b32(d, b, a); break; // commutative + case IAND: masm.and_b32(d, b, a); break; default: throw GraalInternalError.shouldNotReachHere(); } } else if (isConstant(src2)) { @@ -478,64 +288,6 @@ } } - private static void emitConvertFixup(TargetMethodAssembler tasm, PTXMacroAssembler masm, Value result, Value x) { -// ConvertSlowPath slowPath = new ConvertSlowPath(result, x); -// tasm.stubs.add(slowPath); -// switch (result.getKind()) { -// case Int: masm.cmpl(asIntReg(result), Integer.MIN_VALUE); break; -// case Long: masm.cmpq(asLongReg(result), tasm.asLongConstRef(Constant.forLong(java.lang.Long.MIN_VALUE))); break; -// default: throw GraalInternalError.shouldNotReachHere(); -// } -// masm.jcc(ConditionFlag.equal, slowPath.start); -// masm.bind(slowPath.continuation); - throw new InternalError("NYI"); - } - -// private static class ConvertSlowPath extends PTXCode { -// public final Label start = new Label(); -// public final Label continuation = new Label(); -// private final Value result; -// private final Value x; -// -// public ConvertSlowPath(Value result, Value x) { -// this.result = result; -// this.x = x; -// } -// -// @Override -// public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { -// masm.bind(start); -// switch (x.getKind()) { -// case Float: masm.ucomiss(asFloatReg(x), tasm.asFloatConstRef(Constant.FLOAT_0)); break; -// case Double: masm.ucomisd(asDoubleReg(x), tasm.asDoubleConstRef(Constant.DOUBLE_0)); break; -// default: throw GraalInternalError.shouldNotReachHere(); -// } -// Label nan = new Label(); -// masm.jcc(ConditionFlag.parity, nan); -// masm.jcc(ConditionFlag.below, continuation); -// -// // input is > 0 -> return maxInt -// // result register already contains 0x80000000, so subtracting 1 gives 0x7fffffff -// switch (result.getKind()) { -// case Int: masm.decrementl(asIntReg(result), 1); break; -// case Long: masm.decrementq(asLongReg(result), 1); break; -// default: throw GraalInternalError.shouldNotReachHere(); -// } -// masm.jmp(continuation); -// -// // input is NaN -> return 0 -// masm.bind(nan); -// masm.xorptr(asRegister(result), asRegister(result)); -// masm.jmp(continuation); -// } -// -// @Override -// public String description() { -// return "convert " + x + " to " + result; -// } -// } - - private static void verifyKind(PTXArithmetic opcode, Value result, Value x, Value y) { assert (opcode.name().startsWith("I") && result.getKind() == Kind.Int && x.getKind().getStackKind() == Kind.Int && y.getKind().getStackKind() == Kind.Int) || (opcode.name().startsWith("L") && result.getKind() == Kind.Long && x.getKind() == Kind.Long && y.getKind() == Kind.Long)
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXBitManipulationOp.java Mon Feb 18 14:47:54 2013 -0800 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXBitManipulationOp.java Mon Feb 18 16:04:25 2013 -0800 @@ -25,6 +25,7 @@ import com.oracle.graal.api.code.*; import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.ptx.*; +import com.oracle.graal.graph.*; import com.oracle.graal.lir.asm.*; public class PTXBitManipulationOp extends PTXLIRInstruction { @@ -44,47 +45,18 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { Register dst = ValueUtil.asIntReg(result); -// if (ValueUtil.isAddress(input)) { -// Address src = ValueUtil.asAddress(input); -// switch (opcode) { -// case IPOPCNT: -// masm.popcntl(dst, src); -// break; -// case LPOPCNT: -// masm.popcntq(dst, src); -// break; -// case BSF: -// masm.bsfq(dst, src); -// break; -// case IBSR: -// masm.bsrl(dst, src); -// break; -// case LBSR: -// masm.bsrq(dst, src); -// break; -// } -// } else { - Register src = ValueUtil.asRegister(input); - switch (opcode) { - case IPOPCNT: - masm.popc_b32(dst, src); - break; - case LPOPCNT: - masm.popc_b64(dst, src); - break; -// case BSF: -// masm.bsfq(dst, src); -// break; -// case IBSR: -// masm.bsrl(dst, src); -// break; -// case LBSR: -// masm.bsrq(dst, src); -// break; - } -// } + Register src = ValueUtil.asRegister(input); + switch (opcode) { + case IPOPCNT: + masm.popc_b32(dst, src); + break; + case LPOPCNT: + masm.popc_b64(dst, src); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } } - }
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXCode.java Mon Feb 18 14:47:54 2013 -0800 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXCode.java Mon Feb 18 16:04:25 2013 -0800 @@ -27,14 +27,14 @@ import com.oracle.graal.lir.asm.*; /** - * Convenience class to provide PTXMacroAssembler for the {@link #emitCode} method. + * Convenience class to provide PTXAssembler for the {@link #emitCode} method. */ public abstract class PTXCode implements LIR.Code { @Override public final void emitCode(TargetMethodAssembler tasm) { - emitCode(tasm, (PTXMacroAssembler) tasm.asm); + emitCode(tasm, (PTXAssembler) tasm.asm); } - public abstract void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm); + public abstract void emitCode(TargetMethodAssembler tasm, PTXAssembler masm); }
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXCompare.java Mon Feb 18 14:47:54 2013 -0800 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXCompare.java Mon Feb 18 16:04:25 2013 -0800 @@ -32,11 +32,11 @@ import com.oracle.graal.lir.asm.*; import com.oracle.graal.nodes.calc.*; -// @formatter:off public enum PTXCompare { ICMP, LCMP, ACMP, FCMP, DCMP; public static class CompareOp extends PTXLIRInstruction { + @Opcode private final PTXCompare opcode; @Use({REG}) protected Value x; @Use({REG, STACK, CONST}) protected Value y; @@ -50,91 +50,177 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { emit(tasm, masm, opcode, condition, x, y); } @Override protected void verify() { super.verify(); - assert (name().startsWith("I") && x.getKind() == Kind.Int && y.getKind().getStackKind() == Kind.Int) - || (name().startsWith("L") && x.getKind() == Kind.Long && y.getKind() == Kind.Long) - || (name().startsWith("A") && x.getKind() == Kind.Object && y.getKind() == Kind.Object) - || (name().startsWith("F") && x.getKind() == Kind.Float && y.getKind() == Kind.Float) - || (name().startsWith("D") && x.getKind() == Kind.Double && y.getKind() == Kind.Double); + assert (name().startsWith("I") && x.getKind() == Kind.Int && y.getKind().getStackKind() == Kind.Int) || (name().startsWith("L") && x.getKind() == Kind.Long && y.getKind() == Kind.Long) || + (name().startsWith("A") && x.getKind() == Kind.Object && y.getKind() == Kind.Object) || + (name().startsWith("F") && x.getKind() == Kind.Float && y.getKind() == Kind.Float) || (name().startsWith("D") && x.getKind() == Kind.Double && y.getKind() == Kind.Double); } } - public static void emit(TargetMethodAssembler tasm, PTXMacroAssembler masm, PTXCompare opcode, Condition condition, Value x, Value y) { + public static void emit(TargetMethodAssembler tasm, PTXAssembler masm, PTXCompare opcode, Condition condition, Value x, Value y) { if (isConstant(x)) { - int a = tasm.asIntConst(x); + int a = tasm.asIntConst(x); Register b = asIntReg(y); switch (opcode) { case ICMP: - switch (condition) { - case EQ: masm.setp_eq_s32(a, b); break; - case NE: masm.setp_ne_s32(a, b); break; - case AE: masm.setp_ge_u32(a, b); break; - case BT: masm.setp_lt_u32(a, b); break; - default: throw GraalInternalError.shouldNotReachHere(); - } + emitCompareConstReg(masm, condition, a, b); break; -// case LCMP: masm.cmpq(asLongReg(x), asLongReg(y)); break; -// case ACMP: masm.cmpptr(asObjectReg(x), asObjectReg(y)); break; -// case FCMP: masm.ucomiss(asFloatReg(x), asFloatReg(y)); break; -// case DCMP: masm.ucomisd(asDoubleReg(x), asDoubleReg(y)); break; - default: throw GraalInternalError.shouldNotReachHere(); + default: + throw GraalInternalError.shouldNotReachHere(); } } else if (isConstant(y)) { Register a = asIntReg(x); - int b = tasm.asIntConst(y); + int b = tasm.asIntConst(y); switch (opcode) { case ICMP: - switch (condition) { - case EQ: masm.setp_eq_s32(a, b); break; - case NE: masm.setp_ne_s32(a, b); break; - default: throw GraalInternalError.shouldNotReachHere(); - } + emitCompareRegConst(masm, condition, a, b); break; -// case ICMP: masm.cmpl(asIntReg(x), tasm.asIntConst(y)); break; -// case LCMP: masm.cmpq(asLongReg(x), tasm.asIntConst(y)); break; case ACMP: if (((Constant) y).isNull()) { switch (condition) { - case EQ: masm.setp_eq_s32(a, b); break; - case NE: masm.setp_ne_s32(a, b); break; - default: throw GraalInternalError.shouldNotReachHere(); + case EQ: + masm.setp_eq_s32(a, b); + break; + case NE: + masm.setp_ne_s32(a, b); + break; + default: + throw GraalInternalError.shouldNotReachHere(); } } else { throw GraalInternalError.shouldNotReachHere("Only null object constants are allowed in comparisons"); } - break; -// case FCMP: masm.ucomiss(asFloatReg(x), tasm.asFloatConstRef(y)); break; -// case DCMP: masm.ucomisd(asDoubleReg(x), tasm.asDoubleConstRef(y)); break; - default: throw GraalInternalError.shouldNotReachHere(); + break; + default: + throw GraalInternalError.shouldNotReachHere(); } } else { Register a = asIntReg(x); Register b = asIntReg(y); switch (opcode) { case ICMP: - switch (condition) { - case EQ: masm.setp_eq_s32(a, b); break; - case NE: masm.setp_ne_s32(a, b); break; - case GE: masm.setp_ge_s32(a, b); break; - case LT: masm.setp_lt_s32(a, b); break; - case AE: masm.setp_ge_u32(a, b); break; - case BT: masm.setp_lt_u32(a, b); break; - default: throw GraalInternalError.shouldNotReachHere(); - } - break; -// case ICMP: masm.cmpl(asIntReg(x), tasm.asIntAddr(y)); break; -// case LCMP: masm.cmpq(asLongReg(x), tasm.asLongAddr(y)); break; -// case ACMP: masm.cmpptr(asObjectReg(x), tasm.asObjectAddr(y)); break; -// case FCMP: masm.ucomiss(asFloatReg(x), tasm.asFloatAddr(y)); break; -// case DCMP: masm.ucomisd(asDoubleReg(x), tasm.asDoubleAddr(y)); break; - default: throw GraalInternalError.shouldNotReachHere(); + emitCompareRegReg(masm, condition, a, b); + break; + default: + throw GraalInternalError.shouldNotReachHere(); } } } + + private static void emitCompareConstReg(PTXAssembler masm, Condition condition, int a, Register b) { + switch (condition) { + case EQ: + masm.setp_eq_s32(a, b); + break; + case NE: + masm.setp_ne_s32(a, b); + break; + case LT: + masm.setp_lt_s32(a, b); + break; + case LE: + masm.setp_le_s32(a, b); + break; + case GT: + masm.setp_gt_s32(a, b); + break; + case GE: + masm.setp_ge_s32(a, b); + break; + case AT: + masm.setp_gt_u32(a, b); + break; + case AE: + masm.setp_ge_u32(a, b); + break; + case BT: + masm.setp_lt_u32(a, b); + break; + case BE: + masm.setp_le_u32(a, b); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + + private static void emitCompareRegConst(PTXAssembler masm, Condition condition, Register a, int b) { + switch (condition) { + case EQ: + masm.setp_eq_s32(a, b); + break; + case NE: + masm.setp_ne_s32(a, b); + break; + case LT: + masm.setp_lt_s32(a, b); + break; + case LE: + masm.setp_le_s32(a, b); + break; + case GT: + masm.setp_gt_s32(a, b); + break; + case GE: + masm.setp_ge_s32(a, b); + break; + case AT: + masm.setp_gt_u32(a, b); + break; + case AE: + masm.setp_ge_u32(a, b); + break; + case BT: + masm.setp_lt_u32(a, b); + break; + case BE: + masm.setp_le_u32(a, b); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + + private static void emitCompareRegReg(PTXAssembler masm, Condition condition, Register a, Register b) { + switch (condition) { + case EQ: + masm.setp_eq_s32(a, b); + break; + case NE: + masm.setp_ne_s32(a, b); + break; + case LT: + masm.setp_lt_s32(a, b); + break; + case LE: + masm.setp_le_s32(a, b); + break; + case GT: + masm.setp_gt_s32(a, b); + break; + case GE: + masm.setp_ge_s32(a, b); + break; + case AT: + masm.setp_gt_u32(a, b); + break; + case AE: + masm.setp_ge_u32(a, b); + break; + case BT: + masm.setp_lt_u32(a, b); + break; + case BE: + masm.setp_le_u32(a, b); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } }
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java Mon Feb 18 14:47:54 2013 -0800 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java Mon Feb 18 16:04:25 2013 -0800 @@ -22,27 +22,19 @@ */ package com.oracle.graal.lir.ptx; -import static com.oracle.graal.api.code.ValueUtil.*; import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; -import com.oracle.graal.ptx.*; -import com.oracle.graal.api.code.*; -import com.oracle.graal.api.code.Address.Scale; -import com.oracle.graal.api.code.CompilationResult.JumpTable; import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.*; import com.oracle.graal.asm.ptx.*; -import com.oracle.graal.graph.*; import com.oracle.graal.lir.*; -import com.oracle.graal.lir.LIRInstruction.Opcode; -import com.oracle.graal.lir.StandardOp.FallThroughOp; import com.oracle.graal.lir.asm.*; import com.oracle.graal.nodes.calc.*; -// @formatter:off public class PTXControlFlow { public static class ReturnOp extends PTXLIRInstruction { + @Use({REG, ILLEGAL}) protected Value x; public ReturnOp(Value x) { @@ -50,7 +42,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { if (tasm.frameContext != null) { tasm.frameContext.leave(tasm); } @@ -58,8 +50,8 @@ } } + public static class BranchOp extends PTXLIRInstruction implements StandardOp.BranchOp { - public static class BranchOp extends PTXLIRInstruction implements StandardOp.BranchOp { protected Condition condition; protected LabelRef destination; @State protected LIRFrameState state; @@ -70,14 +62,12 @@ this.state = state; } - public static final String UNBOUND_TARGET = "L"+Integer.MAX_VALUE; - @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { masm.at(); Label l = destination.label(); l.addPatchAt(tasm.asm.codeBuffer.position()); - String target = l.isBound() ? "L"+l.toString() : UNBOUND_TARGET; + String target = l.isBound() ? "L" + l.toString() : AbstractPTXAssembler.UNBOUND_TARGET; masm.bra(target); } @@ -92,380 +82,4 @@ condition = condition.negate(); } } - - -// public static class FloatBranchOp extends BranchOp { -// protected boolean unorderedIsTrue; -// -// public FloatBranchOp(Condition condition, boolean unorderedIsTrue, LabelRef destination, LIRFrameState info) { -// super(floatCond(condition), destination, info); -// this.unorderedIsTrue = unorderedIsTrue; -// } -// -// @Override -// public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) { -// floatJcc(masm, condition, unorderedIsTrue, destination.label()); -// } -// -// @Override -// public void negate(LabelRef newDestination) { -// super.negate(newDestination); -// unorderedIsTrue = !unorderedIsTrue; -// } -// } -// -// -// public static class TableSwitchOp extends AMD64LIRInstruction { -// private final int lowKey; -// private final LabelRef defaultTarget; -// private final LabelRef[] targets; -// @Alive protected Value index; -// @Temp protected Value scratch; -// -// public TableSwitchOp(final int lowKey, final LabelRef defaultTarget, final LabelRef[] targets, Variable index, Variable scratch) { -// this.lowKey = lowKey; -// this.defaultTarget = defaultTarget; -// this.targets = targets; -// this.index = index; -// this.scratch = scratch; -// } -// -// @Override -// public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) { -// tableswitch(tasm, masm, lowKey, defaultTarget, targets, asIntReg(index), asLongReg(scratch)); -// } -// } -// -// public static class SequentialSwitchOp extends AMD64LIRInstruction implements FallThroughOp { -// @Use({CONST}) protected Constant[] keyConstants; -// private final LabelRef[] keyTargets; -// private LabelRef defaultTarget; -// @Alive({REG}) protected Value key; -// @Temp({REG, ILLEGAL}) protected Value scratch; -// -// public SequentialSwitchOp(Constant[] keyConstants, LabelRef[] keyTargets, LabelRef defaultTarget, Value key, Value scratch) { -// assert keyConstants.length == keyTargets.length; -// this.keyConstants = keyConstants; -// this.keyTargets = keyTargets; -// this.defaultTarget = defaultTarget; -// this.key = key; -// this.scratch = scratch; -// } -// -// @Override -// public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) { -// if (key.getKind() == Kind.Int) { -// Register intKey = asIntReg(key); -// for (int i = 0; i < keyConstants.length; i++) { -// if (tasm.runtime.needsDataPatch(keyConstants[i])) { -// tasm.recordDataReferenceInCode(keyConstants[i], 0, true); -// } -// long lc = keyConstants[i].asLong(); -// assert NumUtil.isInt(lc); -// masm.cmpl(intKey, (int) lc); -// masm.jcc(ConditionFlag.equal, keyTargets[i].label()); -// } -// } else if (key.getKind() == Kind.Long) { -// Register longKey = asLongReg(key); -// for (int i = 0; i < keyConstants.length; i++) { -// masm.cmpq(longKey, tasm.asLongConstRef(keyConstants[i])); -// masm.jcc(ConditionFlag.equal, keyTargets[i].label()); -// } -// } else if (key.getKind() == Kind.Object) { -// Register intKey = asObjectReg(key); -// Register temp = asObjectReg(scratch); -// for (int i = 0; i < keyConstants.length; i++) { -// AMD64Move.move(tasm, masm, temp.asValue(Kind.Object), keyConstants[i]); -// masm.cmpptr(intKey, temp); -// masm.jcc(ConditionFlag.equal, keyTargets[i].label()); -// } -// } else { -// throw new GraalInternalError("sequential switch only supported for int, long and object"); -// } -// if (defaultTarget != null) { -// masm.jmp(defaultTarget.label()); -// } else { -// masm.hlt(); -// } -// } -// -// @Override -// public LabelRef fallThroughTarget() { -// return defaultTarget; -// } -// -// @Override -// public void setFallThroughTarget(LabelRef target) { -// defaultTarget = target; -// } -// } -// -// public static class SwitchRangesOp extends AMD64LIRInstruction implements FallThroughOp { -// private final LabelRef[] keyTargets; -// private LabelRef defaultTarget; -// private final int[] lowKeys; -// private final int[] highKeys; -// @Alive protected Value key; -// -// public SwitchRangesOp(int[] lowKeys, int[] highKeys, LabelRef[] keyTargets, LabelRef defaultTarget, Value key) { -// this.lowKeys = lowKeys; -// this.highKeys = highKeys; -// this.keyTargets = keyTargets; -// this.defaultTarget = defaultTarget; -// this.key = key; -// } -// -// @Override -// public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) { -// assert isSorted(lowKeys) && isSorted(highKeys); -// -// Label actualDefaultTarget = defaultTarget == null ? new Label() : defaultTarget.label(); -// int prevHighKey = 0; -// boolean skipLowCheck = false; -// for (int i = 0; i < lowKeys.length; i++) { -// int lowKey = lowKeys[i]; -// int highKey = highKeys[i]; -// if (lowKey == highKey) { -// masm.cmpl(asIntReg(key), lowKey); -// masm.jcc(ConditionFlag.equal, keyTargets[i].label()); -// skipLowCheck = false; -// } else { -// if (!skipLowCheck || (prevHighKey + 1) != lowKey) { -// masm.cmpl(asIntReg(key), lowKey); -// masm.jcc(ConditionFlag.less, actualDefaultTarget); -// } -// masm.cmpl(asIntReg(key), highKey); -// masm.jcc(ConditionFlag.lessEqual, keyTargets[i].label()); -// skipLowCheck = true; -// } -// prevHighKey = highKey; -// } -// if (defaultTarget != null) { -// masm.jmp(defaultTarget.label()); -// } else { -// masm.bind(actualDefaultTarget); -// masm.hlt(); -// } -// } -// -// @Override -// protected void verify() { -// super.verify(); -// assert lowKeys.length == keyTargets.length; -// assert highKeys.length == keyTargets.length; -// assert key.getKind() == Kind.Int; -// } -// -// @Override -// public LabelRef fallThroughTarget() { -// return defaultTarget; -// } -// -// @Override -// public void setFallThroughTarget(LabelRef target) { -// defaultTarget = target; -// } -// -// private static boolean isSorted(int[] values) { -// for (int i = 1; i < values.length; i++) { -// if (values[i - 1] >= values[i]) { -// return false; -// } -// } -// return true; -// } -// } -// -// -// @Opcode("CMOVE") -// public static class CondMoveOp extends AMD64LIRInstruction { -// @Def({REG, HINT}) protected Value result; -// @Alive({REG}) protected Value trueValue; -// @Use({REG, STACK, CONST}) protected Value falseValue; -// private final ConditionFlag condition; -// -// public CondMoveOp(Variable result, Condition condition, Variable trueValue, Value falseValue) { -// this.result = result; -// this.condition = intCond(condition); -// this.trueValue = trueValue; -// this.falseValue = falseValue; -// } -// -// @Override -// public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) { -// cmove(tasm, masm, result, false, condition, false, trueValue, falseValue); -// } -// } -// -// -// @Opcode("CMOVE") -// public static class FloatCondMoveOp extends AMD64LIRInstruction { -// @Def({REG}) protected Value result; -// @Alive({REG}) protected Value trueValue; -// @Alive({REG}) protected Value falseValue; -// private final ConditionFlag condition; -// private final boolean unorderedIsTrue; -// -// public FloatCondMoveOp(Variable result, Condition condition, boolean unorderedIsTrue, Variable trueValue, Variable falseValue) { -// this.result = result; -// this.condition = floatCond(condition); -// this.unorderedIsTrue = unorderedIsTrue; -// this.trueValue = trueValue; -// this.falseValue = falseValue; -// } -// -// @Override -// public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) { -// cmove(tasm, masm, result, true, condition, unorderedIsTrue, trueValue, falseValue); -// } -// } -// -// private static void tableswitch(TargetMethodAssembler tasm, AMD64MacroAssembler masm, int lowKey, LabelRef defaultTarget, LabelRef[] targets, Register value, Register scratch) { -// Buffer buf = masm.codeBuffer; -// // Compare index against jump table bounds -// int highKey = lowKey + targets.length - 1; -// if (lowKey != 0) { -// // subtract the low value from the switch value -// masm.subl(value, lowKey); -// masm.cmpl(value, highKey - lowKey); -// } else { -// masm.cmpl(value, highKey); -// } -// -// // Jump to default target if index is not within the jump table -// if (defaultTarget != null) { -// masm.jcc(ConditionFlag.above, defaultTarget.label()); -// } -// -// // Set scratch to address of jump table -// int leaPos = buf.position(); -// masm.leaq(scratch, new Address(tasm.target.wordKind, AMD64.rip.asValue(), 0)); -// int afterLea = buf.position(); -// -// // Load jump table entry into scratch and jump to it -// masm.movslq(value, new Address(Kind.Int, scratch.asValue(), value.asValue(), Scale.Times4, 0)); -// masm.addq(scratch, value); -// masm.jmp(scratch); -// -// // Inserting padding so that jump table address is 4-byte aligned -// if ((buf.position() & 0x3) != 0) { -// masm.nop(4 - (buf.position() & 0x3)); -// } -// -// // Patch LEA instruction above now that we know the position of the jump table -// int jumpTablePos = buf.position(); -// buf.setPosition(leaPos); -// masm.leaq(scratch, new Address(tasm.target.wordKind, AMD64.rip.asValue(), jumpTablePos - afterLea)); -// buf.setPosition(jumpTablePos); -// -// // Emit jump table entries -// for (LabelRef target : targets) { -// Label label = target.label(); -// int offsetToJumpTableBase = buf.position() - jumpTablePos; -// if (label.isBound()) { -// int imm32 = label.position() - jumpTablePos; -// buf.emitInt(imm32); -// } else { -// label.addPatchAt(buf.position()); -// -// buf.emitByte(0); // pseudo-opcode for jump table entry -// buf.emitShort(offsetToJumpTableBase); -// buf.emitByte(0); // padding to make jump table entry 4 bytes wide -// } -// } -// -// JumpTable jt = new JumpTable(jumpTablePos, lowKey, highKey, 4); -// tasm.targetMethod.addAnnotation(jt); -// } -// -// private static void floatJcc(AMD64MacroAssembler masm, ConditionFlag condition, boolean unorderedIsTrue, Label label) { -// Label endLabel = new Label(); -// if (unorderedIsTrue && !trueOnUnordered(condition)) { -// masm.jcc(ConditionFlag.parity, label); -// } else if (!unorderedIsTrue && trueOnUnordered(condition)) { -// masm.jcc(ConditionFlag.parity, endLabel); -// } -// masm.jcc(condition, label); -// masm.bind(endLabel); -// } -// -// private static void cmove(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Value result, boolean isFloat, ConditionFlag condition, boolean unorderedIsTrue, Value trueValue, Value falseValue) { -// // check that we don't overwrite an input operand before it is used. -// assert !result.equals(trueValue); -// -// AMD64Move.move(tasm, masm, result, falseValue); -// cmove(tasm, masm, result, condition, trueValue); -// -// if (isFloat) { -// if (unorderedIsTrue && !trueOnUnordered(condition)) { -// cmove(tasm, masm, result, ConditionFlag.parity, trueValue); -// } else if (!unorderedIsTrue && trueOnUnordered(condition)) { -// cmove(tasm, masm, result, ConditionFlag.parity, falseValue); -// } -// } -// } -// -// private static void cmove(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Value result, ConditionFlag cond, Value other) { -// if (isRegister(other)) { -// assert asRegister(other) != asRegister(result) : "other already overwritten by previous move"; -// switch (other.getKind()) { -// case Int: masm.cmovl(cond, asRegister(result), asRegister(other)); break; -// case Long: masm.cmovq(cond, asRegister(result), asRegister(other)); break; -// default: throw GraalInternalError.shouldNotReachHere(); -// } -// } else { -// switch (other.getKind()) { -// case Int: masm.cmovl(cond, asRegister(result), tasm.asAddress(other)); break; -// case Long: masm.cmovq(cond, asRegister(result), tasm.asAddress(other)); break; -// default: throw GraalInternalError.shouldNotReachHere(); -// } -// } -// } -// -// private static ConditionFlag intCond(Condition cond) { -// switch (cond) { -// case EQ: return ConditionFlag.equal; -// case NE: return ConditionFlag.notEqual; -// case LT: return ConditionFlag.less; -// case LE: return ConditionFlag.lessEqual; -// case GE: return ConditionFlag.greaterEqual; -// case GT: return ConditionFlag.greater; -// case BE: return ConditionFlag.belowEqual; -// case AE: return ConditionFlag.aboveEqual; -// case AT: return ConditionFlag.above; -// case BT: return ConditionFlag.below; -// default: throw GraalInternalError.shouldNotReachHere(); -// } -// } -// -// private static ConditionFlag floatCond(Condition cond) { -// switch (cond) { -// case EQ: return ConditionFlag.equal; -// case NE: return ConditionFlag.notEqual; -// case LT: return ConditionFlag.below; -// case LE: return ConditionFlag.belowEqual; -// case GE: return ConditionFlag.aboveEqual; -// case GT: return ConditionFlag.above; -// default: throw GraalInternalError.shouldNotReachHere(); -// } -// } -// -// private static boolean trueOnUnordered(ConditionFlag condition) { -// switch(condition) { -// case aboveEqual: -// case notEqual: -// case above: -// case less: -// case overflow: -// return false; -// case equal: -// case belowEqual: -// case below: -// case greaterEqual: -// case noOverflow: -// return true; -// default: -// throw GraalInternalError.shouldNotReachHere(); -// } -// } }
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXLIRInstruction.java Mon Feb 18 14:47:54 2013 -0800 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXLIRInstruction.java Mon Feb 18 16:04:25 2013 -0800 @@ -27,14 +27,14 @@ import com.oracle.graal.lir.asm.*; /** - * Convenience class to provide PTXMacroAssembler for the {@link #emitCode} method. + * Convenience class to provide PTXAssembler for the {@link #emitCode} method. */ public abstract class PTXLIRInstruction extends LIRInstruction { @Override public final void emitCode(TargetMethodAssembler tasm) { - emitCode(tasm, (PTXMacroAssembler) tasm.asm); + emitCode(tasm, (PTXAssembler) tasm.asm); } - public abstract void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm); + public abstract void emitCode(TargetMethodAssembler tasm, PTXAssembler masm); }
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java Mon Feb 18 14:47:54 2013 -0800 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java Mon Feb 18 16:04:25 2013 -0800 @@ -24,13 +24,9 @@ import static com.oracle.graal.api.code.ValueUtil.*; import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; -import static java.lang.Double.*; -import static java.lang.Float.*; -import com.oracle.graal.ptx.*; import com.oracle.graal.api.code.*; import com.oracle.graal.api.meta.*; -import com.oracle.graal.asm.*; import com.oracle.graal.asm.ptx.*; import com.oracle.graal.graph.*; import com.oracle.graal.lir.*; @@ -38,11 +34,11 @@ import com.oracle.graal.lir.StandardOp.MoveOp; import com.oracle.graal.lir.asm.*; -// @formatter:off public class PTXMove { @Opcode("MOVE") public static class SpillMoveOp extends PTXLIRInstruction implements MoveOp { + @Def({REG, STACK}) protected Value result; @Use({REG, STACK, CONST}) protected Value input; @@ -52,7 +48,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { move(tasm, masm, getResult(), getInput()); } @@ -60,15 +56,16 @@ public Value getInput() { return input; } + @Override public Value getResult() { return result; } } - @Opcode("MOVE") public static class MoveToRegOp extends PTXLIRInstruction implements MoveOp { + @Def({REG, HINT}) protected Value result; @Use({REG, STACK, CONST}) protected Value input; @@ -78,7 +75,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { move(tasm, masm, getResult(), getInput()); } @@ -86,15 +83,16 @@ public Value getInput() { return input; } + @Override public Value getResult() { return result; } } - @Opcode("MOVE") public static class MoveFromRegOp extends PTXLIRInstruction implements MoveOp { + @Def({REG, STACK}) protected Value result; @Use({REG, CONST, HINT}) protected Value input; @@ -104,7 +102,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { move(tasm, masm, getResult(), getInput()); } @@ -112,14 +110,15 @@ public Value getInput() { return input; } + @Override public Value getResult() { return result; } } + public static class LoadOp extends PTXLIRInstruction { - public static class LoadOp extends PTXLIRInstruction { @Def({REG}) protected Value result; @Use({ADDR}) protected Value address; @State protected LIRFrameState state; @@ -131,13 +130,13 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { load(tasm, masm, result, (Address) address, state); } } + public static class StoreOp extends PTXLIRInstruction { - public static class StoreOp extends PTXLIRInstruction { @Use({ADDR}) protected Value address; @Use({REG, CONST}) protected Value input; @State protected LIRFrameState state; @@ -149,13 +148,13 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { store(tasm, masm, (Address) address, input, state); } } + public static class LeaOp extends PTXLIRInstruction { - public static class LeaOp extends PTXLIRInstruction { @Def({REG}) protected Value result; @Use({ADDR, STACK, UNINITIALIZED}) protected Value address; @@ -165,48 +164,14 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { -// masm.leaq(asLongReg(result), tasm.asAddress(address)); - throw new InternalError("NYI"); - } - } - - - public static class MembarOp extends PTXLIRInstruction { - private final int barriers; - - public MembarOp(final int barriers) { - this.barriers = barriers; - } - - @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { -// masm.membar(barriers); + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { throw new InternalError("NYI"); } } - - public static class NullCheckOp extends PTXLIRInstruction { - @Use protected Value input; - @State protected LIRFrameState state; - - public NullCheckOp(Variable input, LIRFrameState state) { - this.input = input; - this.state = state; - } - - @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { -// tasm.recordImplicitException(masm.codeBuffer.position(), state); -// masm.nullCheck(asRegister(input)); - throw new InternalError("NYI"); - } - } - - @Opcode("CAS") public static class CompareAndSwapOp extends PTXLIRInstruction { + @Def protected Value result; @Use({ADDR}) protected Value address; @Use protected Value cmpValue; @@ -220,32 +185,21 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { compareAndSwap(tasm, masm, result, (Address) address, cmpValue, newValue); } } - - public static void move(TargetMethodAssembler tasm, PTXMacroAssembler masm, Value result, Value input) { + public static void move(TargetMethodAssembler tasm, PTXAssembler masm, Value result, Value input) { if (isRegister(input)) { if (isRegister(result)) { reg2reg(masm, result, input); - } else if (isStackSlot(result)) { - reg2stack(tasm, masm, result, input); - } else { - throw GraalInternalError.shouldNotReachHere(); - } - } else if (isStackSlot(input)) { - if (isRegister(result)) { - stack2reg(tasm, masm, result, input); } else { throw GraalInternalError.shouldNotReachHere(); } } else if (isConstant(input)) { if (isRegister(result)) { const2reg(tasm, masm, result, (Constant) input); - } else if (isStackSlot(result)) { - const2stack(tasm, masm, result, (Constant) input); } else { throw GraalInternalError.shouldNotReachHere(); } @@ -254,50 +208,23 @@ } } - private static void reg2reg(PTXMacroAssembler masm, Value result, Value input) { + private static void reg2reg(PTXAssembler masm, Value result, Value input) { if (asRegister(input).equals(asRegister(result))) { return; } switch (input.getKind()) { - case Int: masm.mov_s32(asRegister(result), asRegister(input)); break; -// case Long: masm.movq(asRegister(result), asRegister(input)); break; -// case Float: masm.movflt(asFloatReg(result), asFloatReg(input)); break; -// case Double: masm.movdbl(asDoubleReg(result), asDoubleReg(input)); break; - case Object: masm.mov_u64(asRegister(result), asRegister(input)); break; - default: throw GraalInternalError.shouldNotReachHere("kind=" + result.getKind()); + case Int: + masm.mov_s32(asRegister(result), asRegister(input)); + break; + case Object: + masm.mov_u64(asRegister(result), asRegister(input)); + break; + default: + throw GraalInternalError.shouldNotReachHere("kind=" + result.getKind()); } } - private static void reg2stack(TargetMethodAssembler tasm, PTXMacroAssembler masm, Value result, Value input) { -// switch (input.getKind()) { -// case Jsr: -// case Int: masm.movl(tasm.asAddress(result), asRegister(input)); break; -// case Long: masm.movq(tasm.asAddress(result), asRegister(input)); break; -// case Float: masm.movflt(tasm.asAddress(result), asFloatReg(input)); break; -// case Double: masm.movsd(tasm.asAddress(result), asDoubleReg(input)); break; -// case Object: masm.movq(tasm.asAddress(result), asRegister(input)); break; -// default: throw GraalInternalError.shouldNotReachHere(); -// } - throw new InternalError("NYI"); - } - - private static void stack2reg(TargetMethodAssembler tasm, PTXMacroAssembler masm, Value result, Value input) { -// switch (input.getKind()) { -// case Jsr: -// case Int: masm.movl(asRegister(result), tasm.asAddress(input)); break; -// case Long: masm.movq(asRegister(result), tasm.asAddress(input)); break; -// case Float: masm.movflt(asFloatReg(result), tasm.asAddress(input)); break; -// case Double: masm.movdbl(asDoubleReg(result), tasm.asAddress(input)); break; -// case Object: masm.movq(asRegister(result), tasm.asAddress(input)); break; -// default: throw GraalInternalError.shouldNotReachHere(); -// } - throw new InternalError("NYI"); - } - - private static void const2reg(TargetMethodAssembler tasm, PTXMacroAssembler masm, Value result, Constant input) { - // Note: we use the kind of the input operand (and not the kind of the result operand) because they don't match - // in all cases. For example, an object constant can be loaded to a long register when unsafe casts occurred (e.g., - // for a write barrier where arithmetic operations are then performed on the pointer). + private static void const2reg(TargetMethodAssembler tasm, PTXAssembler masm, Value result, Constant input) { switch (input.getKind().getStackKind()) { case Jsr: case Int: @@ -306,139 +233,42 @@ } masm.mov_s32(asRegister(result), input.asInt()); break; -// case Long: -// if (tasm.runtime.needsDataPatch(input)) { -// tasm.recordDataReferenceInCode(input, 0, true); -// } -// // Do not optimize with an XOR as this instruction may be between -// // a CMP and a Jcc in which case the XOR will modify the condition -// // flags and interfere with the Jcc. -// masm.movq(asRegister(result), input.asLong()); -// break; -// case Float: -// // This is *not* the same as 'constant == 0.0f' in the case where constant is -0.0f -// if (Float.floatToRawIntBits(input.asFloat()) == Float.floatToRawIntBits(0.0f)) { -// assert !tasm.runtime.needsDataPatch(input); -// masm.xorps(asFloatReg(result), asFloatReg(result)); -// } else { -// masm.movflt(asFloatReg(result), tasm.asFloatConstRef(input)); -// } -// break; -// case Double: -// // This is *not* the same as 'constant == 0.0d' in the case where constant is -0.0d -// if (Double.doubleToRawLongBits(input.asDouble()) == Double.doubleToRawLongBits(0.0d)) { -// assert !tasm.runtime.needsDataPatch(input); -// masm.xorpd(asDoubleReg(result), asDoubleReg(result)); -// } else { -// masm.movdbl(asDoubleReg(result), tasm.asDoubleConstRef(input)); -// } -// break; -// case Object: -// // Do not optimize with an XOR as this instruction may be between -// // a CMP and a Jcc in which case the XOR will modify the condition -// // flags and interfere with the Jcc. -// if (input.isNull()) { -// masm.movq(asRegister(result), 0x0L); -// } else if (tasm.target.inlineObjects) { -// tasm.recordDataReferenceInCode(input, 0, true); -// masm.movq(asRegister(result), 0xDEADDEADDEADDEADL); -// } else { -// masm.movq(asRegister(result), tasm.recordDataReferenceInCode(input, 0, false)); -// } -// break; default: throw GraalInternalError.shouldNotReachHere(); } } - private static void const2stack(TargetMethodAssembler tasm, PTXMacroAssembler masm, Value result, Constant input) { -// assert !tasm.runtime.needsDataPatch(input); -// switch (input.getKind().getStackKind()) { -// case Jsr: -// case Int: masm.movl(tasm.asAddress(result), input.asInt()); break; -// case Long: masm.movlong(tasm.asAddress(result), input.asLong()); break; -// case Float: masm.movl(tasm.asAddress(result), floatToRawIntBits(input.asFloat())); break; -// case Double: masm.movlong(tasm.asAddress(result), doubleToRawLongBits(input.asDouble())); break; -// case Object: -// if (input.isNull()) { -// masm.movlong(tasm.asAddress(result), 0L); -// } else { -// throw GraalInternalError.shouldNotReachHere("Non-null object constants must be in register"); -// } -// break; -// default: -// throw GraalInternalError.shouldNotReachHere(); -// } - throw new InternalError("NYI"); - } - - - public static void load(TargetMethodAssembler tasm, PTXMacroAssembler masm, Value result, Address loadAddr, LIRFrameState info) { -// if (info != null) { -// tasm.recordImplicitException(masm.codeBuffer.position(), info); -// } + @SuppressWarnings("unused") + public static void load(TargetMethodAssembler tasm, PTXAssembler masm, Value result, Address loadAddr, LIRFrameState info) { Register a = asRegister(loadAddr.getBase()); int immOff = loadAddr.getDisplacement(); switch (loadAddr.getKind()) { -// case Boolean: -// case Byte: masm.movsxb(asRegister(result), loadAddr); break; -// case Char: masm.movzxl(asRegister(result), loadAddr); break; -// case Short: masm.movswl(asRegister(result), loadAddr); break; - case Int: masm.ld_global_s32(asRegister(result), a, immOff); break; -// case Long: masm.movq(asRegister(result), loadAddr); break; -// case Float: masm.movflt(asFloatReg(result), loadAddr); break; -// case Double: masm.movdbl(asDoubleReg(result), loadAddr); break; - case Object: masm.ld_global_u32(asRegister(result), a, immOff); break; - default: throw GraalInternalError.shouldNotReachHere(); + case Int: + masm.ld_global_s32(asRegister(result), a, immOff); + break; + case Object: + masm.ld_global_u32(asRegister(result), a, immOff); + break; + default: + throw GraalInternalError.shouldNotReachHere(); } } - public static void store(TargetMethodAssembler tasm, PTXMacroAssembler masm, Address storeAddr, Value input, LIRFrameState info) { -// if (info != null) { -// tasm.recordImplicitException(masm.codeBuffer.position(), info); -// } -// + @SuppressWarnings("unused") + public static void store(TargetMethodAssembler tasm, PTXAssembler masm, Address storeAddr, Value input, LIRFrameState info) { Register a = asRegister(storeAddr.getBase()); int immOff = storeAddr.getDisplacement(); - System.out.println("store: "+storeAddr); if (isRegister(input)) { switch (storeAddr.getKind()) { -// case Boolean: -// case Byte: masm.movb(storeAddr, asRegister(input)); break; -// case Char: -// case Short: masm.movw(storeAddr, asRegister(input)); break; - case Int: masm.st_global_s32(a, immOff, asRegister(input)); break; -// case Long: masm.movq(storeAddr, asRegister(input)); break; -// case Float: masm.movflt(storeAddr, asFloatReg(input)); break; -// case Double: masm.movsd(storeAddr, asDoubleReg(input)); break; -// case Object: masm.movq(storeAddr, asRegister(input)); break; - default: throw GraalInternalError.shouldNotReachHere(); + case Int: + masm.st_global_s32(a, immOff, asRegister(input)); + break; + default: + throw GraalInternalError.shouldNotReachHere(); } } else if (isConstant(input)) { Constant c = (Constant) input; switch (storeAddr.getKind()) { -// case Boolean: -// case Byte: masm.movb(storeAddr, c.asInt() & 0xFF); break; -// case Char: -// case Short: masm.movw(storeAddr, c.asInt() & 0xFFFF); break; -// case Jsr: -// case Int: masm.st_st32(a, immOff, c.asInt()); break; -// case Long: -// if (NumUtil.isInt(c.asLong())) { -// masm.movslq(storeAddr, (int) c.asLong()); -// } else { -// throw GraalInternalError.shouldNotReachHere("Cannot store 64-bit constants to memory"); -// } -// break; -// case Float: masm.movl(storeAddr, floatToRawIntBits(c.asFloat())); break; -// case Double: throw GraalInternalError.shouldNotReachHere("Cannot store 64-bit constants to memory"); -// case Object: -// if (c.isNull()) { -// masm.movptr(storeAddr, 0); -// } else { -// throw GraalInternalError.shouldNotReachHere("Cannot store 64-bit constants to memory"); -// } -// break; default: throw GraalInternalError.shouldNotReachHere(); } @@ -448,18 +278,8 @@ } } - protected static void compareAndSwap(TargetMethodAssembler tasm, PTXMacroAssembler masm, Value result, Address address, Value cmpValue, Value newValue) { -// assert asRegister(cmpValue) == PTX.rax && asRegister(result) == PTX.rax; -// -// if (tasm.target.isMP) { -// masm.lock(); -// } -// switch (cmpValue.getKind()) { -// case Int: masm.cmpxchgl(asRegister(newValue), address); break; -// case Long: -// case Object: masm.cmpxchgq(asRegister(newValue), address); break; -// default: throw GraalInternalError.shouldNotReachHere(); -// } + @SuppressWarnings("unused") + protected static void compareAndSwap(TargetMethodAssembler tasm, PTXAssembler masm, Value result, Address address, Value cmpValue, Value newValue) { throw new InternalError("NYI"); } }