Merge branch 'master' of git+ssh://scm.gforge.inria.fr/gitroot/starpu/starpu into ft_checkpoint

julia/examples/cholesky/cholesky_codelets.jl

@@ -31,8 +31,7 @@ chol_model22 = starpu_perfmodel(
 
 cl_11 = starpu_codelet(
     cpu_func = "u11",
-    # This kernel cannot be translated to CUDA yet.
-    # cuda_func = "u11",
+    cuda_func = "u11",
     modes = [STARPU_RW],
     color = 0xffff00,
     perfmodel = chol_model11

julia/src/StarPU.jl

@@ -26,8 +26,8 @@ const starpu_wrapper_library_name=fstarpu_task_library_name()
 
 include("translate_headers.jl")
 
-if !isfile((@__DIR__)*"/../gen/libstarpu_common.jl") || !isfile((@__DIR__)*"/../gen/libstarpu_api.jl") ||
-    mtime(@__FILE__) > mtime((@__DIR__)*"/../gen/libstarpu_common.jl")
+if !isfile(joinpath(fstarpu_build_dir(), "julia/gen/libstarpu_common.jl")) || !isfile(joinpath(fstarpu_build_dir(), "julia/gen/libstarpu_api.jl")) ||
+    mtime(joinpath(@__FILE__, "translate_headers.jl")) > mtime(joinpath(fstarpu_build_dir(), "julia/gen/libstarpu_api.jl"))
     starpu_translate_headers()
 end
 

julia/src/compiler/cuda.jl

@@ -258,6 +258,150 @@ function translate_cublas(expr :: StarpuExpr)
     return apply(func_to_run, expr)
 end
 
+function get_all_assignments(cpu_instr)
+    ret = StarpuExpr[]
+
+    function func_to_run(x :: StarpuExpr)
+        if isa(x, StarpuExprAffect)
+            push!(ret, x)
+        end
+
+        return x
+    end
+
+    apply(func_to_run, cpu_instr)
+    return ret
+end
+
+function get_all_buffer_vars(cpu_instr)
+    ret = StarpuExprTypedVar[]
+    assignments = get_all_assignments(cpu_instr)
+    for x in assignments
+        var = x.var
+        expr = x.expr
+        if isa(expr, StarpuExprCall) && expr.func in [:STARPU_MATRIX_GET_PTR, :STARPU_VECTOR_GET_PTR]
+            push!(ret, var)
+        end
+    end
+
+    return ret
+end
+
+function get_all_buffer_stores(cpu_instr, vars)
+    ret = StarpuExprAffect[]
+
+    function func_to_run(x :: StarpuExpr)
+        if isa(x, StarpuExprAffect) && isa(x.var, StarpuExprRef) && isa(x.var.ref, StarpuExprVar) &&
+            x.var.ref.name in map(x -> x.name, vars)
+            push!(ret, x)
+        end
+
+        return x
+    end
+
+    apply(func_to_run, cpu_instr)
+    return ret
+end
+
+function get_all_buffer_refs(cpu_instr, vars)
+    ret = []
+
+    current_instr = nothing
+    InstrTy = Union{StarpuExprAffect,
+                    StarpuExprCall,
+                    StarpuExprCudaCall,
+                    StarpuExprFor,
+                    StarpuExprIf,
+                    StarpuExprIfElse,
+                    StarpuExprReturn,
+                    StarpuExprBreak,
+                    StarpuExprWhile}
+    parent = nothing
+
+    function func_to_run(x :: StarpuExpr)
+        if isa(x, InstrTy) && !(isa(x, StarpuExprCall) && x.func in [:(+), :(-), :(*), :(/), :(%), :(<), :(<=), :(==), :(!=), :(>=), :(>), :sqrt])
+            current_instr = x
+        end
+
+        if isa(x, StarpuExprRef) && isa(x.ref, StarpuExprVar) && x.ref.name in map(x -> x.name, vars) && # var[...]
+            !isa(parent, StarpuExprAddress) && # filter &var[..]
+            !(isa(current_instr, StarpuExprAffect) && current_instr.var == x) # filter lhs ref
+            push!(ret, (current_instr, x))
+        end
+
+        parent = x
+        return x
+    end
+
+    visit_preorder(func_to_run, cpu_instr)
+    return ret
+end
+
+function transform_cuda_device_loadstore(cpu_instr :: StarpuExprBlock)
+    # Get all CUDA buffer pointers
+    buffer_vars = get_all_buffer_vars(cpu_instr)
+
+    buffer_types = Dict{Symbol, Type}()
+    for var in buffer_vars
+        buffer_types[var.name] = var.typ
+    end
+
+    # Get all store to a CUDA buffer
+    stores = get_all_buffer_stores(cpu_instr, buffer_vars)
+
+    # Get all load from CUDA buffer
+    loads = get_all_buffer_refs(cpu_instr, buffer_vars)
+
+    # Replace each load L:
+    # L: ... buffer[id]
+    # With the following instruction block:
+    # Type varX
+    # cudaMemcpy(&varX, &buffer[id], sizeof(Type), cudaMemcpyDeviceToHost)
+    # L: ... varX
+    for l in loads
+        (instr, ref) = l
+        block = []
+        buffer = ref.ref.name
+        varX = "var"*rand_string()
+        type = buffer_types[Symbol(buffer)]
+        ctype = starpu_type_traduction(eltype(type))
+        push!(block, StarpuExprTypedVar(Symbol(varX), eltype(type)))
+        push!(block, StarpuExprCall(:cudaMemcpy,
+                                    [StarpuExprAddress(StarpuExprVar(Symbol(varX))),
+                                     StarpuExprAddress(ref),
+                                     StarpuExprVar(Symbol("sizeof($ctype)")),
+                                     StarpuExprVar(:cudaMemcpyDeviceToHost)]))
+        push!(block, substitute(instr, ref, StarpuExprVar(Symbol("$varX"))))
+
+        cpu_instr = substitute(cpu_instr, instr, StarpuExprBlock(block))
+    end
+
+    # Replace each Store S:
+    # S: buffer[id] = expr
+    # With the following instruction block:
+    # Type varX
+    # varX = expr
+    # cudaMemcpy(&buffer[id], &varX, sizeof(Type), cudaMemcpyHostToDevice)
+    for s in stores
+        block = []
+        buffer = s.var.ref.name
+        varX = "var"*rand_string()
+        type = buffer_types[Symbol(buffer)]
+        ctype = starpu_type_traduction(eltype(type))
+        push!(block, StarpuExprTypedVar(Symbol(varX), eltype(type)))
+        push!(block, StarpuExprAffect(StarpuExprVar(Symbol("$varX")), s.expr))
+        push!(block, StarpuExprCall(:cudaMemcpy,
+                                    [StarpuExprAddress(s.var),
+                                     StarpuExprAddress(StarpuExprVar(Symbol(varX))),
+                                     StarpuExprVar(Symbol("sizeof($ctype)")),
+                                     StarpuExprVar(:cudaMemcpyHostToDevice)]))
+
+        cpu_instr = substitute(cpu_instr, s, StarpuExprBlock(block))
+    end
+
+    return cpu_instr
+end
+
 function transform_to_cuda_kernel(func :: StarpuExprFunction)
 
     cpu_func = transform_to_cpu_kernel(func)
@@ -300,9 +444,11 @@ function transform_to_cuda_kernel(func :: StarpuExprFunction)
     end
 
     cpu_instr = vcat(cpu_instr, finish)
+    cpu_instr = StarpuExprBlock(cpu_instr)
+    cpu_instr = transform_cuda_device_loadstore(cpu_instr)
 
     prekernel_name = Symbol("CUDA_", func.func)
-    prekernel = StarpuExprFunction(Nothing, prekernel_name, cpu_func.args, StarpuExprBlock(cpu_instr))
+    prekernel = StarpuExprFunction(Nothing, prekernel_name, cpu_func.args, cpu_instr)
     prekernel = translate_cublas(prekernel)
     prekernel = flatten_blocks(prekernel)
 

julia/src/compiler/expression_manipulation.jl

@@ -1,4 +1,28 @@
 
+"""
+    Lenient comparison operator for structures and arrays.
+"""
+@generated function ≂(x, y)
+    if x != y || x <: Type
+        :(x == y)
+    elseif !isempty(fieldnames(x))
+        mapreduce(n -> :(x.$n ≂ y.$n), (a,b)->:($a && $b), fieldnames(x))
+    elseif x <: Array
+        quote
+            if length(x) != length(y)
+                return false
+            end
+            for i in 1:length(x)
+                if !(x[i] ≂ y[i])
+                    return false
+                end
+            end
+            return true
+        end
+    else
+        :(x == y)
+    end
+end
 
 """
     Returns a new expression where every occurrence of expr_to_replace into expr
@@ -7,8 +31,7 @@
 function substitute(expr :: StarpuExpr, expr_to_replace :: StarpuExpr, new_expr :: StarpuExpr)
 
     function func_to_apply(x :: StarpuExpr)
-
-        if (x == expr_to_replace)
+        if (x ≂ expr_to_replace)
             return new_expr
         end
 
@@ -18,7 +41,6 @@ function substitute(expr :: StarpuExpr, expr_to_replace :: StarpuExpr, new_expr
     return apply(func_to_apply, expr)
 end
 
-
 """
     Returns an expression where "€" symbols  in expr were replaced
     by the following expression list.
@@ -110,3 +132,329 @@ import Base.all
 function all(cond :: Function, expr :: StarpuExpr)
     return !any(!cond, expr)
 end
+
+function visit_preorder(func :: Function, expr :: StarpuExprAffect)
+    func(expr)
+    visit_preorder(func, expr.var)
+    visit_preorder(func, expr.expr)
+    return expr
+end
+
+function visit_preorder(func :: Function, expr :: StarpuExprBlock)
+    func(expr)
+    for e in expr.exprs
+        visit_preorder(func, e)
+    end
+    return expr
+end
+
+function visit_preorder(func :: Function, expr :: StarpuExprCall)
+    func(expr)
+    for a in expr.args
+        visit_preorder(func, a)
+    end
+    return expr
+end
+
+function visit_preorder(func :: Function, expr :: StarpuExprCudaCall)
+    func(expr)
+    func(expr.nblocks)
+    func(expr.threads_per_block)
+    for a in expr.args
+        visit_preorder(func, a)
+    end
+    return expr
+end
+
+function visit_preorder(func :: Function, expr :: StarpuExprField)
+    func(expr)
+    func(expr.left)
+    func(expr.field)
+    func(expr.is_an_arrow)
+    return expr
+end
+
+function visit_preorder(func :: Function, expr :: StarpuExprFor)
+    func(expr)
+    for d in expr.set_declarations
+        visit_preorder(func, d)
+    end
+    visit_preorder(func, expr.set)
+    visit_preorder(func, expr.body)
+    return expr
+end
+
+function visit_preorder(func :: Function, expr :: StarpuExprFunction)
+    func(expr)
+    for a in expr.args
+        visit_preorder(func, a)
+    end
+    visit_preorder(func, e.body)
+    return expr
+end
+
+function visit_preorder(func :: Function, expr :: StarpuExprIf)
+    func(expr)
+    visit_preorder(func, expr.cond)
+    visit_preorder(func, expr.then_statement)
+    return expr
+end
+
+
+
+function visit_preorder(func :: Function, expr :: StarpuExprIfElse)
+    func(expr)
+    visit_preorder(func, expr.cond)
+    visit_preorder(func, expr.then_statement)
+    visit_preorder(func, expr.else_statement)
+    return expr
+end
+
+function visit_preorder(func :: Function, expr :: StarpuExprInterval)
+    func(expr)
+    visit_preorder(func, expr.start)
+    visit_preorder(func, expr.step)
+    visit_preorder(func, expr.stop)
+    return expr
+end
+
+function visit_preorder(func :: Function, expr :: StarpuExprRef)
+    func(expr)
+    visit_preorder(func, expr.ref)
+    for i in expr.indexes
+        visit_preorder(func, i)
+    end
+    return expr
+end
+
+function visit_preorder(func :: Function, expr :: StarpuExprAddress)
+    func(expr)
+    visit_preorder(func, expr.ref)
+    return expr
+end
+
+function visit_preorder(func :: Function, expr :: StarpuExprBreak)
+    func(expr)
+    return expr
+end
+
+function visit_preorder(func :: Function, expr :: StarpuExprReturn)
+    func(expr)
+    visit_preorder(func, expr.value)
+    return expr
+end
+
+function visit_preorder(func :: Function, expr :: StarpuExpr)
+    func(expr)
+    return expr
+end
+
+function visit_preorder(func :: Function, expr :: StarpuExprTypedExpr)
+    func(expr)
+    visit_preorder(func, expr.expr)
+    return expr
+end
+
+function visit_preorder(func :: Function, expr :: StarpuExprWhile)
+    func(expr)
+    visit_preorder(func, expr.cond)
+    visit_preorder(func, expr.body)
+    return expr
+end
+
+# function substitute_preorder(expr :: StarpuExprAffect, match :: StarpuExpr, replace :: StarpuExpr)
+#     if expr == match
+#         return replace
+#     end
+#     var = substitute_preorder(func, expr.var)
+#     expr = substitute_preorder(func, expr.expr)
+
+#     if var != expr.var || expr != expr.expr
+#         return StarpuExprAffect(var, expr)
+#     end
+#     return expr
+# end
+
+# function substitute_preorder(expr :: StarpuExprBlock, match :: StarpuExpr, replace :: StarpuExpr)
+#     if expr == match
+#         return replace
+#     end
+
+#     modified = false
+#     new_exprs = Vector{StarpuExpr}()
+#     for e in expr.exprs
+#         push!(new_exprs, substitute_preorder(func, e))
+#     end
+#     if new_exprs != expr.exprs
+#         return StarpuExprBlock(new_exprs)
+#     end
+#     return expr
+# end
+
+# function substitute_preorder(expr :: StarpuExprCall, match :: StarpuExpr, replace :: StarpuExpr)
+#     if expr == match
+#         return replace
+#     end
+
+#     new_args = Vector{StarpuExpr}()
+#     for a in expr.args
+#         push!(new_args, substitute_preorder(func, a))
+#     end
+#     if new_args != expr.args
+#         return StarpuExprCall(expr.func, new_args)
+#     end
+#     return expr
+# end
+
+# function substitute_preorder(expr :: StarpuExprCudaCall, match :: StarpuExpr, replace :: StarpuExpr)
+#     if expr == match
+#         return replace
+#     end
+
+#     new_args = Vector{StarpuExpr}()
+#     for a in expr.args
+#         push!(new_args, substitute_preorder(func, a))
+#     end
+#     if new_args != expr.args
+#         return new StarpuExprCudaCall(expr.ker_name, expr.nblocks, expr.threads_per_block, new_args)
+#     end
+#     return expr
+# end
+
+# function substitute_preorder(expr :: StarpuExprField, match :: StarpuExpr, replace :: StarpuExpr)
+#     if expr == match
+#         return replace
+#     end
+
+#     left = substitute_preorder(expr.left, match, replace)
+#     if left != expr.left
+#         return StarpuExprField(left, expr.field, expr.is_an_arrow)
+#     end
+#     return expr
+# end
+
+# function substitute_preorder(expr :: StarpuExprFor, match :: StarpuExpr, replace :: StarpuExpr)
+#     if expr == match
+#         return replace
+#     end
+
+#     new_set_declarations = Vector{StarpuExpr}()
+    
+#     for d in expr.set_declarations
+#         substitute_preorder(func, d)
+#     end
+#     substitute_preorder(expr.set, match :: StarpuExpr, replace :: StarpuExpr)
+#     substitute_preorder(func, expr.body)
+#     return expr
+# end
+
+# function substitute_preorder(expr :: StarpuExprFunction, match :: StarpuExpr, replace :: StarpuExpr)
+#         if expr == match
+#         return replace
+#     end
+
+#     for a in expr.args
+#         substitute_preorder(func, a)
+#     end
+#     substitute_preorder(e.body, match :: StarpuExpr, replace :: StarpuExpr)
+#     return expr
+# end
+
+# function substitute_preorder(expr :: StarpuExprIf, match :: StarpuExpr, replace :: StarpuExpr)
+#         if expr == match
+#         return replace
+#     end
+
+#     substitute_preorder(func, expr.cond)
+#     substitute_preorder(func, expr.then_statement)
+#     return expr
+# end
+
+
+
+# function substitute_preorder(expr :: StarpuExprIfElse, match :: StarpuExpr, replace :: StarpuExpr)
+#         if expr == match
+#         return replace
+#     end
+
+#     substitute_preorder(func, expr.cond)
+#     substitute_preorder(func, expr.then_statement)
+#     substitute_preorder(func, expr.else_statement)
+#     return expr
+# end
+
+# function substitute_preorder(expr :: StarpuExprInterval, match :: StarpuExpr, replace :: StarpuExpr)
+#         if expr == match
+#         return replace
+#     end
+
+#     substitute_preorder(func, expr.start)
+#     substitute_preorder(func, expr.step)
+#     substitute_preorder(func, expr.stop)
+#     return expr
+# end
+
+# function substitute_preorder(expr :: StarpuExprRef, match :: StarpuExpr, replace :: StarpuExpr)
+#         if expr == match
+#         return replace
+#     end
+
+#     substitute_preorder(func, expr.ref)
+#     for i in expr.indexes
+#         substitute_preorder(func, i)
+#     end
+#     return expr
+# end
+
+# function substitute_preorder(expr :: StarpuExprAddress, match :: StarpuExpr, replace :: StarpuExpr)
+#         if expr == match
+#         return replace
+#     end
+
+#     substitute_preorder(func, expr.ref)
+#     return expr
+# end
+
+# function substitute_preorder(expr :: StarpuExprBreak, match :: StarpuExpr, replace :: StarpuExpr)
+#         if expr == match
+#         return replace
+#     end
+
+#     return expr
+# end
+
+# function substitute_preorder(expr :: StarpuExprReturn, match :: StarpuExpr, replace :: StarpuExpr)
+#         if expr == match
+#         return replace
+#     end
+
+#     substitute_preorder(func, expr.value)
+#     return expr
+# end
+
+# function substitute_preorder(expr :: StarpuExpr, match :: StarpuExpr, replace :: StarpuExpr)
+#         if expr == match
+#         return replace
+#     end
+
+#     return expr
+# end
+
+# function substitute_preorder(expr :: StarpuExprTypedExpr, match :: StarpuExpr, replace :: StarpuExpr)
+#         if expr == match
+#         return replace
+#     end
+
+#     substitute_preorder(func, expr.expr)
+#     return expr
+# end
+
+# function substitute_preorder(expr :: StarpuExprWhile, match :: StarpuExpr, replace :: StarpuExpr)
+#         if expr == match
+#         return replace
+#     end
+
+#     substitute_preorder(func, expr.cond)
+#     substitute_preorder(func, expr.body)
+#     return expr
+# end