Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Issue #506 has been fixed. #513

Merged
merged 3 commits into from
Jul 30, 2024
Merged

Issue #506 has been fixed. #513

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
aa244f9
Issue #506 has been fixed.
andrii0lomakin Jul 27, 2024
aa5271d
Fixed issue that caused printing of zero values in `printf` instead o…
andrii0lomakin Jul 28, 2024
b0683c7
Additional tests were added for Debug.printf
andrii0lomakin Jul 30, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 8 additions & 0 deletions pom.xml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2088,6 +2088,14 @@
<artifactId>collections</artifactId>
<version>${graalvm.version}</version>
</dependency>

<dependency>
<groupId>org.graalvm.compiler</groupId>
<artifactId>compiler</artifactId>
<version>23.1.0</version>
<scope>provided</scope>
</dependency>

</dependencies>
<modules>
<module>tornado-runtime</module>
Expand Down
120 changes: 14 additions & 106 deletions ...k/ac/manchester/tornado/drivers/opencl/graal/compiler/plugins/OCLGraphBuilderPlugins.java
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,6 +48,8 @@
import org.graalvm.compiler.core.common.memory.MemoryOrderMode;
import org.graalvm.compiler.core.common.type.StampFactory;
import org.graalvm.compiler.graph.Node;
import org.graalvm.compiler.lir.Variable;
import org.graalvm.compiler.lir.gen.LIRGeneratorTool;
import org.graalvm.compiler.nodes.ConstantNode;
import org.graalvm.compiler.nodes.FixedWithNextNode;
import org.graalvm.compiler.nodes.PiNode;
Expand Down Expand Up @@ -79,10 +81,14 @@
import uk.ac.manchester.tornado.api.exceptions.Debug;
import uk.ac.manchester.tornado.api.exceptions.TornadoRuntimeException;
import uk.ac.manchester.tornado.drivers.opencl.graal.OCLArchitecture;
import uk.ac.manchester.tornado.drivers.opencl.graal.asm.OCLAssembler.OCLUnaryIntrinsic;
import uk.ac.manchester.tornado.drivers.opencl.graal.lir.OCLKind;
import uk.ac.manchester.tornado.drivers.opencl.graal.lir.OCLLIRStmt.AssignStmt;
import uk.ac.manchester.tornado.drivers.opencl.graal.lir.OCLUnary;
import uk.ac.manchester.tornado.drivers.opencl.graal.nodes.AtomicAddNodeTemplate;
import uk.ac.manchester.tornado.drivers.opencl.graal.nodes.DecAtomicNode;
import uk.ac.manchester.tornado.drivers.opencl.graal.nodes.GetAtomicNode;
import uk.ac.manchester.tornado.drivers.opencl.graal.nodes.GlobalThreadIdNode;
import uk.ac.manchester.tornado.drivers.opencl.graal.nodes.IncAtomicNode;
import uk.ac.manchester.tornado.drivers.opencl.graal.nodes.LocalArrayNode;
import uk.ac.manchester.tornado.drivers.opencl.graal.nodes.OCLBarrierNode;
Expand All @@ -91,7 +97,6 @@
import uk.ac.manchester.tornado.drivers.opencl.graal.nodes.OCLIntBinaryIntrinsicNode;
import uk.ac.manchester.tornado.drivers.opencl.graal.nodes.OCLIntUnaryIntrinsicNode;
import uk.ac.manchester.tornado.drivers.opencl.graal.nodes.PrintfNode;
import uk.ac.manchester.tornado.drivers.opencl.graal.nodes.TPrintfNode;
import uk.ac.manchester.tornado.drivers.opencl.graal.nodes.TornadoAtomicIntegerNode;
import uk.ac.manchester.tornado.runtime.common.TornadoOptions;

Expand Down Expand Up @@ -309,75 +314,6 @@ private static void registerKernelContextPlugins(InvocationPlugins plugins) {
localArraysPlugins(r);
}

