parse kernel global variable

src/gputask/gputask.py

@@ -20,6 +20,7 @@ class GPUTaskParams:
     # kernel can have different parameters types
     # use of variadic template here to be able to handle all of them
     FUNCTION            = "\tvoid (*f)(Types...);\n";
+    # @TODO add here FUNCTION ARGUMENTS
 
    ### IMPL SECTION
     # RT params
@@ -36,6 +37,7 @@ class GPUTaskParams:
     # GENERAL params
     TASK_ID_IMPL        = "\t.task_id    = {task_id},\n"
     FUNCTION_IMPL       = "\t.{name}=&{kernel_name},\n"
+    # @TODO add here FUNCTION ARGUMENTS
     NAME_IMPL           = "\t.kernel_name = std::string(\"{name}\")\n"
 
     #Parse Index
@@ -113,7 +115,8 @@ class GPUTask:
         #specific parameters to retrieve the cuda kernel declarations
         self.kfile      = kfile
 
-        
+        self.kernel_parameters_variable = []
+        self.init_function_name = ""
 
     def dump_str(self) -> str:
         dump = "struct {struct_name} task_{no} {{\n".format(struct_name=GPUTaskParams.STRUCT_NAME, no=self.task_id)
@@ -137,12 +140,15 @@ class GPUTask:
         return dump
 
     def get_kernel_declaration(self) -> str:
+        self.get_variables_declaration()
+        print(self.kernel_parameters_variable)
+        print(self.init_function_name)
         with open(self.kfile, 'r') as inFile:
 
             inFileLines = inFile.readlines()
 
             for line in inFileLines:
-                if self.name in line:
+                if "__global__" in line and self.name in line:
                     kernel_arguments = line.split(self.name)[1].strip()
                     if(kernel_arguments[-1] == ')'):
                         return kernel_arguments
@@ -154,3 +160,34 @@ class GPUTask:
                 if line == '-1\n':
                     break
             return kernel_arguments
+    def get_variables_declaration(self):
+        parameters = []
+        kernel_detected = False
+        parameters_detected = False
+        init_name = ""
+        with open(self.kfile, 'r') as inFile:
+            inFileLines = inFile.readlines()
+            for idx,line in enumerate(inFileLines):
+                if self.name in line and "Start" in line:
+                    kernel_detected = True
+                if self.name in line and "End" in line:
+                    flag_kernel_detected = False
+                    break
+
+                if kernel_detected and "Start Parameters" in line:
+                    parameters_detected = True
+                elif "End Parameters" in line:
+                    parameters_detected = False
+                elif parameters_detected and not line.isspace():
+                    parameters.append(line)
+
+                if kernel_detected and "Start Init function" in line:
+                    init_name = inFileLines[idx+1]
+                    break
+        for parameter in parameters:
+            parameter = (parameter.split(" ")[-1]).rstrip(';\n')
+            self.kernel_parameters_variable.append(parameter)
+        self.init_function_name = init_name.split("void")[1].lstrip().rstrip(" {\n")
+
+
+

src/main.py

@@ -31,13 +31,13 @@ def generate_main_unroll_step():
         for job in range(0, currStep.number_kernels):
             relatedTask = int(relation_job_id_task_id[currStep.list_jobs[job]],10)
             currTask    = list_gpu_tasks_object[relatedTask]
-            unroll +="\t{kernel_name}<<<{block},{thread},{smem},{stream}>>>({params});\n".format(kernel_name=currTask.name, block = currTask.blocks, thread = currTask.threads, smem = currTask.smem,stream = "stream_{no}".format(no=job),params="test")
+            unroll +="\t{kernel_name}<<<{block},{thread},{smem},{stream}>>>({params});\n".format(kernel_name=currTask.name, block = currTask.blocks, thread = currTask.threads, smem = currTask.smem,stream = "stream_{no}".format(no=job),params=",".join(currTask.kernel_parameters_variable))
         unroll += "\tcudaDeviceSynchronize();\n\n"
         if i != len(list_gpusched_step_object) - 1:
             unroll += "\tstep_end = high_resolution_clock::now();\n"
             unroll += "\telapsed_time = step_end - base;\n"
             unroll += "\telapsed_time_ms = duration_cast<ms>(elasped_time);\n"
-            unroll += "\tidle_time = milliseconds({next_step_min_time}) - elapsed_time_ms);\n".format(next_step_min_time=list_gpusched_step_object[i+1].min_time_start)
+            unroll += "\tidle_time = milliseconds({next_step_min_time}) - elapsed_time_ms;\n".format(next_step_min_time=list_gpusched_step_object[i+1].min_time_start)
             unroll += "\tif (idle_time.count() > 0)\n\t\tstd::this_thread::sleep_for(idle_time);\n"