Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
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
28 changes: 17 additions & 11 deletions src/main/java/com/aparapi/internal/writer/BlockWriter.java
Original file line number Diff line number Diff line change
Expand Up @@ -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();
Expand Down Expand Up @@ -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();
Expand Down Expand Up @@ -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()) {
Expand All @@ -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()) + " = ");
}
}

Expand Down Expand Up @@ -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));
}
}

Expand All @@ -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) {

Expand All @@ -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;

Expand All @@ -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<?>) {
Expand Down Expand Up @@ -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) {
Expand Down Expand Up @@ -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()) + " = ");
}

}
Expand All @@ -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)) {
Expand Down
62 changes: 48 additions & 14 deletions src/main/java/com/aparapi/internal/writer/KernelWriter.java
Original file line number Diff line number Diff line change
Expand Up @@ -117,6 +117,38 @@ public abstract class KernelWriter extends BlockWriter{

private Entrypoint entryPoint = null;

private static final Set<String> openCLReservedIdentifiers = createOpenCLReservedIdentifiers();

private static Set<String> createOpenCLReservedIdentifiers() {
final Set<String> identifiers = new HashSet<String>();
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<String, String> javaToCLIdentifierMap = new HashMap<String, String>();
{
javaToCLIdentifierMap.put("getGlobalId()I", "get_global_id(0)");
Expand Down Expand Up @@ -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()
Expand Down Expand Up @@ -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("])");
Expand Down Expand Up @@ -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();

Expand All @@ -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);
}
Expand Down Expand Up @@ -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());
}
Expand All @@ -441,15 +475,15 @@ 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();
final StringBuilder lenArgLine = new StringBuilder();
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);

Expand All @@ -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);

Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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++;
Expand Down
Original file line number Diff line number Diff line change
@@ -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);
}
}
Original file line number Diff line number Diff line change
@@ -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<? extends com.aparapi.internal.exception.AparapiException> 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);
}
}