test2 :: (Input (Exp a), 
         Input (Exp b),
         Sources' (GArr (Exp a)),
         Sources' (GArr (Exp b))) => 
        (GArr (Exp a,Exp b) -> GPU (GArr (Exp c,Exp d))) -> [(Exp a,Exp b)] -> String
test2 inp inputs = 
    let 
        -- generate C style array
        len = length inputs
        (i1,i2) = unzip inputs      
        cInputs1 = strListToCArrayStr $ map render i1
        cInputs2 = strListToCArrayStr $ map render i2
        -- names of the result arrays, used to read back result 
        -- the results are also alive when leaving the program
        --nms = names $ getResult inp
        [nm1,nm2] = names $ getResult inp len
        res = Set.fromList [nm1,nm2]   
        -- perform liveness analysis. 
        icl = liveness (getIC inp len) res 

        -- evaluate the symbolic table
        env = [("n0",len),("n1",len)]
        st  = symtab $ getState inp len
        concreteT  = st -- evalSymTab env st 
        ((ict,memMap),mstate) = 
            runState (genMemoryMap icl Set.empty) 
                         (MState sharedMem globalMem concreteT Map.empty)

        nthreads = maximum [snd x | x <- ict] -- number of threads needed 
        ic = mm_IC ict memMap 
    
        --ic' = (map fst ic) -- threadsIC ic nthreads
        ic' = threadsIC ic nthreads
        (kernel,mstate') = runState (genCudaKernel "generated" env ic') mstate
        inputType1  = elemType $ typeof "source0" concreteT
        inputType2  = elemType $ typeof "source1" concreteT
        outputType1 = elemType $ typeof nm1 st
        outputType2 = elemType $ typeof nm2 st
        olen       = threads nm1 concreteT
                     
        sharedMemoryNeeded = size (shared mstate') 
        globalMemoryNeeded = size (global mstate')

        result_pos1 = case Map.lookup nm1 memMap of 
                       Nothing -> error "error"
                       Just a -> "(gbase + " ++ show (fst a) ++ ")"
        result_pos2 = case Map.lookup nm2 memMap of 
                       Nothing -> error "error"
                       Just a -> "(gbase + " ++ show (fst a) ++ ")"


        
        --generate code that launches kernel     
    in "/*\n" ++ "number of threads:" ++ show nthreads ++ "\n" ++ 
       cInputs1 ++ "\n" ++
       cInputs2 ++ "\n" ++
       "*/\n\n\n" ++ includes  ++ kernel ++  
       cmain (
             "  "++ typeStr inputType1 ++ " values1[" ++ show len ++ "] = " ++ cInputs1 ++ ";\n" ++
             "  "++ typeStr inputType2 ++ " values2[" ++ show len ++ "] = " ++ cInputs2 ++ ";\n" ++
             "  "++ typeStr outputType1 ++ " result1[" ++ show olen ++ "];\n" ++
             "  "++ typeStr outputType2 ++ " result2[" ++ show olen ++ "];\n" ++
             "  char *gbase;\n" ++ 
             "  " ++typeStr inputType1 ++" * dvalues1;\n" ++
             "  " ++typeStr inputType2 ++" * dvalues2;\n" ++
             "  cudaMalloc((void**)&dvalues1, sizeof("++typeStr inputType1++") * "++ show len ++" ); \n" ++
             "  cudaMemcpy(dvalues1, values1, sizeof("++typeStr inputType1++") * "++ show len ++", cudaMemcpyHostToDevice);\n" ++
             "  cudaMalloc((void**)&dvalues2, sizeof("++typeStr inputType2++") * "++ show len ++" ); \n" ++
             "  cudaMemcpy(dvalues2, values2, sizeof("++typeStr inputType2++") * "++ show len ++", cudaMemcpyHostToDevice);\n" ++
             "  cudaMalloc((void**)&gbase," ++ show globalMemoryNeeded ++ "); \n" ++
             "  " ++ runKernelp "generated" nthreads sharedMemoryNeeded ++ 
             "  cudaMemcpy(result1," ++ result_pos1 ++", sizeof("++typeStr outputType1++") * "++ show olen ++" , cudaMemcpyDeviceToHost);\n" ++
             "  cudaMemcpy(result2," ++ result_pos2 ++", sizeof("++typeStr outputType2++") * "++ show olen ++" , cudaMemcpyDeviceToHost);\n" ++
             "  cudaFree(dvalues1);\n" ++
             "  cudaFree(dvalues2);\n" ++
             "  cudaFree(gbase);\n" ++
             "  for(int i = 0; i < " ++ show olen ++ "; i++){\n" ++ 
             "    printf(\"" ++ pfSpecStr outputType1 ++ " \",result1[i]);\n" ++
             "  }\n" ++
             "  printf(\"\\nSEPARATOR\\n\");\n" ++ 
             "  for(int i = 0; i < " ++ show olen ++ "; i++){\n" ++ 
             "    printf(\"" ++ pfSpecStr outputType1 ++ " \",result2[i]);\n" ++
             "  }\n" 

            )
test2 :: (Input (Exp a), 
         Input (Exp b),
         Sources' (GArr (Exp a)),
         Sources' (GArr (Exp b))) => 
        (GArr (Exp a,Exp b) -> GPU (GArr (Exp c,Exp d))) -> [(Exp a,Exp b)] -> String
test2 inp inputs = 
    let 
        -- generate C style array
        len = length inputs
        (i1,i2) = unzip inputs      
        cInputs1 = strListToCArrayStr $ map render i1
        cInputs2 = strListToCArrayStr $ map render i2
        -- names of the result arrays, used to read back result 
        -- the results are also alive when leaving the program
        --nms = names $ getResult inp
        [nm1,nm2] = names $ getResult inp len
        res = Set.fromList [nm1,nm2]   
        -- perform liveness analysis. 
        icl = liveness (getIC inp len) res 

        -- evaluate the symbolic table
        env = [("n0",len),("n1",len)]
        st  = symtab $ getState inp len
        concreteT  = st -- evalSymTab env st 
        ((ict,memMap),mstate) = 
            runState (genMemoryMap icl Set.empty) 
                         (MState sharedMem globalMem concreteT Map.empty)

        nthreads = maximum [snd x | x <- ict] -- number of threads needed 
        ic = mm_IC ict memMap 
    
        --ic' = (map fst ic) -- threadsIC ic nthreads
        ic' = threadsIC ic nthreads
        (kernel,mstate') = runState (genCudaKernel "generated" env ic') mstate
        inputType1  = elemType $ typeof "source0" concreteT
        inputType2  = elemType $ typeof "source1" concreteT
        outputType1 = elemType $ typeof nm1 st
        outputType2 = elemType $ typeof nm2 st
        olen       = threads nm1 concreteT
                     
        sharedMemoryNeeded = size (shared mstate') 
        globalMemoryNeeded = size (global mstate')

        result_pos1 = case Map.lookup nm1 memMap of 
                       Nothing -> error "error"
                       Just a -> "(gbase + " ++ show (fst a) ++ ")"
        result_pos2 = case Map.lookup nm2 memMap of 
                       Nothing -> error "error"
                       Just a -> "(gbase + " ++ show (fst a) ++ ")"


        
        --generate code that launches kernel     
    in "/*\n" ++ "number of threads:" ++ show nthreads ++ "\n" ++ 
       cInputs1 ++ "\n" ++
       cInputs2 ++ "\n" ++
       "*/\n\n\n" ++ includes  ++ kernel ++  
       cmain (
             "  "++ typeStr inputType1 ++ " values1[" ++ show len ++ "] = " ++ cInputs1 ++ ";\n" ++
             "  "++ typeStr inputType2 ++ " values2[" ++ show len ++ "] = " ++ cInputs2 ++ ";\n" ++
             "  "++ typeStr outputType1 ++ " result1[" ++ show olen ++ "];\n" ++
             "  "++ typeStr outputType2 ++ " result2[" ++ show olen ++ "];\n" ++
             "  char *gbase;\n" ++ 
             "  " ++typeStr inputType1 ++" * dvalues1;\n" ++
             "  " ++typeStr inputType2 ++" * dvalues2;\n" ++
             "  cudaMalloc((void**)&dvalues1, sizeof("++typeStr inputType1++") * "++ show len ++" ); \n" ++
             "  cudaMemcpy(dvalues1, values1, sizeof("++typeStr inputType1++") * "++ show len ++", cudaMemcpyHostToDevice);\n" ++
             "  cudaMalloc((void**)&dvalues2, sizeof("++typeStr inputType2++") * "++ show len ++" ); \n" ++
             "  cudaMemcpy(dvalues2, values2, sizeof("++typeStr inputType2++") * "++ show len ++", cudaMemcpyHostToDevice);\n" ++
             "  cudaMalloc((void**)&gbase," ++ show globalMemoryNeeded ++ "); \n" ++
             "  " ++ runKernelp "generated" nthreads sharedMemoryNeeded ++ 
             "  cudaMemcpy(result1," ++ result_pos1 ++", sizeof("++typeStr outputType1++") * "++ show olen ++" , cudaMemcpyDeviceToHost);\n" ++
             "  cudaMemcpy(result2," ++ result_pos2 ++", sizeof("++typeStr outputType2++") * "++ show olen ++" , cudaMemcpyDeviceToHost);\n" ++
             "  cudaFree(dvalues1);\n" ++
             "  cudaFree(dvalues2);\n" ++
             "  cudaFree(gbase);\n" ++
             "  for(int i = 0; i < " ++ show olen ++ "; i++){\n" ++ 
             "    printf(\"" ++ pfSpecStr outputType1 ++ " \",result1[i]);\n" ++
             "  }\n" ++
             "  printf(\"\\nSEPARATOR\\n\");\n" ++ 
             "  for(int i = 0; i < " ++ show olen ++ "; i++){\n" ++ 
             "    printf(\"" ++ pfSpecStr outputType1 ++ " \",result2[i]);\n" ++
             "  }\n" 

            )
test2 inp inputs = 
    let 
        -- generate C style array
        len = length inputs
        (i1,i2) = unzip inputs      
        cInputs1 = strListToCArrayStr $ map render i1
        cInputs2 = strListToCArrayStr $ map render i2
        -- names of the result arrays, used to read back result 
        -- the results are also alive when leaving the program
        --nms = names $ getResult inp
        [nm1,nm2] = names $ getResult inp len
        res = Set.fromList [nm1,nm2]   
        -- perform liveness analysis. 
        icl = liveness (getIC inp len) res 

        -- evaluate the symbolic table
        env = [("n0",len),("n1",len)]
        st  = symtab $ getState inp len
        concreteT  = st -- evalSymTab env st 
        ((ict,memMap),mstate) = 
            runState (genMemoryMap icl Set.empty) 
                         (MState sharedMem globalMem concreteT Map.empty)

        nthreads = maximum [snd x | x <- ict] -- number of threads needed 
        ic = mm_IC ict memMap 
    
        --ic' = (map fst ic) -- threadsIC ic nthreads
        ic' = threadsIC ic nthreads
        (kernel,mstate') = runState (genCudaKernel "generated" env ic') mstate
        inputType1  = elemType $ typeof "source0" concreteT
        inputType2  = elemType $ typeof "source1" concreteT
        outputType1 = elemType $ typeof nm1 st
        outputType2 = elemType $ typeof nm2 st
        olen       = threads nm1 concreteT
                     
        sharedMemoryNeeded = size (shared mstate') 
        globalMemoryNeeded = size (global mstate')

        result_pos1 = case Map.lookup nm1 memMap of 
                       Nothing -> error "error"
                       Just a -> "(gbase + " ++ show (fst a) ++ ")"
        result_pos2 = case Map.lookup nm2 memMap of 
                       Nothing -> error "error"
                       Just a -> "(gbase + " ++ show (fst a) ++ ")"


        
        --generate code that launches kernel     
    in "/*\n" ++ "number of threads:" ++ show nthreads ++ "\n" ++ 
       cInputs1 ++ "\n" ++
       cInputs2 ++ "\n" ++
       "*/\n\n\n" ++ includes  ++ kernel ++  
       cmain (
             "  "++ typeStr inputType1 ++ " values1[" ++ show len ++ "] = " ++ cInputs1 ++ ";\n" ++
             "  "++ typeStr inputType2 ++ " values2[" ++ show len ++ "] = " ++ cInputs2 ++ ";\n" ++
             "  "++ typeStr outputType1 ++ " result1[" ++ show olen ++ "];\n" ++
             "  "++ typeStr outputType2 ++ " result2[" ++ show olen ++ "];\n" ++
             "  char *gbase;\n" ++ 
             "  " ++typeStr inputType1 ++" * dvalues1;\n" ++
             "  " ++typeStr inputType2 ++" * dvalues2;\n" ++
             "  cudaMalloc((void**)&dvalues1, sizeof("++typeStr inputType1++") * "++ show len ++" ); \n" ++
             "  cudaMemcpy(dvalues1, values1, sizeof("++typeStr inputType1++") * "++ show len ++", cudaMemcpyHostToDevice);\n" ++
             "  cudaMalloc((void**)&dvalues2, sizeof("++typeStr inputType2++") * "++ show len ++" ); \n" ++
             "  cudaMemcpy(dvalues2, values2, sizeof("++typeStr inputType2++") * "++ show len ++", cudaMemcpyHostToDevice);\n" ++
             "  cudaMalloc((void**)&gbase," ++ show globalMemoryNeeded ++ "); \n" ++
             "  " ++ runKernelp "generated" nthreads sharedMemoryNeeded ++ 
             "  cudaMemcpy(result1," ++ result_pos1 ++", sizeof("++typeStr outputType1++") * "++ show olen ++" , cudaMemcpyDeviceToHost);\n" ++
             "  cudaMemcpy(result2," ++ result_pos2 ++", sizeof("++typeStr outputType2++") * "++ show olen ++" , cudaMemcpyDeviceToHost);\n" ++
             "  cudaFree(dvalues1);\n" ++
             "  cudaFree(dvalues2);\n" ++
             "  cudaFree(gbase);\n" ++
             "  for(int i = 0; i < " ++ show olen ++ "; i++){\n" ++ 
             "    printf(\"" ++ pfSpecStr outputType1 ++ " \",result1[i]);\n" ++
             "  }\n" ++
             "  printf(\"\\nSEPARATOR\\n\");\n" ++ 
             "  for(int i = 0; i < " ++ show olen ++ "; i++){\n" ++ 
             "    printf(\"" ++ pfSpecStr outputType1 ++ " \",result2[i]);\n" ++
             "  }\n" 

            )