Mercurial > hg > truffle
changeset 7804:447f9ba1962b
Experimental PTX backend. Contribution by Christian Thalinger.
line wrap: on
line diff
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/Register.java Mon Feb 18 18:58:39 2013 +0100 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/Register.java Mon Feb 18 14:47:54 2013 -0800 @@ -65,6 +65,13 @@ public final int encoding; /** + * The assembler calls this method to get the register's encoding. + */ + public final int encoding() { + return encoding; + } + + /** * The size of the stack slot used to spill the value of this register. */ public final int spillSlotSize;
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/AbstractPTXAssembler.java Mon Feb 18 14:47:54 2013 -0800 @@ -0,0 +1,67 @@ +/* + * 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.*; +import com.oracle.graal.asm.*; +import com.oracle.graal.lir.ptx.*; + +/** + * The platform-dependent base class for the PTX assembler. + */ +public abstract class AbstractPTXAssembler extends AbstractAssembler { + + public AbstractPTXAssembler(TargetDescription target) { + super(target); + } + + @Override + public final void bind(Label l) { + super.bind(l); + emitString0("L"+l.toString() + ":\n"); + } + + @Override + public void align(int modulus) { + // TODO Auto-generated method stub + } + + @Override + public void jmp(Label l) { + // TODO Auto-generated method stub + } + + @Override + protected void patchJumpTarget(int branch, int jumpTarget) { + final int spaces = PTXControlFlow.BranchOp.UNBOUND_TARGET.length(); + String target = 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); + } + + @Override + public void bangStack(int disp) { + // TODO Auto-generated method stub + } + +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAsmOptions.java Mon Feb 18 14:47:54 2013 -0800 @@ -0,0 +1,27 @@ +/* + * 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; + +public class PTXAsmOptions { + // Nothing for now +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Mon Feb 18 14:47:54 2013 -0800 @@ -0,0 +1,551 @@ +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 { + 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() + ";" + ""); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXMacroAssembler.java Mon Feb 18 14:47:54 2013 -0800 @@ -0,0 +1,36 @@ +/* + * 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 18:58:39 2013 +0100 +++ b/graal/com.oracle.graal.asm/src/com/oracle/graal/asm/AbstractAssembler.java Mon Feb 18 14:47:54 2013 -0800 @@ -44,7 +44,7 @@ } } - public final void bind(Label l) { + public void bind(Label l) { assert !l.isBound() : "can bind label only once"; l.bind(codeBuffer.position()); l.patchInstructions(this); @@ -79,4 +79,16 @@ protected final void emitLong(long x) { codeBuffer.emitLong(x); } + + /** + * Some GPU architectures have a text based encoding. + */ + 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.asm/src/com/oracle/graal/asm/Buffer.java Mon Feb 18 18:58:39 2013 +0100 +++ b/graal/com.oracle.graal.asm/src/com/oracle/graal/asm/Buffer.java Mon Feb 18 14:47:54 2013 -0800 @@ -106,6 +106,27 @@ position = emitLong(b, position); } + private static final String NEWLINE = System.getProperty("line.separator"); + + public void emitString(String s) { + position = emitString("\t", position); // XXX REMOVE ME pretty-printing + position = emitString(s, position); + position = emitString(NEWLINE, position); + } + + // XXX for pretty-printing + public void emitString0(String s) { + emitBytes(s.getBytes(), 0, s.length()); + } + + public int emitBytes(byte[] arr, int pos) { + final int len = arr.length; + final int newPos = pos + len; + ensureSize(newPos); + System.arraycopy(arr, 0, data, pos, len); + return newPos; + } + public int emitByte(int b, int pos) { assert NumUtil.isUByte(b); int newPos = pos + 1; @@ -120,6 +141,10 @@ public abstract int emitLong(long b, int pos); + public int emitString(String s, int pos) { + return emitBytes(s.getBytes(), pos); + } + public int getByte(int pos) { return data[pos] & 0xff; }
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java Mon Feb 18 14:47:54 2013 -0800 @@ -0,0 +1,278 @@ +/* + * 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.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.lir.*; +import com.oracle.graal.lir.asm.*; +import com.oracle.graal.nodes.*; +import com.oracle.graal.nodes.java.*; +import com.oracle.graal.phases.GraalOptions; + +/** + * PTX specific backend. + */ +public class PTXBackend extends Backend { + + public PTXBackend(CodeCacheProvider runtime, TargetDescription target) { + super(runtime, target); + } + + 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"); + } + } + + 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); +// } + } + + @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)); +// } +// } + } + } + + @Override + public TargetMethodAssembler newAssembler(FrameMap frameMap, LIR lir) { + // Omit the frame if the method: + // - has no spill slots or other slots allocated during register allocation + // - has no callee-saved registers + // - has no incoming arguments passed on the stack + // - has no instructions with debug info + 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); + HotSpotFrameContext frameContext = omitFrame ? null : new HotSpotFrameContext(); + TargetMethodAssembler tasm = new TargetMethodAssembler(target, runtime(), frameMap, masm, frameContext, lir.stubs); + tasm.setFrameSize(frameMap.frameSize()); + tasm.targetMethod.setCustomStackAreaOffset(frameMap.offsetToCustomArea()); + return tasm; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, ResolvedJavaMethod method, LIR lir) { + // 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; + 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)); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXDeoptimizationStub.java Mon Feb 18 14:47:54 2013 -0800 @@ -0,0 +1,80 @@ +/* + * 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.compiler.ptx; + +import java.util.*; + +import com.oracle.graal.api.code.*; +import com.oracle.graal.api.code.RuntimeCallTarget.*; +import com.oracle.graal.api.meta.*; +import com.oracle.graal.asm.*; +import com.oracle.graal.asm.ptx.*; +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 { + + public static final Descriptor DEOPTIMIZE = new Descriptor("deoptimize", true, void.class); + public static final Descriptor SET_DEOPT_INFO = new Descriptor("setDeoptInfo", true, void.class, Object.class); + + public final Label label = new Label(); + public final LIRFrameState info; + public final DeoptimizationAction action; + public final DeoptimizationReason reason; + public final Object deoptInfo; + + public PTXDeoptimizationStub(DeoptimizationAction action, DeoptimizationReason reason, LIRFrameState info, Object deoptInfo) { + this.action = action; + this.reason = reason; + this.info = info; + 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); + masm.exit(); + } + + @Override + public String description() { + return "deopt stub[reason=" + reason + ", action=" + action + "]"; + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java Mon Feb 18 14:47:54 2013 -0800 @@ -0,0 +1,896 @@ +/* + * 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.compiler.ptx; + +import static com.oracle.graal.api.code.ValueUtil.*; +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.*; +import com.oracle.graal.compiler.target.*; +import com.oracle.graal.graph.*; +import com.oracle.graal.lir.*; +import com.oracle.graal.lir.StandardOp.JumpOp; +import com.oracle.graal.lir.ptx.PTXBitManipulationOp; +import com.oracle.graal.lir.ptx.PTXCompare.CompareOp; +import com.oracle.graal.lir.ptx.PTXControlFlow.BranchOp; +import com.oracle.graal.lir.ptx.PTXControlFlow.ReturnOp; +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.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; + +/** + * This class implements the PTX specific portion of the LIR generator. + */ +public abstract 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"); + } + } + + public PTXLIRGenerator(StructuredGraph graph, CodeCacheProvider runtime, TargetDescription target, FrameMap frameMap, ResolvedJavaMethod method, LIR lir) { + super(graph, runtime, target, frameMap, method, lir); + lir.spillMoveFactory = new PTXSpillMoveFactory(); + } + + @Override + protected void emitNode(ValueNode node) { + if (node instanceof LIRGenLowerable) { + ((LIRGenLowerable) node).generate(this); + } else { + super.emitNode(node); + } + } + + @Override + public boolean canStoreConstant(Constant c) { + // Operand b must be in the .reg state space. + return false; + } + + @Override + public boolean canInlineConstant(Constant c) { + switch (c.getKind()) { + case Long: + return NumUtil.isInt(c.asLong()) && !runtime.needsDataPatch(c); + case Object: + return c.isNull(); + default: + return true; + } + } + + @Override + public Address makeAddress(LocationNode location, ValueNode object) { + Value base = operand(object); + Value index = Value.ILLEGAL; + int scale = 1; + int displacement = location.displacement(); + + if (isConstant(base)) { + if (asConstant(base).isNull()) { + base = Value.ILLEGAL; + } else if (asConstant(base).getKind() != Kind.Object) { + long newDisplacement = displacement + asConstant(base).asLong(); + if (NumUtil.isInt(newDisplacement)) { + assert !runtime.needsDataPatch(asConstant(base)); + displacement = (int) newDisplacement; + base = Value.ILLEGAL; + } else { + Value newBase = newVariable(Kind.Long); + emitMove(base, newBase); + base = newBase; + } + } + } + + if (location instanceof IndexedLocationNode) { + IndexedLocationNode indexedLoc = (IndexedLocationNode) location; + + index = operand(indexedLoc.index()); + if (indexedLoc.indexScalingEnabled()) { + scale = target().sizeInBytes(location.getValueKind()); + } + if (isConstant(index)) { + long newDisplacement = displacement + asConstant(index).asLong() * scale; + // only use the constant index if the resulting displacement fits into a 32 bit + // offset + if (NumUtil.isInt(newDisplacement)) { + displacement = (int) newDisplacement; + index = Value.ILLEGAL; + } else { + // create a temporary variable for the index, the pointer load cannot handle a + // constant index + Value newIndex = newVariable(Kind.Long); + emitMove(index, newIndex); + index = newIndex; + } + } + } + + return new Address(location.getValueKind(), base, index, Address.Scale.fromInt(scale), displacement); + } + + @Override + public Variable emitMove(Value input) { + Variable result = newVariable(input.getKind()); + emitMove(input, result); + return result; + } + + @Override + public void emitMove(Value src, Value dst) { + if (isRegister(src) || isStackSlot(dst)) { + append(new MoveFromRegOp(dst, src)); + } else { + append(new MoveToRegOp(dst, src)); + } + } + + @Override + public Variable emitLoad(Value loadAddress, boolean canTrap) { + Variable result = newVariable(loadAddress.getKind()); + append(new LoadOp(result, loadAddress, canTrap ? state() : null)); + return result; + } + + @Override + public void emitStore(Value storeAddress, Value inputVal, boolean canTrap) { + Value input = loadForStore(inputVal, storeAddress.getKind()); + append(new StoreOp(storeAddress, input, canTrap ? state() : null)); + } + + @Override + public Variable emitLea(Value address) { +// Variable result = newVariable(target().wordKind); +// append(new LeaOp(result, address)); +// return result; + throw new InternalError("NYI"); + } + + @Override + public void emitJump(LabelRef label, LIRFrameState info) { + append(new JumpOp(label, info)); + } + + @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()); + } + } + + @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()); + switch (input.getKind()) { + 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(); + } + return result; + } + + @Override + public Variable emitAdd(Value a, Value b) { + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + 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(); + } + return result; + } + + @Override + public Variable emitSub(Value a, Value b) { + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + 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(); + } + return result; + } + + @Override + public Variable emitMul(Value a, Value b) { + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + 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(); + } + return result; + } + + @Override + protected boolean peephole(ValueNode valueNode) { + // No peephole optimizations for now + 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"); + } + + @Override + public Variable emitAnd(Value a, Value b) { + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + 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(); + } + return result; + } + + @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"); + } + + @Override + public Variable emitUShr(Value a, Value b) { + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + case Int: + append(new ShiftOp(IUSHR, result, a, b)); + break; +// case Long: +// append(new ShiftOp(LUSHR, result, a, b)); +// break; + default: + GraalInternalError.shouldNotReachHere(); + } + return result; + } + + @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"); + } + + @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"); + } + + @Override + public void emitBitCount(Variable result, Value value) { + if (value.getKind().getStackKind() == Kind.Int) { + append(new PTXBitManipulationOp(IPOPCNT, result, value)); + } else { + append(new PTXBitManipulationOp(LPOPCNT, result, value)); + } + } + + @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"); + } + + @Override + protected void emitReturn(Value input) { + append(new ReturnOp(input)); + } + + @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"); + } + + @Override + protected LabelRef createDeoptStub(DeoptimizationAction action, DeoptimizationReason reason, LIRFrameState info, Object deoptInfo) { + assert info.topFrame.getBCI() >= 0 : "invalid bci for deopt framestate"; + PTXDeoptimizationStub stub = new PTXDeoptimizationStub(action, reason, info, deoptInfo); + lir.stubs.add(stub); + return LabelRef.forLabel(stub.label); + } + + @Override + protected void emitNullCheckGuard(ValueNode object) { + Variable value = load(operand(object)); + LIRFrameState info = state(); + append(new NullCheckOp(value, info)); + } + + @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"); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXArithmetic.java Mon Feb 18 14:47:54 2013 -0800 @@ -0,0 +1,545 @@ +/* + * 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.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.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.*; + +// @formatter:off +public enum PTXArithmetic { + IADD, ISUB, IMUL, IDIV, IDIVREM, IREM, IUDIV, IUREM, IAND, IOR, IXOR, ISHL, ISHR, IUSHR, + LADD, LSUB, LMUL, LDIV, LDIVREM, LREM, LUDIV, LUREM, LAND, LOR, LXOR, LSHL, LSHR, LUSHR, + FADD, FSUB, FMUL, FDIV, FAND, FOR, FXOR, + DADD, DSUB, DMUL, DDIV, DAND, DOR, DXOR, + INEG, LNEG, + I2L, L2I, I2B, I2C, I2S, + F2D, D2F, + I2F, I2D, F2I, D2I, + L2F, L2D, F2L, D2L, + MOV_I2F, MOV_L2D, MOV_F2I, MOV_D2L; + + + public static class Op1Reg extends PTXLIRInstruction { + @Opcode private final PTXArithmetic opcode; + @Def({REG, HINT}) protected Value result; + @Use({REG}) protected Value x; + + public Op1Reg(PTXArithmetic opcode, Value result, Value x) { + this.opcode = opcode; + this.result = result; + this.x = x; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + emit(tasm, masm, opcode, result, x, null); + } + } + + public static class Op1Stack extends PTXLIRInstruction { + @Opcode private final PTXArithmetic opcode; + @Def({REG, HINT}) protected Value result; + @Use({REG, STACK, CONST}) protected Value x; + + public Op1Stack(PTXArithmetic opcode, Value result, Value x) { + this.opcode = opcode; + this.result = result; + this.x = x; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + emit(tasm, masm, opcode, result, x, null); + } + } + + public static class Op2Stack extends PTXLIRInstruction { + @Opcode private final PTXArithmetic opcode; + @Def({REG, HINT}) protected Value result; + @Use({REG, STACK, CONST}) protected Value x; + @Alive({REG, STACK, CONST}) protected Value y; + + public Op2Stack(PTXArithmetic opcode, Value result, Value x, Value y) { + this.opcode = opcode; + this.result = result; + this.x = x; + this.y = y; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + emit(tasm, masm, opcode, result, x, y, null); + } + + @Override + public void verify() { + super.verify(); + verifyKind(opcode, result, x, y); + } + } + + public static class Op2Reg extends PTXLIRInstruction { + @Opcode private final PTXArithmetic opcode; + @Def({REG, HINT}) protected Value result; + @Use({REG, STACK, CONST}) protected Value x; + @Alive({REG, CONST}) protected Value y; + + public Op2Reg(PTXArithmetic opcode, Value result, Value x, Value y) { + this.opcode = opcode; + this.result = result; + this.x = x; + this.y = y; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + emit(tasm, masm, opcode, result, x, y, null); + } + + @Override + public void verify() { + super.verify(); + verifyKind(opcode, result, x, y); + } + } + + public static class Op2RegCommutative extends PTXLIRInstruction { + @Opcode private final PTXArithmetic opcode; + @Def({REG, HINT}) protected Value result; + @Use({REG, STACK, CONST}) protected Value x; + @Use({REG, CONST}) protected Value y; + + public Op2RegCommutative(PTXArithmetic opcode, Value result, Value x, Value y) { + this.opcode = opcode; + this.result = result; + this.x = x; + this.y = y; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + if (sameRegister(result, y)) { + emit(tasm, masm, opcode, result, x, null); + } else { + PTXMove.move(tasm, masm, result, x); + emit(tasm, masm, opcode, result, y, null); + } + } + + @Override + protected void verify() { + super.verify(); + verifyKind(opcode, result, x, y); + } + } + + public static class ShiftOp extends PTXLIRInstruction { + @Opcode private final PTXArithmetic opcode; + @Def({REG, HINT}) protected Value result; + @Use({REG, STACK, CONST}) protected Value x; + @Alive({REG, CONST}) protected Value y; + + public ShiftOp(PTXArithmetic opcode, Value result, Value x, Value y) { + this.opcode = opcode; + this.result = result; + this.x = x; + this.y = y; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + emit(tasm, masm, opcode, result, x, y, null); + } + + @Override + public void verify() { + super.verify(); + verifyKind(opcode, result, x, x); + assert y.getKind().getStackKind() == Kind.Int; + } + } + + 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) { + this.opcode = opcode; + 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) { + 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) { + 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) { + 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(); + } + } + + if (info != null) { + assert exceptionOffset != -1; + tasm.recordImplicitException(exceptionOffset, info); + } + } + + public static void emit(TargetMethodAssembler tasm, PTXMacroAssembler masm, PTXArithmetic opcode, Value dst, Value src1, Value src2, LIRFrameState info) { + int exceptionOffset = -1; + if (isConstant(src1)) { + int a = tasm.asIntConst(src1); + Register b = asIntReg(src2); + 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 + default: throw GraalInternalError.shouldNotReachHere(); + } + } else if (isConstant(src2)) { + Register a = asIntReg(src1); + int b = tasm.asIntConst(src2); + Register d = asIntReg(dst); + switch (opcode) { + case IADD: masm.add_s32(d, a, b); break; + case IAND: masm.and_b32(d, a, b); break; + case IUSHR: masm.shr_u32(d, a, b); break; + default: throw GraalInternalError.shouldNotReachHere(); + } + } else { + Register a = asIntReg(src1); + Register b = asIntReg(src2); + Register d = asIntReg(dst); + switch (opcode) { + case IADD: masm.add_s32(d, a, b); break; + case ISUB: masm.sub_s32(d, a, b); break; + case IMUL: masm.mul_s32(d, a, b); break; + default: throw GraalInternalError.shouldNotReachHere(); + } + } + + if (info != null) { + assert exceptionOffset != -1; + tasm.recordImplicitException(exceptionOffset, info); + } + } + + 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) + || (opcode.name().startsWith("F") && result.getKind() == Kind.Float && x.getKind() == Kind.Float && y.getKind() == Kind.Float) + || (opcode.name().startsWith("D") && result.getKind() == Kind.Double && x.getKind() == Kind.Double && y.getKind() == Kind.Double); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXBitManipulationOp.java Mon Feb 18 14:47:54 2013 -0800 @@ -0,0 +1,90 @@ +/* + * 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.lir.ptx; + +import com.oracle.graal.api.code.*; +import com.oracle.graal.api.meta.*; +import com.oracle.graal.asm.ptx.*; +import com.oracle.graal.lir.asm.*; + +public class PTXBitManipulationOp extends PTXLIRInstruction { + + public enum IntrinsicOpcode { + IPOPCNT, LPOPCNT, IBSR, LBSR, BSF; + } + + @Opcode private final IntrinsicOpcode opcode; + @Def protected Value result; + @Use({OperandFlag.REG, OperandFlag.ADDR}) protected Value input; + + public PTXBitManipulationOp(IntrinsicOpcode opcode, Value result, Value input) { + this.opcode = opcode; + this.result = result; + this.input = input; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler 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; + } +// } + } + +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXCode.java Mon Feb 18 14:47:54 2013 -0800 @@ -0,0 +1,40 @@ +/* + * 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.lir.ptx; + +import com.oracle.graal.asm.ptx.*; +import com.oracle.graal.lir.*; +import com.oracle.graal.lir.asm.*; + +/** + * Convenience class to provide PTXMacroAssembler for the {@link #emitCode} method. + */ +public abstract class PTXCode implements LIR.Code { + + @Override + public final void emitCode(TargetMethodAssembler tasm) { + emitCode(tasm, (PTXMacroAssembler) tasm.asm); + } + + public abstract void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm); +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXCompare.java Mon Feb 18 14:47:54 2013 -0800 @@ -0,0 +1,140 @@ +/* + * Copyright (c) 2011, 2012, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ +package com.oracle.graal.lir.ptx; + +import static com.oracle.graal.api.code.ValueUtil.*; +import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; + +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.*; +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; + private final Condition condition; + + public CompareOp(PTXCompare opcode, Condition condition, Value x, Value y) { + this.opcode = opcode; + this.condition = condition; + this.x = x; + this.y = y; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler 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); + } + } + + public static void emit(TargetMethodAssembler tasm, PTXMacroAssembler masm, PTXCompare opcode, Condition condition, Value x, Value y) { + if (isConstant(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(); + } + 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(); + } + } else if (isConstant(y)) { + Register a = asIntReg(x); + 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(); + } + 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(); + } + } 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(); + } + } 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(); + } + } + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java Mon Feb 18 14:47:54 2013 -0800 @@ -0,0 +1,471 @@ +/* + * 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.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) { + this.x = x; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + if (tasm.frameContext != null) { + tasm.frameContext.leave(tasm); + } + masm.exit(); + } + } + + + public static class BranchOp extends PTXLIRInstruction implements StandardOp.BranchOp { + protected Condition condition; + protected LabelRef destination; + @State protected LIRFrameState state; + + public BranchOp(Condition condition, LabelRef destination, LIRFrameState state) { + this.condition = condition; + this.destination = destination; + this.state = state; + } + + public static final String UNBOUND_TARGET = "L"+Integer.MAX_VALUE; + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + masm.at(); + Label l = destination.label(); + l.addPatchAt(tasm.asm.codeBuffer.position()); + String target = l.isBound() ? "L"+l.toString() : UNBOUND_TARGET; + masm.bra(target); + } + + @Override + public LabelRef destination() { + return destination; + } + + @Override + public void negate(LabelRef newDestination) { + destination = newDestination; + 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(); +// } +// } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXLIRInstruction.java Mon Feb 18 14:47:54 2013 -0800 @@ -0,0 +1,40 @@ +/* + * 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.lir.ptx; + +import com.oracle.graal.asm.ptx.*; +import com.oracle.graal.lir.*; +import com.oracle.graal.lir.asm.*; + +/** + * Convenience class to provide PTXMacroAssembler for the {@link #emitCode} method. + */ +public abstract class PTXLIRInstruction extends LIRInstruction { + + @Override + public final void emitCode(TargetMethodAssembler tasm) { + emitCode(tasm, (PTXMacroAssembler) tasm.asm); + } + + public abstract void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm); +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java Mon Feb 18 14:47:54 2013 -0800 @@ -0,0 +1,465 @@ +/* + * 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.lir.ptx; + +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.*; +import com.oracle.graal.lir.LIRInstruction.Opcode; +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; + + public SpillMoveOp(Value result, Value input) { + this.result = result; + this.input = input; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + move(tasm, masm, getResult(), getInput()); + } + + @Override + 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; + + public MoveToRegOp(Value result, Value input) { + this.result = result; + this.input = input; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + move(tasm, masm, getResult(), getInput()); + } + + @Override + 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; + + public MoveFromRegOp(Value result, Value input) { + this.result = result; + this.input = input; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + move(tasm, masm, getResult(), getInput()); + } + + @Override + public Value getInput() { + return input; + } + @Override + public Value getResult() { + return result; + } + } + + + public static class LoadOp extends PTXLIRInstruction { + @Def({REG}) protected Value result; + @Use({ADDR}) protected Value address; + @State protected LIRFrameState state; + + public LoadOp(Value result, Value address, LIRFrameState state) { + this.result = result; + this.address = address; + this.state = state; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + load(tasm, masm, result, (Address) address, state); + } + } + + + public static class StoreOp extends PTXLIRInstruction { + @Use({ADDR}) protected Value address; + @Use({REG, CONST}) protected Value input; + @State protected LIRFrameState state; + + public StoreOp(Value address, Value input, LIRFrameState state) { + this.address = address; + this.input = input; + this.state = state; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + store(tasm, masm, (Address) address, input, state); + } + } + + + public static class LeaOp extends PTXLIRInstruction { + @Def({REG}) protected Value result; + @Use({ADDR, STACK, UNINITIALIZED}) protected Value address; + + public LeaOp(Value result, Value address) { + this.result = result; + this.address = address; + } + + @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); + 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; + @Use protected Value newValue; + + public CompareAndSwapOp(Value result, Address address, Value cmpValue, Value newValue) { + this.result = result; + this.address = address; + this.cmpValue = cmpValue; + this.newValue = newValue; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { + compareAndSwap(tasm, masm, result, (Address) address, cmpValue, newValue); + } + } + + + public static void move(TargetMethodAssembler tasm, PTXMacroAssembler 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(); + } + } else { + throw GraalInternalError.shouldNotReachHere(); + } + } + + private static void reg2reg(PTXMacroAssembler 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()); + } + } + + 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). + switch (input.getKind().getStackKind()) { + case Jsr: + case Int: + if (tasm.runtime.needsDataPatch(input)) { + tasm.recordDataReferenceInCode(input, 0, true); + } + 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); +// } + 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(); + } + } + + public static void store(TargetMethodAssembler tasm, PTXMacroAssembler masm, Address storeAddr, Value input, LIRFrameState info) { +// if (info != null) { +// tasm.recordImplicitException(masm.codeBuffer.position(), 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(); + } + } 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(); + } + + } else { + throw GraalInternalError.shouldNotReachHere(); + } + } + + 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(); +// } + throw new InternalError("NYI"); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.ptx/src/com/oracle/graal/ptx/PTX.java Mon Feb 18 14:47:54 2013 -0800 @@ -0,0 +1,119 @@ +/* + * 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.ptx; + +import static com.oracle.graal.api.code.MemoryBarriers.*; +import static com.oracle.graal.api.code.Register.RegisterFlag.*; + +import java.nio.*; + +import com.oracle.graal.api.code.*; +import com.oracle.graal.api.code.Register.*; + +/** + * Represents the PTX architecture. + */ +public class PTX extends Architecture { + + // @formatter:off + + /* + * Register State Space + * + * Registers (.reg state space) are fast storage locations. The number of + * registers is limited, and will vary from platform to platform. When the + * limit is exceeded, register variables will be spilled to memory, causing + * changes in performance. For each architecture, there is a recommended + * maximum number of registers to use (see the "CUDA Programming Guide" for + * details). + */ + + // General purpose registers + public static final Register r0 = new Register(0, 0, 8, "r0", CPU, RegisterFlag.Byte); + public static final Register r1 = new Register(1, 1, 8, "r1", CPU, RegisterFlag.Byte); + public static final Register r2 = new Register(2, 2, 8, "r2", CPU, RegisterFlag.Byte); + public static final Register r3 = new Register(3, 3, 8, "r3", CPU, RegisterFlag.Byte); + public static final Register r4 = new Register(4, 4, 8, "r4", CPU, RegisterFlag.Byte); + public static final Register r5 = new Register(5, 5, 8, "r5", CPU, RegisterFlag.Byte); + public static final Register r6 = new Register(6, 6, 8, "r6", CPU, RegisterFlag.Byte); + public static final Register r7 = new Register(7, 7, 8, "r7", CPU, RegisterFlag.Byte); + + public static final Register r8 = new Register(8, 8, 8, "r8", CPU, RegisterFlag.Byte); + public static final Register r9 = new Register(9, 9, 8, "r9", CPU, RegisterFlag.Byte); + public static final Register r10 = new Register(10, 10, 8, "r10", CPU, RegisterFlag.Byte); + public static final Register r11 = new Register(11, 11, 8, "r11", CPU, RegisterFlag.Byte); + public static final Register r12 = new Register(12, 12, 8, "r12", CPU, RegisterFlag.Byte); + public static final Register r13 = new Register(13, 13, 8, "r13", CPU, RegisterFlag.Byte); + public static final Register r14 = new Register(14, 14, 8, "r14", CPU, RegisterFlag.Byte); + public static final Register r15 = new Register(15, 15, 8, "r15", CPU, RegisterFlag.Byte); + + public static final Register[] gprRegisters = { + r0, r1, r2, r3, r4, r5, r6, r7, + r8, r9, r10, r11, r12, r13, r14, r15 + }; + + // Floating point registers + public static final Register f0 = new Register(16, 0, 8, "f0", FPU); + public static final Register f1 = new Register(17, 1, 8, "f1", FPU); + public static final Register f2 = new Register(18, 2, 8, "f2", FPU); + public static final Register f3 = new Register(19, 3, 8, "f3", FPU); + public static final Register f4 = new Register(20, 4, 8, "f4", FPU); + public static final Register f5 = new Register(21, 5, 8, "f5", FPU); + public static final Register f6 = new Register(22, 6, 8, "f6", FPU); + public static final Register f7 = new Register(23, 7, 8, "f7", FPU); + + public static final Register f8 = new Register(24, 8, 8, "f8", FPU); + public static final Register f9 = new Register(25, 9, 8, "f9", FPU); + public static final Register f10 = new Register(26, 10, 8, "f10", FPU); + public static final Register f11 = new Register(27, 11, 8, "f11", FPU); + public static final Register f12 = new Register(28, 12, 8, "f12", FPU); + public static final Register f13 = new Register(29, 13, 8, "f13", FPU); + public static final Register f14 = new Register(30, 14, 8, "f14", FPU); + public static final Register f15 = new Register(31, 15, 8, "f15", FPU); + + public static final Register[] fpuRegisters = { + f0, f1, f2, f3, f4, f5, f6, f7, + f8, f9, f10, f11, f12, f13, f14, f15 + }; + + public static final Register[] allRegisters = { + // GPR + r0, r1, r2, r3, r4, r5, r6, r7, + r8, r9, r10, r11, r12, r13, r14, r15, + // FPU + f0, f1, f2, f3, f4, f5, f6, f7, + f8, f9, f10, f11, f12, f13, f14, f15 + }; + + public PTX() { + super("PTX", + 8, + ByteOrder.LITTLE_ENDIAN, + allRegisters, + LOAD_STORE | STORE_STORE, + 0, + r15.encoding + 1, + 8); + } + // @formatter:on +}
--- a/mx/projects Mon Feb 18 18:58:39 2013 +0100 +++ b/mx/projects Mon Feb 18 14:47:54 2013 -0800 @@ -62,6 +62,13 @@ project@com.oracle.graal.amd64@checkstyle=com.oracle.graal.graph project@com.oracle.graal.amd64@javaCompliance=1.7 +# graal.ptx +project@com.oracle.graal.ptx@subDir=graal +project@com.oracle.graal.ptx@sourceDirs=src +project@com.oracle.graal.ptx@dependencies=com.oracle.graal.api.code +project@com.oracle.graal.ptx@checkstyle=com.oracle.graal.graph +project@com.oracle.graal.ptx@javaCompliance=1.7 + # graal.sparc project@com.oracle.graal.sparc@subDir=graal project@com.oracle.graal.sparc@sourceDirs=src @@ -83,6 +90,13 @@ project@com.oracle.graal.hotspot.amd64@checkstyle=com.oracle.graal.graph project@com.oracle.graal.hotspot.amd64@javaCompliance=1.7 +# graal.hotspot.ptx +project@com.oracle.graal.hotspot.ptx@subDir=graal +project@com.oracle.graal.hotspot.ptx@sourceDirs=src +project@com.oracle.graal.hotspot.ptx@dependencies=com.oracle.graal.hotspot,com.oracle.graal.compiler.ptx +project@com.oracle.graal.hotspot.ptx@checkstyle=com.oracle.graal.graph +project@com.oracle.graal.hotspot.ptx@javaCompliance=1.7 + # graal.hotspot.sparc project@com.oracle.graal.hotspot.sparc@subDir=graal project@com.oracle.graal.hotspot.sparc@sourceDirs=src @@ -137,6 +151,13 @@ project@com.oracle.graal.lir.amd64@checkstyle=com.oracle.graal.graph project@com.oracle.graal.lir.amd64@javaCompliance=1.7 +# graal.lir.ptx +project@com.oracle.graal.lir.ptx@subDir=graal +project@com.oracle.graal.lir.ptx@sourceDirs=src +project@com.oracle.graal.lir.ptx@dependencies=com.oracle.graal.lir,com.oracle.graal.asm.ptx +project@com.oracle.graal.lir.ptx@checkstyle=com.oracle.graal.graph +project@com.oracle.graal.lir.ptx@javaCompliance=1.7 + # graal.lir.sparc project@com.oracle.graal.lir.sparc@subDir=graal project@com.oracle.graal.lir.sparc@sourceDirs=src @@ -221,6 +242,13 @@ project@com.oracle.graal.compiler.amd64@checkstyle=com.oracle.graal.graph project@com.oracle.graal.compiler.amd64@javaCompliance=1.7 +# graal.compiler.ptx +project@com.oracle.graal.compiler.ptx@subDir=graal +project@com.oracle.graal.compiler.ptx@sourceDirs=src +project@com.oracle.graal.compiler.ptx@dependencies=com.oracle.graal.compiler,com.oracle.graal.lir.ptx +project@com.oracle.graal.compiler.ptx@checkstyle=com.oracle.graal.graph +project@com.oracle.graal.compiler.ptx@javaCompliance=1.7 + # graal.compiler.sparc project@com.oracle.graal.compiler.sparc@subDir=graal project@com.oracle.graal.compiler.sparc@sourceDirs=src @@ -297,6 +325,13 @@ project@com.oracle.graal.asm.amd64.test@checkstyle=com.oracle.graal.graph project@com.oracle.graal.asm.amd64.test@javaCompliance=1.7 +# graal.asm.ptx +project@com.oracle.graal.asm.ptx@subDir=graal +project@com.oracle.graal.asm.ptx@sourceDirs=src +project@com.oracle.graal.asm.ptx@dependencies=com.oracle.graal.asm,com.oracle.graal.ptx +project@com.oracle.graal.asm.ptx@checkstyle=com.oracle.graal.graph +project@com.oracle.graal.asm.ptx@javaCompliance=1.7 + # graal.asm.sparc project@com.oracle.graal.asm.sparc@subDir=graal project@com.oracle.graal.asm.sparc@sourceDirs=src