ADD: add cudaGetParameterBufferV2 and add cudaLaunchDeviceV2 implementation. Kernel launch to stream not yet implemented

7 files changed, 219 insertions, 6 deletions

diff --git a/src/cuda-sim/Makefile b/src/cuda-sim/Makefile
index 166e256..f479294 100644
--- a/src/cuda-sim/Makefile
+++ b/src/cuda-sim/Makefile
@@ -62,7 +62,7 @@ ifeq ($(GNUC_CPP0X),1)
 endif
 endif
 
-OBJS	:= $(OUTPUT_DIR)/ptx_parser.o $(OUTPUT_DIR)/ptx_loader.o $(OUTPUT_DIR)/cuda_device_printf.o $(OUTPUT_DIR)/instructions.o $(OUTPUT_DIR)/cuda-sim.o $(OUTPUT_DIR)/ptx_ir.o $(OUTPUT_DIR)/ptx_sim.o  $(OUTPUT_DIR)/memory.o $(OUTPUT_DIR)/ptx-stats.o $(OUTPUT_DIR)/decuda_pred_table/decuda_pred_table.o $(OUTPUT_DIR)/ptx.tab.o $(OUTPUT_DIR)/lex.ptx_.o $(OUTPUT_DIR)/ptxinfo.tab.o $(OUTPUT_DIR)/lex.ptxinfo_.o
+OBJS	:= $(OUTPUT_DIR)/ptx_parser.o $(OUTPUT_DIR)/ptx_loader.o $(OUTPUT_DIR)/cuda_device_printf.o $(OUTPUT_DIR)/instructions.o $(OUTPUT_DIR)/cuda-sim.o $(OUTPUT_DIR)/ptx_ir.o $(OUTPUT_DIR)/ptx_sim.o  $(OUTPUT_DIR)/memory.o $(OUTPUT_DIR)/ptx-stats.o $(OUTPUT_DIR)/decuda_pred_table/decuda_pred_table.o $(OUTPUT_DIR)/ptx.tab.o $(OUTPUT_DIR)/lex.ptx_.o $(OUTPUT_DIR)/ptxinfo.tab.o $(OUTPUT_DIR)/lex.ptxinfo_.o $(OUTPUT_DIR)/cuda_device_runtime.o
 
 
 OPT += -DCUDART_VERSION=$(CUDART_VERSION)
@@ -145,5 +145,6 @@ $(OUTPUT_DIR)/ptx_sim.o: $(OUTPUT_DIR)/ptx.tab.c
 $(OUTPUT_DIR)/cuda-sim.o: $(OUTPUT_DIR)/ptx.tab.c
 $(OUTPUT_DIR)/lex.ptxinfo_.o: $(OUTPUT_DIR)/ptx.tab.c
 $(OUTPUT_DIR)/lex.ptx_.o: $(OUTPUT_DIR)/ptx.tab.c
+$(OUTPUT_DIR)/cuda_device_runtime.o: $(OUTPUT_DIR)/ptx.tab.c
 
 include $(OUTPUT_DIR)/Makefile.makedepend
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 715be98..4933029 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1065,6 +1065,16 @@ void function_info::add_param_data( unsigned argn, struct gpgpu_ptx_sim_arg *arg
    } 
 }
 
