From 0bbbb4cd8cd52342c6615481e66fe39881f29654 Mon Sep 17 00:00:00 2001
From: Eric Caspole <eric.caspole@amd.com>
Date: Wed, 23 Jan 2013 16:09:04 +0000
Subject: [PATCH] Fix for issue 92, static arrays in kernels
---
.../src/java/com/amd/aparapi/Config.java | 3 +-
.../src/java/com/amd/aparapi/Entrypoint.java | 10 +++--
.../test/ClassHasStaticFieldAccess.java | 31 +++++++++++++-
.../aparapi/test/runtime/UseStaticArray.java | 40 +++++++++++++++++++
4 files changed, 78 insertions(+), 6 deletions(-)
create mode 100644 test/runtime/src/java/com/amd/aparapi/test/runtime/UseStaticArray.java
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 f45c3f8e..699be065 100644
--- a/com.amd.aparapi/src/java/com/amd/aparapi/Config.java
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/Config.java
@@ -147,7 +147,8 @@ class Config{
static final boolean enablePUTSTATIC = Boolean.getBoolean(propPkgName + ".enable.PUTSTATIC");
- static final boolean enableGETSTATIC = Boolean.getBoolean(propPkgName + ".enable.GETSTATIC");
+ // Allow static array accesses
+ static final boolean enableGETSTATIC = true; //Boolean.getBoolean(propPkgName + ".enable.GETSTATIC");
static final boolean enableINVOKEINTERFACE = Boolean.getBoolean(propPkgName + ".enable.INVOKEINTERFACE");
diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/Entrypoint.java b/com.amd.aparapi/src/java/com/amd/aparapi/Entrypoint.java
index 34868697..a77b2ae4 100644
--- a/com.amd.aparapi/src/java/com/amd/aparapi/Entrypoint.java
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/Entrypoint.java
@@ -558,8 +558,9 @@ class Entrypoint{
AssignToArrayElement assignment = (AssignToArrayElement) instruction;
Instruction arrayRef = assignment.getArrayRef();
- if (arrayRef instanceof I_GETFIELD) {
- I_GETFIELD getField = (I_GETFIELD) arrayRef;
+ // AccessField here allows instance and static array refs
+ if (arrayRef instanceof AccessField) {
+ AccessField getField = (AccessField) arrayRef;
FieldEntry field = getField.getConstantPoolFieldEntry();
String assignedArrayFieldName = field.getNameAndTypeEntry().getNameUTF8Entry().getUTF8();
arrayFieldAssignments.add(assignedArrayFieldName);
@@ -570,8 +571,9 @@ class Entrypoint{
AccessArrayElement access = (AccessArrayElement) instruction;
Instruction arrayRef = access.getArrayRef();
- if (arrayRef instanceof I_GETFIELD) {
- I_GETFIELD getField = (I_GETFIELD) arrayRef;
+ // AccessField here allows instance and static array refs
+ if (arrayRef instanceof AccessField) {
+ AccessField getField = (AccessField) arrayRef;
FieldEntry field = getField.getConstantPoolFieldEntry();
String accessedArrayFieldName = field.getNameAndTypeEntry().getNameUTF8Entry().getUTF8();
arrayFieldAccesses.add(accessedArrayFieldName);
diff --git a/test/codegen/src/java/com/amd/aparapi/test/ClassHasStaticFieldAccess.java b/test/codegen/src/java/com/amd/aparapi/test/ClassHasStaticFieldAccess.java
index ea061b7e..b04651b2 100644
--- a/test/codegen/src/java/com/amd/aparapi/test/ClassHasStaticFieldAccess.java
+++ b/test/codegen/src/java/com/amd/aparapi/test/ClassHasStaticFieldAccess.java
@@ -13,4 +13,33 @@ public class ClassHasStaticFieldAccess{
}
}
}
-/**{Throws{ClassParseException}Throws}**/
+/**{OpenCL{
+typedef struct This_s{
+ __global int *ints;
+ int foo;
+ int passid;
+}This;
+int get_pass_id(This *this){
+ return this->passid;
+}
+__kernel void run(
+ __global int *ints,
+ int foo,
+ int passid
+){
+ This thisStruct;
+ This* this=&thisStruct;
+ this->ints = ints;
+ this->foo = foo;
+ this->passid = passid;
+ {
+ for (int i = 0; i<1024; i++){
+ if ((i % 2)==0){
+ this->ints[i] = foo;
+ }
+ }
+ return;
+ }
+}
+
+}OpenCL}**/
diff --git a/test/runtime/src/java/com/amd/aparapi/test/runtime/UseStaticArray.java b/test/runtime/src/java/com/amd/aparapi/test/runtime/UseStaticArray.java
new file mode 100644
index 00000000..2564518f
--- /dev/null
+++ b/test/runtime/src/java/com/amd/aparapi/test/runtime/UseStaticArray.java
@@ -0,0 +1,40 @@
+package com.amd.aparapi.test.runtime;
+
+import static org.junit.Assert.assertTrue;
+import org.junit.Test;
+import com.amd.aparapi.Kernel;
+
+
+public class UseStaticArray extends Kernel {
+
+ static final int size = 256;
+
+ static final int[] values = new int[size];
+ static final int[] results = new int[size];
+
+ @Override public void run() {
+ int gid = getGlobalId();
+ results[gid] = values[gid];
+ }
+
+ @Test public void test() {
+
+ for(int i=0; i<size; i++) {
+ values[i] = i;
+ results[i] = 0;
+ }
+
+ execute(size);
+
+ assertTrue("ran on GPU", getExecutionMode()==Kernel.EXECUTION_MODE.GPU);
+
+ for(int i=0; i<size; i++) {
+ assertTrue( "results == fooBar", results[i] == values[i] );
+ }
+ }
+
+ public static void main(String args[]) {
+ UseStaticArray k = new UseStaticArray();
+ k.test();
+ }
+}
--
GitLab