From 1b2241c3bfb4f93d46108566db269e285a7c0616 Mon Sep 17 00:00:00 2001 From: gshaowei6 Date: Sat, 16 May 2026 18:10:38 +0800 Subject: [PATCH] Mangle OpenCL reserved Java identifiers --- .../aparapi/internal/writer/BlockWriter.java | 44 +++++++++---- .../aparapi/internal/writer/KernelWriter.java | 30 ++++----- .../test/OpenCLReservedIdentifier.java | 33 ++++++++++ .../test/OpenCLReservedIdentifierTest.java | 61 +++++++++++++++++++ 4 files changed, 143 insertions(+), 25 deletions(-) create mode 100644 src/test/java/com/aparapi/codegen/test/OpenCLReservedIdentifier.java create mode 100644 src/test/java/com/aparapi/codegen/test/OpenCLReservedIdentifierTest.java diff --git a/src/main/java/com/aparapi/internal/writer/BlockWriter.java b/src/main/java/com/aparapi/internal/writer/BlockWriter.java index 013239d0..bc636bd3 100644 --- a/src/main/java/com/aparapi/internal/writer/BlockWriter.java +++ b/src/main/java/com/aparapi/internal/writer/BlockWriter.java @@ -80,8 +80,30 @@ public abstract class BlockWriter{ public final static String arrayDimMangleSuffix = "__javaArrayDimension"; + private static final Set openCLReservedIdentifiers = new HashSet(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(); @@ -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(); @@ -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()) { @@ -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()) + " = "); } } @@ -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); } } @@ -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) { @@ -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) { @@ -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) { @@ -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) { @@ -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()) + " = "); } } @@ -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)) { diff --git a/src/main/java/com/aparapi/internal/writer/KernelWriter.java b/src/main/java/com/aparapi/internal/writer/KernelWriter.java index 0a72a33f..10b531a2 100644 --- a/src/main/java/com/aparapi/internal/writer/KernelWriter.java +++ b/src/main/java/com/aparapi/internal/writer/KernelWriter.java @@ -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; @@ -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("])"); @@ -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; @@ -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); } @@ -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()); } @@ -441,7 +443,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(rawFieldName) || isPointer && numDimensions > 1) { for (int i = 0; i < numDimensions; i++) { final StringBuilder lenStructLine = new StringBuilder(); @@ -449,7 +451,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 = fieldName + BlockWriter.arrayLengthMangleSuffix + suffix; lenStructLine.append("int " + lenName); @@ -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); @@ -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 @@ -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++; diff --git a/src/test/java/com/aparapi/codegen/test/OpenCLReservedIdentifier.java b/src/test/java/com/aparapi/codegen/test/OpenCLReservedIdentifier.java new file mode 100644 index 00000000..18141733 --- /dev/null +++ b/src/test/java/com/aparapi/codegen/test/OpenCLReservedIdentifier.java @@ -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; + } +} diff --git a/src/test/java/com/aparapi/codegen/test/OpenCLReservedIdentifierTest.java b/src/test/java/com/aparapi/codegen/test/OpenCLReservedIdentifierTest.java new file mode 100644 index 00000000..73dc19cf --- /dev/null +++ b/src/test/java/com/aparapi/codegen/test/OpenCLReservedIdentifierTest.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 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 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); + } +}