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 b74c59c501b9798a2f4df8170d53b7d013516dba..57f259350c762d55c085b4bcf416c3f566e2d6a3 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 72812b94265004265d558faffb75261adf72792a..68f76233b21030617a2adc8b5b0579bf2ff2c499 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 57b36ff33b7a2bbad726668795165eb07c127f9c..46d2266a191eadbf983dd9540ba7c711dede66de 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 a3fc293f5df609bd9c67ba343c25466b9f0aa7bc..9751d14564e9e017a79891d07ed92cb9fe5b3a54 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 9f82ffa02dd0adcf66ba900df8e4362cbab61f60..1af511ed29ab58700aec01553b03ab2b74c7e503 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 6631510bc97daf0f671e06eae6aca557c05de80b..d8f93a58556e55979efe6fa13066b33b953766ec 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 017ae676444bfad7f4ad8a9a75fbe737a58382f7..34bc6efe6cbbbabea3e612cbff59b18bfa9678b6 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 8be389b5128681c67e992a0dc339132d68ab44f1..3398d3fee9a5590c8a7bb06edc089073acd1bea8 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 0000000000000000000000000000000000000000..5c09b9e92fc308b94cb05b19595a4b738b0c11bc --- /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}**/