From 76e042adf1a0f2d988b208708d1c89fc59fc4715 Mon Sep 17 00:00:00 2001 From: Thanh Nguyen Date: Sun, 17 May 2026 17:47:10 +0700 Subject: [PATCH] fix: escape OpenCL reserved local variable names --- .../aparapi/internal/writer/BlockWriter.java | 26 +++++++++--- .../aparapi/internal/writer/KernelWriter.java | 2 +- .../test/OpenCLReservedLocalVariable.java | 15 +++++++ .../test/OpenCLReservedLocalVariableTest.java | 41 +++++++++++++++++++ 4 files changed, 77 insertions(+), 7 deletions(-) create mode 100644 src/test/java/com/aparapi/codegen/test/OpenCLReservedLocalVariable.java create mode 100644 src/test/java/com/aparapi/codegen/test/OpenCLReservedLocalVariableTest.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..5a23dd2e 100644 --- a/src/main/java/com/aparapi/internal/writer/BlockWriter.java +++ b/src/main/java/com/aparapi/internal/writer/BlockWriter.java @@ -80,8 +80,22 @@ public abstract class BlockWriter{ public final static String arrayDimMangleSuffix = "__javaArrayDimension"; + private static final Set openCLReservedWords = new HashSet(Arrays.asList( + "__kernel", "kernel", "__global", "global", "__local", "local", "__constant", "constant", "__private", + "private", "read_only", "write_only", "read_write", "image2d_t", "image3d_t", "sampler_t", "event_t", + "bool", "char", "uchar", "short", "ushort", "int", "uint", "long", "ulong", "float", "double", "half", + "void", "struct", "union", "enum", "typedef", "const", "volatile", "restrict", "unsigned", "signed", "if", + "else", "while", "do", "for", "switch", "case", "default", "break", "continue", "return", "goto", "sizeof")); + public abstract void write(String _string); + public String escapeOpenCLIdentifier(String identifier) { + if (openCLReservedWords.contains(identifier)) { + return identifier + "$"; + } + return identifier; + } + public void writeln(String _string) { write(_string); newLine(); @@ -417,7 +431,7 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException { } else { final String descriptor = localVariableInfo.getVariableDescriptor(); write(convertType(descriptor, true, true)); - write(localVariableInfo.getVariableName()); + write(escapeOpenCLIdentifier(localVariableInfo.getVariableName())); } } else { if (assignToLocalVariable.isDeclaration()) { @@ -431,7 +445,7 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException { if (localVariableInfo == null) { throw new CodeGenException("outOfScope" + _instruction.getThisPC() + " = "); } else { - write(localVariableInfo.getVariableName() + " = "); + write(escapeOpenCLIdentifier(localVariableInfo.getVariableName()) + " = "); } } @@ -581,13 +595,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(escapeOpenCLIdentifier(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(escapeOpenCLIdentifier(localVariable.getVariableName())); if (adjust == 1) { write("++"); } else if (adjust == -1) { @@ -688,7 +702,7 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException { if (localVariableInfo == null) { throw new CodeGenException("outOfScope" + _instruction.getThisPC() + " = "); } else { - write(localVariableInfo.getVariableName() + " = "); + write(escapeOpenCLIdentifier(localVariableInfo.getVariableName()) + " = "); } } @@ -703,7 +717,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(escapeOpenCLIdentifier(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..6ef24420 100644 --- a/src/main/java/com/aparapi/internal/writer/KernelWriter.java +++ b/src/main/java/com/aparapi/internal/writer/KernelWriter.java @@ -687,7 +687,7 @@ public void writePragma(String _name, boolean _enable) { } write(convertType(descriptor, true, false)); - write(lvi.getVariableName()); + write(escapeOpenCLIdentifier(lvi.getVariableName())); alreadyHasFirstArg = true; localVariableIndex++; diff --git a/src/test/java/com/aparapi/codegen/test/OpenCLReservedLocalVariable.java b/src/test/java/com/aparapi/codegen/test/OpenCLReservedLocalVariable.java new file mode 100644 index 00000000..f5fe51b7 --- /dev/null +++ b/src/test/java/com/aparapi/codegen/test/OpenCLReservedLocalVariable.java @@ -0,0 +1,15 @@ +/** + * Copyright (c) 2016 - 2018 Syncleus, Inc. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + */ +package com.aparapi.codegen.test; + +public class OpenCLReservedLocalVariable { + + public void run() { + int kernel = 1; + int global = kernel + 1; + kernel = global; + } +} diff --git a/src/test/java/com/aparapi/codegen/test/OpenCLReservedLocalVariableTest.java b/src/test/java/com/aparapi/codegen/test/OpenCLReservedLocalVariableTest.java new file mode 100644 index 00000000..40718dcf --- /dev/null +++ b/src/test/java/com/aparapi/codegen/test/OpenCLReservedLocalVariableTest.java @@ -0,0 +1,41 @@ +/** + * Copyright (c) 2016 - 2018 Syncleus, Inc. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + */ +package com.aparapi.codegen.test; + +import org.junit.Test; + +public class OpenCLReservedLocalVariableTest extends com.aparapi.codegen.CodeGenJUnitBase { + + private static final String[] expectedOpenCL = { + "typedef struct This_s{\n" + + "\n" + + " int passid;\n" + + " }This;\n" + + " int get_pass_id(This *this){\n" + + " return this->passid;\n" + + " }\n" + + "\n" + + " __kernel void run(\n" + + " int passid\n" + + " ){\n" + + " This thisStruct;\n" + + " This* this=&thisStruct;\n" + + " this->passid = passid;\n" + + " {\n" + + " int kernel$ = 1;\n" + + " int global$ = kernel$ + 1;\n" + + " kernel$ = global$;\n" + + " return;\n" + + " }\n" + + " }\n" + + " "}; + private static final Class expectedException = null; + + @Test + public void OpenCLReservedLocalVariableTest() { + test(com.aparapi.codegen.test.OpenCLReservedLocalVariable.class, expectedException, expectedOpenCL); + } +}