+unsigned function_info::get_args_aligned_size() {
+   unsigned int align_size = 4; // a word
+   unsigned int total_size = 0;
+   for(unsigned int i = 0; i < num_args(); i++) {
+       total_size += ((m_args[i]->get_size_in_bytes() + align_size - 1) / align_size) * align_size;
+   }
+   return total_size;
+}
+
+
 void function_info::finalize( memory_space *param_mem ) 
 {
    unsigned param_address = 0;
diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc
new file mode 100644
index 0000000..937eec8
--- /dev/null
+++ b/src/cuda-sim/cuda_device_runtime.cc
@@ -0,0 +1,175 @@
+//Jin: cuda_device_runtime.cc
+//Defines CUDA device runtime APIs for CDP support
+
+#include <iostream>
+#include <map>
+
+#define __CUDA_RUNTIME_API_H__
+
+#include <builtin_types.h>
+#include <driver_types.h>
+#include "../gpgpu-sim/gpu-sim.h"
+#include "cuda-sim.h"
+#include "ptx_ir.h"
+#include "../stream_manager.h"
+#include "cuda_device_runtime.h"
+
+#define DEV_RUNTIME_REPORT(a) \
+   if( g_debug_execution ) { \
+      std::cout << __FILE__ << ", " << __LINE__ << ": " << a << "\n"; \
+      std::cout.flush(); \
+   }
+
+std::map<void *, kernel_info_t *> g_cuda_device_launch_map;
+struct CUstream_st * g_device_default_stream = NULL;
+extern stream_manager *g_stream_manager;
+
+//Handling device runtime api:
+//void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize)
+void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func)
+{
+    DEV_RUNTIME_REPORT("Calling cudaGetParameterBufferV2");
+      
+    unsigned n_return = target_func->has_return();
+    assert(n_return);
+    unsigned n_args = target_func->num_args();
+    assert( n_args == 4 );
+
+    function_info * child_kernel_entry;
+    struct dim3 gridDim, blockDim;
+    unsigned int sharedMem;
+
+    for( unsigned arg=0; arg < n_args; arg ++ ) {
+        const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param#
+        const symbol *formal_param = target_func->get_arg(arg); //cudaGetParameterBufferV2_param_#
+        unsigned size=formal_param->get_size_in_bytes();
+        assert( formal_param->is_param_local() );
+        assert( actual_param_op.is_param_local() );
+        addr_t from_addr = actual_param_op.get_symbol()->get_address();
+
+        if(arg == 0) {//function_info* for the child kernel
+            unsigned long long buf;
+			assert(size == sizeof(function_info *));
+            thread->m_local_mem->read(from_addr, size, &buf);
+            child_kernel_entry = (function_info *)buf;
+            assert(child_kernel_entry);
+            DEV_RUNTIME_REPORT("child kernel name " << child_kernel_entry->get_name());
+        }
+        else if(arg == 1) { //dim3 gridDim for the child kernel
+			assert(size == sizeof(struct dim3));
+            thread->m_local_mem->read(from_addr, size, & gridDim);
+            DEV_RUNTIME_REPORT("grid (" << gridDim.x << ", " << gridDim.y << ", " << gridDim.z << ")");
+        }
+        else if(arg == 2) { //dim3 blockDim for the child kernel
+			assert(size == sizeof(struct dim3));
+            thread->m_local_mem->read(from_addr, size, & blockDim);
+            DEV_RUNTIME_REPORT("block (" << blockDim.x << ", " << blockDim.y << ", " << blockDim.z << ")");
+        }
+        else if(arg == 3) { //unsigned int sharedMem
+			assert(size == sizeof(unsigned int));
+            thread->m_local_mem->read(from_addr, size, & sharedMem);
+            DEV_RUNTIME_REPORT("shared memory " << sharedMem);
+        }
+    }
+
+	//get total child kernel argument size and malloc buffer in global memory
+	unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size();
+	void * param_buffer = thread->get_gpu()->gpu_malloc(child_kernel_arg_size);
+	DEV_RUNTIME_REPORT("child kernel arg size total " << child_kernel_arg_size << ", parameter buffer allocated at " << param_buffer);
+	
+	//create child kernel_info_t and index it with parameter_buffer address
+	kernel_info_t * child_grid = new kernel_info_t(gridDim, blockDim, child_kernel_entry); 
+    assert(g_cuda_device_launch_map.find(param_buffer) == g_cuda_device_launch_map.end());
+    g_cuda_device_launch_map[param_buffer] = child_grid;
+
+	//copy the buffer address to retval0
+    const operand_info &actual_return_op = pI->operand_lookup(0); //retval0
+    const symbol *formal_return = target_func->get_return_var(); //void *
+	unsigned int return_size = formal_return->get_size_in_bytes();
+	DEV_RUNTIME_REPORT("cudaGetParameterBufferV2 return value has size of " << return_size);
+	assert(actual_return_op.is_param_local());
+	assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size && return_size == sizeof(void *));
+    addr_t ret_param_addr = actual_return_op.get_symbol()->get_address();
+	thread->m_local_mem->write(ret_param_addr, return_size, &param_buffer, NULL, NULL);
+
+}
+
+//Handling device runtime api:
+//cudaError_t cudaLaunchDeviceV2(void *parameterBuffer, cudaStream_t stream)
+void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func) {
+    DEV_RUNTIME_REPORT("Calling cudaLaunchDeviceV2");
+
+    unsigned n_return = target_func->has_return();
+    assert(n_return);
+    unsigned n_args = target_func->num_args();
+    assert( n_args == 2 );
+
+    kernel_info_t * child_grid = NULL;
+    function_info * child_kernel_entry = NULL;
+    void * parameter_buffer;
+    struct CUstream_st * child_stream;
+    for( unsigned arg=0; arg < n_args; arg ++ ) {
+        const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param#
+        const symbol *formal_param = target_func->get_arg(arg); //cudaLaunchDeviceV2_param_#
+        unsigned size=formal_param->get_size_in_bytes();
+        assert( formal_param->is_param_local() );
+        assert( actual_param_op.is_param_local() );
+        addr_t from_addr = actual_param_op.get_symbol()->get_address();
+
+        if(arg == 0) {//paramter buffer for child kernel (in global memory)
+            //get parameter_buffer from the cudaDeviceLaunchV2_param0
+			assert(size == sizeof(void *));
+            thread->m_local_mem->read(from_addr, size, &parameter_buffer);
+            assert((size_t)parameter_buffer >= GLOBAL_HEAP_START);
+            DEV_RUNTIME_REPORT("Parameter buffer locating at global memory " << parameter_buffer);
+
+            //get child grid info through parameter_buffer address
+            assert(g_cuda_device_launch_map.find(parameter_buffer) != g_cuda_device_launch_map.end());
+            child_grid = g_cuda_device_launch_map[parameter_buffer];
+            child_kernel_entry = child_grid->entry();
+            DEV_RUNTIME_REPORT("find child kernel " << child_kernel_entry->get_name());
+
+            //copy data in parameter_buffer to child kernel param memory
+	        unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size();
+            DEV_RUNTIME_REPORT("child_kernel_arg_size " << child_kernel_arg_size);
+            memory_space *child_kernel_param_mem = child_grid->get_param_memory();
+            size_t param_start_address = 0;
+            for(unsigned n = 0; n < child_kernel_arg_size; n++) {
+                unsigned char one_byte;
+                thread->get_gpu()->get_global_memory()->read((size_t)parameter_buffer + n, 1, &one_byte);
+                child_kernel_param_mem->write(param_start_address + n, 1, &one_byte, NULL, NULL); 
+            }
+        }
+        else if(arg == 1) { //cudaStream for the child kernel
+			assert(size == sizeof(cudaStream_t));
+            thread->m_local_mem->read(from_addr, size, &child_stream);
+            if(child_stream == 0) { //default stream on device
+                if(!g_device_default_stream) {
+                    //g_device_default_stream = new struct CUstream_st();
+                    //g_stream_manager->add_stream(g_device_default_stream);
+                }
+                child_stream = g_device_default_stream;
+            }
+//            DEV_RUNTIME_REPORT("launching child kernel to stream " << child_stream->get_uid());
+        }
+        
+    }
+
+    //launch child kernel
+    stream_operation op(child_grid, g_ptx_sim_mode, child_stream);
+//    g_stream_manager->push(op);
+//    g_cuda_device_launch_map.erase(parameter_buffer);
+
+    //set retval0
+    const operand_info &actual_return_op = pI->operand_lookup(0); //retval0
+    const symbol *formal_return = target_func->get_return_var(); //cudaError_t
+	unsigned int return_size = formal_return->get_size_in_bytes();
+	DEV_RUNTIME_REPORT("cudaLaunchDeviceV2 return value has size of " << return_size);
+	assert(actual_return_op.is_param_local());
+	assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size 
+        && return_size == sizeof(cudaError_t));
+    cudaError_t error = cudaSuccess;
+    addr_t ret_param_addr = actual_return_op.get_symbol()->get_address();
+	thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL);
+
+}
diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h
new file mode 100644
index 0000000..1b10407
--- /dev/null
+++ b/src/cuda-sim/cuda_device_runtime.h
@@ -0,0 +1,7 @@
+//Jin: cuda_device_runtime.h
+//Defines CUDA device runtime APIs for CDP support
+
+#pragma once
+
+void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func);
+void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func);
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 254427b..1c47ad3 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -41,6 +41,9 @@
 #include "../gpgpu-sim/gpu-sim.h"
 #include "../gpgpu-sim/shader.h"
 
