Skip to content

Fix OpenCL reserved identifier code generation #214

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

Open
gshaowei6 wants to merge 1 commit into
base: master
Choose a base branch
from
Open

Fix OpenCL reserved identifier code generation #214

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
44 changes: 33 additions & 11 deletions src/main/java/com/aparapi/internal/writer/BlockWriter.java
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -80,8 +80,30 @@ public abstract class BlockWriter{

public final static String arrayDimMangleSuffix = "__javaArrayDimension";

private static final Set<String> openCLReservedIdentifiers = new HashSet<String>(Arrays.asList(new String[]{
"__constant", "__global", "__kernel", "__local", "__private", "__read_only", "__read_write",
"__write_only", "atomic_double", "atomic_float", "atomic_int", "atomic_long", "atomic_uint",
"atomic_ulong", "auto", "bool", "bool16", "bool2", "bool3", "bool4", "bool8", "break", "case",
"char", "char16", "char2", "char3", "char4", "char8", "complex", "const", "constant", "continue",
"default", "do", "double", "double16", "double2", "double3", "double4", "double8", "else", "enum",
"event_t", "extern", "float", "float16", "float2", "float3", "float4", "float8", "for", "generic",
"global", "goto", "half", "half16", "half2", "half3", "half4", "half8", "if", "image1d_array_t",
"image1d_buffer_t", "image1d_t", "image2d_array_t", "image2d_t", "image3d_t", "imaginary", "inline",
"int", "int16", "int2", "int3", "int4", "int8", "intptr_t", "kernel", "local", "long", "long16",
"long2", "long3", "long4", "long8", "ndrange_t", "pipe", "private", "ptrdiff_t", "queue_t",
"read_only", "read_write", "register", "reserve_id_t", "restrict", "return", "sampler_t", "short",
"short16", "short2", "short3", "short4", "short8", "signed", "size_t", "sizeof", "static",
"struct", "switch", "typedef", "uchar", "uchar16", "uchar2", "uchar3", "uchar4", "uchar8", "uint",
"uint16", "uint2", "uint3", "uint4", "uint8", "uintptr_t", "ulong", "ulong16", "ulong2", "ulong3",
"ulong4", "ulong8", "union", "unsigned", "ushort", "ushort16", "ushort2", "ushort3", "ushort4",
"ushort8", "void", "volatile", "while", "write_only"}));

public abstract void write(String _string);

protected String mangleIdentifier(String identifier) {
return openCLReservedIdentifiers.contains(identifier) ? identifier + "__javaKeyword" : identifier;
}

public void writeln(String _string) {
write(_string);
newLine();
Expand Down Expand Up @@ -309,7 +331,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 +439,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 +453,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 @@ -481,7 +503,7 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException {

NameAndTypeEntry nameAndTypeEntry = ((AccessField) load).getConstantPoolFieldEntry().getNameAndTypeEntry();
if (isMultiDimensionalArray(nameAndTypeEntry)) {
String arrayName = nameAndTypeEntry.getNameUTF8Entry().getUTF8();
String arrayName = mangleIdentifier(nameAndTypeEntry.getNameUTF8Entry().getUTF8());
write(" * this->" + arrayName + arrayDimMangleSuffix + dim);
}
}
Expand All @@ -506,7 +528,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 @@ -521,7 +543,7 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException {
dim++;
}
NameAndTypeEntry nameAndTypeEntry = ((AccessInstanceField) load).getConstantPoolFieldEntry().getNameAndTypeEntry();
final String arrayName = nameAndTypeEntry.getNameUTF8Entry().getUTF8();
final String arrayName = mangleIdentifier(nameAndTypeEntry.getNameUTF8Entry().getUTF8());
String dimSuffix = isMultiDimensionalArray(nameAndTypeEntry) ? Integer.toString(dim) : "";
write("this->" + arrayName + arrayLengthMangleSuffix + dimSuffix);
} else if (_instruction instanceof AssignToField) {
Expand All @@ -537,7 +559,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 +603,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 +710,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 +725,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
30 changes: 16 additions & 14 deletions src/main/java/com/aparapi/internal/writer/KernelWriter.java
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -241,7 +241,7 @@ public abstract class KernelWriter extends BlockWriter{
getterField = m.getAccessorVariableFieldEntry();
}
if (getterField != null && isThis(_methodCall.getArg(0))) {
String fieldName = getterField.getNameAndTypeEntry().getNameUTF8Entry().getUTF8();
String fieldName = mangleIdentifier(getterField.getNameAndTypeEntry().getNameUTF8Entry().getUTF8());
write("this->");
write(fieldName);
return;
Expand Down Expand Up @@ -285,7 +285,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 @@ -344,6 +344,8 @@ public void writePragma(String _name, boolean _enable) {
final StringBuilder argLine = new StringBuilder();
final StringBuilder assignLine = new StringBuilder();

final String rawFieldName = field.getName();
final String fieldName = mangleIdentifier(rawFieldName);
String signature = field.getDescriptor();

boolean isPointer = false;
Expand All @@ -352,11 +354,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 = rawFieldName.endsWith(Kernel.LOCAL_SUFFIX) ? __local
: (rawFieldName.endsWith(Kernel.CONSTANT_SUFFIX) ? __constant : __global);
Integer privateMemorySize = null;
try {
privateMemorySize = _entryPoint.getClassModel().getPrivateMemorySize(field.getName());
privateMemorySize = _entryPoint.getClassModel().getPrivateMemorySize(rawFieldName);
} catch (ClassParseException e) {
throw new CodeGenException(e);
}
Expand Down Expand Up @@ -423,13 +425,13 @@ public void writePragma(String _name, boolean _enable) {

if (privateMemorySize == null) {
assignLine.append("this->");
assignLine.append(field.getName());
assignLine.append(fieldName);
assignLine.append(" = ");
assignLine.append(field.getName());
assignLine.append(fieldName);
}

argLine.append(field.getName());
thisStructLine.append(field.getName());
argLine.append(fieldName);
thisStructLine.append(fieldName);
if (privateMemorySize == null) {
assigns.add(assignLine.toString());
}
Expand All @@ -441,15 +443,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(rawFieldName) || 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 = fieldName + BlockWriter.arrayLengthMangleSuffix + suffix;

lenStructLine.append("int " + lenName);

Expand All @@ -468,7 +470,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 = fieldName + BlockWriter.arrayDimMangleSuffix + suffix;

dimStructLine.append("int " + dimName);

Expand Down Expand Up @@ -562,7 +564,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 +689,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
33 changes: 33 additions & 0 deletions src/test/java/com/aparapi/codegen/test/OpenCLReservedIdentifier.java
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,33 @@
/**
* 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 com.aparapi.Kernel;

public class OpenCLReservedIdentifier extends Kernel {
int[] global = new int[1];

int add(int kernel) {
int local = kernel + 1;
return local;
}

public void run() {
int event_t = getGlobalId();
int queue_t = add(event_t);
global[event_t] = queue_t;
}
}
61 changes: 61 additions & 0 deletions src/test/java/com/aparapi/codegen/test/OpenCLReservedIdentifierTest.java
View file
Open in desktop
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 OpenCLReservedIdentifierTest extends com.aparapi.codegen.CodeGenJUnitBase {

private static final String[] expectedOpenCL = {
"typedef struct This_s{\n" +
" __global int *global__javaKeyword;\n" +
" int passid;\n" +
" }This;\n" +
" int get_pass_id(This *this){\n" +
" return this->passid;\n" +
" }\n" +
" int com_aparapi_codegen_test_OpenCLReservedIdentifier__add(This *this, int kernel__javaKeyword){\n" +
" int local__javaKeyword = kernel__javaKeyword + 1;\n" +
" return(local__javaKeyword);\n" +
" }\n" +
" __kernel void run(\n" +
" __global int *global__javaKeyword, \n" +
" int passid\n" +
" ){\n" +
" This thisStruct;\n" +
" This* this=&thisStruct;\n" +
" this->global__javaKeyword = global__javaKeyword;\n" +
" this->passid = passid;\n" +
" {\n" +
" int event_t__javaKeyword = get_global_id(0);\n" +
" int queue_t__javaKeyword = com_aparapi_codegen_test_OpenCLReservedIdentifier__add(this, event_t__javaKeyword);\n" +
" this->global__javaKeyword[event_t__javaKeyword] = queue_t__javaKeyword;\n" +
" return;\n" +
" }\n" +
" }"
};
private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;

@Test
public void OpenCLReservedIdentifierTest() {
test(com.aparapi.codegen.test.OpenCLReservedIdentifier.class, expectedException, expectedOpenCL);
}

@Test
public void OpenCLReservedIdentifierTestWorksWithCaching() {
test(com.aparapi.codegen.test.OpenCLReservedIdentifier.class, expectedException, expectedOpenCL);
}
}