Mercurial > hg > graal-jvmci-8

diff graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotBackend.java @ 15249:54dd156e5f09

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
[PTX] Fix regression in generation of predicate register declaration
author bharadwaj
date Fri, 18 Apr 2014 16:11:53 -0400
parents 96bb07a5d667
children a38d791982e1
line wrap: on
line diff
--- a/graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotBackend.java	Fri Apr 18 09:38:29 2014 -1000
+++ b/graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotBackend.java	Fri Apr 18 16:11:53 2014 -0400
@@ -51,6 +51,7 @@
 import com.oracle.graal.lir.StandardOp.LabelOp;
 import com.oracle.graal.lir.asm.*;
 import com.oracle.graal.lir.ptx.*;
+import com.oracle.graal.lir.ptx.PTXControlFlow.PTXPredicatedLIRInstruction;
 import com.oracle.graal.lir.ptx.PTXMemOp.LoadReturnAddrOp;
 import com.oracle.graal.nodes.*;
 import com.oracle.graal.phases.*;
@@ -409,12 +410,20 @@
         assert codeCacheOwner != null : lir + " is not associated with a method";
 
         RegisterAnalysis registerAnalysis = new RegisterAnalysis();
+        // Assume no predicate registers are used
+        int maxPredRegNum = -1;
 
         for (AbstractBlock<?> b : lir.codeEmittingOrder()) {
             for (LIRInstruction op : lir.getLIRforBlock(b)) {
                 if (op instanceof LabelOp) {
                     // Don't consider this as a definition
                 } else {
+                    if (op instanceof PTXPredicatedLIRInstruction) {
+                        // Update maximum predicate register number if op uses a larger number
+                        int opPredRegNum = ((PTXPredicatedLIRInstruction) op).getPredRegNum();
+                        maxPredRegNum = (opPredRegNum > maxPredRegNum) ? opPredRegNum : maxPredRegNum;
+                    }
+                    // Record registers used in the kernel
                     registerAnalysis.op = op;
                     op.forEachTemp(registerAnalysis);
                     op.forEachOutput(registerAnalysis);
@@ -422,13 +431,13 @@
             }
         }
 
+        // Emit register declarations
         Assembler asm = crb.asm;
         registerAnalysis.emitDeclarations(asm);
 
         // emit predicate register declaration
-        int maxPredRegNum = lir.numVariables();
-        if (maxPredRegNum > 0) {
-            asm.emitString(".reg .pred %p<" + maxPredRegNum + ">;");
+        if (maxPredRegNum > -1) {
+            asm.emitString(".reg .pred %p<" + ++maxPredRegNum + ">;");
         }
     }