+//Jin: include device runtime for CDP
+#include "cuda_device_runtime.h"
+
 #include <stdarg.h>
 
 unsigned ptx_instruction::g_num_ptx_inst_uid=0;
@@ -149,8 +152,7 @@ ptx_reg_t ptx_thread_info::get_operand_value( const operand_info &op, operand_in
          } else if ( op.is_local() ) {
             result.u64 = op.get_symbol()->get_address();
          } else if ( op.is_function_address() ) {
-		 	result.u64 = op.get_symbol()->get_pc()->get_start_PC();
-		 	printf("Get pc for kernel function %u\n", op.get_symbol()->get_pc()->get_start_PC());
+		 	result.u64 = (size_t)op.get_symbol()->get_pc();
          } else {
             const char *name = op.name().c_str();
             printf("GPGPU-Sim PTX: ERROR ** get_operand_value : unknown operand type for %s\n", name );
@@ -1411,12 +1413,14 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread )
       gpgpusim_cuda_vprintf(pI, thread, target_func);
       return;
    }
+
+   //Jin: handle device runtime apis for CDP
    else if(fname == "cudaGetParameterBufferV2") {
-      printf("calling cudaGetParameterBufferV2\n");
+      gpgpusim_cuda_getParameterBufferV2(pI, thread, target_func);
 	  return;
    }
    else if(fname == "cudaLaunchDeviceV2") {
-      printf("calling cudaLaunchDeviceV2\n");
+      gpgpusim_cuda_launchDeviceV2(pI, thread, target_func);
 	  return;
    }
 
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index 7325e5f..a7ca27e 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -699,7 +699,7 @@ public:
    }
 
    bool is_immediate_address() const {
-	   return   m_immediate_address;
+       return   m_immediate_address;
    }
 
    bool is_literal() const { return m_type == int_t ||
@@ -1209,6 +1209,8 @@ public:
    {
       return m_args.size();
    }
+   unsigned get_args_aligned_size();
+
    const symbol* get_arg( unsigned n ) const
    {
       assert( n < m_args.size() );
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 39257da..baa3bcd 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -120,6 +120,20 @@ symbol_table *init_parser( const char *ptx_filename )
 #define DEF(X,Y) g_ptx_token_decode[X] = Y;
 #include "ptx_parser_decode.def"
 #undef DEF
+   g_ptx_token_decode[undefined_space] = "undefined_space";
+   g_ptx_token_decode[undefined_space] = "undefined_space=0";
+   g_ptx_token_decode[reg_space] = "reg_space";
+   g_ptx_token_decode[local_space] = "local_space";
+   g_ptx_token_decode[shared_space] = "shared_space";
+   g_ptx_token_decode[param_space_unclassified] = "param_space_unclassified";
+   g_ptx_token_decode[param_space_kernel] = "param_space_kernel";
+   g_ptx_token_decode[param_space_local] = "param_space_local";
+   g_ptx_token_decode[const_space] = "const_space";
+   g_ptx_token_decode[tex_space] = "tex_space";
+   g_ptx_token_decode[surf_space] = "surf_space";
+   g_ptx_token_decode[global_space] = "global_space";
+   g_ptx_token_decode[generic_space] = "generic_space";
+   g_ptx_token_decode[instruction_space] = "instruction_space";
 
    return g_global_symbol_table;
 }