diff --git a/src/main/java/com/aparapi/internal/writer/BlockWriter.java b/src/main/java/com/aparapi/internal/writer/BlockWriter.java
index 013239d0..93c7faa6 100644
--- a/src/main/java/com/aparapi/internal/writer/BlockWriter.java
+++ b/src/main/java/com/aparapi/internal/writer/BlockWriter.java
@@ -304,6 +304,14 @@ public void writeSequence(Instruction _first, Instruction _last) throws CodeGenE
 
    }
 
+   private void writeNewArrayDimensions(I_NEWARRAY _newArray) throws CodeGenException {
+      for (Instruction operand = _newArray.getFirstChild(); operand != null; operand = operand.getNextExpr()) {
+         write("[");
+         writeInstruction(operand);
+         write("]");
+      }
+   }
+
    protected void writeGetterBlock(FieldEntry accessorVariableFieldEntry) {
       write("{");
       in();
@@ -396,15 +404,7 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException {
           writeComposite((CompositeInstruction) _instruction);
 
       } else if (_instruction instanceof I_NEWARRAY) {
-          if (_instruction.getParentExpr() instanceof Return) {
-              throw new CodeGenException("'newarray' is not allowed after 'return'");
-          }
-
-          for (Instruction operand = _instruction.getFirstChild(); operand != null; operand = operand.getNextExpr()) {
-              write("[");
-              writeInstruction(operand);
-              write("]");
-          }
+          writeNewArrayDimensions((I_NEWARRAY) _instruction);
 
       } else if (_instruction instanceof AssignToLocalVariable) {
          final AssignToLocalVariable assignToLocalVariable = (AssignToLocalVariable) _instruction;
@@ -632,6 +632,19 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException {
       } else if (_instruction instanceof Return) {
 
          final Return ret = (Return) _instruction;
+         if (ret.getFirstChild() instanceof I_NEWARRAY) {
+            final I_NEWARRAY newArray = (I_NEWARRAY) ret.getFirstChild();
+            final String returnArrayName = "returnArray" + ret.getThisPC();
+            write(convertType(ret.getMethod().getReturnType(), true, true));
+            write(returnArrayName);
+            writeNewArrayDimensions(newArray);
+            write(";");
+            newLine();
+            write("return(");
+            write(returnArrayName);
+            write(")");
+            return;
+         }
          write("return");
          if (ret.getStackConsumeCount() > 0) {
             write("(");
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnArrayNewTestSupport.java b/src/test/java/com/aparapi/codegen/test/ReturnArrayNewTestSupport.java
new file mode 100644
index 00000000..6934f08b
--- /dev/null
+++ b/src/test/java/com/aparapi/codegen/test/ReturnArrayNewTestSupport.java
@@ -0,0 +1,49 @@
+/**
+ * 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;
+
+final class ReturnArrayNewTestSupport {
+
+    private ReturnArrayNewTestSupport() {
+    }
+
+    static String[] expectedOpenCL(String openCLType, String className, String methodName) {
+        String qualifiedMethodName = "com_aparapi_codegen_test_" + className + "__" + methodName;
+        return new String[] {
+            "typedef struct This_s{\n"
+                + "   int passid;\n"
+                + "}This;\n"
+                + "int get_pass_id(This *this){\n"
+                + "   return this->passid;\n"
+                + "}\n"
+                + " __global " + openCLType + "* " + qualifiedMethodName + "(This *this){\n"
+                + "   " + openCLType + " returnArray5[1024];\n"
+                + "   return(returnArray5);\n"
+                + "}\n"
+                + "__kernel void run(\n"
+                + "   int passid\n"
+                + "){\n"
+                + "   This thisStruct;\n"
+                + "   This* this=&thisStruct;\n"
+                + "   this->passid = passid;\n"
+                + "   {\n"
+                + "      " + qualifiedMethodName + "(this);\n"
+                + "      return;\n"
+                + "   }\n"
+                + "}"
+        };
+    }
+}
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnBooleanNewArray.java b/src/test/java/com/aparapi/codegen/test/ReturnBooleanNewArray.java
index 2befe672..ed0277fe 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnBooleanNewArray.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnBooleanNewArray.java
@@ -26,4 +26,27 @@ public void run() {
         returnBooleanNewArray();
     }
 }
-/**{Throws{CodeGenException}Throws}**/
+/**{OpenCL{
+ typedef struct This_s{
+ int passid;
+ }This;
+ int get_pass_id(This *this){
+ return this->passid;
+ }
+ __global char* com_aparapi_codegen_test_ReturnBooleanNewArray__returnBooleanNewArray(This *this){
+ char returnArray5[1024];
+ return(returnArray5);
+ }
+ __kernel void run(
+ int passid
+ ){
+ This thisStruct;
+ This* this=&thisStruct;
+ this->passid = passid;
+ {
+ com_aparapi_codegen_test_ReturnBooleanNewArray__returnBooleanNewArray(this);
+ return;
+ }
+ }
+
+ }OpenCL}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnBooleanNewArrayTest.java b/src/test/java/com/aparapi/codegen/test/ReturnBooleanNewArrayTest.java
index 85a81499..606dad80 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnBooleanNewArrayTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnBooleanNewArrayTest.java
@@ -15,13 +15,13 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.CodeGenException;
 import org.junit.Test;
 
 public class ReturnBooleanNewArrayTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
-    private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = CodeGenException.class;
+    private static final String[] expectedOpenCL = ReturnArrayNewTestSupport.expectedOpenCL("char",
+            "ReturnBooleanNewArray", "returnBooleanNewArray");
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     @Test
     public void ReturnBooleanNewArrayTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnByteArrayNew.java b/src/test/java/com/aparapi/codegen/test/ReturnByteArrayNew.java
index ea6db60f..f2db8f2e 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnByteArrayNew.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnByteArrayNew.java
@@ -25,4 +25,27 @@ public void run() {
         returnByteArrayNew();
     }
 }
-/**{Throws{CodeGenException}Throws}**/
+/**{OpenCL{
+ typedef struct This_s{
+ int passid;
+ }This;
+ int get_pass_id(This *this){
+ return this->passid;
+ }
+ __global char* com_aparapi_codegen_test_ReturnByteArrayNew__returnByteArrayNew(This *this){
+ char returnArray5[1024];
+ return(returnArray5);
+ }
+ __kernel void run(
+ int passid
+ ){
+ This thisStruct;
+ This* this=&thisStruct;
+ this->passid = passid;
+ {
+ com_aparapi_codegen_test_ReturnByteArrayNew__returnByteArrayNew(this);
+ return;
+ }
+ }
+
+ }OpenCL}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnByteArrayNewTest.java b/src/test/java/com/aparapi/codegen/test/ReturnByteArrayNewTest.java
index 24ab1fb8..2e8ab5af 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnByteArrayNewTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnByteArrayNewTest.java
@@ -15,12 +15,12 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.CodeGenException;
 import org.junit.Test;
 
 public class ReturnByteArrayNewTest extends com.aparapi.codegen.CodeGenJUnitBase {
-    private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = CodeGenException.class;
+    private static final String[] expectedOpenCL = ReturnArrayNewTestSupport.expectedOpenCL("char",
+            "ReturnByteArrayNew", "returnByteArrayNew");
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     
     @Test
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayNew.java b/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayNew.java
index 93f1800d..5620f8be 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayNew.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayNew.java
@@ -25,4 +25,27 @@ public void run() {
         returnDoubleArrayNew();
     }
 }
-/**{Throws{CodeGenException}Throws}**/
+/**{OpenCL{
+ typedef struct This_s{
+ int passid;
+ }This;
+ int get_pass_id(This *this){
+ return this->passid;
+ }
+ __global double* com_aparapi_codegen_test_ReturnDoubleArrayNew__returnDoubleArrayNew(This *this){
+ double returnArray5[1024];
+ return(returnArray5);
+ }
+ __kernel void run(
+ int passid
+ ){
+ This thisStruct;
+ This* this=&thisStruct;
+ this->passid = passid;
+ {
+ com_aparapi_codegen_test_ReturnDoubleArrayNew__returnDoubleArrayNew(this);
+ return;
+ }
+ }
+
+ }OpenCL}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayNewTest.java b/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayNewTest.java
index 95496478..f851c922 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayNewTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayNewTest.java
@@ -15,13 +15,13 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.CodeGenException;
 import org.junit.Test;
 
 public class ReturnDoubleArrayNewTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
-    private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = CodeGenException.class;
+    private static final String[] expectedOpenCL = ReturnArrayNewTestSupport.expectedOpenCL("double",
+            "ReturnDoubleArrayNew", "returnDoubleArrayNew");
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     @Test
     public void ReturnDoubleArrayNewTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayNew.java b/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayNew.java
index c15fd367..f3a7bd48 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayNew.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayNew.java
@@ -25,4 +25,27 @@ public void run() {
         returnFloatArrayNew();
     }
 }
-/**{Throws{CodeGenException}Throws}**/
+/**{OpenCL{
+ typedef struct This_s{
+ int passid;
+ }This;
+ int get_pass_id(This *this){
+ return this->passid;
+ }
+ __global float* com_aparapi_codegen_test_ReturnFloatArrayNew__returnFloatArrayNew(This *this){
+ float returnArray5[1024];
+ return(returnArray5);
+ }
+ __kernel void run(
+ int passid
+ ){
+ This thisStruct;
+ This* this=&thisStruct;
+ this->passid = passid;
+ {
+ com_aparapi_codegen_test_ReturnFloatArrayNew__returnFloatArrayNew(this);
+ return;
+ }
+ }
+
+ }OpenCL}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayNewTest.java b/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayNewTest.java
index 59fd675b..15dd4b1b 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayNewTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayNewTest.java
@@ -15,13 +15,13 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.CodeGenException;
 import org.junit.Test;
 
 public class ReturnFloatArrayNewTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
-    private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = CodeGenException.class;
+    private static final String[] expectedOpenCL = ReturnArrayNewTestSupport.expectedOpenCL("float",
+            "ReturnFloatArrayNew", "returnFloatArrayNew");
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     @Test
     public void ReturnFloatArrayNewTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnIntArrayNew.java b/src/test/java/com/aparapi/codegen/test/ReturnIntArrayNew.java
index 114e8504..986dcd51 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnIntArrayNew.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnIntArrayNew.java
@@ -25,4 +25,27 @@ public void run() {
         returnIntArrayNew();
     }
 }
-/**{Throws{CodeGenException}Throws}**/
+/**{OpenCL{
+ typedef struct This_s{
+ int passid;
+ }This;
+ int get_pass_id(This *this){
+ return this->passid;
+ }
+ __global int* com_aparapi_codegen_test_ReturnIntArrayNew__returnIntArrayNew(This *this){
+ int returnArray5[1024];
+ return(returnArray5);
+ }
+ __kernel void run(
+ int passid
+ ){
+ This thisStruct;
+ This* this=&thisStruct;
+ this->passid = passid;
+ {
+ com_aparapi_codegen_test_ReturnIntArrayNew__returnIntArrayNew(this);
+ return;
+ }
+ }
+
+ }OpenCL}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnIntArrayNewTest.java b/src/test/java/com/aparapi/codegen/test/ReturnIntArrayNewTest.java
index a24fd3c2..9b916e39 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnIntArrayNewTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnIntArrayNewTest.java
@@ -15,13 +15,13 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.CodeGenException;
 import org.junit.Test;
 
 public class ReturnIntArrayNewTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
-    private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = CodeGenException.class;
+    private static final String[] expectedOpenCL = ReturnArrayNewTestSupport.expectedOpenCL("int",
+            "ReturnIntArrayNew", "returnIntArrayNew");
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     @Test
     public void ReturnIntArrayNewTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnLongArrayNew.java b/src/test/java/com/aparapi/codegen/test/ReturnLongArrayNew.java
index 0a45fe1a..1a18722b 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnLongArrayNew.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnLongArrayNew.java
@@ -25,4 +25,27 @@ public void run() {
         returnLongArrayNew();
     }
 }
-/**{Throws{CodeGenException}Throws}**/
+/**{OpenCL{
+ typedef struct This_s{
+ int passid;
+ }This;
+ int get_pass_id(This *this){
+ return this->passid;
+ }
+ __global long* com_aparapi_codegen_test_ReturnLongArrayNew__returnLongArrayNew(This *this){
+ long returnArray5[1024];
+ return(returnArray5);
+ }
+ __kernel void run(
+ int passid
+ ){
+ This thisStruct;
+ This* this=&thisStruct;
+ this->passid = passid;
+ {
+ com_aparapi_codegen_test_ReturnLongArrayNew__returnLongArrayNew(this);
+ return;
+ }
+ }
+
+ }OpenCL}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnLongArrayNewTest.java b/src/test/java/com/aparapi/codegen/test/ReturnLongArrayNewTest.java
index 15b68c5d..d491c8f3 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnLongArrayNewTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnLongArrayNewTest.java
@@ -15,13 +15,13 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.CodeGenException;
 import org.junit.Test;
 
 public class ReturnLongArrayNewTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
-    private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = CodeGenException.class;
+    private static final String[] expectedOpenCL = ReturnArrayNewTestSupport.expectedOpenCL("long",
+            "ReturnLongArrayNew", "returnLongArrayNew");
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     @Test
     public void ReturnLongArrayNewTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnShortArrayNew.java b/src/test/java/com/aparapi/codegen/test/ReturnShortArrayNew.java
index 8d167073..7691c89e 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnShortArrayNew.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnShortArrayNew.java
@@ -25,4 +25,27 @@ public void run() {
         returnShortArrayNew();
     }
 }
-/**{Throws{CodeGenException}Throws}**/
+/**{OpenCL{
+ typedef struct This_s{
+ int passid;
+ }This;
+ int get_pass_id(This *this){
+ return this->passid;
+ }
+ __global short* com_aparapi_codegen_test_ReturnShortArrayNew__returnShortArrayNew(This *this){
+ short returnArray5[1024];
+ return(returnArray5);
+ }
+ __kernel void run(
+ int passid
+ ){
+ This thisStruct;
+ This* this=&thisStruct;
+ this->passid = passid;
+ {
+ com_aparapi_codegen_test_ReturnShortArrayNew__returnShortArrayNew(this);
+ return;
+ }
+ }
+
+ }OpenCL}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnShortArrayNewTest.java b/src/test/java/com/aparapi/codegen/test/ReturnShortArrayNewTest.java
index c5f26823..5a71b712 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnShortArrayNewTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnShortArrayNewTest.java
@@ -15,13 +15,13 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.CodeGenException;
 import org.junit.Test;
 
 public class ReturnShortArrayNewTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
-    private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = CodeGenException.class;
+    private static final String[] expectedOpenCL = ReturnArrayNewTestSupport.expectedOpenCL("short",
+            "ReturnShortArrayNew", "returnShortArrayNew");
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     @Test
     public void ReturnShortArrayNewTest() {
diff --git a/src/test/java/com/aparapi/runtime/ReturnInstantiatedArrayDirectlyTest.java b/src/test/java/com/aparapi/runtime/ReturnInstantiatedArrayDirectlyTest.java
index 5ae9c182..106842b2 100644
--- a/src/test/java/com/aparapi/runtime/ReturnInstantiatedArrayDirectlyTest.java
+++ b/src/test/java/com/aparapi/runtime/ReturnInstantiatedArrayDirectlyTest.java
@@ -16,7 +16,6 @@
 package com.aparapi.runtime;
 
 import com.aparapi.Kernel;
-import org.junit.Ignore;
 import org.junit.Test;
 
 import java.util.Arrays;
@@ -26,7 +25,6 @@
 
 public class ReturnInstantiatedArrayDirectlyTest {
     @Test
-    @Ignore("Knon bug, ignoring until fixed.")
     public void test() {
         ReturnDoubleArrayNew kernel = new ReturnDoubleArrayNew();
         kernel.execute(1);