private static boolean printfHandler(GraphBuilderContext b, ResolvedJavaMethod targetMethod, Receiver receiver, ValueNode... args) {
int idCount = 0;
int index = 0;
for (; index < 3; index++) {
ValueNode arg = args[index];
if (arg instanceof ConstantNode && arg.getStackKind().isObject()) {
break;
}
idCount++;
}

NewArrayNode newArrayNode = (NewArrayNode) args[index + 1];
ConstantNode lengthNode = (ConstantNode) newArrayNode.dimension(0);
int length = ((JavaConstant) lengthNode.getValue()).asInt();

ValueNode[] actualArgs = new ValueNode[4 + length];
if (idCount >= 0) {
System.arraycopy(args, 0, actualArgs, 0, idCount);
}

for (int i = idCount; i < 3; i++) {
actualArgs[i] = ConstantNode.forInt(0);
}

actualArgs[3] = args[index];

int argIndex = 0;
for (Node n : newArrayNode.usages()) {
if (n instanceof StoreIndexedNode) {
StoreIndexedNode storeNode = (StoreIndexedNode) n;
ValueNode value = storeNode.value();
if (value instanceof BoxNode) {
BoxNode box = (BoxNode) value;
value = box.getValue();
GraphUtil.unlinkFixedNode(box);
box.safeDelete();
}
actualArgs[4 + argIndex] = value;
argIndex++;
}

}

TPrintfNode printfNode = new TPrintfNode(actualArgs);

b.add(b.append(printfNode));
while (newArrayNode.hasUsages()) {
Node n = newArrayNode.usages().first();
// need to remove all nodes from the graph that operate on
// the new array,
// however, we cannot remove all inputs as they may be used
// by the currently
// unbuilt part of the graph. We also need to ensure that we
// do not leave any
// gaps inbetween fixed nodes
if (n instanceof FixedWithNextNode) {
GraphUtil.unlinkFixedNode((FixedWithNextNode) n);
}
n.clearInputs();
n.safeDelete();
}

GraphUtil.unlinkFixedNode(newArrayNode);
newArrayNode.clearInputs();
newArrayNode.safeDelete();

return true;
}

