From a272d8e7dd68bc9de532e38a7cd24cf0ff413534 Mon Sep 17 00:00:00 2001 From: Gary Frost <frost.gary@gmail.com> Date: Sun, 20 Nov 2011 17:18:05 +0000 Subject: [PATCH] Submitted Witold Bolt's patch to allow static methods. Added Witold's junit test case and modified previous test cases which trapped the use of static methods. --- .../src/java/com/amd/aparapi/ClassModel.java | 4 ++ .../com/amd/aparapi/ClassParseException.java | 1 - .../src/java/com/amd/aparapi/Config.java | 2 - .../java/com/amd/aparapi/KernelWriter.java | 22 ++++++---- .../src/java/com/amd/aparapi/MethodModel.java | 12 ++++-- .../com/amd/aparapi/CodeGenJUnitBase.java | 2 +- .../amd/aparapi/test/CallObjectStatic.java | 23 +++++++++- .../aparapi/test/ClassHasStaticMethod.java | 34 ++++++++++++++- .../amd/aparapi/test/StaticMethodCall.java | 42 +++++++++++++++++++ 9 files changed, 126 insertions(+), 16 deletions(-) create mode 100644 test/codegen/src/java/com/amd/aparapi/test/StaticMethodCall.java diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/ClassModel.java b/com.amd.aparapi/src/java/com/amd/aparapi/ClassModel.java index b74c59c5..57f25935 100644 --- a/com.amd.aparapi/src/java/com/amd/aparapi/ClassModel.java +++ b/com.amd.aparapi/src/java/com/amd/aparapi/ClassModel.java @@ -2089,6 +2089,10 @@ class ClassModel{ return (methodAccessFlags); } + public boolean isStatic() { + return (Access.STATIC.bitIsSet(methodAccessFlags)); + } + AttributePool getAttributePool() { return (methodAttributePool); } diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/ClassParseException.java b/com.amd.aparapi/src/java/com/amd/aparapi/ClassParseException.java index 72812b94..68f76233 100644 --- a/com.amd.aparapi/src/java/com/amd/aparapi/ClassParseException.java +++ b/com.amd.aparapi/src/java/com/amd/aparapi/ClassParseException.java @@ -51,7 +51,6 @@ package com.amd.aparapi; NONE("none"), // ARRAY_RETURN("We don't support areturn instructions"), // PUTFIELD("We don't support putstatic instructions"), // - INVOKESTATIC("We don't support invokestatic instructions"), // INVOKEINTERFACE("We don't support invokeinterface instructions"), // GETSTATIC("We don't support getstatic instructions"), // ATHROW("We don't support athrow instructions"), // diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/Config.java b/com.amd.aparapi/src/java/com/amd/aparapi/Config.java index 57b36ff3..46d2266a 100644 --- a/com.amd.aparapi/src/java/com/amd/aparapi/Config.java +++ b/com.amd.aparapi/src/java/com/amd/aparapi/Config.java @@ -120,8 +120,6 @@ class Config{ static final boolean enableGETSTATIC = Boolean.getBoolean(propPkgName + ".enable.GETSTATIC"); - static final boolean enableINVOKESTATIC = Boolean.getBoolean(propPkgName + ".enable.INVOKESTATIC"); - static final boolean enableINVOKEINTERFACE = Boolean.getBoolean(propPkgName + ".enable.INVOKEINTERFACE"); static final boolean enableMONITOR = Boolean.getBoolean(propPkgName + ".enable.MONITOR"); diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/KernelWriter.java b/com.amd.aparapi/src/java/com/amd/aparapi/KernelWriter.java index a3fc293f..9751d145 100644 --- a/com.amd.aparapi/src/java/com/amd/aparapi/KernelWriter.java +++ b/com.amd.aparapi/src/java/com/amd/aparapi/KernelWriter.java @@ -429,30 +429,37 @@ abstract class KernelWriter extends BlockWriter{ } write(convertType(returnType, true)); - write(mm.getName()); + write(mm.getName()+"("); + if(!mm.getMethod().isStatic()) { if (mm.getMethod().getClassModel() == _entryPoint.getClassModel() || mm.getMethod().getClassModel().isSuperClass(_entryPoint.getClassModel().getClassWeAreModelling())) { - write("(This *this"); + write("This *this"); } else { // Call to an object member or superclass of member for (ClassModel c : _entryPoint.getObjectArrayFieldsClasses().values()) { if (mm.getMethod().getClassModel() == c) { - write("( __global " + mm.getMethod().getClassModel().getClassWeAreModelling().getName().replace(".", "_") + write("__global " + mm.getMethod().getClassModel().getClassWeAreModelling().getName().replace(".", "_") + " *this"); break; } else if (mm.getMethod().getClassModel().isSuperClass(c.getClassWeAreModelling())) { - write("( __global " + c.getClassWeAreModelling().getName().replace(".", "_") + " *this"); + write("__global " + c.getClassWeAreModelling().getName().replace(".", "_") + " *this"); break; } } } + } + + boolean alreadyHasFirstArg = !mm.getMethod().isStatic(); + LocalVariableTableEntry lvte = mm.getLocalVariableTableEntry(); for (LocalVariableInfo lvi : lvte) { - if (lvi.getStart() == 0 && lvi.getVariableIndex() != 0) { // full scope but skip this + if (lvi.getStart() == 0 && (lvi.getVariableIndex() != 0 || mm.getMethod().isStatic())) { // full scope but skip this String descriptor = lvi.getVariableDescriptor(); - write(", "); - + if(alreadyHasFirstArg) { + write(", "); + } + // Arrays always map to __global arrays if (descriptor.startsWith("[")) { write(" __global "); @@ -460,6 +467,7 @@ abstract class KernelWriter extends BlockWriter{ write(convertType(descriptor, true)); write(lvi.getVariableName()); + alreadyHasFirstArg=true; } } write(")"); diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/MethodModel.java b/com.amd.aparapi/src/java/com/amd/aparapi/MethodModel.java index 9f82ffa0..1af511ed 100644 --- a/com.amd.aparapi/src/java/com/amd/aparapi/MethodModel.java +++ b/com.amd.aparapi/src/java/com/amd/aparapi/MethodModel.java @@ -120,6 +120,8 @@ class MethodModel{ private boolean methodIsGetter; private boolean methodIsSetter; + + // Only setters can use putfield private boolean usesPutfield; @@ -246,27 +248,31 @@ class MethodModel{ if ((!Config.enablePUTSTATIC) && instruction instanceof I_PUTSTATIC) { throw new ClassParseException(instruction, ClassParseException.TYPE.PUTFIELD); } - if ((!Config.enableINVOKESTATIC) && instruction instanceof I_INVOKESTATIC) { - throw new ClassParseException(instruction, ClassParseException.TYPE.INVOKESTATIC); - } + if ((!Config.enableINVOKEINTERFACE) && instruction instanceof I_INVOKEINTERFACE) { throw new ClassParseException(instruction, ClassParseException.TYPE.INVOKEINTERFACE); } + if ((!Config.enableGETSTATIC) && instruction instanceof I_GETSTATIC) { throw new ClassParseException(instruction, ClassParseException.TYPE.GETSTATIC); } + if ((!Config.enableATHROW) && instruction instanceof I_ATHROW) { throw new ClassParseException(instruction, ClassParseException.TYPE.ATHROW); } + if ((!Config.enableMONITOR) && ((instruction instanceof I_MONITORENTER) || (instruction instanceof I_MONITOREXIT))) { throw new ClassParseException(instruction, ClassParseException.TYPE.SYNCHRONIZE); } + if ((!Config.enableNEW) && instruction instanceof New) { throw new ClassParseException(instruction, ClassParseException.TYPE.NEW); } + if (instruction instanceof I_CALOAD || instruction instanceof I_CASTORE) { throw new ClassParseException(instruction, ClassParseException.TYPE.CHARARRAY); } + if (instruction instanceof I_AASTORE) { throw new ClassParseException(instruction, ClassParseException.TYPE.ARRAYALIAS); } diff --git a/test/codegen/src/java/com/amd/aparapi/CodeGenJUnitBase.java b/test/codegen/src/java/com/amd/aparapi/CodeGenJUnitBase.java index 6631510b..d8f93a58 100644 --- a/test/codegen/src/java/com/amd/aparapi/CodeGenJUnitBase.java +++ b/test/codegen/src/java/com/amd/aparapi/CodeGenJUnitBase.java @@ -91,7 +91,7 @@ public class CodeGenJUnitBase{ } assertTrue(_class.getSimpleName(), same); } else { - assertTrue("Expected exception " + _expectedExceptionType, false); + assertTrue("Expected exception " + _expectedExceptionType +" Instead we got {\n"+actual+"\n}", false); } } catch (Throwable t) { diff --git a/test/codegen/src/java/com/amd/aparapi/test/CallObjectStatic.java b/test/codegen/src/java/com/amd/aparapi/test/CallObjectStatic.java index 017ae676..34bc6efe 100644 --- a/test/codegen/src/java/com/amd/aparapi/test/CallObjectStatic.java +++ b/test/codegen/src/java/com/amd/aparapi/test/CallObjectStatic.java @@ -16,4 +16,25 @@ public class CallObjectStatic extends Kernel{ int out[] = new int[2]; } -/**{Throws{ClassParseException}Throws}**/ +/**{OpenCL{ +typedef struct This_s{ +__global int *out; +int passid; +}This; +int get_pass_id(This *this){ +return this->passid; +} +__kernel void run( +__global int *out, +int passid +){ +This thisStruct; +This* this=&thisStruct; +this->out = out; +this->passid = passid; +{ +this->out[0] = foo(); +return; +} +} +}OpenCL}**/ diff --git a/test/codegen/src/java/com/amd/aparapi/test/ClassHasStaticMethod.java b/test/codegen/src/java/com/amd/aparapi/test/ClassHasStaticMethod.java index 8be389b5..3398d3fe 100644 --- a/test/codegen/src/java/com/amd/aparapi/test/ClassHasStaticMethod.java +++ b/test/codegen/src/java/com/amd/aparapi/test/ClassHasStaticMethod.java @@ -19,4 +19,36 @@ public class ClassHasStaticMethod{ } } } -/**{Throws{ClassParseException}Throws}**/ +/**{OpenCL{ +typedef struct This_s{ +__global int *ints; +int passid; +}This; +int get_pass_id(This *this){ +return this->passid; +} +int com_amd_aparapi_test_ClassHasStaticMethod__getIntAndReturnIt(int a){ +return((1 - a)); +} +__kernel void run( +__global int *ints, +int passid +){ +This thisStruct; +This* this=&thisStruct; +this->ints = ints; +this->passid = passid; +{ +int foo = 1; +for (int i = 0; i<1024; i++){ +if ((i % 2)==0){ +this->ints[i] = foo; +} else { +this->ints[i] = com_amd_aparapi_test_ClassHasStaticMethod__getIntAndReturnIt(foo); +} +} +return; +} +} +}OpenCL}**/ + diff --git a/test/codegen/src/java/com/amd/aparapi/test/StaticMethodCall.java b/test/codegen/src/java/com/amd/aparapi/test/StaticMethodCall.java new file mode 100644 index 00000000..5c09b9e9 --- /dev/null +++ b/test/codegen/src/java/com/amd/aparapi/test/StaticMethodCall.java @@ -0,0 +1,42 @@ +package com.amd.aparapi.test; + +import com.amd.aparapi.Kernel; + +public class StaticMethodCall extends Kernel{ + public static int add(int i, int j) { + return i+j; + } + + public void run() { + out[0] = add(1,2); + } + + int out[] = new int[1]; +} + +/**{OpenCL{ +typedef struct This_s{ + __global int *out; + int passid; +}This; +int get_pass_id(This *this){ + return this->passid; +} +int com_amd_aparapi_test_StaticMethodCall__add(int i, int j){ + return((i + j)); +} +__kernel void run( + __global int *out, + int passid +){ + This thisStruct; + This* this=&thisStruct; + this->out = out; + this->passid = passid; + { + this->out[0] = com_amd_aparapi_test_StaticMethodCall__add(1, 2); + return; + } +} + +}OpenCL}**/ -- GitLab