diff options
Diffstat (limited to 'libcuda/cuda_runtime_api.cc')
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 2102 |
1 files changed, 1192 insertions, 910 deletions
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 7a651bb..c0d7be3 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -109,11 +109,11 @@ #include <stdarg.h> #ifdef OPENGL_SUPPORT #define GL_GLEXT_PROTOTYPES - #ifdef __APPLE__ - #include <GLUT/glut.h> // Apple's version of GLUT is here - #else - #include <GL/gl.h> - #endif +#ifdef __APPLE__ +#include <GLUT/glut.h> // Apple's version of GLUT is here +#else +#include <GL/gl.h> +#endif #endif #define __CUDA_RUNTIME_API_H__ @@ -140,27 +140,27 @@ static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsign static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu ); static kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, - gpgpu_ptx_sim_arg_list_t args, - struct dim3 gridDim, - struct dim3 blockDim, - struct CUctx_st* context ); + gpgpu_ptx_sim_arg_list_t args, + struct dim3 gridDim, + struct dim3 blockDim, + struct CUctx_st* context ); /*DEVICE_BUILTIN*/ struct cudaArray { - void *devPtr; - int devPtr32; - struct cudaChannelFormatDesc desc; - int width; - int height; - int size; //in bytes - unsigned dimensions; + void *devPtr; + int devPtr32; + struct cudaChannelFormatDesc desc; + int width; + int height; + int size; //in bytes + unsigned dimensions; }; #if !defined(__dv) #if defined(__cplusplus) #define __dv(v) \ - = v + = v #else /* __cplusplus */ #define __dv(v) #endif /* __cplusplus */ @@ -172,7 +172,7 @@ extern stream_manager *g_stream_manager; void register_ptx_function( const char *name, function_info *impl ) { - // no longer need this + // no longer need this } #if defined __APPLE__ @@ -190,174 +190,174 @@ void register_ptx_function( const char *name, function_info *impl ) #endif struct _cuda_device_id { - _cuda_device_id(gpgpu_sim* gpu) {m_id = 0; m_next = NULL; m_gpgpu=gpu;} - struct _cuda_device_id *next() { return m_next; } - unsigned num_shader() const { return m_gpgpu->get_config().num_shader(); } - int num_devices() const { - if( m_next == NULL ) return 1; - else return 1 + m_next->num_devices(); - } - struct _cuda_device_id *get_device( unsigned n ) - { - assert( n < (unsigned)num_devices() ); - struct _cuda_device_id *p=this; - for(unsigned i=0; i<n; i++) - p = p->m_next; - return p; - } - const struct cudaDeviceProp *get_prop() const - { - return m_gpgpu->get_prop(); - } - unsigned get_id() const { return m_id; } + _cuda_device_id(gpgpu_sim* gpu) {m_id = 0; m_next = NULL; m_gpgpu=gpu;} + struct _cuda_device_id *next() { return m_next; } + unsigned num_shader() const { return m_gpgpu->get_config().num_shader(); } + int num_devices() const { + if( m_next == NULL ) return 1; + else return 1 + m_next->num_devices(); + } + struct _cuda_device_id *get_device( unsigned n ) + { + assert( n < (unsigned)num_devices() ); + struct _cuda_device_id *p=this; + for(unsigned i=0; i<n; i++) + p = p->m_next; + return p; + } + const struct cudaDeviceProp *get_prop() const + { + return m_gpgpu->get_prop(); + } + unsigned get_id() const { return m_id; } - gpgpu_sim *get_gpgpu() { return m_gpgpu; } + gpgpu_sim *get_gpgpu() { return m_gpgpu; } private: - unsigned m_id; - class gpgpu_sim *m_gpgpu; - struct _cuda_device_id *m_next; + unsigned m_id; + class gpgpu_sim *m_gpgpu; + struct _cuda_device_id *m_next; }; struct CUctx_st { - CUctx_st( _cuda_device_id *gpu ) { m_gpu = gpu; } + CUctx_st( _cuda_device_id *gpu ) { m_gpu = gpu; } - _cuda_device_id *get_device() { return m_gpu; } + _cuda_device_id *get_device() { return m_gpu; } - void add_binary( symbol_table *symtab, unsigned fat_cubin_handle ) - { - m_code[fat_cubin_handle] = symtab; - m_last_fat_cubin_handle = fat_cubin_handle; - } + void add_binary( symbol_table *symtab, unsigned fat_cubin_handle ) + { + m_code[fat_cubin_handle] = symtab; + m_last_fat_cubin_handle = fat_cubin_handle; + } - void add_ptxinfo( const char *deviceFun, const struct gpgpu_ptx_sim_kernel_info &info ) - { - symbol *s = m_code[m_last_fat_cubin_handle]->lookup(deviceFun); - assert( s != NULL ); - function_info *f = s->get_pc(); - assert( f != NULL ); - f->set_kernel_info(info); - } + void add_ptxinfo( const char *deviceFun, const struct gpgpu_ptx_sim_kernel_info &info ) + { + symbol *s = m_code[m_last_fat_cubin_handle]->lookup(deviceFun); + assert( s != NULL ); + function_info *f = s->get_pc(); + assert( f != NULL ); + f->set_kernel_info(info); + } - void register_function( unsigned fat_cubin_handle, const char *hostFun, const char *deviceFun ) - { - if( m_code.find(fat_cubin_handle) != m_code.end() ) { - symbol *s = m_code[fat_cubin_handle]->lookup(deviceFun); - assert( s != NULL ); - function_info *f = s->get_pc(); - assert( f != NULL ); - m_kernel_lookup[hostFun] = f; - } else { - m_kernel_lookup[hostFun] = NULL; - } - } + void register_function( unsigned fat_cubin_handle, const char *hostFun, const char *deviceFun ) + { + if( m_code.find(fat_cubin_handle) != m_code.end() ) { + symbol *s = m_code[fat_cubin_handle]->lookup(deviceFun); + assert( s != NULL ); + function_info *f = s->get_pc(); + assert( f != NULL ); + m_kernel_lookup[hostFun] = f; + } else { + m_kernel_lookup[hostFun] = NULL; + } + } - function_info *get_kernel(const char *hostFun) - { - std::map<const void*,function_info*>::iterator i=m_kernel_lookup.find(hostFun); - assert( i != m_kernel_lookup.end() ); - return i->second; - } + function_info *get_kernel(const char *hostFun) + { + std::map<const void*,function_info*>::iterator i=m_kernel_lookup.find(hostFun); + assert( i != m_kernel_lookup.end() ); + return i->second; + } private: - _cuda_device_id *m_gpu; // selected gpu - std::map<unsigned,symbol_table*> m_code; // fat binary handle => global symbol table - unsigned m_last_fat_cubin_handle; - std::map<const void*,function_info*> m_kernel_lookup; // unique id (CUDA app function address) => kernel entry point + _cuda_device_id *m_gpu; // selected gpu + std::map<unsigned,symbol_table*> m_code; // fat binary handle => global symbol table + unsigned m_last_fat_cubin_handle; + std::map<const void*,function_info*> m_kernel_lookup; // unique id (CUDA app function address) => kernel entry point }; class kernel_config { public: - kernel_config( dim3 GridDim, dim3 BlockDim, size_t sharedMem, struct CUstream_st *stream ) - { - m_GridDim=GridDim; - m_BlockDim=BlockDim; - m_sharedMem=sharedMem; - m_stream = stream; - } - void set_arg( const void *arg, size_t size, size_t offset ) - { - m_args.push_front( gpgpu_ptx_sim_arg(arg,size,offset) ); - } - dim3 grid_dim() const { return m_GridDim; } - dim3 block_dim() const { return m_BlockDim; } - gpgpu_ptx_sim_arg_list_t get_args() { return m_args; } - struct CUstream_st *get_stream() { return m_stream; } - + kernel_config( dim3 GridDim, dim3 BlockDim, size_t sharedMem, struct CUstream_st *stream ) + { + m_GridDim=GridDim; + m_BlockDim=BlockDim; + m_sharedMem=sharedMem; + m_stream = stream; + } + void set_arg( const void *arg, size_t size, size_t offset ) + { + m_args.push_front( gpgpu_ptx_sim_arg(arg,size,offset) ); + } + dim3 grid_dim() const { return m_GridDim; } + dim3 block_dim() const { return m_BlockDim; } + gpgpu_ptx_sim_arg_list_t get_args() { return m_args; } + struct CUstream_st *get_stream() { return m_stream; } + private: - dim3 m_GridDim; - dim3 m_BlockDim; - size_t m_sharedMem; - struct CUstream_st *m_stream; - gpgpu_ptx_sim_arg_list_t m_args; + dim3 m_GridDim; + dim3 m_BlockDim; + size_t m_sharedMem; + struct CUstream_st *m_stream; + gpgpu_ptx_sim_arg_list_t m_args; }; class _cuda_device_id *GPGPUSim_Init() { - static _cuda_device_id *the_device = NULL; - if( !the_device ) { - gpgpu_sim *the_gpu = gpgpu_ptx_sim_init_perf(); + static _cuda_device_id *the_device = NULL; + if( !the_device ) { + gpgpu_sim *the_gpu = gpgpu_ptx_sim_init_perf(); - cudaDeviceProp *prop = (cudaDeviceProp *) calloc(sizeof(cudaDeviceProp),1); - snprintf(prop->name,256,"GPGPU-Sim_v%s", g_gpgpusim_version_string ); - prop->major = 2; - prop->minor = 0; - prop->totalGlobalMem = 0x40000000 /* 1 GB */; - prop->memPitch = 0; - prop->maxThreadsPerBlock = 512; - prop->maxThreadsDim[0] = 512; - prop->maxThreadsDim[1] = 512; - prop->maxThreadsDim[2] = 512; - prop->maxGridSize[0] = 0x40000000; - prop->maxGridSize[1] = 0x40000000; - prop->maxGridSize[2] = 0x40000000; - prop->totalConstMem = 0x40000000; - prop->textureAlignment = 0; - prop->sharedMemPerBlock = the_gpu->shared_mem_size(); - prop->regsPerBlock = the_gpu->num_registers_per_core(); - prop->warpSize = the_gpu->wrp_size(); - prop->clockRate = the_gpu->shader_clock(); + cudaDeviceProp *prop = (cudaDeviceProp *) calloc(sizeof(cudaDeviceProp),1); + snprintf(prop->name,256,"GPGPU-Sim_v%s", g_gpgpusim_version_string ); + prop->major = 2; + prop->minor = 0; + prop->totalGlobalMem = 0x40000000 /* 1 GB */; + prop->memPitch = 0; + prop->maxThreadsPerBlock = 512; + prop->maxThreadsDim[0] = 512; + prop->maxThreadsDim[1] = 512; + prop->maxThreadsDim[2] = 512; + prop->maxGridSize[0] = 0x40000000; + prop->maxGridSize[1] = 0x40000000; + prop->maxGridSize[2] = 0x40000000; + prop->totalConstMem = 0x40000000; + prop->textureAlignment = 0; + prop->sharedMemPerBlock = the_gpu->shared_mem_size(); + prop->regsPerBlock = the_gpu->num_registers_per_core(); + prop->warpSize = the_gpu->wrp_size(); + prop->clockRate = the_gpu->shader_clock(); #if (CUDART_VERSION >= 2010) - prop->multiProcessorCount = the_gpu->get_config().num_shader(); + prop->multiProcessorCount = the_gpu->get_config().num_shader(); #endif - the_gpu->set_prop(prop); - the_device = new _cuda_device_id(the_gpu); - } - start_sim_thread(1); - return the_device; + the_gpu->set_prop(prop); + the_device = new _cuda_device_id(the_gpu); + } + start_sim_thread(1); + return the_device; } static CUctx_st* GPGPUSim_Context() { - static CUctx_st *the_context = NULL; - if( the_context == NULL ) { - _cuda_device_id *the_gpu = GPGPUSim_Init(); - the_context = new CUctx_st(the_gpu); - } - return the_context; + static CUctx_st *the_context = NULL; + if( the_context == NULL ) { + _cuda_device_id *the_gpu = GPGPUSim_Init(); + the_context = new CUctx_st(the_gpu); + } + return the_context; } extern "C" void ptxinfo_addinfo() { - if( !strcmp("__cuda_dummy_entry__",get_ptxinfo_kname()) ) { - // this string produced by ptxas for empty ptx files (e.g., bandwidth test) - clear_ptxinfo(); - return; - } - CUctx_st *context = GPGPUSim_Context(); - print_ptxinfo(); - context->add_ptxinfo( get_ptxinfo_kname(), get_ptxinfo_kinfo() ); - clear_ptxinfo(); + if( !strcmp("__cuda_dummy_entry__",get_ptxinfo_kname()) ) { + // this string produced by ptxas for empty ptx files (e.g., bandwidth test) + clear_ptxinfo(); + return; + } + CUctx_st *context = GPGPUSim_Context(); + print_ptxinfo(); + context->add_ptxinfo( get_ptxinfo_kname(), get_ptxinfo_kinfo() ); + clear_ptxinfo(); } void cuda_not_implemented( const char* func, unsigned line ) { - fflush(stdout); - fflush(stderr); - printf("\n\nGPGPU-Sim PTX: Execution error: CUDA API function \"%s()\" has not been implemented yet.\n" - " [$GPGPUSIM_ROOT/libcuda/%s around line %u]\n\n\n", - func,__FILE__, line ); - fflush(stdout); - abort(); + fflush(stdout); + fflush(stderr); + printf("\n\nGPGPU-Sim PTX: Execution error: CUDA API function \"%s()\" has not been implemented yet.\n" + " [$GPGPUSIM_ROOT/libcuda/%s around line %u]\n\n\n", + func,__FILE__, line ); + fflush(stdout); + abort(); } @@ -366,27 +366,27 @@ void cuda_not_implemented( const char* func, unsigned line ) void gpgpusim_ptx_error_impl( const char *func, const char *file, unsigned line, const char *msg, ... ) { - va_list ap; - char buf[1024]; - va_start(ap,msg); - vsnprintf(buf,1024,msg,ap); - va_end(ap); + va_list ap; + char buf[1024]; + va_start(ap,msg); + vsnprintf(buf,1024,msg,ap); + va_end(ap); - printf("GPGPU-Sim CUDA API: %s\n", buf); - printf(" [%s:%u : %s]\n", file, line, func ); - abort(); + printf("GPGPU-Sim CUDA API: %s\n", buf); + printf(" [%s:%u : %s]\n", file, line, func ); + abort(); } void gpgpusim_ptx_assert_impl( int test_value, const char *func, const char *file, unsigned line, const char *msg, ... ) { - va_list ap; - char buf[1024]; - va_start(ap,msg); - vsnprintf(buf,1024,msg,ap); - va_end(ap); + va_list ap; + char buf[1024]; + va_start(ap,msg); + vsnprintf(buf,1024,msg,ap); + va_end(ap); - if ( test_value == 0 ) - gpgpusim_ptx_error_impl(func, file, line, msg); + if ( test_value == 0 ) + gpgpusim_ptx_error_impl(func, file, line, msg); } @@ -398,860 +398,1142 @@ int g_active_device = 0; //active gpu that runs the code std::list<kernel_config> g_cuda_launch_stack; /******************************************************************************* -* * -* * -* * -*******************************************************************************/ + * * + * * + * * + *******************************************************************************/ extern "C" { /******************************************************************************* -* * -* * -* * -*******************************************************************************/ + * * + * * + * * + *******************************************************************************/ __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) { - CUctx_st* context = GPGPUSim_Context(); - *devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size); - if(g_debug_execution >= 3) - printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr); - if ( *devPtr ) { - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorMemoryAllocation; - } + CUctx_st* context = GPGPUSim_Context(); + *devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size); + if(g_debug_execution >= 3) + printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr); + if ( *devPtr ) { + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorMemoryAllocation; + } } - + __host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size) { - GPGPUSim_Context(); - *ptr = malloc(size); - if ( *ptr ) { - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorMemoryAllocation; - } - } - __host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height) + GPGPUSim_Context(); + *ptr = malloc(size); + if ( *ptr ) { + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorMemoryAllocation; + } +} +__host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height) { - unsigned malloc_width_inbytes = width; - printf("GPGPU-Sim PTX: cudaMallocPitch (width = %d)\n", malloc_width_inbytes); - CUctx_st* ctx = GPGPUSim_Context(); - *devPtr = ctx->get_device()->get_gpgpu()->gpu_malloc(malloc_width_inbytes*height); - pitch[0] = malloc_width_inbytes; - if ( *devPtr ) { - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorMemoryAllocation; - } + unsigned malloc_width_inbytes = width; + printf("GPGPU-Sim PTX: cudaMallocPitch (width = %d)\n", malloc_width_inbytes); + CUctx_st* ctx = GPGPUSim_Context(); + *devPtr = ctx->get_device()->get_gpgpu()->gpu_malloc(malloc_width_inbytes*height); + pitch[0] = malloc_width_inbytes; + if ( *devPtr ) { + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorMemoryAllocation; + } } __host__ cudaError_t CUDARTAPI cudaMallocArray(struct cudaArray **array, const struct cudaChannelFormatDesc *desc, size_t width, size_t height __dv(1)) { - unsigned size = width * height * ((desc->x + desc->y + desc->z + desc->w)/8); - CUctx_st* context = GPGPUSim_Context(); - (*array) = (struct cudaArray*) malloc(sizeof(struct cudaArray)); - (*array)->desc = *desc; - (*array)->width = width; - (*array)->height = height; - (*array)->size = size; - (*array)->dimensions = 2; - ((*array)->devPtr32)= (int) (long long)context->get_device()->get_gpgpu()->gpu_mallocarray(size); - printf("GPGPU-Sim PTX: cudaMallocArray: devPtr32 = %d\n", ((*array)->devPtr32)); - ((*array)->devPtr) = (void*) (long long) ((*array)->devPtr32); - if ( ((*array)->devPtr) ) { - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorMemoryAllocation; - } + unsigned size = width * height * ((desc->x + desc->y + desc->z + desc->w)/8); + CUctx_st* context = GPGPUSim_Context(); + (*array) = (struct cudaArray*) malloc(sizeof(struct cudaArray)); + (*array)->desc = *desc; + (*array)->width = width; + (*array)->height = height; + (*array)->size = size; + (*array)->dimensions = 2; + ((*array)->devPtr32)= (int) (long long)context->get_device()->get_gpgpu()->gpu_mallocarray(size); + printf("GPGPU-Sim PTX: cudaMallocArray: devPtr32 = %d\n", ((*array)->devPtr32)); + ((*array)->devPtr) = (void*) (long long) ((*array)->devPtr32); + if ( ((*array)->devPtr) ) { + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorMemoryAllocation; + } } __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr) { - // TODO... manage g_global_mem space? - return g_last_cudaError = cudaSuccess; + // TODO... manage g_global_mem space? + return g_last_cudaError = cudaSuccess; } - __host__ cudaError_t CUDARTAPI cudaFreeHost(void *ptr) +__host__ cudaError_t CUDARTAPI cudaFreeHost(void *ptr) { - free (ptr); // this will crash the system if called twice - return g_last_cudaError = cudaSuccess; + free (ptr); // this will crash the system if called twice + return g_last_cudaError = cudaSuccess; } - __host__ cudaError_t CUDARTAPI cudaFreeArray(struct cudaArray *array) +__host__ cudaError_t CUDARTAPI cudaFreeArray(struct cudaArray *array) { - // TODO... manage g_global_mem space? - return g_last_cudaError = cudaSuccess; + // TODO... manage g_global_mem space? + return g_last_cudaError = cudaSuccess; }; /******************************************************************************* -* * -* * -* * -*******************************************************************************/ + * * + * * + * * + *******************************************************************************/ __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind) { - //CUctx_st *context = GPGPUSim_Context(); - //gpgpu_t *gpu = context->get_device()->get_gpgpu(); - if(g_debug_execution >= 3) - printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst); - if( kind == cudaMemcpyHostToDevice ) - g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) ); - else if( kind == cudaMemcpyDeviceToHost ) - g_stream_manager->push( stream_operation((size_t)src,dst,count,0) ); - else if( kind == cudaMemcpyDeviceToDevice ) - g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) ); - else { - printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported cudaMemcpyKind\n"); - abort(); - } - return g_last_cudaError = cudaSuccess; + //CUctx_st *context = GPGPUSim_Context(); + //gpgpu_t *gpu = context->get_device()->get_gpgpu(); + if(g_debug_execution >= 3) + printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst); + if( kind == cudaMemcpyHostToDevice ) + g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) ); + else if( kind == cudaMemcpyDeviceToHost ) + g_stream_manager->push( stream_operation((size_t)src,dst,count,0) ); + else if( kind == cudaMemcpyDeviceToDevice ) + g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) ); + else { + printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported cudaMemcpyKind\n"); + abort(); + } + return g_last_cudaError = cudaSuccess; } - __host__ cudaError_t CUDARTAPI cudaMemcpyToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind) +__host__ cudaError_t CUDARTAPI cudaMemcpyToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind) { - CUctx_st *context = GPGPUSim_Context(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); - size_t size = count; - printf("GPGPU-Sim PTX: cudaMemcpyToArray\n"); - if( kind == cudaMemcpyHostToDevice ) - gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size); - else if( kind == cudaMemcpyDeviceToHost ) - gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size); - else if( kind == cudaMemcpyDeviceToDevice ) - gpu->memcpy_gpu_to_gpu( (size_t)(dst->devPtr), (size_t)src, size); - else { - printf("GPGPU-Sim PTX: cudaMemcpyToArray - ERROR : unsupported cudaMemcpyKind\n"); - abort(); - } - dst->devPtr32 = (unsigned) (size_t)(dst->devPtr); - return g_last_cudaError = cudaSuccess; + CUctx_st *context = GPGPUSim_Context(); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + size_t size = count; + printf("GPGPU-Sim PTX: cudaMemcpyToArray\n"); + if( kind == cudaMemcpyHostToDevice ) + gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size); + else if( kind == cudaMemcpyDeviceToHost ) + gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size); + else if( kind == cudaMemcpyDeviceToDevice ) + gpu->memcpy_gpu_to_gpu( (size_t)(dst->devPtr), (size_t)src, size); + else { + printf("GPGPU-Sim PTX: cudaMemcpyToArray - ERROR : unsupported cudaMemcpyKind\n"); + abort(); + } + dst->devPtr32 = (unsigned) (size_t)(dst->devPtr); + return g_last_cudaError = cudaSuccess; } - __host__ cudaError_t CUDARTAPI cudaMemcpyFromArray(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind) +__host__ cudaError_t CUDARTAPI cudaMemcpyFromArray(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } - __host__ cudaError_t CUDARTAPI cudaMemcpyArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t count, enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToDevice)) +__host__ cudaError_t CUDARTAPI cudaMemcpyArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t count, enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToDevice)) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } - __host__ cudaError_t CUDARTAPI cudaMemcpy2D(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind) +__host__ cudaError_t CUDARTAPI cudaMemcpy2D(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind) { - CUctx_st *context = GPGPUSim_Context(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); - struct cudaArray *cuArray_ptr; - size_t size = spitch*height; - cuArray_ptr = (cudaArray*)dst; - gpgpusim_ptx_assert( (dpitch==spitch), "different src and dst pitch not supported yet" ); - if( kind == cudaMemcpyHostToDevice ) - gpu->memcpy_to_gpu( (size_t)dst, src, size ); - else if( kind == cudaMemcpyDeviceToHost ) - gpu->memcpy_from_gpu( dst, (size_t)src, size ); - else if( kind == cudaMemcpyDeviceToDevice ) - gpu->memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, size); - else { - printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n"); - abort(); - } - return g_last_cudaError = cudaSuccess; + CUctx_st *context = GPGPUSim_Context(); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + struct cudaArray *cuArray_ptr; + size_t size = spitch*height; + cuArray_ptr = (cudaArray*)dst; + gpgpusim_ptx_assert( (dpitch==spitch), "different src and dst pitch not supported yet" ); + if( kind == cudaMemcpyHostToDevice ) + gpu->memcpy_to_gpu( (size_t)dst, src, size ); + else if( kind == cudaMemcpyDeviceToHost ) + gpu->memcpy_from_gpu( dst, (size_t)src, size ); + else if( kind == cudaMemcpyDeviceToDevice ) + gpu->memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, size); + else { + printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n"); + abort(); + } + return g_last_cudaError = cudaSuccess; } - __host__ cudaError_t CUDARTAPI cudaMemcpy2DToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind) +__host__ cudaError_t CUDARTAPI cudaMemcpy2DToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind) { - CUctx_st *context = GPGPUSim_Context(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); - size_t size = spitch*height; - size_t channel_size = dst->desc.w+dst->desc.x+dst->desc.y+dst->desc.z; - gpgpusim_ptx_assert( ((channel_size%8) == 0), "none byte multiple destination channel size not supported (sz=%u)", channel_size ); - unsigned elem_size = channel_size/8; - gpgpusim_ptx_assert( (dst->dimensions==2), "copy to none 2D array not supported" ); - gpgpusim_ptx_assert( (wOffset==0), "non-zero wOffset not yet supported" ); - gpgpusim_ptx_assert( (hOffset==0), "non-zero hOffset not yet supported" ); - gpgpusim_ptx_assert( (dst->height == (int)height), "partial copy not supported" ); - gpgpusim_ptx_assert( (elem_size*dst->width == width), "partial copy not supported" ); - gpgpusim_ptx_assert( (spitch == width), "spitch != width not supported" ); - if( kind == cudaMemcpyHostToDevice ) - gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size); - else if( kind == cudaMemcpyDeviceToHost ) - gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size); - else if( kind == cudaMemcpyDeviceToDevice ) - gpu->memcpy_gpu_to_gpu( (size_t)dst->devPtr, (size_t)src, size); - else { - printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n"); - abort(); - } - dst->devPtr32 = (unsigned) (size_t)(dst->devPtr); - return g_last_cudaError = cudaSuccess; + CUctx_st *context = GPGPUSim_Context(); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + size_t size = spitch*height; + size_t channel_size = dst->desc.w+dst->desc.x+dst->desc.y+dst->desc.z; + gpgpusim_ptx_assert( ((channel_size%8) == 0), "none byte multiple destination channel size not supported (sz=%u)", channel_size ); + unsigned elem_size = channel_size/8; + gpgpusim_ptx_assert( (dst->dimensions==2), "copy to none 2D array not supported" ); + gpgpusim_ptx_assert( (wOffset==0), "non-zero wOffset not yet supported" ); + gpgpusim_ptx_assert( (hOffset==0), "non-zero hOffset not yet supported" ); + gpgpusim_ptx_assert( (dst->height == (int)height), "partial copy not supported" ); + gpgpusim_ptx_assert( (elem_size*dst->width == width), "partial copy not supported" ); + gpgpusim_ptx_assert( (spitch == width), "spitch != width not supported" ); + if( kind == cudaMemcpyHostToDevice ) + gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size); + else if( kind == cudaMemcpyDeviceToHost ) + gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size); + else if( kind == cudaMemcpyDeviceToDevice ) + gpu->memcpy_gpu_to_gpu( (size_t)dst->devPtr, (size_t)src, size); + else { + printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n"); + abort(); + } + dst->devPtr32 = (unsigned) (size_t)(dst->devPtr); + return g_last_cudaError = cudaSuccess; } - __host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArray(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind) +__host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArray(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } - __host__ cudaError_t CUDARTAPI cudaMemcpy2DArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToDevice)) +__host__ cudaError_t CUDARTAPI cudaMemcpy2DArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToDevice)) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } - __host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const char *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice)) +__host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const char *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice)) { - //CUctx_st *context = GPGPUSim_Context(); - assert(kind == cudaMemcpyHostToDevice); - printf("GPGPU-Sim PTX: cudaMemcpyToSymbol: symbol = %p\n", symbol); - //stream_operation( const char *symbol, const void *src, size_t count, size_t offset ) - g_stream_manager->push( stream_operation(src,symbol,count,offset,0) ); - //gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu()); - return g_last_cudaError = cudaSuccess; + //CUctx_st *context = GPGPUSim_Context(); + assert(kind == cudaMemcpyHostToDevice); + printf("GPGPU-Sim PTX: cudaMemcpyToSymbol: symbol = %p\n", symbol); + //stream_operation( const char *symbol, const void *src, size_t count, size_t offset ) + g_stream_manager->push( stream_operation(src,symbol,count,offset,0) ); + //gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu()); + return g_last_cudaError = cudaSuccess; } - __host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const char *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost)) +__host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const char *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost)) { - //CUctx_st *context = GPGPUSim_Context(); - assert(kind == cudaMemcpyDeviceToHost); - printf("GPGPU-Sim PTX: cudaMemcpyFromSymbol: symbol = %p\n", symbol); - g_stream_manager->push( stream_operation(symbol,dst,count,offset,0) ); - //gpgpu_ptx_sim_memcpy_symbol(symbol,dst,count,offset,0,context->get_device()->get_gpgpu()); - return g_last_cudaError = cudaSuccess; + //CUctx_st *context = GPGPUSim_Context(); + assert(kind == cudaMemcpyDeviceToHost); + printf("GPGPU-Sim PTX: cudaMemcpyFromSymbol: symbol = %p\n", symbol); + g_stream_manager->push( stream_operation(symbol,dst,count,offset,0) ); + //gpgpu_ptx_sim_memcpy_symbol(symbol,dst,count,offset,0,context->get_device()->get_gpgpu()); + return g_last_cudaError = cudaSuccess; } /******************************************************************************* -* * -* * -* * -*******************************************************************************/ + * * + * * + * * + *******************************************************************************/ - __host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream) +__host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream) { - struct CUstream_st *s = (struct CUstream_st *)stream; - switch( kind ) { - case cudaMemcpyHostToDevice: g_stream_manager->push( stream_operation(src,(size_t)dst,count,s) ); break; - case cudaMemcpyDeviceToHost: g_stream_manager->push( stream_operation((size_t)src,dst,count,s) ); break; - case cudaMemcpyDeviceToDevice: g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,s) ); break; - default: - abort(); - } - return g_last_cudaError = cudaSuccess; + struct CUstream_st *s = (struct CUstream_st *)stream; + switch( kind ) { + case cudaMemcpyHostToDevice: g_stream_manager->push( stream_operation(src,(size_t)dst,count,s) ); break; + case cudaMemcpyDeviceToHost: g_stream_manager->push( stream_operation((size_t)src,dst,count,s) ); break; + case cudaMemcpyDeviceToDevice: g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,s) ); break; + default: + abort(); + } + return g_last_cudaError = cudaSuccess; } - __host__ cudaError_t CUDARTAPI cudaMemcpyToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream) +__host__ cudaError_t CUDARTAPI cudaMemcpyToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } - __host__ cudaError_t CUDARTAPI cudaMemcpyFromArrayAsync(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream) +__host__ cudaError_t CUDARTAPI cudaMemcpyFromArrayAsync(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } - __host__ cudaError_t CUDARTAPI cudaMemcpy2DAsync(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream) +__host__ cudaError_t CUDARTAPI cudaMemcpy2DAsync(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } - __host__ cudaError_t CUDARTAPI cudaMemcpy2DToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream) +__host__ cudaError_t CUDARTAPI cudaMemcpy2DToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } - __host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArrayAsync(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream) +__host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArrayAsync(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } /******************************************************************************* -* * -* * -* * -*******************************************************************************/ + * * + * * + * * + *******************************************************************************/ __host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count) { - CUctx_st *context = GPGPUSim_Context(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); - gpu->gpu_memset((size_t)mem, c, count); - return g_last_cudaError = cudaSuccess; + CUctx_st *context = GPGPUSim_Context(); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + gpu->gpu_memset((size_t)mem, c, count); + return g_last_cudaError = cudaSuccess; } - __host__ cudaError_t CUDARTAPI cudaMemset2D(void *mem, size_t pitch, int c, size_t width, size_t height) +__host__ cudaError_t CUDARTAPI cudaMemset2D(void *mem, size_t pitch, int c, size_t width, size_t height) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } /******************************************************************************* -* * -* * -* * -*******************************************************************************/ + * * + * * + * * + *******************************************************************************/ - __host__ cudaError_t CUDARTAPI cudaGetSymbolAddress(void **devPtr, const char *symbol) +__host__ cudaError_t CUDARTAPI cudaGetSymbolAddress(void **devPtr, const char *symbol) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } - __host__ cudaError_t CUDARTAPI cudaGetSymbolSize(size_t *size, const char *symbol) +__host__ cudaError_t CUDARTAPI cudaGetSymbolSize(size_t *size, const char *symbol) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } /******************************************************************************* -* * -* * -* * -*******************************************************************************/ - __host__ cudaError_t CUDARTAPI cudaGetDeviceCount(int *count) + * * + * * + * * + *******************************************************************************/ +__host__ cudaError_t CUDARTAPI cudaGetDeviceCount(int *count) { - _cuda_device_id *dev = GPGPUSim_Init(); - *count = dev->num_devices(); - return g_last_cudaError = cudaSuccess; + _cuda_device_id *dev = GPGPUSim_Init(); + *count = dev->num_devices(); + return g_last_cudaError = cudaSuccess; } - __host__ cudaError_t CUDARTAPI cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device) +__host__ cudaError_t CUDARTAPI cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device) { - _cuda_device_id *dev = GPGPUSim_Init(); - if (device <= dev->num_devices() ) { - *prop= *dev->get_prop(); - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorInvalidDevice; - } + _cuda_device_id *dev = GPGPUSim_Init(); + if (device <= dev->num_devices() ) { + *prop= *dev->get_prop(); + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorInvalidDevice; + } } - __host__ cudaError_t CUDARTAPI cudaChooseDevice(int *device, const struct cudaDeviceProp *prop) +__host__ cudaError_t CUDARTAPI cudaChooseDevice(int *device, const struct cudaDeviceProp *prop) { - _cuda_device_id *dev = GPGPUSim_Init(); - *device = dev->get_id(); - return g_last_cudaError = cudaSuccess; + _cuda_device_id *dev = GPGPUSim_Init(); + *device = dev->get_id(); + return g_last_cudaError = cudaSuccess; } - - __host__ cudaError_t CUDARTAPI cudaSetDevice(int device) + +__host__ cudaError_t CUDARTAPI cudaSetDevice(int device) { - //set the active device to run cuda - if ( device <= GPGPUSim_Init()->num_devices() ) { - g_active_device = device; - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorInvalidDevice; - } + //set the active device to run cuda + if ( device <= GPGPUSim_Init()->num_devices() ) { + g_active_device = device; + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorInvalidDevice; + } } - - __host__ cudaError_t CUDARTAPI cudaGetDevice(int *device) + +__host__ cudaError_t CUDARTAPI cudaGetDevice(int *device) { - *device = g_active_device; - return g_last_cudaError = cudaSuccess; + *device = g_active_device; + return g_last_cudaError = cudaSuccess; } /******************************************************************************* -* * -* * -* * -*******************************************************************************/ + * * + * * + * * + *******************************************************************************/ - __host__ cudaError_t CUDARTAPI cudaBindTexture(size_t *offset, - const struct textureReference *texref, - const void *devPtr, - const struct cudaChannelFormatDesc *desc, - size_t size __dv(UINT_MAX)) +__host__ cudaError_t CUDARTAPI cudaBindTexture(size_t *offset, + const struct textureReference *texref, + const void *devPtr, + const struct cudaChannelFormatDesc *desc, + size_t size __dv(UINT_MAX)) { - CUctx_st *context = GPGPUSim_Context(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); - printf("GPGPU-Sim PTX: in cudaBindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference)); - struct cudaArray *array; - array = (struct cudaArray*) malloc(sizeof(struct cudaArray)); - array->desc = *desc; - array->size = size; - array->width = size; - array->height = 1; - array->dimensions = 1; - array->devPtr = (void*)devPtr; - array->devPtr32 = (int)(long long)devPtr; - offset = 0; - printf("GPGPU-Sim PTX: size = %zu\n", size); - printf("GPGPU-Sim PTX: texref = %p, array = %p\n", texref, array); - printf("GPGPU-Sim PTX: devPtr32 = %x\n", array->devPtr32); - printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref)); - printf("GPGPU-Sim PTX: ChannelFormatDesc: x=%d, y=%d, z=%d, w=%d\n", desc->x, desc->y, desc->z, desc->w); - printf("GPGPU-Sim PTX: Texture Normalized? = %d\n", texref->normalized); - gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array); - devPtr = (void*)(long long)array->devPtr32; - printf("GPGPU-Sim PTX: devPtr = %p\n", devPtr); - return g_last_cudaError = cudaSuccess; + CUctx_st *context = GPGPUSim_Context(); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + printf("GPGPU-Sim PTX: in cudaBindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference)); + struct cudaArray *array; + array = (struct cudaArray*) malloc(sizeof(struct cudaArray)); + array->desc = *desc; + array->size = size; + array->width = size; + array->height = 1; + array->dimensions = 1; + array->devPtr = (void*)devPtr; + array->devPtr32 = (int)(long long)devPtr; + offset = 0; + printf("GPGPU-Sim PTX: size = %zu\n", size); + printf("GPGPU-Sim PTX: texref = %p, array = %p\n", texref, array); + printf("GPGPU-Sim PTX: devPtr32 = %x\n", array->devPtr32); + printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref)); + printf("GPGPU-Sim PTX: ChannelFormatDesc: x=%d, y=%d, z=%d, w=%d\n", desc->x, desc->y, desc->z, desc->w); + printf("GPGPU-Sim PTX: Texture Normalized? = %d\n", texref->normalized); + gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array); + devPtr = (void*)(long long)array->devPtr32; + printf("GPGPU-Sim PTX: devPtr = %p\n", devPtr); + return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaBindTextureToArray(const struct textureReference *texref, const struct cudaArray *array, const struct cudaChannelFormatDesc *desc) { - CUctx_st *context = GPGPUSim_Context(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); - printf("GPGPU-Sim PTX: in cudaBindTextureToArray: %p %p\n", texref, array); - printf("GPGPU-Sim PTX: devPtr32 = %x\n", array->devPtr32); - printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref)); - printf("GPGPU-Sim PTX: Texture Normalized? = %d\n", texref->normalized); - gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array); - return g_last_cudaError = cudaSuccess; + CUctx_st *context = GPGPUSim_Context(); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + printf("GPGPU-Sim PTX: in cudaBindTextureToArray: %p %p\n", texref, array); + printf("GPGPU-Sim PTX: devPtr32 = %x\n", array->devPtr32); + printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref)); + printf("GPGPU-Sim PTX: Texture Normalized? = %d\n", texref->normalized); + gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array); + return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaUnbindTexture(const struct textureReference *texref) { - return g_last_cudaError = cudaSuccess; + return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaGetTextureAlignmentOffset(size_t *offset, const struct textureReference *texref) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } __host__ cudaError_t CUDARTAPI cudaGetTextureReference(const struct textureReference **texref, const char *symbol) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } __host__ cudaError_t CUDARTAPI cudaGetChannelDesc(struct cudaChannelFormatDesc *desc, const struct cudaArray *array) { - *desc = array->desc; - return g_last_cudaError = cudaSuccess; + *desc = array->desc; + return g_last_cudaError = cudaSuccess; } __host__ struct cudaChannelFormatDesc CUDARTAPI cudaCreateChannelDesc(int x, int y, int z, int w, enum cudaChannelFormatKind f) { - struct cudaChannelFormatDesc dummy; - dummy.x = x; - dummy.y = y; - dummy.z = z; - dummy.w = w; - dummy.f = f; - return dummy; + struct cudaChannelFormatDesc dummy; + dummy.x = x; + dummy.y = y; + dummy.z = z; + dummy.w = w; + dummy.f = f; + return dummy; } __host__ cudaError_t CUDARTAPI cudaGetLastError(void) { - return g_last_cudaError; + return g_last_cudaError; } __host__ const char* CUDARTAPI cudaGetErrorString(cudaError_t error) { - if( g_last_cudaError == cudaSuccess ) - return "no error"; - char buf[1024]; - snprintf(buf,1024,"<<GPGPU-Sim PTX: there was an error (code = %d)>>", g_last_cudaError); - return strdup(buf); + if( g_last_cudaError == cudaSuccess ) + return "no error"; + char buf[1024]; + snprintf(buf,1024,"<<GPGPU-Sim PTX: there was an error (code = %d)>>", g_last_cudaError); + return strdup(buf); } __host__ cudaError_t CUDARTAPI cudaConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, cudaStream_t stream) { - struct CUstream_st *s = (struct CUstream_st *)stream; - g_cuda_launch_stack.push_back( kernel_config(gridDim,blockDim,sharedMem,s) ); - return g_last_cudaError = cudaSuccess; + struct CUstream_st *s = (struct CUstream_st *)stream; + g_cuda_launch_stack.push_back( kernel_config(gridDim,blockDim,sharedMem,s) ); + return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaSetupArgument(const void *arg, size_t size, size_t offset) { - gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); - kernel_config &config = g_cuda_launch_stack.back(); - config.set_arg(arg,size,offset); + gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); + kernel_config &config = g_cuda_launch_stack.back(); + config.set_arg(arg,size,offset); + + struct gpgpu_ptx_sim_arg *param = (gpgpu_ptx_sim_arg*) calloc(1,sizeof(struct gpgpu_ptx_sim_arg)); + param->m_start = arg; + param->m_nbytes = size; + param->m_offset = offset; - struct gpgpu_ptx_sim_arg *param = (gpgpu_ptx_sim_arg*) calloc(1,sizeof(struct gpgpu_ptx_sim_arg)); - param->m_start = arg; - param->m_nbytes = size; - param->m_offset = offset; - - return g_last_cudaError = cudaSuccess; + return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) { - CUctx_st* context = GPGPUSim_Context(); - char *mode = getenv("PTX_SIM_MODE_FUNC"); - if( mode ) - sscanf(mode,"%u", &g_ptx_sim_mode); - gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); - kernel_config config = g_cuda_launch_stack.back(); - struct CUstream_st *stream = config.get_stream(); - printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun, - g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 ); - kernel_info_t *grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context); - std::string kname = grid->name(); - dim3 gridDim = config.grid_dim(); - dim3 blockDim = config.block_dim(); - printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n", - kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z ); - stream_operation op(grid,g_ptx_sim_mode,stream); - g_stream_manager->push(op); - g_cuda_launch_stack.pop_back(); - return g_last_cudaError = cudaSuccess; + CUctx_st* context = GPGPUSim_Context(); + char *mode = getenv("PTX_SIM_MODE_FUNC"); + if( mode ) + sscanf(mode,"%u", &g_ptx_sim_mode); + gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); + kernel_config config = g_cuda_launch_stack.back(); + struct CUstream_st *stream = config.get_stream(); + printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun, + g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 ); + kernel_info_t *grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context); + std::string kname = grid->name(); + dim3 gridDim = config.grid_dim(); + dim3 blockDim = config.block_dim(); + printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n", + kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z ); + stream_operation op(grid,g_ptx_sim_mode,stream); + g_stream_manager->push(op); + g_cuda_launch_stack.pop_back(); + return g_last_cudaError = cudaSuccess; } /******************************************************************************* -* * -* * -* * -*******************************************************************************/ + * * + * * + * * + *******************************************************************************/ __host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream) { - printf("GPGPU-Sim PTX: cudaStreamCreate\n"); + printf("GPGPU-Sim PTX: cudaStreamCreate\n"); #if (CUDART_VERSION >= 3000) - *stream = new struct CUstream_st(); - g_stream_manager->add_stream(*stream); + *stream = new struct CUstream_st(); + g_stream_manager->add_stream(*stream); #else - *stream = 0; - printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); + *stream = 0; + printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); #endif - return g_last_cudaError = cudaSuccess; + return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) { #if (CUDART_VERSION >= 3000) - g_stream_manager->destroy_stream(stream); + g_stream_manager->destroy_stream(stream); #endif - return g_last_cudaError = cudaSuccess; + return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaStreamSynchronize(cudaStream_t stream) { #if (CUDART_VERSION >= 3000) - if( stream == NULL ) - return g_last_cudaError = cudaErrorInvalidResourceHandle; - stream->synchronize(); + if( stream == NULL ) + return g_last_cudaError = cudaErrorInvalidResourceHandle; + stream->synchronize(); #else - printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); + printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); #endif - return g_last_cudaError = cudaSuccess; + return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaStreamQuery(cudaStream_t stream) { #if (CUDART_VERSION >= 3000) - if( stream == NULL ) - return g_last_cudaError = cudaErrorInvalidResourceHandle; - return g_last_cudaError = stream->empty()?cudaSuccess:cudaErrorNotReady; + if( stream == NULL ) + return g_last_cudaError = cudaErrorInvalidResourceHandle; + return g_last_cudaError = stream->empty()?cudaSuccess:cudaErrorNotReady; #else - printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); - return g_last_cudaError = cudaSuccess; // it is always success because all cuda calls are synchronous + printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); + return g_last_cudaError = cudaSuccess; // it is always success because all cuda calls are synchronous #endif } /******************************************************************************* -* * -* * -* * -*******************************************************************************/ + * * + * * + * * + *******************************************************************************/ __host__ cudaError_t CUDARTAPI cudaEventCreate(cudaEvent_t *event) { - CUevent_st *e = new CUevent_st(false); - g_timer_events[e->get_uid()] = e; + CUevent_st *e = new CUevent_st(false); + g_timer_events[e->get_uid()] = e; #if CUDART_VERSION >= 3000 - *event = e; + *event = e; #else - *event = e->get_uid(); + *event = e->get_uid(); #endif - return g_last_cudaError = cudaSuccess; + return g_last_cudaError = cudaSuccess; } CUevent_st *get_event(cudaEvent_t event) { - unsigned event_uid; + unsigned event_uid; #if CUDART_VERSION >= 3000 - event_uid = event->get_uid(); + event_uid = event->get_uid(); #else - event_uid = event; + event_uid = event; #endif - event_tracker_t::iterator e = g_timer_events.find(event_uid); - if( e == g_timer_events.end() ) - return NULL; - return e->second; + event_tracker_t::iterator e = g_timer_events.find(event_uid); + if( e == g_timer_events.end() ) + return NULL; + return e->second; } __host__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t stream) { - CUevent_st *e = get_event(event); - if( !e ) return g_last_cudaError = cudaErrorUnknown; - struct CUstream_st *s = (struct CUstream_st *)stream; - stream_operation op(e,s); - g_stream_manager->push(op); - return g_last_cudaError = cudaSuccess; + CUevent_st *e = get_event(event); + if( !e ) return g_last_cudaError = cudaErrorUnknown; + struct CUstream_st *s = (struct CUstream_st *)stream; + stream_operation op(e,s); + g_stream_manager->push(op); + return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaEventQuery(cudaEvent_t event) { - CUevent_st *e = get_event(event); - if( e == NULL ) { - return g_last_cudaError = cudaErrorInvalidValue; - } else if( e->done() ) { - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorNotReady; - } + CUevent_st *e = get_event(event); + if( e == NULL ) { + return g_last_cudaError = cudaErrorInvalidValue; + } else if( e->done() ) { + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorNotReady; + } } __host__ cudaError_t CUDARTAPI cudaEventSynchronize(cudaEvent_t event) { - printf("GPGPU-Sim API: cudaEventSynchronize ** waiting for event\n"); - fflush(stdout); - CUevent_st *e = (CUevent_st*) event; - while( !e->done() ) - ; - printf("GPGPU-Sim API: cudaEventSynchronize ** event detected\n"); - fflush(stdout); - return g_last_cudaError = cudaSuccess; + printf("GPGPU-Sim API: cudaEventSynchronize ** waiting for event\n"); + fflush(stdout); + CUevent_st *e = (CUevent_st*) event; + while( !e->done() ) + ; + printf("GPGPU-Sim API: cudaEventSynchronize ** event detected\n"); + fflush(stdout); + return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaEventDestroy(cudaEvent_t event) { - CUevent_st *e = get_event(event); - unsigned event_uid = e->get_uid(); - event_tracker_t::iterator pe = g_timer_events.find(event_uid); - if( pe == g_timer_events.end() ) - return g_last_cudaError = cudaErrorInvalidValue; - g_timer_events.erase(pe); - return g_last_cudaError = cudaSuccess; + CUevent_st *e = get_event(event); + unsigned event_uid = e->get_uid(); + event_tracker_t::iterator pe = g_timer_events.find(event_uid); + if( pe == g_timer_events.end() ) + return g_last_cudaError = cudaErrorInvalidValue; + g_timer_events.erase(pe); + return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaEventElapsedTime(float *ms, cudaEvent_t start, cudaEvent_t end) { - time_t elapsed_time; - CUevent_st *s = get_event(start); - CUevent_st *e = get_event(end); - if( s==NULL || e==NULL ) - return g_last_cudaError = cudaErrorUnknown; - elapsed_time = e->clock() - s->clock(); - *ms = 1000*elapsed_time; - return g_last_cudaError = cudaSuccess; + time_t elapsed_time; + CUevent_st *s = get_event(start); + CUevent_st *e = get_event(end); + if( s==NULL || e==NULL ) + return g_last_cudaError = cudaErrorUnknown; + elapsed_time = e->clock() - s->clock(); + *ms = 1000*elapsed_time; + return g_last_cudaError = cudaSuccess; } /******************************************************************************* -* * -* * -* * -*******************************************************************************/ + * * + * * + * * + *******************************************************************************/ __host__ cudaError_t CUDARTAPI cudaThreadExit(void) { - exit_simulation(); - return g_last_cudaError = cudaSuccess; + exit_simulation(); + return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaThreadSynchronize(void) { - //Called on host side - synchronize(); - return g_last_cudaError = cudaSuccess; + //Called on host side + synchronize(); + return g_last_cudaError = cudaSuccess; }; int CUDARTAPI __cudaSynchronizeThreads(void**, void*) { - return cudaThreadExit(); + return cudaThreadExit(); } + + + /******************************************************************************* -* * -* * -* * -*******************************************************************************/ + * * + * * + * * + *******************************************************************************/ + +//#include "../../cuobjdump_to_ptxplus/cuobjdump_parser.h" + +enum cuobjdumpSectionType { + PTXSECTION=0, + ELFSECTION +}; + +typedef struct cuobjdumpSectionRec { + enum cuobjdumpSectionType type; + char* arch; + char* identifier; + char* ptxfilename; + char* elffilename; + char* sassfilename; +}cuobjdumpSection; + + + +std::list<cuobjdumpSection> cuobjdump; + +// sectiontype: 0 for ptx, 1 for elf +void addCuobjdumpSection(int sectiontype){ + cuobjdumpSection x; + x.type = sectiontype? ELFSECTION: PTXSECTION; + cuobjdump.push_front(x); + printf("## Adding new section %s\n", x.type==PTXSECTION?"PTX":"ELF"); +} + +void setCuobjdumparch(const char* arch){ + printf("Adding arch: %s\n", arch); + cuobjdump.front().arch = strdup(arch); +} + +void setCuobjdumpidentifier(const char* identifier){ + printf("Adding identifier: %s\n", identifier); + cuobjdump.front().identifier = strdup(identifier); +} + +void setCuobjdumpptxfilename(const char* filename){ + printf("Adding ptx filename: %s\n", filename); + cuobjdump.front().ptxfilename = strdup(filename); +} + +void setCuobjdumpelffilename(const char* filename){ + cuobjdump.front().elffilename = strdup(filename); +} + +void setCuobjdumpsassfilename(const char* filename){ + cuobjdump.front().sassfilename = strdup(filename); +} +extern "C" int cuobjdump_parse(); +extern "C" FILE *cuobjdump_in; +//! Return PTX file +/*! + * This Function extract the whole PTX (for all the files) using cuobjdump + * */ +void extract_ptx(){ + char command[1000]; + char* whole_code; + char fname[1024]="_ptx_whole_code_WESWW"; + //! Running CUobjdump using dynamic link to current process + snprintf(command,1000,"cuobjdump -ptx -elf -sass /proc/%d/exe > %s",getpid(),fname); + printf("Running cuobjdump using \"%s\"\n", command); + int result = system(command); + if(result) {printf("ERROR: Failed to execute: %s\n", command); exit(1);} + + printf("Parsing file %s\n", fname); + cuobjdump_in = fopen(fname, "r"); + + cuobjdump_parse(); + fclose(cuobjdump_in); + printf("Done parsing!!!\n"); + // Just to stop things here + + /* Hadi's old code + FILE* fp = fopen(fname,"r"); + //! finding size of the file + int lSize= 0; + fseek (fp , 0 , SEEK_END); + lSize = ftell (fp); + rewind (fp); + //! allocate and copy the entire ptx + whole_code = (char*)malloc(lSize * sizeof(char)); + fread(whole_code,1,lSize,fp); + fclose(fp); + snprintf(command,1000,"rm %s",fname); + //system(command); + return whole_code; + */ +} + +//! Return proper ptx version +/*! + * This function return newest ptx version inside the argument + * which is is not newer than forced_max_capability + */ +unsigned get_best_version(std::list<cuobjdumpSection> sectionlist, CUctx_st *context){ + + unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability(); + unsigned max_capability=0; -void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) + std::list<cuobjdumpSection>::iterator iter; + + for ( iter = sectionlist.begin(); + iter != sectionlist.end(); + iter++ + ){ + unsigned capability = 0; + sscanf(iter->arch,"sm_%u", &capability); + if (capability > max_capability && + (forced_max_capability == 0 || capability <= forced_max_capability)){ + max_capability = capability; + } + } + return max_capability; + + /* Hadi's old code + char* start_entry; + char* arch_pointer; + start_entry = strstr(ptx,"Fatbin ptx code:"); + while(start_entry){ + arch_pointer = strstr (start_entry,"arch = "); + unsigned int capability = 0; + sscanf(arch_pointer,"arch = sm_%u",&capability); + if(capability > max_capability && (forced_max_capability == 0 || capability <= forced_max_capability)){ + max_capability = capability; + } + start_entry = strstr(arch_pointer,"Fatbin ptx code:"); + } + return max_capability; + */ +} +/*! + * Return number of diffrent ptx files generated in first argument + * which have sm version equal to selected_capability + */ +unsigned get_number_of_ptx(std::list<cuobjdumpSection> sectionlist, unsigned selected_capability){ + int result = 0; + + for ( std::list<cuobjdumpSection>::iterator iter = sectionlist.begin(); + iter != sectionlist.end(); + iter++ + ){ + unsigned capability = 0; + sscanf(iter->arch, "sm_%u", &capability); + if(capability <= selected_capability && + iter->type == PTXSECTION){ + printf("Found compatible PTX section with capability %s\n", iter->arch); + result++; + } + } + return result; + /* Hadi's old code + char* start_entry; + char* arch_pointer; + start_entry = strstr(ptx,"Fatbin ptx code:"); + while(start_entry){ + arch_pointer = strstr (start_entry,"arch = "); + unsigned int capability = 0; + sscanf(arch_pointer,"arch = sm_%u",&capability); + if(capability == selected_capability) + result++; + start_entry = strstr(arch_pointer,"Fatbin ptx code:"); + } + */ + +} + +char* readfile (const char* filename){ + assert (filename != NULL); + char str[128]; + FILE* fp = fopen(filename,"r"); + if (!fp) { + printf("ERROR: Could not open file %s for reading\n", filename); + assert (0); + } + //! finding size of the file + int filesize= 0; + fseek (fp , 0 , SEEK_END); + + filesize = ftell (fp); + fseek (fp, 0, SEEK_SET); + //! allocate and copy the entire ptx + char* ret = (char*)malloc((filesize +1)* sizeof(char)); + fread(ret,1,filesize,fp); + ret[filesize]='\0'; + fclose(fp); + return ret; +} + +cuobjdumpSection* findelfsection(std::list<cuobjdumpSection> sectionlist, unsigned selected_capability, const char* identifier){ + + std::list<cuobjdumpSection>::iterator iter; + for ( iter = sectionlist.begin(); + iter != sectionlist.end(); + iter++ + ){ + unsigned capability = 0; + sscanf(iter->arch,"sm_%u", &capability); + if(capability <= selected_capability && + strcmp(identifier, iter->identifier)==0 && + iter->type == ELFSECTION) + return &(*iter); + } + return NULL; +} + +void useCuobjdump() { + + CUctx_st *context = GPGPUSim_Context(); + unsigned source_num=1; + char *sass, *elf; + extract_ptx(); //extract all the output of cuobjdump to _cuobjdump_*.* + unsigned selected_capability = get_best_version(cuobjdump, context); //Find max capability less than forced_max_capability + unsigned total_ptx_files = get_number_of_ptx(cuobjdump,selected_capability); //Count ptx files for the given capability + + for ( std::list<cuobjdumpSection>::iterator iter = cuobjdump.begin(); + iter != cuobjdump.end(); + iter++ + ){ + unsigned capability = 0; + sscanf(iter->arch,"sm_%u", &capability); + if((capability <= selected_capability) && (iter->type == PTXSECTION)){ + symbol_table *symtab; + char *ptxcode = readfile(iter->ptxfilename); + if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) { + cuobjdumpSection* elfsection = findelfsection(cuobjdump, selected_capability, iter->identifier); + assert (elfsection!= NULL); + //char *elfcode = readfile(elfsection->elffilename); + //char *sasscode = readfile(elfsection->sassfilename); + char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus( + iter->ptxfilename, + elfsection->elffilename, + elfsection->sassfilename); + symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxplus_str,source_num); + printf("Adding %s with cubin handle %u\n", iter->ptxfilename, source_num); + //context->add_binary(symtab,total_ptx_files/*fat_cubin_handle*/); + //gpgpu_ptxinfo_load_from_string(ptxcode,total_ptx_files/*source_num*/); + context->add_binary(symtab, source_num/*fat_cubin_handle*/); + gpgpu_ptxinfo_load_from_string( ptxcode,total_ptx_files-source_num/*source_num*/); + delete[] ptxplus_str; + } else { + symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxcode, source_num/*total_ptx_files*//*source_num*/); + context->add_binary(symtab,source_num/*fat_cubin_handle*/); + gpgpu_ptxinfo_load_from_string( ptxcode, total_ptx_files-source_num/*source_num */); + } + source_num++; + /*! + * The order of files in output of cuobjdump is reverse + */ + load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); + load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); + } + } +} + +void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) { #if (CUDART_VERSION < 2010) - printf("GPGPU-Sim PTX: ERROR ** this version of GPGPU-Sim requires CUDA 2.1 or higher\n"); - exit(1); + printf("GPGPU-Sim PTX: ERROR ** this version of GPGPU-Sim requires CUDA 2.1 or higher\n"); + exit(1); #endif + CUctx_st *context = GPGPUSim_Context(); + static unsigned next_fat_bin_handle = 1; + if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump()) { + unsigned fat_cubin_handle = next_fat_bin_handle; + next_fat_bin_handle++; + printf("GPGPU-Sim PTX: __cudaRegisterFatBinary, fat_cubin_handle = %u\n", fat_cubin_handle); + /*! + * This function extracting all data from all files in first call + * then for next calls, only return the appropriate number + */ + assert(fat_cubin_handle >= 1); + if(fat_cubin_handle == 1) + useCuobjdump(); - CUctx_st *context = GPGPUSim_Context(); - static unsigned next_fat_bin_handle = 1; - static unsigned source_num=1; - unsigned fat_cubin_handle = next_fat_bin_handle++; - __cudaFatCudaBinary *info = (__cudaFatCudaBinary *)fatCubin; - assert( info->version >= 3 ); - unsigned num_ptx_versions=0; - unsigned max_capability=0; - unsigned selected_capability=0; - bool found=false; - unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability(); - while( info->ptx[num_ptx_versions].gpuProfileName != NULL ) { - unsigned capability=0; - sscanf(info->ptx[num_ptx_versions].gpuProfileName,"compute_%u",&capability); - printf("GPGPU-Sim PTX: __cudaRegisterFatBinary found PTX versions for '%s', ", info->ident); - printf("capability = %s\n", info->ptx[num_ptx_versions].gpuProfileName ); - if( forced_max_capability ) { - if( capability > max_capability && capability <= forced_max_capability ) { - found = true; - max_capability=capability; - selected_capability = num_ptx_versions; - } - } else { - if( capability > max_capability ) { - found = true; - max_capability=capability; - selected_capability = num_ptx_versions; - } - } - num_ptx_versions++; - } - if( found ) { - printf("GPGPU-Sim PTX: Loading PTX for %s, capability = %s\n", - info->ident, info->ptx[selected_capability].gpuProfileName ); - symbol_table *symtab; - const char *ptx = info->ptx[selected_capability].ptx; - if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) { - if (info->cubin[selected_capability].cubin ==NULL) { - printf("GPGPU-Sim PTX: Cannot convert to ptxplus no cubin found, probably because it was compiled using newer version of cuda (>=3.0)\nGPGPU-Sim PTX: Exiting ...\n"); - exit(1); - } - char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_to_ptxplus(ptx, info->cubin[selected_capability].cubin, source_num++, - context->get_device()->get_gpgpu()->get_config().saved_converted_ptxplus()); - symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxplus_str,source_num); - context->add_binary(symtab,fat_cubin_handle); - gpgpu_ptxinfo_load_from_string(ptx,source_num); - delete[] ptxplus_str; - } else { - symtab=gpgpu_ptx_sim_load_ptx_from_string(ptx,source_num); - context->add_binary(symtab,fat_cubin_handle); - gpgpu_ptxinfo_load_from_string( ptx, source_num ); - } - source_num++; - load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); - load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); - } else { - printf("GPGPU-Sim PTX: warning -- did not find an appropriate PTX in cubin\n"); - } - return (void**)fat_cubin_handle; + return (void**)fat_cubin_handle; + } else { + static unsigned source_num=1; + unsigned fat_cubin_handle = next_fat_bin_handle++; + __cudaFatCudaBinary *info = (__cudaFatCudaBinary *)fatCubin; + assert( info->version >= 3 ); + unsigned num_ptx_versions=0; + unsigned max_capability=0; + unsigned selected_capability=0; + bool found=false; + unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability(); + while( info->ptx[num_ptx_versions].gpuProfileName != NULL ) { + unsigned capability=0; + sscanf(info->ptx[num_ptx_versions].gpuProfileName,"compute_%u",&capability); + printf("GPGPU-Sim PTX: __cudaRegisterFatBinary found PTX versions for '%s', ", info->ident); + printf("capability = %s\n", info->ptx[num_ptx_versions].gpuProfileName ); + if( forced_max_capability ) { + if( capability > max_capability && capability <= forced_max_capability ) { + found = true; + max_capability=capability; + selected_capability = num_ptx_versions; + } + } else { + if( capability > max_capability ) { + found = true; + max_capability=capability; + selected_capability = num_ptx_versions; + } + } + num_ptx_versions++; + } + if( found ) { + printf("GPGPU-Sim PTX: Loading PTX for %s, capability = %s\n", + info->ident, info->ptx[selected_capability].gpuProfileName ); + symbol_table *symtab; + const char *ptx = info->ptx[selected_capability].ptx; + if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) { + if (info->cubin[selected_capability].cubin ==NULL) { + printf("GPGPU-Sim PTX: Cannot convert to ptxplus no cubin found, probably because it was compiled using newer version of cuda (>=3.0)\nGPGPU-Sim PTX: Exiting ...\n"); + exit(1); + } + char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_to_ptxplus(ptx, info->cubin[selected_capability].cubin, source_num++, + context->get_device()->get_gpgpu()->get_config().saved_converted_ptxplus(), false); + symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxplus_str,source_num); + context->add_binary(symtab,fat_cubin_handle); + gpgpu_ptxinfo_load_from_string(ptx,source_num); + delete[] ptxplus_str; + } else { + symtab=gpgpu_ptx_sim_load_ptx_from_string(ptx,source_num); + context->add_binary(symtab,fat_cubin_handle); + gpgpu_ptxinfo_load_from_string( ptx, source_num ); + } + source_num++; + load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); + load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); + } else { + printf("GPGPU-Sim PTX: warning -- did not find an appropriate PTX in cubin\n"); + } + return (void**)fat_cubin_handle; + } } void __cudaUnregisterFatBinary(void **fatCubinHandle) { - ; + ; +} + +cudaError_t cudaDeviceReset ( void ) { + // Should reset the simulated GPU + return g_last_cudaError = cudaSuccess; +} +cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ + // I don't know what this should do + return g_last_cudaError = cudaSuccess; } void CUDARTAPI __cudaRegisterFunction( - void **fatCubinHandle, - const char *hostFun, - char *deviceFun, - const char *deviceName, - int thread_limit, - uint3 *tid, - uint3 *bid, - dim3 *bDim, - dim3 *gDim - ) + void **fatCubinHandle, + const char *hostFun, + char *deviceFun, + const char *deviceName, + int thread_limit, + uint3 *tid, + uint3 *bid, + dim3 *bDim, + dim3 *gDim +) { - CUctx_st *context = GPGPUSim_Context(); - unsigned fat_cubin_handle = (unsigned)(unsigned long long)fatCubinHandle; - printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : hostFun 0x%p, fat_cubin_handle = %u\n", - deviceFun, hostFun, fat_cubin_handle); - context->register_function( fat_cubin_handle, hostFun, deviceFun ); + CUctx_st *context = GPGPUSim_Context(); + unsigned fat_cubin_handle = (unsigned)(unsigned long long)fatCubinHandle; + printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : hostFun 0x%p, fat_cubin_handle = %u\n", + deviceFun, hostFun, fat_cubin_handle); + context->register_function( fat_cubin_handle, hostFun, deviceFun ); } extern void __cudaRegisterVar( - void **fatCubinHandle, - char *hostVar, //pointer to...something - char *deviceAddress, //name of variable - const char *deviceName, //name of variable (same as above) - int ext, - int size, - int constant, - int global ) + void **fatCubinHandle, + char *hostVar, //pointer to...something + char *deviceAddress, //name of variable + const char *deviceName, //name of variable (same as above) + int ext, + int size, + int constant, + int global ) { - printf("GPGPU-Sim PTX: __cudaRegisterVar: hostVar = %p; deviceAddress = %s; deviceName = %s\n", hostVar, deviceAddress, deviceName); - printf("GPGPU-Sim PTX: __cudaRegisterVar: Registering const memory space of %d bytes\n", size); - fflush(stdout); - if ( constant && !global && !ext ) { - gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size); - } else if ( !constant && !global && !ext ) { - gpgpu_ptx_sim_register_global_variable(hostVar,deviceName,size); - } else cuda_not_implemented(__my_func__,__LINE__); + printf("GPGPU-Sim PTX: __cudaRegisterVar: hostVar = %p; deviceAddress = %s; deviceName = %s\n", hostVar, deviceAddress, deviceName); + printf("GPGPU-Sim PTX: __cudaRegisterVar: Registering const memory space of %d bytes\n", size); + fflush(stdout); + if ( constant && !global && !ext ) { + gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size); + } else if ( !constant && !global && !ext ) { + gpgpu_ptx_sim_register_global_variable(hostVar,deviceName,size); + } else cuda_not_implemented(__my_func__,__LINE__); } void __cudaRegisterShared( - void **fatCubinHandle, - void **devicePtr - ) + void **fatCubinHandle, + void **devicePtr +) { - // we don't do anything here - printf("GPGPU-Sim PTX: __cudaRegisterShared\n" ); + // we don't do anything here + printf("GPGPU-Sim PTX: __cudaRegisterShared\n" ); } void CUDARTAPI __cudaRegisterSharedVar( - void **fatCubinHandle, - void **devicePtr, - size_t size, - size_t alignment, - int storage - ) + void **fatCubinHandle, + void **devicePtr, + size_t size, + size_t alignment, + int storage +) { - // we don't do anything here - printf("GPGPU-Sim PTX: __cudaRegisterSharedVar\n" ); + // we don't do anything here + printf("GPGPU-Sim PTX: __cudaRegisterSharedVar\n" ); } void __cudaRegisterTexture( - void **fatCubinHandle, - const struct textureReference *hostVar, - const void **deviceAddress, - const char *deviceName, - int dim, - int norm, - int ext - ) //passes in a newly created textureReference + void **fatCubinHandle, + const struct textureReference *hostVar, + const void **deviceAddress, + const char *deviceName, + int dim, + int norm, + int ext +) //passes in a newly created textureReference { - CUctx_st *context = GPGPUSim_Context(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); - printf("GPGPU-Sim PTX: in __cudaRegisterTexture:\n"); - gpu->gpgpu_ptx_sim_bindNameToTexture(deviceName, hostVar); - printf("GPGPU-Sim PTX: int dim = %d\n", dim); - printf("GPGPU-Sim PTX: int norm = %d\n", norm); - printf("GPGPU-Sim PTX: int ext = %d\n", ext); - printf("GPGPU-Sim PTX: Execution warning: Not finished implementing \"%s\"\n", __my_func__ ); + CUctx_st *context = GPGPUSim_Context(); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + printf("GPGPU-Sim PTX: in __cudaRegisterTexture:\n"); + gpu->gpgpu_ptx_sim_bindNameToTexture(deviceName, hostVar); + printf("GPGPU-Sim PTX: int dim = %d\n", dim); + printf("GPGPU-Sim PTX: int norm = %d\n", norm); + printf("GPGPU-Sim PTX: int ext = %d\n", ext); + printf("GPGPU-Sim PTX: Execution warning: Not finished implementing \"%s\"\n", __my_func__ ); } #ifndef OPENGL_SUPPORT @@ -1260,15 +1542,15 @@ typedef unsigned long GLuint; cudaError_t cudaGLRegisterBufferObject(GLuint bufferObj) { - printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ ); - return g_last_cudaError = cudaSuccess; + printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ ); + return g_last_cudaError = cudaSuccess; } struct glbmap_entry { - GLuint m_bufferObj; - void *m_devPtr; - size_t m_size; - struct glbmap_entry *m_next; + GLuint m_bufferObj; + void *m_devPtr; + size_t m_size; + struct glbmap_entry *m_next; }; typedef struct glbmap_entry glbmap_entry_t; @@ -1277,182 +1559,182 @@ glbmap_entry_t* g_glbmap = NULL; cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj) { #ifdef OPENGL_SUPPORT - GLint buffer_size=0; - CUctx_st* ctx = GPGPUSim_Context(); + GLint buffer_size=0; + CUctx_st* ctx = GPGPUSim_Context(); - glbmap_entry_t *p = g_glbmap; - while ( p && p->m_bufferObj != bufferObj ) - p = p->m_next; - if ( p == NULL ) { - glBindBuffer(GL_ARRAY_BUFFER,bufferObj); - glGetBufferParameteriv(GL_ARRAY_BUFFER,GL_BUFFER_SIZE,&buffer_size); - assert( buffer_size != 0 ); - *devPtr = ctx->get_device()->get_gpgpu()->gpu_malloc(buffer_size); + glbmap_entry_t *p = g_glbmap; + while ( p && p->m_bufferObj != bufferObj ) + p = p->m_next; + if ( p == NULL ) { + glBindBuffer(GL_ARRAY_BUFFER,bufferObj); + glGetBufferParameteriv(GL_ARRAY_BUFFER,GL_BUFFER_SIZE,&buffer_size); + assert( buffer_size != 0 ); + *devPtr = ctx->get_device()->get_gpgpu()->gpu_malloc(buffer_size); - // create entry and insert to front of list - glbmap_entry_t *n = (glbmap_entry_t *) calloc(1,sizeof(glbmap_entry_t)); - n->m_next = g_glbmap; - g_glbmap = n; + // create entry and insert to front of list + glbmap_entry_t *n = (glbmap_entry_t *) calloc(1,sizeof(glbmap_entry_t)); + n->m_next = g_glbmap; + g_glbmap = n; - // initialize entry - n->m_bufferObj = bufferObj; - n->m_devPtr = *devPtr; - n->m_size = buffer_size; + // initialize entry + n->m_bufferObj = bufferObj; + n->m_devPtr = *devPtr; + n->m_size = buffer_size; - p = n; - } else { - buffer_size = p->m_size; - *devPtr = p->m_devPtr; - } + p = n; + } else { + buffer_size = p->m_size; + *devPtr = p->m_devPtr; + } - if ( *devPtr ) { - char *data = (char *) calloc(p->m_size,1); - glGetBufferSubData(GL_ARRAY_BUFFER,0,buffer_size,data); - memcpy_to_gpu( (size_t) *devPtr, data, buffer_size ); - free(data); - printf("GPGPU-Sim PTX: cudaGLMapBufferObject %zu bytes starting at 0x%llx..\n", (size_t)buffer_size, - (unsigned long long) *devPtr); - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorMemoryAllocation; - } + if ( *devPtr ) { + char *data = (char *) calloc(p->m_size,1); + glGetBufferSubData(GL_ARRAY_BUFFER,0,buffer_size,data); + memcpy_to_gpu( (size_t) *devPtr, data, buffer_size ); + free(data); + printf("GPGPU-Sim PTX: cudaGLMapBufferObject %zu bytes starting at 0x%llx..\n", (size_t)buffer_size, + (unsigned long long) *devPtr); + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorMemoryAllocation; + } - return g_last_cudaError = cudaSuccess; + return g_last_cudaError = cudaSuccess; #else - fflush(stdout); - fflush(stderr); - printf("GPGPU-Sim PTX: GPGPU-Sim support for OpenGL integration disabled -- exiting\n"); - fflush(stdout); - exit(50); + fflush(stdout); + fflush(stderr); + printf("GPGPU-Sim PTX: GPGPU-Sim support for OpenGL integration disabled -- exiting\n"); + fflush(stdout); + exit(50); #endif } cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj) { #ifdef OPENGL_SUPPORT - glbmap_entry_t *p = g_glbmap; - while ( p && p->m_bufferObj != bufferObj ) - p = p->m_next; - if ( p == NULL ) - return g_last_cudaError = cudaErrorUnknown; + glbmap_entry_t *p = g_glbmap; + while ( p && p->m_bufferObj != bufferObj ) + p = p->m_next; + if ( p == NULL ) + return g_last_cudaError = cudaErrorUnknown; - char *data = (char *) calloc(p->m_size,1); - memcpy_from_gpu( data,(size_t)p->m_devPtr,p->m_size ); - glBufferSubData(GL_ARRAY_BUFFER,0,p->m_size,data); - free(data); + char *data = (char *) calloc(p->m_size,1); + memcpy_from_gpu( data,(size_t)p->m_devPtr,p->m_size ); + glBufferSubData(GL_ARRAY_BUFFER,0,p->m_size,data); + free(data); - return g_last_cudaError = cudaSuccess; + return g_last_cudaError = cudaSuccess; #else - fflush(stdout); - fflush(stderr); - printf("GPGPU-Sim PTX: support for OpenGL integration disabled -- exiting\n"); - fflush(stdout); - exit(50); + fflush(stdout); + fflush(stderr); + printf("GPGPU-Sim PTX: support for OpenGL integration disabled -- exiting\n"); + fflush(stdout); + exit(50); #endif } cudaError_t cudaGLUnregisterBufferObject(GLuint bufferObj) { - printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ ); - return g_last_cudaError = cudaSuccess; + printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ ); + return g_last_cudaError = cudaSuccess; } #if (CUDART_VERSION >= 2010) cudaError_t CUDARTAPI cudaHostAlloc(void **pHost, size_t bytes, unsigned int flags) { - *pHost = malloc(bytes); - if( *pHost ) - return g_last_cudaError = cudaSuccess; - else - return g_last_cudaError = cudaErrorMemoryAllocation; + *pHost = malloc(bytes); + if( *pHost ) + return g_last_cudaError = cudaSuccess; + else + return g_last_cudaError = cudaErrorMemoryAllocation; } cudaError_t CUDARTAPI cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } cudaError_t CUDARTAPI cudaSetValidDevices(int *device_arr, int len) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } cudaError_t CUDARTAPI cudaSetDeviceFlags( int flags ) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, const char *hostFun ) { - CUctx_st *context = GPGPUSim_Context(); - function_info *entry = context->get_kernel(hostFun); - if( entry ) { - const struct gpgpu_ptx_sim_kernel_info *kinfo = entry->get_kernel_info(); - attr->sharedSizeBytes = kinfo->smem; - attr->constSizeBytes = kinfo->cmem; - attr->localSizeBytes = kinfo->lmem; - attr->numRegs = kinfo->regs; - attr->maxThreadsPerBlock = 0; // from pragmas? + CUctx_st *context = GPGPUSim_Context(); + function_info *entry = context->get_kernel(hostFun); + if( entry ) { + const struct gpgpu_ptx_sim_kernel_info *kinfo = entry->get_kernel_info(); + attr->sharedSizeBytes = kinfo->smem; + attr->constSizeBytes = kinfo->cmem; + attr->localSizeBytes = kinfo->lmem; + attr->numRegs = kinfo->regs; + attr->maxThreadsPerBlock = 0; // from pragmas? #if CUDART_VERSION >= 3000 - attr->ptxVersion = kinfo->ptx_version; - attr->binaryVersion = kinfo->sm_target; + attr->ptxVersion = kinfo->ptx_version; + attr->binaryVersion = kinfo->sm_target; #endif - } - return g_last_cudaError = cudaSuccess; + } + return g_last_cudaError = cudaSuccess; } cudaError_t CUDARTAPI cudaEventCreateWithFlags(cudaEvent_t *event, int flags) { - CUevent_st *e = new CUevent_st(flags==cudaEventBlockingSync); - g_timer_events[e->get_uid()] = e; + CUevent_st *e = new CUevent_st(flags==cudaEventBlockingSync); + g_timer_events[e->get_uid()] = e; #if CUDART_VERSION >= 3000 - *event = e; + *event = e; #else - *event = e->get_uid(); + *event = e->get_uid(); #endif - return g_last_cudaError = cudaSuccess; + return g_last_cudaError = cudaSuccess; } cudaError_t CUDARTAPI cudaDriverGetVersion(int *driverVersion) { - *driverVersion = CUDART_VERSION; - return g_last_cudaError = cudaErrorUnknown; + *driverVersion = CUDART_VERSION; + return g_last_cudaError = cudaErrorUnknown; } cudaError_t CUDARTAPI cudaRuntimeGetVersion(int *runtimeVersion) { - *runtimeVersion = CUDART_VERSION; - return g_last_cudaError = cudaErrorUnknown; + *runtimeVersion = CUDART_VERSION; + return g_last_cudaError = cudaErrorUnknown; } #endif cudaError_t CUDARTAPI cudaGLSetGLDevice(int device) { - printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ ); - return g_last_cudaError = cudaErrorUnknown; + printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ ); + return g_last_cudaError = cudaErrorUnknown; } typedef void* HGPUNV; cudaError_t CUDARTAPI cudaWGLGetDevice(int *device, HGPUNV hGpu) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } void CUDARTAPI __cudaMutexOperation(int lock) { - cuda_not_implemented(__my_func__,__LINE__); + cuda_not_implemented(__my_func__,__LINE__); } void CUDARTAPI __cudaTextureFetch(const void *tex, void *index, int integer, void *val) { - cuda_not_implemented(__my_func__,__LINE__); + cuda_not_implemented(__my_func__,__LINE__); } } @@ -1461,18 +1743,18 @@ namespace cuda_math { void CUDARTAPI __cudaMutexOperation(int lock) { - cuda_not_implemented(__my_func__,__LINE__); + cuda_not_implemented(__my_func__,__LINE__); } void CUDARTAPI __cudaTextureFetch(const void *tex, void *index, int integer, void *val) { - cuda_not_implemented(__my_func__,__LINE__); + cuda_not_implemented(__my_func__,__LINE__); } int CUDARTAPI __cudaSynchronizeThreads(void**, void*) { - //TODO This function should syncronize if we support Asyn kernel calls - return g_last_cudaError = cudaSuccess; + //TODO This function should syncronize if we support Asyn kernel calls + return g_last_cudaError = cudaSuccess; } } @@ -1492,106 +1774,106 @@ extern "C" FILE *ptxinfo_in; static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu ) { - printf( "GPGPU-Sim PTX: loading globals with explicit initializers... \n" ); - fflush(stdout); - int ng_bytes=0; - symbol_table::iterator g=symtab->global_iterator_begin(); + printf( "GPGPU-Sim PTX: loading globals with explicit initializers... \n" ); + fflush(stdout); + int ng_bytes=0; + symbol_table::iterator g=symtab->global_iterator_begin(); - for ( ; g!=symtab->global_iterator_end(); g++) { - symbol *global = *g; - if ( global->has_initializer() ) { - printf( "GPGPU-Sim PTX: initializing '%s' ... ", global->name().c_str() ); - unsigned addr=global->get_address(); - const type_info *type = global->type(); - type_info_key ti=type->get_key(); - size_t size; - int t; - ti.type_decode(size,t); - int nbytes = size/8; - int offset=0; - std::list<operand_info> init_list = global->get_initializer(); - for ( std::list<operand_info>::iterator i=init_list.begin(); i!=init_list.end(); i++ ) { - operand_info op = *i; - ptx_reg_t value = op.get_literal_value(); - assert( (addr+offset+nbytes) < min_gaddr ); // min_gaddr is start of "heap" for cudaMalloc - gpu->get_global_memory()->write(addr+offset,nbytes,&value,NULL,NULL); // assuming little endian here - offset+=nbytes; - ng_bytes+=nbytes; - } - printf(" wrote %u bytes\n", offset ); - } - } - printf( "GPGPU-Sim PTX: finished loading globals (%u bytes total).\n", ng_bytes ); - fflush(stdout); - return ng_bytes; + for ( ; g!=symtab->global_iterator_end(); g++) { + symbol *global = *g; + if ( global->has_initializer() ) { + printf( "GPGPU-Sim PTX: initializing '%s' ... ", global->name().c_str() ); + unsigned addr=global->get_address(); + const type_info *type = global->type(); + type_info_key ti=type->get_key(); + size_t size; + int t; + ti.type_decode(size,t); + int nbytes = size/8; + int offset=0; + std::list<operand_info> init_list = global->get_initializer(); + for ( std::list<operand_info>::iterator i=init_list.begin(); i!=init_list.end(); i++ ) { + operand_info op = *i; + ptx_reg_t value = op.get_literal_value(); + assert( (addr+offset+nbytes) < min_gaddr ); // min_gaddr is start of "heap" for cudaMalloc + gpu->get_global_memory()->write(addr+offset,nbytes,&value,NULL,NULL); // assuming little endian here + offset+=nbytes; + ng_bytes+=nbytes; + } + printf(" wrote %u bytes\n", offset ); + } + } + printf( "GPGPU-Sim PTX: finished loading globals (%u bytes total).\n", ng_bytes ); + fflush(stdout); + return ng_bytes; } static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu ) { - printf( "GPGPU-Sim PTX: loading constants with explicit initializers... " ); - fflush(stdout); - int nc_bytes = 0; - symbol_table::iterator g=symtab->const_iterator_begin(); + printf( "GPGPU-Sim PTX: loading constants with explicit initializers... " ); + fflush(stdout); + int nc_bytes = 0; + symbol_table::iterator g=symtab->const_iterator_begin(); - for ( ; g!=symtab->const_iterator_end(); g++) { - symbol *constant = *g; - if ( constant->is_const() && constant->has_initializer() ) { + for ( ; g!=symtab->const_iterator_end(); g++) { + symbol *constant = *g; + if ( constant->is_const() && constant->has_initializer() ) { - // get the constant element data size - int basic_type; - size_t num_bits; - constant->type()->get_key().type_decode(num_bits,basic_type); + // get the constant element data size + int basic_type; + size_t num_bits; + constant->type()->get_key().type_decode(num_bits,basic_type); - std::list<operand_info> init_list = constant->get_initializer(); - int nbytes_written = 0; - for ( std::list<operand_info>::iterator i=init_list.begin(); i!=init_list.end(); i++ ) { - operand_info op = *i; - ptx_reg_t value = op.get_literal_value(); - int nbytes = num_bits/8; - switch ( op.get_type() ) { - case int_t: assert(nbytes >= 1); break; - case float_op_t: assert(nbytes == 4); break; - case double_op_t: assert(nbytes >= 4); break; // account for double DEMOTING - default: - abort(); - } - unsigned addr=constant->get_address() + nbytes_written; - assert( addr+nbytes < min_gaddr ); + std::list<operand_info> init_list = constant->get_initializer(); + int nbytes_written = 0; + for ( std::list<operand_info>::iterator i=init_list.begin(); i!=init_list.end(); i++ ) { + operand_info op = *i; + ptx_reg_t value = op.get_literal_value(); + int nbytes = num_bits/8; + switch ( op.get_type() ) { + case int_t: assert(nbytes >= 1); break; + case float_op_t: assert(nbytes == 4); break; + case double_op_t: assert(nbytes >= 4); break; // account for double DEMOTING + default: + abort(); + } + unsigned addr=constant->get_address() + nbytes_written; + assert( addr+nbytes < min_gaddr ); - gpu->get_global_memory()->write(addr,nbytes,&value,NULL,NULL); // assume little endian (so u8 is the first byte in u32) - nc_bytes+=nbytes; - nbytes_written += nbytes; - } - } - } - printf( " done.\n"); - fflush(stdout); - return nc_bytes; + gpu->get_global_memory()->write(addr,nbytes,&value,NULL,NULL); // assume little endian (so u8 is the first byte in u32) + nc_bytes+=nbytes; + nbytes_written += nbytes; + } + } + } + printf( " done.\n"); + fflush(stdout); + return nc_bytes; } kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, - gpgpu_ptx_sim_arg_list_t args, - struct dim3 gridDim, - struct dim3 blockDim, - CUctx_st* context ) + gpgpu_ptx_sim_arg_list_t args, + struct dim3 gridDim, + struct dim3 blockDim, + CUctx_st* context ) { - function_info *entry = context->get_kernel(hostFun); - kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry); - if( entry == NULL ) { - printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found\n"); - abort(); - } - unsigned argcount=args.size(); - unsigned argn=1; - for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) { - entry->add_param_data(argcount-argn,&(*a)); - argn++; - } + function_info *entry = context->get_kernel(hostFun); + kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry); + if( entry == NULL ) { + printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found for %p\n", hostFun); + abort(); + } + unsigned argcount=args.size(); + unsigned argn=1; + for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) { + entry->add_param_data(argcount-argn,&(*a)); + argn++; + } - entry->finalize(result->get_param_memory()); - g_ptx_kernel_count++; - fflush(stdout); + entry->finalize(result->get_param_memory()); + g_ptx_kernel_count++; + fflush(stdout); - return result; + return result; } |