diff --git a/src/main/java/com/aparapi/internal/writer/BlockWriter.java b/src/main/java/com/aparapi/internal/writer/BlockWriter.java
index 013239d0..35c37da4 100644
--- a/src/main/java/com/aparapi/internal/writer/BlockWriter.java
+++ b/src/main/java/com/aparapi/internal/writer/BlockWriter.java
@@ -296,6 +296,7 @@ public void writeSequence(Instruction _first, Instruction _last) throws CodeGenE
             writeComposite((CompositeInstruction) instruction);
          } else if (!instruction.getByteCode().equals(ByteCode.NONE)) {
             newLine();
+            writeInlineLocalVariableDeclarations(instruction);
             writeInstruction(instruction);
             write(";");
 
@@ -304,6 +305,50 @@ public void writeSequence(Instruction _first, Instruction _last) throws CodeGenE
 
    }
 
+   private void writeInlineLocalVariableDeclarations(Instruction instruction) throws CodeGenException {
+      final Map<Integer, AssignToLocalVariable> declarations = new LinkedHashMap<Integer, AssignToLocalVariable>();
+      collectInlineLocalVariableDeclarations(instruction, declarations);
+
+      for (final AssignToLocalVariable declaration : declarations.values()) {
+         final LocalVariableInfo localVariableInfo = declaration.getLocalVariableInfo();
+         if (localVariableInfo == null) {
+            throw new CodeGenException("outOfScope" + ((Instruction) declaration).getThisPC() + " = ");
+         }
+
+         final String descriptor = localVariableInfo.getVariableDescriptor();
+         if (descriptor.startsWith("[")) {
+            write(" __global ");
+         }
+         write(convertType(descriptor, true, false));
+         write(localVariableInfo.getVariableName());
+         write(";");
+         newLine();
+      }
+   }
+
+   private void collectInlineLocalVariableDeclarations(Instruction instruction, Map<Integer, AssignToLocalVariable> declarations) {
+      if (instruction == null) {
+         return;
+      }
+
+      if (instruction instanceof InlineAssignInstruction) {
+         final InlineAssignInstruction inlineAssignInstruction = (InlineAssignInstruction) instruction;
+         final AssignToLocalVariable assignToLocalVariable = inlineAssignInstruction.getAssignToLocalVariable();
+         if (assignToLocalVariable.isDeclaration()) {
+            declarations.put(assignToLocalVariable.getLocalVariableTableIndex(), assignToLocalVariable);
+         }
+         collectInlineLocalVariableDeclarations(inlineAssignInstruction.getRhs(), declarations);
+      }
+
+      final Instruction lastChild = instruction.getLastChild();
+      for (Instruction child = instruction.getFirstChild(); child != null; child = child.getNextExpr()) {
+         collectInlineLocalVariableDeclarations(child, declarations);
+         if (child == lastChild) {
+            break;
+         }
+      }
+   }
+
    protected void writeGetterBlock(FieldEntry accessorVariableFieldEntry) {
       write("{");
       in();
@@ -698,11 +743,6 @@ public void writeInstruction(Instruction _instruction) throws CodeGenException {
          final AssignToLocalVariable assignToLocalVariable = inlineAssignInstruction.getAssignToLocalVariable();
 
          final LocalVariableInfo localVariableInfo = assignToLocalVariable.getLocalVariableInfo();
-         if (assignToLocalVariable.isDeclaration()) {
-            // this is bad! we need a general way to hoist up a required declaration
-            throw new CodeGenException("/* we can't declare this " + convertType(localVariableInfo.getVariableDescriptor(), true, false)
-                  + " here */");
-         }
          write(localVariableInfo.getVariableName());
          write("=");
          writeInstruction(inlineAssignInstruction.getRhs());
diff --git a/src/test/java/com/aparapi/codegen/test/AssignAndPassAsParameterSimpleTest.java b/src/test/java/com/aparapi/codegen/test/AssignAndPassAsParameterSimpleTest.java
index efff6cd5..cfc11c79 100644
--- a/src/test/java/com/aparapi/codegen/test/AssignAndPassAsParameterSimpleTest.java
+++ b/src/test/java/com/aparapi/codegen/test/AssignAndPassAsParameterSimpleTest.java
@@ -15,15 +15,35 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.ClassParseException;
-import com.aparapi.internal.exception.CodeGenException;
-import org.junit.Ignore;
 import org.junit.Test;
 
 public class AssignAndPassAsParameterSimpleTest 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 = {
+        "typedef struct This_s{\n"
+        + " int passid;\n"
+        + " }This;\n"
+        + " int get_pass_id(This *this){\n"
+        + " return this->passid;\n"
+        + " }\n"
+        + " void com_aparapi_codegen_test_AssignAndPassAsParameterSimple__actuallyDoIt(This *this, int a){\n"
+        + " return;\n"
+        + " }\n"
+        + " __kernel void run(\n"
+        + " int passid\n"
+        + " ){\n"
+        + " This thisStruct;\n"
+        + " This* this=&thisStruct;\n"
+        + " this->passid = passid;\n"
+        + " {\n"
+        + " int z;\n"
+        + " com_aparapi_codegen_test_AssignAndPassAsParameterSimple__actuallyDoIt(this, z=1);\n"
+        + " return;\n"
+        + " }\n"
+        + " }\n"
+        + "\n"
+        + " "};
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     @Test
     public void AssignAndPassAsParameterSimpleTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/FirstAssignInExpressionTest.java b/src/test/java/com/aparapi/codegen/test/FirstAssignInExpressionTest.java
index ecc5b660..27822efe 100644
--- a/src/test/java/com/aparapi/codegen/test/FirstAssignInExpressionTest.java
+++ b/src/test/java/com/aparapi/codegen/test/FirstAssignInExpressionTest.java
@@ -15,13 +15,35 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.CodeGenException;
-import org.junit.Ignore;
 import org.junit.Test;
 
 public class FirstAssignInExpressionTest 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 = {
+        "typedef struct This_s{\n"
+        + " int passid;\n"
+        + " }This;\n"
+        + " int get_pass_id(This *this){\n"
+        + " return this->passid;\n"
+        + " }\n"
+        + " void com_aparapi_codegen_test_FirstAssignInExpression__func(This *this, int _arg){\n"
+        + " return;\n"
+        + " }\n"
+        + " __kernel void run(\n"
+        + " int passid\n"
+        + " ){\n"
+        + " This thisStruct;\n"
+        + " This* this=&thisStruct;\n"
+        + " this->passid = passid;\n"
+        + " {\n"
+        + " int value = 1;\n"
+        + " int result;\n"
+        + " com_aparapi_codegen_test_FirstAssignInExpression__func(this, result=value);\n"
+        + " return;\n"
+        + " }\n"
+        + " }\n"
+        + "\n"
+        + " "};
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     @Test
     public void FirstAssignInExpressionTest() {