julia: Handle scalar parameters for CUDA kernels.

julia/StarPU.jl/src/StarPU.jl

@@ -496,13 +496,26 @@ mutable struct StarpuTask
 
         # handle scalar_parameters
         codelet_name = cl.cpu_func
-        scalar_parameters = CODELETS_SCALARS[codelet_name]
-        nb_scalar_required = length(scalar_parameters)
-        nb_scalar_provided = length(cl_arg)
-        if (nb_scalar_provided != nb_scalar_required)
-            error("$nb_scalar_provided scalar parameters provided but $nb_scalar_required are required by $codelet_name.")
+        if isempty(codelet_name)
+            codelet_name = cl.cuda_func
+        end
+        if isempty(codelet_name)
+            codelet_name = cl.opencl_func
+        end
+        if isempty(codelet_name)
+            error("No function provided with codelet.")
+        end
+        scalar_parameters = get(CODELETS_SCALARS, codelet_name, nothing)
+        if scalar_parameters != nothing
+            nb_scalar_required = length(scalar_parameters)
+            nb_scalar_provided = length(cl_arg)
+            if (nb_scalar_provided != nb_scalar_required)
+                error("$nb_scalar_provided scalar parameters provided but $nb_scalar_required are required by $codelet_name.")
+            end
+            output.cl_arg = create_param_struct_from_clarg(codelet_name, cl_arg)
+        else
+            output.cl_arg = nothing
         end
-        output.cl_arg = create_param_struct_from_clarg(codelet_name, cl_arg)
 
         output.synchronous = false
         output.handle_pointers = StarpuDataHandlePointer[]
@@ -522,8 +535,13 @@ mutable struct StarpuTask
 
 end
 
-function create_param_struct_from_clarg(codelet_name, cl_arg)
-    struct_params_name = CODELETS_PARAMS_STRUCT[codelet_name]
+function create_param_struct_from_clarg(name, cl_arg)
+    struct_params_name = CODELETS_PARAMS_STRUCT[name]
+
+    if struct_params_name == false
+        error("structure name not found in CODELET_PARAMS_STRUCT")
+    end
+
     nb_scalar_provided = length(cl_arg)
     create_struct_param_str = "output = $struct_params_name("
     for i in 1:nb_scalar_provided-1
@@ -856,7 +874,7 @@ macro starpu_async_cl(expr,modes,cl_arg=[])
     println(CPU_CODELETS[string(expr.args[1])])
     cl = StarpuCodelet(
         cpu_func = CPU_CODELETS[string(expr.args[1])],
-        #cuda_func = "matrix_mult",
+        # cuda_func = CUDA_CODELETS[string(expr.args[1])],
         #opencl_func="ocl_matrix_mult",
         ### TODO: CORRECT !
         modes = map((x -> starpu_modes(x)),modes.args),
@@ -865,7 +883,7 @@ macro starpu_async_cl(expr,modes,cl_arg=[])
     handles = Expr(:vect, expr.args[2:end]...)
     #dump(handles)
     quote
-        task = StarpuTask(cl = $(esc(cl)), handles = $(esc(handles)), cl_arg=cl_arg)
+        task = StarpuTask(cl = $(esc(cl)), handles = $(esc(handles)), cl_arg=$(esc(cl_arg)))
         starpu_task_submit(task)
     end
 end

julia/StarPU.jl/src/compiler/file_generation.jl

@@ -106,8 +106,11 @@ global CODELETS_PARAMS_STRUCT=Dict{String,Any}()
 macro codelet(x)
     parsed = starpu_parse(x)
     name=string(x.args[1].args[1].args[1]);
+    cpu_name = name
+    cuda_name = "CUDA_"*name
     dump(name)
-    parse_scalar_parameters(parsed, name)
+    parse_scalar_parameters(parsed, cpu_name, cuda_name)
+    c_struct_param_decl = generate_c_struct_param_declaration(name)
     cpu_expr = transform_to_cpu_kernel(parsed)
     prekernel, kernel = transform_to_cuda_kernel(parsed)
     generated_cpu_kernel_file_name=string("genc_",string(x.args[1].args[1].args[1]),".c")
@@ -119,10 +122,10 @@ macro codelet(x)
             kernel_file = open($(esc(generated_cpu_kernel_file_name)), "w")
             @debugprint "generating " $(generated_cpu_kernel_file_name)
             print(kernel_file, $(esc(cpu_kernel_file_start)))
-            print(kernel_file, generate_c_struct_param_declaration($name))
+            print(kernel_file, $c_struct_param_decl)
             print(kernel_file, $cpu_expr)
             close(kernel_file)
-            CPU_CODELETS[$name]=$name
+            CPU_CODELETS[$name]=$cpu_name
         end
         
         if ($targets&$STARPU_CUDA!=0)
@@ -130,9 +133,10 @@ macro codelet(x)
             @debugprint "generating " $(generated_cuda_kernel_file_name)
             print(kernel_file, $(esc(cuda_kernel_file_start)))
             print(kernel_file, "__global__ ", $kernel)
+            print(kernel_file, $c_struct_param_decl) # TODO: extern C ?
             print(kernel_file, "\nextern \"C\" ", $prekernel)
             close(kernel_file)
-            CUDA_CODELETS[$name]="CUDA_"*$name
+            CUDA_CODELETS[$name]=$cuda_name
         end
         print("end generation")
         #starpu_task_library_name="generated_tasks"
@@ -140,7 +144,7 @@ macro codelet(x)
     end
 end
 
-function parse_scalar_parameters(expr :: StarpuExprFunction, name::String)
+function parse_scalar_parameters(expr :: StarpuExprFunction, cpu_name::String, cuda_name::String)
     scalar_parameters = []
     for i in (1 : length(expr.args))
         type = expr.args[i].typ
@@ -149,7 +153,8 @@ function parse_scalar_parameters(expr :: StarpuExprFunction, name::String)
         end
     end
 
-    CODELETS_SCALARS[name] = scalar_parameters
+    CODELETS_SCALARS[cpu_name] = scalar_parameters
+    CODELETS_SCALARS[cuda_name] = scalar_parameters
 
     # declare structure carrying scalar parameters
     struct_params_name = Symbol("params_", rand_string())
@@ -165,5 +170,6 @@ function parse_scalar_parameters(expr :: StarpuExprFunction, name::String)
     eval(Meta.parse(add_to_dict_str))
 
     # save structure name
-    CODELETS_PARAMS_STRUCT[name] = struct_params_name
+    CODELETS_PARAMS_STRUCT[cpu_name] = struct_params_name
+    CODELETS_PARAMS_STRUCT[cuda_name] = struct_params_name
 end