diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/BlockWriter.java b/com.amd.aparapi/src/java/com/amd/aparapi/BlockWriter.java
index 7a9a63d3d2314f9e80c51478779ecc0654c53e7c..6ca9d5b1bdc9c235c9fd3081763559db00b1b759 100644
--- a/com.amd.aparapi/src/java/com/amd/aparapi/BlockWriter.java
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/BlockWriter.java
@@ -75,6 +75,8 @@ import com.amd.aparapi.InstructionSet.I_IINC;
 import com.amd.aparapi.InstructionSet.I_POP;
 import com.amd.aparapi.InstructionSet.If;
 import com.amd.aparapi.InstructionSet.IfUnary;
+import com.amd.aparapi.InstructionSet.I_IFNULL;
+import com.amd.aparapi.InstructionSet.I_IFNONNULL;
 import com.amd.aparapi.InstructionSet.IncrementInstruction;
 import com.amd.aparapi.InstructionSet.InlineAssignInstruction;
 import com.amd.aparapi.InstructionSet.MethodCall;
@@ -127,7 +129,25 @@ abstract class BlockWriter{
          writeInstruction(iff.getLhs());
          write(_branch16.getOperator().getText(_invert));
          writeInstruction(iff.getRhs());
+      } else if (_branch16 instanceof I_IFNULL) {
+         I_IFNULL iff = (I_IFNULL) _branch16;
+         writeInstruction(iff.getFirstChild());
 
+         if (_invert) {
+            write(" != NULL");
+         } else {
+            write(" == NULL");
+         }
+
+      } else if (_branch16 instanceof I_IFNONNULL) {
+         I_IFNONNULL iff = (I_IFNONNULL) _branch16;
+         writeInstruction(iff.getFirstChild());
+
+         if (_invert) {
+            write(" == NULL");
+         } else {
+            write(" != NULL");
+         }
       } else if (_branch16 instanceof IfUnary) {
          IfUnary branch16 = (IfUnary) _branch16;
          Instruction comparison = branch16.getUnary();
@@ -154,7 +174,6 @@ abstract class BlockWriter{
          }
 
       }
-
    }
 
    protected void writeComposite(CompositeInstruction instruction) throws CodeGenException {
diff --git a/test/codegen/src/java/com/amd/aparapi/test/NonNullCheck.java b/test/codegen/src/java/com/amd/aparapi/test/NonNullCheck.java
new file mode 100644
index 0000000000000000000000000000000000000000..5c6580fef0f6262c85a52e8821add537aa8f5015
--- /dev/null
+++ b/test/codegen/src/java/com/amd/aparapi/test/NonNullCheck.java
@@ -0,0 +1,37 @@
+package com.amd.aparapi.test;
+
+public class NonNullCheck{
+   int[] ints = new int[1024];
+
+   public void run() {
+      if (ints != null){
+         int value = ints[0];
+      }
+    
+   }
+}
+/**{OpenCL{
+typedef struct This_s{
+   __global int *ints;
+    int passid;
+}This;
+int get_pass_id(This *this){
+   return this->passid;
+}
+
+__kernel void run(
+   __global int *ints,
+   int passid
+){
+   This thisStruct;
+   This* this=&thisStruct;
+   this->ints = ints;
+   this->passid = passid;
+   {
+      if (this->ints != NULL){
+         int value = this->ints[0];
+      }
+      return;
+   }
+}
+}OpenCL}**/
diff --git a/test/codegen/src/java/com/amd/aparapi/test/NullCheck.java b/test/codegen/src/java/com/amd/aparapi/test/NullCheck.java
new file mode 100644
index 0000000000000000000000000000000000000000..319815a1fbc1f3399491b4ef1c965e6becd80253
--- /dev/null
+++ b/test/codegen/src/java/com/amd/aparapi/test/NullCheck.java
@@ -0,0 +1,38 @@
+package com.amd.aparapi.test;
+
+public class NullCheck{
+   int[] ints = new int[1024];
+
+   public void run() {
+      if (ints == null){
+         return;
+      }
+      int value = ints[0];
+   }
+}
+/**{OpenCL{
+typedef struct This_s{
+   __global int *ints;
+    int passid;
+}This;
+int get_pass_id(This *this){
+   return this->passid;
+}
+
+__kernel void run(
+   __global int *ints,
+   int passid
+){
+   This thisStruct;
+   This* this=&thisStruct;
+   this->ints = ints;
+   this->passid = passid;
+   {
+      if (this->ints == NULL){
+         return;
+      }
+      int value = this->ints[0];
+      return;
+   }
+}
+}OpenCL}**/