public static Class getValueLayoutClass(Class k) {
if (k == int.class) {
return ValueLayout.OfInt.class;
Expand Down Expand Up @@ -430,39 +366,6 @@ public boolean apply(GraphBuilderContext b, ResolvedJavaMethod targetMethod, Rec
}

private static void registerTornadoVMIntrinsicsPlugins(InvocationPlugins plugins) {
final InvocationPlugin tprintfPlugin2 = new InvocationPlugin("tprintf", String.class, Object[].class) {
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I see the tprinf intrinsic is removed, which makes sense. We can debug directly using the Debug.printf method.

@Override
public boolean defaultHandler(GraphBuilderContext b, ResolvedJavaMethod targetMethod, Receiver receiver, ValueNode... args) {
return printfHandler(b, targetMethod, receiver, args);
}
};

final InvocationPlugin tprintfPlugin3 = new InvocationPlugin("tprintf", int.class, String.class, Object[].class) {
@Override
public boolean defaultHandler(GraphBuilderContext b, ResolvedJavaMethod targetMethod, Receiver receiver, ValueNode... args) {
return printfHandler(b, targetMethod, receiver, args);
}
};

final InvocationPlugin tprintfPlugin4 = new InvocationPlugin("tprintf", int.class, int.class, String.class, Object[].class) {
@Override
public boolean defaultHandler(GraphBuilderContext b, ResolvedJavaMethod targetMethod, Receiver receiver, ValueNode... args) {
return printfHandler(b, targetMethod, receiver, args);
}
};

final InvocationPlugin tprintfPlugin5 = new InvocationPlugin("tprintf", int.class, int.class, int.class, String.class, Object[].class) {
@Override
public boolean defaultHandler(GraphBuilderContext b, ResolvedJavaMethod targetMethod, Receiver receiver, ValueNode... args) {
return printfHandler(b, targetMethod, receiver, args);
}
};

plugins.register(Debug.class, tprintfPlugin2);
plugins.register(Debug.class, tprintfPlugin3);
plugins.register(Debug.class, tprintfPlugin4);
plugins.register(Debug.class, tprintfPlugin5);

final InvocationPlugin printfPlugin = new InvocationPlugin("printf", String.class, Object[].class) {

@Override
Expand All @@ -472,9 +375,13 @@ public boolean defaultHandler(GraphBuilderContext b, ResolvedJavaMethod targetMe
ConstantNode lengthNode = (ConstantNode) newArrayNode.dimension(0);
int length = ((JavaConstant) lengthNode.getValue()).asInt();

ValueNode[] actualArgs = new ValueNode[length + 1];
ValueNode[] actualArgs = new ValueNode[length + 4];
actualArgs[0] = args[0];

actualArgs[1] = b.append(new GlobalThreadIdNode(ConstantNode.forInt(0)));
actualArgs[2] = b.append(new GlobalThreadIdNode(ConstantNode.forInt(1)));
actualArgs[3] = b.append(new GlobalThreadIdNode(ConstantNode.forInt(2)));

int argIndex = 0;
for (Node n : newArrayNode.usages()) {
if (n instanceof StoreIndexedNode) {
Expand All @@ -486,14 +393,15 @@ public boolean defaultHandler(GraphBuilderContext b, ResolvedJavaMethod targetMe
GraphUtil.unlinkFixedNode(box);
box.safeDelete();
}
actualArgs[argIndex + 1] = value;
actualArgs[argIndex + 4] = value;
argIndex++;
}

}


PrintfNode printfNode = new PrintfNode(actualArgs);
b.add(b.append(printfNode));
b.append(printfNode);

while (newArrayNode.hasUsages()) {
Node n = newArrayNode.usages().first();
Expand Down
28 changes: 16 additions & 12 deletions ...ers/opencl/src/main/java/uk/ac/manchester/tornado/drivers/opencl/graal/lir/OCLPrintf.java
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,39 +43,43 @@ public OCLPrintf(Value[] inputs) {

@Override
public void emit(OCLCompilationResultBuilder crb, OCLAssembler asm) {
asm.emit("printf( \"tornado[%%3d,%%3d,%%3d]> %s", asm.formatConstant((ConstantValue) inputs[0]));
asm.emit("printf( \"tornado[%%3d,%%3d,%%3d]> %s\"",
asm.formatConstant((ConstantValue) inputs[0]));

asm.emit("\", ");
for (int i = 0; i < 2; i++) {
asm.emit("get_global_id(%d), ", i);
for (int i = 1; i < 4; i++) {
asm.emit(", ");
asm.emitValue(crb, inputs[i]);
}
asm.emit("get_global_id(%d) ", 2);
if (inputs.length > 1) {

if (inputs.length > 4) {
asm.emit(", ");
}
for (int i = 1; i < inputs.length - 1; i++) {

for (int i = 4; i < inputs.length - 1; i++) {
asm.emitValue(crb, inputs[i]);
asm.emit(", ");
}

if (inputs.length > 1) {
if (inputs.length > 4) {
asm.emitValue(crb, inputs[inputs.length - 1]);
}

asm.emit(")");
}

@Override
public String toString() {
StringBuilder sb = new StringBuilder();
sb.append(String.format("printf( %s", inputs[0]));
if (inputs.length > 1) {
sb.append(String.format("printf( \"%s\"", inputs[0]));

if (inputs.length > 4) {
sb.append(", ");
}
for (int i = 1; i < inputs.length - 1; i++) {
for (int i = 4; i < inputs.length - 1; i++) {
sb.append(inputs[i]);
sb.append(", ");
}
if (inputs.length > 1) {
if (inputs.length > 4) {
sb.append(inputs[inputs.length - 1]);
}
sb.append(" )");
Expand Down
97 changes: 0 additions & 97 deletions ...rs/opencl/src/main/java/uk/ac/manchester/tornado/drivers/opencl/graal/lir/OCLTPrintf.java
View file
Open in desktop

This file was deleted.

5 changes: 1 addition & 4 deletions .../opencl/src/main/java/uk/ac/manchester/tornado/drivers/opencl/graal/nodes/PrintfNode.java
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,10 +44,7 @@ public class PrintfNode extends FixedWithNextNode implements LIRLowerable {

public PrintfNode(ValueNode... values) {
super(TYPE, StampFactory.forVoid());
this.inputs = new NodeInputList<>(this, values.length);
for (int i = 0; i < values.length; i++) {
inputs.set(i, values[i]);
}
this.inputs = new NodeInputList<>(this, values);
}

@Override
Expand Down
Loading