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

        -- evaluate the symbolic table
        env = [("n0",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
        inputType  = elemType $ typeof "source0" concreteT
        outputType = elemType $ typeof (head nms) st
        olen       = threads (head nms) concreteT
                     
        sharedMemoryNeeded = size (shared mstate') 
        globalMemoryNeeded = size (global mstate')

        result_pos = case Map.lookup (head nms) memMap of 
                       Nothing -> error "error"
                       Just a -> "(gbase + " ++ show (fst a) ++ ")"
        
        --generate code that launches kernel     
    in "/*\n" ++ "number of threads:" ++ show nthreads ++ "\n" ++ cInputs ++ "\n" ++
       "*/\n\n\n" ++ includes  ++ kernel ++  
       cmain (
             "  "++ typeStr inputType ++ " values[" ++ show len ++ "] = " ++ cInputs ++ ";\n" ++
             "  "++ typeStr outputType ++" result[" ++ show olen ++ "];\n" ++
             "  char *gbase;\n" ++ 
             "  " ++ typeStr inputType ++" * dvalues;\n" ++
             "  cudaMalloc((void**)&dvalues, sizeof("++typeStr inputType++") * "++ show len ++" ); \n" ++
             "  cudaMemcpy(dvalues, values, sizeof("++typeStr inputType++") * "++show len ++", cudaMemcpyHostToDevice);\n" ++
             "  cudaMalloc((void**)&gbase," ++ show globalMemoryNeeded ++ "); \n" ++
             "  " ++ runKernel "generated" nthreads sharedMemoryNeeded ++ 
             "  cudaMemcpy(result," ++ result_pos ++", sizeof("++typeStr outputType++") * "++ show olen ++" , cudaMemcpyDeviceToHost);\n" ++
             "  cudaFree(dvalues);\n" ++
             "  cudaFree(gbase);\n" ++
             "  for(int i = 0; i < " ++ show olen ++ "; i++){\n" ++ 
             "    printf(\"" ++ pfSpecStr outputType ++ " \",result[i]);\n" ++
             "  }\n" 
            )
test :: (Input (Exp a), Sources' (GArr (Exp a))) => 
        (GArr (Exp a) -> GPU (GArr (Exp b))) -> [Exp a] -> String
test inp inputs = 
    let 
        -- generate C style array
        len = length inputs
        cInputs = strListToCArrayStr $ map render inputs
        
        -- names of the result arrays, used to read back result 
        -- the results are also alive when leaving the program
        nms = names $ getResult inp len
        res = Set.fromList nms   
        -- perform liveness analysis. 
        icl = liveness (getIC inp len) res 

        -- evaluate the symbolic table
        env = [("n0",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
        inputType  = elemType $ typeof "source0" concreteT
        outputType = elemType $ typeof (head nms) st
        olen       = threads (head nms) concreteT
                     
        sharedMemoryNeeded = size (shared mstate') 
        globalMemoryNeeded = size (global mstate')

        result_pos = case Map.lookup (head nms) memMap of 
                       Nothing -> error "error"
                       Just a -> "(gbase + " ++ show (fst a) ++ ")"
        
        --generate code that launches kernel     
    in "/*\n" ++ "number of threads:" ++ show nthreads ++ "\n" ++ cInputs ++ "\n" ++
       "*/\n\n\n" ++ includes  ++ kernel ++  
       cmain (
             "  "++ typeStr inputType ++ " values[" ++ show len ++ "] = " ++ cInputs ++ ";\n" ++
             "  "++ typeStr outputType ++" result[" ++ show olen ++ "];\n" ++
             "  char *gbase;\n" ++ 
             "  " ++ typeStr inputType ++" * dvalues;\n" ++
             "  cudaMalloc((void**)&dvalues, sizeof("++typeStr inputType++") * "++ show len ++" ); \n" ++
             "  cudaMemcpy(dvalues, values, sizeof("++typeStr inputType++") * "++show len ++", cudaMemcpyHostToDevice);\n" ++
             "  cudaMalloc((void**)&gbase," ++ show globalMemoryNeeded ++ "); \n" ++
             "  " ++ runKernel "generated" nthreads sharedMemoryNeeded ++ 
             "  cudaMemcpy(result," ++ result_pos ++", sizeof("++typeStr outputType++") * "++ show olen ++" , cudaMemcpyDeviceToHost);\n" ++
             "  cudaFree(dvalues);\n" ++
             "  cudaFree(gbase);\n" ++
             "  for(int i = 0; i < " ++ show olen ++ "; i++){\n" ++ 
             "    printf(\"" ++ pfSpecStr outputType ++ " \",result[i]);\n" ++
             "  }\n" 
            )
test inp inputs = 
    let 
        -- generate C style array
        len = length inputs
        cInputs = strListToCArrayStr $ map render inputs
        
        -- names of the result arrays, used to read back result 
        -- the results are also alive when leaving the program
        nms = names $ getResult inp len
        res = Set.fromList nms   
        -- perform liveness analysis. 
        icl = liveness (getIC inp len) res 

        -- evaluate the symbolic table
        env = [("n0",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
        inputType  = elemType $ typeof "source0" concreteT
        outputType = elemType $ typeof (head nms) st
        olen       = threads (head nms) concreteT
                     
        sharedMemoryNeeded = size (shared mstate') 
        globalMemoryNeeded = size (global mstate')

        result_pos = case Map.lookup (head nms) memMap of 
                       Nothing -> error "error"
                       Just a -> "(gbase + " ++ show (fst a) ++ ")"
        
        --generate code that launches kernel     
    in "/*\n" ++ "number of threads:" ++ show nthreads ++ "\n" ++ cInputs ++ "\n" ++
       "*/\n\n\n" ++ includes  ++ kernel ++  
       cmain (
             "  "++ typeStr inputType ++ " values[" ++ show len ++ "] = " ++ cInputs ++ ";\n" ++
             "  "++ typeStr outputType ++" result[" ++ show olen ++ "];\n" ++
             "  char *gbase;\n" ++ 
             "  " ++ typeStr inputType ++" * dvalues;\n" ++
             "  cudaMalloc((void**)&dvalues, sizeof("++typeStr inputType++") * "++ show len ++" ); \n" ++
             "  cudaMemcpy(dvalues, values, sizeof("++typeStr inputType++") * "++show len ++", cudaMemcpyHostToDevice);\n" ++
             "  cudaMalloc((void**)&gbase," ++ show globalMemoryNeeded ++ "); \n" ++
             "  " ++ runKernel "generated" nthreads sharedMemoryNeeded ++ 
             "  cudaMemcpy(result," ++ result_pos ++", sizeof("++typeStr outputType++") * "++ show olen ++" , cudaMemcpyDeviceToHost);\n" ++
             "  cudaFree(dvalues);\n" ++
             "  cudaFree(gbase);\n" ++
             "  for(int i = 0; i < " ++ show olen ++ "; i++){\n" ++ 
             "    printf(\"" ++ pfSpecStr outputType ++ " \",result[i]);\n" ++
             "  }\n" 
            )