diff --git a/src/main/java/com/aparapi/internal/writer/BlockWriter.java b/src/main/java/com/aparapi/internal/writer/BlockWriter.java index 013239d0..b074d630 100644 --- a/src/main/java/com/aparapi/internal/writer/BlockWriter.java +++ b/src/main/java/com/aparapi/internal/writer/BlockWriter.java @@ -80,8 +80,14 @@ public abstract class BlockWriter{ public final static String arrayDimMangleSuffix = "__javaArrayDimension"; + public final static String reservedIdentifierMangleSuffix = "__javaReserved"; + public abstract void write(String _string); + protected String mangleIdentifier(String _identifier) { + return _identifier; + } + public void writeln(String _string) { write(_string); newLine(); @@ -309,7 +315,7 @@ protected void writeGetterBlock(FieldEntry accessorVariableFieldEntry) { in(); newLine(); write("return this->"); - write(accessorVariableFieldEntry.getNameAndTypeEntry().getNameUTF8Entry().getUTF8()); + write(mangleIdentifier(accessorVariableFieldEntry.getNameAndTypeEntry().getNameUTF8Entry().getUTF8())); write(";"); out(); newLine(); @@ -417,7 +423,7 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException { } else { final String descriptor = localVariableInfo.getVariableDescriptor(); write(convertType(descriptor, true, true)); - write(localVariableInfo.getVariableName()); + write(mangleIdentifier(localVariableInfo.getVariableName())); } } else { if (assignToLocalVariable.isDeclaration()) { @@ -431,7 +437,7 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException { if (localVariableInfo == null) { throw new CodeGenException("outOfScope" + _instruction.getThisPC() + " = "); } else { - write(localVariableInfo.getVariableName() + " = "); + write(mangleIdentifier(localVariableInfo.getVariableName()) + " = "); } } @@ -482,7 +488,7 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException { NameAndTypeEntry nameAndTypeEntry = ((AccessField) load).getConstantPoolFieldEntry().getNameAndTypeEntry(); if (isMultiDimensionalArray(nameAndTypeEntry)) { String arrayName = nameAndTypeEntry.getNameUTF8Entry().getUTF8(); - write(" * this->" + arrayName + arrayDimMangleSuffix + dim); + write(" * this->" + mangleIdentifier(arrayName + arrayDimMangleSuffix + dim)); } } @@ -506,7 +512,7 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException { writeThisRef(); } } - write(accessField.getConstantPoolFieldEntry().getNameAndTypeEntry().getNameUTF8Entry().getUTF8()); + write(mangleIdentifier(accessField.getConstantPoolFieldEntry().getNameAndTypeEntry().getNameUTF8Entry().getUTF8())); } else if (_instruction instanceof I_ARRAYLENGTH) { @@ -523,7 +529,7 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException { NameAndTypeEntry nameAndTypeEntry = ((AccessInstanceField) load).getConstantPoolFieldEntry().getNameAndTypeEntry(); final String arrayName = nameAndTypeEntry.getNameUTF8Entry().getUTF8(); String dimSuffix = isMultiDimensionalArray(nameAndTypeEntry) ? Integer.toString(dim) : ""; - write("this->" + arrayName + arrayLengthMangleSuffix + dimSuffix); + write("this->" + mangleIdentifier(arrayName + arrayLengthMangleSuffix + dimSuffix)); } else if (_instruction instanceof AssignToField) { final AssignToField assignedField = (AssignToField) _instruction; @@ -537,7 +543,7 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException { writeThisRef(); } } - write(assignedField.getConstantPoolFieldEntry().getNameAndTypeEntry().getNameUTF8Entry().getUTF8()); + write(mangleIdentifier(assignedField.getConstantPoolFieldEntry().getNameAndTypeEntry().getNameUTF8Entry().getUTF8())); write("="); writeInstruction(assignedField.getValueToAssign()); } else if (_instruction instanceof Constant) { @@ -581,13 +587,13 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException { } else if (_instruction instanceof AccessLocalVariable) { final AccessLocalVariable localVariableLoadInstruction = (AccessLocalVariable) _instruction; final LocalVariableInfo localVariable = localVariableLoadInstruction.getLocalVariableInfo(); - write(localVariable.getVariableName()); + write(mangleIdentifier(localVariable.getVariableName())); } else if (_instruction instanceof I_IINC) { final I_IINC location = (I_IINC) _instruction; final LocalVariableInfo localVariable = location.getLocalVariableInfo(); final int adjust = location.getAdjust(); - write(localVariable.getVariableName()); + write(mangleIdentifier(localVariable.getVariableName())); if (adjust == 1) { write("++"); } else if (adjust == -1) { @@ -688,7 +694,7 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException { if (localVariableInfo == null) { throw new CodeGenException("outOfScope" + _instruction.getThisPC() + " = "); } else { - write(localVariableInfo.getVariableName() + " = "); + write(mangleIdentifier(localVariableInfo.getVariableName()) + " = "); } } @@ -703,7 +709,7 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException { throw new CodeGenException("/* we can't declare this " + convertType(localVariableInfo.getVariableDescriptor(), true, false) + " here */"); } - write(localVariableInfo.getVariableName()); + write(mangleIdentifier(localVariableInfo.getVariableName())); write("="); writeInstruction(inlineAssignInstruction.getRhs()); } else if (_instruction.getByteCode().equals(ByteCode.FIELD_ARRAY_ELEMENT_ASSIGN)) { diff --git a/src/main/java/com/aparapi/internal/writer/KernelWriter.java b/src/main/java/com/aparapi/internal/writer/KernelWriter.java index 0a72a33f..5648d205 100644 --- a/src/main/java/com/aparapi/internal/writer/KernelWriter.java +++ b/src/main/java/com/aparapi/internal/writer/KernelWriter.java @@ -117,6 +117,38 @@ public abstract class KernelWriter extends BlockWriter{ private Entrypoint entryPoint = null; + private static final Set openCLReservedIdentifiers = createOpenCLReservedIdentifiers(); + + private static Set createOpenCLReservedIdentifiers() { + final Set identifiers = new HashSet(); + Collections.addAll(identifiers, + "auto", "bool", "complex", "constant", "event_t", "extern", "global", "half", "imaginary", "inline", + "kernel", "local", "private", "quad", "read_only", "read_write", "register", "restrict", "sampler_t", + "signed", "sizeof", "struct", "typedef", "union", "unsigned", "write_only", "_Bool", "_Complex", + "_Imaginary", "__constant", "__global", "__kernel", "__local", "__private", "__read_only", + "__read_write", "__write_only", "atomic_flag", "atomic_int", "atomic_uint", "atomic_long", + "atomic_ulong", "atomic_float", "atomic_double", "atomic_half", "clk_event_t", "image1d_t", + "image1d_array_t", "image1d_buffer_t", "image2d_t", "image2d_array_t", "image2d_depth_t", + "image2d_array_depth_t", "image2d_msaa_t", "image2d_array_msaa_t", "image2d_msaa_depth_t", + "image2d_array_msaa_depth_t", "image3d_t", "intptr_t", "uintptr_t", "size_t", "ptrdiff_t", + "pipe", "reserve_id_t", "queue_t", "ndrange_t", "uchar", "ushort", "uint", "ulong"); + + final String[] vectorTypes = { + "char", "uchar", "short", "ushort", "int", "uint", "long", "ulong", "float", "double", "half" + }; + final int[] vectorWidths = {2, 3, 4, 8, 16}; + for (final String type : vectorTypes) { + for (final int width : vectorWidths) { + identifiers.add(type + width); + } + } + return Collections.unmodifiableSet(identifiers); + } + + @Override protected String mangleIdentifier(String _identifier) { + return openCLReservedIdentifiers.contains(_identifier) ? _identifier + reservedIdentifierMangleSuffix : _identifier; + } + public final static Map javaToCLIdentifierMap = new HashMap(); { javaToCLIdentifierMap.put("getGlobalId()I", "get_global_id(0)"); @@ -243,7 +275,7 @@ public abstract class KernelWriter extends BlockWriter{ if (getterField != null && isThis(_methodCall.getArg(0))) { String fieldName = getterField.getNameAndTypeEntry().getNameUTF8Entry().getUTF8(); write("this->"); - write(fieldName); + write(mangleIdentifier(fieldName)); return; } boolean noCL = _methodEntry.getOwnerClassModel().getNoCLMethods() @@ -285,7 +317,7 @@ public abstract class KernelWriter extends BlockWriter{ //assert refAccess instanceof I_GETFIELD : "ref should come from getfield"; final String fieldName = ((AccessField) refAccess).getConstantPoolFieldEntry().getNameAndTypeEntry() .getNameUTF8Entry().getUTF8(); - write(" &(this->" + fieldName); + write(" &(this->" + mangleIdentifier(fieldName)); write("["); writeInstruction(arrayAccess.getArrayIndex()); write("])"); @@ -343,6 +375,8 @@ public void writePragma(String _name, boolean _enable) { final StringBuilder thisStructLine = new StringBuilder(); final StringBuilder argLine = new StringBuilder(); final StringBuilder assignLine = new StringBuilder(); + final String fieldName = field.getName(); + final String mangledFieldName = mangleIdentifier(fieldName); String signature = field.getDescriptor(); @@ -352,11 +386,11 @@ public void writePragma(String _name, boolean _enable) { // check the suffix - String type = field.getName().endsWith(Kernel.LOCAL_SUFFIX) ? __local - : (field.getName().endsWith(Kernel.CONSTANT_SUFFIX) ? __constant : __global); + String type = fieldName.endsWith(Kernel.LOCAL_SUFFIX) ? __local + : (fieldName.endsWith(Kernel.CONSTANT_SUFFIX) ? __constant : __global); Integer privateMemorySize = null; try { - privateMemorySize = _entryPoint.getClassModel().getPrivateMemorySize(field.getName()); + privateMemorySize = _entryPoint.getClassModel().getPrivateMemorySize(fieldName); } catch (ClassParseException e) { throw new CodeGenException(e); } @@ -423,13 +457,13 @@ public void writePragma(String _name, boolean _enable) { if (privateMemorySize == null) { assignLine.append("this->"); - assignLine.append(field.getName()); + assignLine.append(mangledFieldName); assignLine.append(" = "); - assignLine.append(field.getName()); + assignLine.append(mangledFieldName); } - argLine.append(field.getName()); - thisStructLine.append(field.getName()); + argLine.append(mangledFieldName); + thisStructLine.append(mangledFieldName); if (privateMemorySize == null) { assigns.add(assignLine.toString()); } @@ -441,7 +475,7 @@ public void writePragma(String _name, boolean _enable) { // Add int field into "this" struct for supporting java arraylength op // named like foo__javaArrayLength - if (isPointer && _entryPoint.getArrayFieldArrayLengthUsed().contains(field.getName()) || isPointer && numDimensions > 1) { + if (isPointer && _entryPoint.getArrayFieldArrayLengthUsed().contains(fieldName) || isPointer && numDimensions > 1) { for (int i = 0; i < numDimensions; i++) { final StringBuilder lenStructLine = new StringBuilder(); @@ -449,7 +483,7 @@ public void writePragma(String _name, boolean _enable) { final StringBuilder lenAssignLine = new StringBuilder(); String suffix = numDimensions == 1 ? "" : Integer.toString(i); - String lenName = field.getName() + BlockWriter.arrayLengthMangleSuffix + suffix; + String lenName = mangleIdentifier(fieldName + BlockWriter.arrayLengthMangleSuffix + suffix); lenStructLine.append("int " + lenName); @@ -468,7 +502,7 @@ public void writePragma(String _name, boolean _enable) { final StringBuilder dimStructLine = new StringBuilder(); final StringBuilder dimArgLine = new StringBuilder(); final StringBuilder dimAssignLine = new StringBuilder(); - String dimName = field.getName() + BlockWriter.arrayDimMangleSuffix + suffix; + String dimName = mangleIdentifier(fieldName + BlockWriter.arrayDimMangleSuffix + suffix); dimStructLine.append("int " + dimName); @@ -562,7 +596,7 @@ public void writePragma(String _name, boolean _enable) { final String cType = convertType(field.getNameAndTypeEntry().getDescriptorUTF8Entry().getUTF8(), true, false); assert cType != null : "could not find type for " + field.getNameAndTypeEntry().getDescriptorUTF8Entry().getUTF8(); - writeln(cType + " " + field.getNameAndTypeEntry().getNameUTF8Entry().getUTF8() + ";"); + writeln(cType + " " + mangleIdentifier(field.getNameAndTypeEntry().getNameUTF8Entry().getUTF8()) + ";"); } // compute total size for OpenCL buffer @@ -687,7 +721,7 @@ public void writePragma(String _name, boolean _enable) { } write(convertType(descriptor, true, false)); - write(lvi.getVariableName()); + write(mangleIdentifier(lvi.getVariableName())); alreadyHasFirstArg = true; localVariableIndex++; diff --git a/src/test/java/com/aparapi/codegen/test/OpenCLKeywordIdentifiers.java b/src/test/java/com/aparapi/codegen/test/OpenCLKeywordIdentifiers.java new file mode 100644 index 00000000..c7f013cd --- /dev/null +++ b/src/test/java/com/aparapi/codegen/test/OpenCLKeywordIdentifiers.java @@ -0,0 +1,31 @@ +/** + * Copyright (c) 2016 - 2018 Syncleus, Inc. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +package com.aparapi.codegen.test; + +public class OpenCLKeywordIdentifiers { + public int global = 1; + public int[] constant = new int[1]; + + int helper(int kernel) { + int local = kernel + global; + int read_only = local + 1; + return read_only; + } + + public void run() { + constant[0] = helper(global); + } +} diff --git a/src/test/java/com/aparapi/codegen/test/OpenCLKeywordIdentifiersTest.java b/src/test/java/com/aparapi/codegen/test/OpenCLKeywordIdentifiersTest.java new file mode 100644 index 00000000..9ac845c0 --- /dev/null +++ b/src/test/java/com/aparapi/codegen/test/OpenCLKeywordIdentifiersTest.java @@ -0,0 +1,61 @@ +/** + * Copyright (c) 2016 - 2018 Syncleus, Inc. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +package com.aparapi.codegen.test; + +import org.junit.Test; + +public class OpenCLKeywordIdentifiersTest extends com.aparapi.codegen.CodeGenJUnitBase { + private static final String[] expectedOpenCL = {"typedef struct This_s{\n" + + " int global__javaReserved;\n" + + " __global int *constant__javaReserved;\n" + + " int passid;\n" + + "}This;\n" + + "int get_pass_id(This *this){\n" + + " return this->passid;\n" + + "}\n" + + "int com_aparapi_codegen_test_OpenCLKeywordIdentifiers__helper(This *this, int kernel__javaReserved){\n" + + " int local__javaReserved = kernel__javaReserved + this->global__javaReserved;\n" + + " int read_only__javaReserved = local__javaReserved + 1;\n" + + " return(read_only__javaReserved);\n" + + "}\n" + + "__kernel void run(\n" + + " int global__javaReserved, \n" + + " __global int *constant__javaReserved, \n" + + " int passid\n" + + "){\n" + + " This thisStruct;\n" + + " This* this=&thisStruct;\n" + + " this->global__javaReserved = global__javaReserved;\n" + + " this->constant__javaReserved = constant__javaReserved;\n" + + " this->passid = passid;\n" + + " {\n" + + " this->constant__javaReserved[0] = com_aparapi_codegen_test_OpenCLKeywordIdentifiers__helper(this, this->global__javaReserved);\n" + + " return;\n" + + " }\n" + + "}\n" + + "\n"}; + private static final Class expectedException = null; + + @Test + public void OpenCLKeywordIdentifiersTest() { + test(com.aparapi.codegen.test.OpenCLKeywordIdentifiers.class, expectedException, expectedOpenCL); + } + + @Test + public void OpenCLKeywordIdentifiersTestWorksWithCaching() { + test(com.aparapi.codegen.test.OpenCLKeywordIdentifiers.class, expectedException, expectedOpenCL); + } +}