diff options
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 113 | ||||
| -rw-r--r-- | libopencl/opencl_runtime_api.cc | 14 | ||||
| -rw-r--r-- | src/abstract_hardware_model.cc | 10 | ||||
| -rw-r--r-- | src/abstract_hardware_model.h | 116 | ||||
| -rw-r--r-- | src/cuda-sim/cuda-sim.cc | 58 | ||||
| -rw-r--r-- | src/cuda-sim/cuda-sim.h | 6 | ||||
| -rw-r--r-- | src/cuda-sim/instructions.cc | 22 | ||||
| -rw-r--r-- | src/cuda-sim/ptx_ir.h | 51 | ||||
| -rw-r--r-- | src/debug.cc | 6 | ||||
| -rw-r--r-- | src/gpgpu-sim/shader.cc | 2 |
10 files changed, 199 insertions, 199 deletions
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index df8ef4f..93f7512 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -439,7 +439,7 @@ extern "C" { __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) { CUctx_st* context = GPGPUSim_Context(); - *devPtr = context->get_device()->get_gpgpu()->gpgpu_ptx_sim_malloc(size); + *devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size); printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr); if ( *devPtr ) { return g_last_cudaError = cudaSuccess; @@ -463,7 +463,7 @@ __host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size) 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()->gpgpu_ptx_sim_malloc(malloc_width_inbytes*height); + *devPtr = ctx->get_device()->get_gpgpu()->gpu_malloc(malloc_width_inbytes*height); pitch[0] = malloc_width_inbytes; if ( *devPtr ) { return g_last_cudaError = cudaSuccess; @@ -482,7 +482,7 @@ __host__ cudaError_t CUDARTAPI cudaMallocArray(struct cudaArray **array, const s (*array)->height = height; (*array)->size = size; (*array)->dimensions = 2; - ((*array)->devPtr32)= (int) (long long)context->get_device()->get_gpgpu()->gpgpu_ptx_sim_mallocarray(size); + ((*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) ) { @@ -522,11 +522,11 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr) gpgpu_t *gpu = context->get_device()->get_gpgpu(); printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst); if( kind == cudaMemcpyHostToDevice ) - gpu->gpgpu_ptx_sim_memcpy_to_gpu( (size_t)dst, src, count ); + gpu->memcpy_to_gpu( (size_t)dst, src, count ); else if( kind == cudaMemcpyDeviceToHost ) - gpu->gpgpu_ptx_sim_memcpy_from_gpu( dst, (size_t)src, count ); + gpu->memcpy_from_gpu( dst, (size_t)src, count ); else if( kind == cudaMemcpyDeviceToDevice ) - gpu->gpgpu_ptx_sim_memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, count ); + gpu->memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, count ); else { printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported cudaMemcpyKind\n"); abort(); @@ -541,11 +541,11 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr) size_t size = count; printf("GPGPU-Sim PTX: cudaMemcpyToArray\n"); if( kind == cudaMemcpyHostToDevice ) - gpu->gpgpu_ptx_sim_memcpy_to_gpu( (size_t)(dst->devPtr), src, size); + gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size); else if( kind == cudaMemcpyDeviceToHost ) - gpu->gpgpu_ptx_sim_memcpy_from_gpu( dst->devPtr, (size_t)src, size); + gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size); else if( kind == cudaMemcpyDeviceToDevice ) - gpu->gpgpu_ptx_sim_memcpy_gpu_to_gpu( (size_t)(dst->devPtr), (size_t)src, size); + gpu->memcpy_gpu_to_gpu( (size_t)(dst->devPtr), (size_t)src, size); else { printf("GPGPU-Sim PTX: cudaMemcpyToArray - ERROR : unsupported cudaMemcpyKind\n"); abort(); @@ -578,11 +578,11 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr) cuArray_ptr = (cudaArray*)dst; gpgpusim_ptx_assert( (dpitch==spitch), "different src and dst pitch not supported yet" ); if( kind == cudaMemcpyHostToDevice ) - gpu->gpgpu_ptx_sim_memcpy_to_gpu( (size_t)dst, src, size ); + gpu->memcpy_to_gpu( (size_t)dst, src, size ); else if( kind == cudaMemcpyDeviceToHost ) - gpu->gpgpu_ptx_sim_memcpy_from_gpu( dst, (size_t)src, size ); + gpu->memcpy_from_gpu( dst, (size_t)src, size ); else if( kind == cudaMemcpyDeviceToDevice ) - gpu->gpgpu_ptx_sim_memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, size); + gpu->memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, size); else { printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n"); abort(); @@ -606,11 +606,11 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr) 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->gpgpu_ptx_sim_memcpy_to_gpu( (size_t)(dst->devPtr), src, size); + gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size); else if( kind == cudaMemcpyDeviceToHost ) - gpu->gpgpu_ptx_sim_memcpy_from_gpu( dst->devPtr, (size_t)src, size); + gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size); else if( kind == cudaMemcpyDeviceToDevice ) - gpu->gpgpu_ptx_sim_memcpy_gpu_to_gpu( (size_t)dst->devPtr, (size_t)src, size); + gpu->memcpy_gpu_to_gpu( (size_t)dst->devPtr, (size_t)src, size); else { printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n"); abort(); @@ -714,7 +714,7 @@ __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->gpgpu_ptx_sim_memset((size_t)mem, c, count); + gpu->gpu_memset((size_t)mem, c, count); return g_last_cudaError = cudaSuccess; } @@ -800,8 +800,14 @@ __host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count) * * *******************************************************************************/ - __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)); @@ -816,62 +822,53 @@ __host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count) 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", gpgpu_ptx_sim_findNamefromTexture(texref)); + 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); - gpgpu_ptx_sim_bindTextureToArray(texref, array); + 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) +__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", gpgpu_ptx_sim_findNamefromTexture(texref)); + 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); - gpgpu_ptx_sim_bindTextureToArray(texref, array); + gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array); return g_last_cudaError = cudaSuccess; } - - __host__ cudaError_t CUDARTAPI cudaUnbindTexture(const struct textureReference *texref) +__host__ cudaError_t CUDARTAPI cudaUnbindTexture(const struct textureReference *texref) { return g_last_cudaError = cudaSuccess; } - - __host__ cudaError_t CUDARTAPI cudaGetTextureAlignmentOffset(size_t *offset, const struct textureReference *texref) +__host__ cudaError_t CUDARTAPI cudaGetTextureAlignmentOffset(size_t *offset, const struct textureReference *texref) { cuda_not_implemented(__my_func__,__LINE__); return g_last_cudaError = cudaErrorUnknown; } - - __host__ cudaError_t CUDARTAPI cudaGetTextureReference(const struct textureReference **texref, const char *symbol) +__host__ cudaError_t CUDARTAPI cudaGetTextureReference(const struct textureReference **texref, const char *symbol) { cuda_not_implemented(__my_func__,__LINE__); return g_last_cudaError = cudaErrorUnknown; } - - -/******************************************************************************* -* * -* * -* * -*******************************************************************************/ - - __host__ cudaError_t CUDARTAPI cudaGetChannelDesc(struct cudaChannelFormatDesc *desc, const struct cudaArray *array) +__host__ cudaError_t CUDARTAPI cudaGetChannelDesc(struct cudaChannelFormatDesc *desc, const struct cudaArray *array) { *desc = array->desc; return g_last_cudaError = cudaSuccess; } - __host__ struct cudaChannelFormatDesc CUDARTAPI cudaCreateChannelDesc(int x, int y, int z, int w, enum cudaChannelFormatKind f) +__host__ struct cudaChannelFormatDesc CUDARTAPI cudaCreateChannelDesc(int x, int y, int z, int w, enum cudaChannelFormatKind f) { struct cudaChannelFormatDesc dummy; dummy.x = x; @@ -882,18 +879,12 @@ __host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count) return dummy; } -/******************************************************************************* -* * -* * -* * -*******************************************************************************/ - - __host__ cudaError_t CUDARTAPI cudaGetLastError(void) +__host__ cudaError_t CUDARTAPI cudaGetLastError(void) { return g_last_cudaError; } - __host__ const char* CUDARTAPI cudaGetErrorString(cudaError_t error) +__host__ const char* CUDARTAPI cudaGetErrorString(cudaError_t error) { if( g_last_cudaError == cudaSuccess ) return "no error"; @@ -902,18 +893,6 @@ __host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count) return strdup(buf); } - - -/******************************************************************************* -* * -* * -* * -*******************************************************************************/ - - - - - __host__ cudaError_t CUDARTAPI cudaConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, cudaStream_t stream) { g_cuda_launch_stack.push_back( kernel_config(gridDim,blockDim,sharedMem,stream) ); @@ -1109,12 +1088,15 @@ int CUDARTAPI __cudaSynchronizeThreads(void**, void*) void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) { - CUctx_st *context = GPGPUSim_Context(); +#if (CUDART_VERSION < 2010) + 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; static unsigned source_num=1; unsigned fat_cubin_handle = next_fat_bin_handle++; -#if (CUDART_VERSION >= 2010) __cudaFatCudaBinary *info = (__cudaFatCudaBinary *)fatCubin; assert( info->version >= 3 ); unsigned num_ptx_versions=0; @@ -1165,7 +1147,6 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) } else { printf("GPGPU-Sim PTX: warning -- did not find an appropriate PTX in cubin\n"); } -#endif return (void**)fat_cubin_handle; } @@ -1246,8 +1227,10 @@ void __cudaRegisterTexture( 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"); - gpgpu_ptx_sim_bindNameToTexture(deviceName, hostVar); + 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); @@ -1287,7 +1270,7 @@ cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj) glBindBuffer(GL_ARRAY_BUFFER,bufferObj); glGetBufferParameteriv(GL_ARRAY_BUFFER,GL_BUFFER_SIZE,&buffer_size); assert( buffer_size != 0 ); - *devPtr = ctx->get_device()->get_gpgpu()->gpgpu_ptx_sim_malloc(buffer_size); + *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)); @@ -1308,7 +1291,7 @@ cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj) if ( *devPtr ) { char *data = (char *) calloc(p->m_size,1); glGetBufferSubData(GL_ARRAY_BUFFER,0,buffer_size,data); - gpgpu_ptx_sim_memcpy_to_gpu( (size_t) *devPtr, data, buffer_size ); + 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); @@ -1337,7 +1320,7 @@ cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj) return g_last_cudaError = cudaErrorUnknown; char *data = (char *) calloc(p->m_size,1); - gpgpu_ptx_sim_memcpy_from_gpu( data,(size_t)p->m_devPtr,p->m_size ); + memcpy_from_gpu( data,(size_t)p->m_devPtr,p->m_size ); glBufferSubData(GL_ARRAY_BUFFER,0,p->m_size,data); free(data); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index b331f92..91505ab 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -330,9 +330,9 @@ _cl_mem::_cl_mem( } if( !(flags & (CL_MEM_USE_HOST_PTR|CL_MEM_ALLOC_HOST_PTR)) ) { // if not allocating on host, then allocate GPU memory and make a copy - m_device_ptr = (size_t) gpu->the_device()->gpgpu_ptx_sim_malloc(size); + m_device_ptr = (size_t) gpu->the_device()->gpu_malloc(size); if( host_ptr ) - gpu->the_device()->gpgpu_ptx_sim_memcpy_to_gpu( m_device_ptr, host_ptr, size ); + gpu->the_device()->memcpy_to_gpu( m_device_ptr, host_ptr, size ); } } @@ -859,7 +859,7 @@ clEnqueueReadBuffer(cl_command_queue command_queue, if( !blocking_read ) gpgpusim_opencl_warning(__my_func__,__LINE__, "non-blocking read treated as blocking read"); gpgpu_t *gpu = command_queue->get_device()->the_device(); - gpu->gpgpu_ptx_sim_memcpy_from_gpu( ptr, (size_t)buffer, cb ); + gpu->memcpy_from_gpu( ptr, (size_t)buffer, cb ); return CL_SUCCESS; } @@ -877,7 +877,7 @@ clEnqueueWriteBuffer(cl_command_queue command_queue, if( !blocking_write ) gpgpusim_opencl_warning(__my_func__,__LINE__, "non-blocking write treated as blocking write"); gpgpu_t *gpu = command_queue->get_device()->the_device(); - gpu->gpgpu_ptx_sim_memcpy_to_gpu( (size_t)buffer, ptr, cb ); + gpu->memcpy_to_gpu( (size_t)buffer, ptr, cb ); return CL_SUCCESS; } @@ -1151,11 +1151,11 @@ clEnqueueCopyBuffer(cl_command_queue command_queue, gpgpu_t *gpu = command_queue->get_device()->the_device(); if( src->is_on_host() && !dst->is_on_host() ) - gpu->gpgpu_ptx_sim_memcpy_to_gpu( ((size_t)dst->device_ptr())+dst_offset, ((char*)src->host_ptr())+src_offset, cb ); + gpu->memcpy_to_gpu( ((size_t)dst->device_ptr())+dst_offset, ((char*)src->host_ptr())+src_offset, cb ); else if( !src->is_on_host() && dst->is_on_host() ) - gpu->gpgpu_ptx_sim_memcpy_from_gpu( ((char*)dst->host_ptr())+dst_offset, ((size_t)src->device_ptr())+src_offset, cb ); + gpu->memcpy_from_gpu( ((char*)dst->host_ptr())+dst_offset, ((size_t)src->device_ptr())+src_offset, cb ); else if( !src->is_on_host() && !dst->is_on_host() ) - gpu->gpgpu_ptx_sim_memcpy_gpu_to_gpu( ((size_t)dst->device_ptr())+dst_offset, ((size_t)src->device_ptr())+src_offset, cb ); + gpu->memcpy_gpu_to_gpu( ((size_t)dst->device_ptr())+dst_offset, ((size_t)src->device_ptr())+src_offset, cb ); else opencl_not_implemented(__my_func__,__LINE__); return CL_SUCCESS; diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index c52be1d..b943a76 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -15,12 +15,12 @@ void move_warp( warp_inst_t *&dst, warp_inst_t *&src ) gpgpu_t::gpgpu_t() { - g_global_mem = new memory_space_impl<8192>("global",64*1024); - g_param_mem = new memory_space_impl<8192>("param",64*1024); - g_tex_mem = new memory_space_impl<8192>("tex",64*1024); - g_surf_mem = new memory_space_impl<8192>("surf",64*1024); + m_global_mem = new memory_space_impl<8192>("global",64*1024); + m_param_mem = new memory_space_impl<8192>("param",64*1024); + m_tex_mem = new memory_space_impl<8192>("tex",64*1024); + m_surf_mem = new memory_space_impl<8192>("surf",64*1024); - g_dev_malloc=GLOBAL_HEAP_START; + m_dev_malloc=GLOBAL_HEAP_START; } void warp_inst_t::sort_accessq( unsigned qbegin ) diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 7162441..62a59e0 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -52,6 +52,7 @@ enum _memory_op_t { #include <vector> #include <assert.h> #include <stdlib.h> +#include <map> #if !defined(__VECTOR_TYPES_H__) struct dim3 { @@ -152,29 +153,110 @@ public: #define LOCAL_GENERIC_START (SHARED_GENERIC_START-TOTAL_LOCAL_MEM) #define STATIC_ALLOC_LIMIT (GLOBAL_HEAP_START - (TOTAL_LOCAL_MEM+TOTAL_SHARED_MEM)) +#if !defined(__CUDA_RUNTIME_API_H__) + +enum cudaChannelFormatKind { + cudaChannelFormatKindSigned, + cudaChannelFormatKindUnsigned, + cudaChannelFormatKindFloat +}; + +struct cudaChannelFormatDesc { + int x; + int y; + int z; + int w; + enum cudaChannelFormatKind f; +}; + +struct cudaArray { + void *devPtr; + int devPtr32; + struct cudaChannelFormatDesc desc; + int width; + int height; + int size; //in bytes + unsigned dimensions; +}; + +enum cudaTextureAddressMode { + cudaAddressModeWrap, + cudaAddressModeClamp +}; + +enum cudaTextureFilterMode { + cudaFilterModePoint, + cudaFilterModeLinear +}; + +enum cudaTextureReadMode { + cudaReadModeElementType, + cudaReadModeNormalizedFloat +}; + +struct textureReference { + int normalized; + enum cudaTextureFilterMode filterMode; + enum cudaTextureAddressMode addressMode[2]; + struct cudaChannelFormatDesc channelDesc; +}; + +#endif + class gpgpu_t { public: - gpgpu_t(); - void* gpgpu_ptx_sim_malloc( size_t size ); - void* gpgpu_ptx_sim_mallocarray( size_t count ); - void gpgpu_ptx_sim_memcpy_to_gpu( size_t dst_start_addr, const void *src, size_t count ); - void gpgpu_ptx_sim_memcpy_from_gpu( void *dst, size_t src_start_addr, size_t count ); - void gpgpu_ptx_sim_memcpy_gpu_to_gpu( size_t dst, size_t src, size_t count ); - void gpgpu_ptx_sim_memset( size_t dst_start_addr, int c, size_t count ); + gpgpu_t(); + void* gpu_malloc( size_t size ); + void* gpu_mallocarray( size_t count ); + void gpu_memset( size_t dst_start_addr, int c, size_t count ); + void memcpy_to_gpu( size_t dst_start_addr, const void *src, size_t count ); + void memcpy_from_gpu( void *dst, size_t src_start_addr, size_t count ); + void memcpy_gpu_to_gpu( size_t dst, size_t src, size_t count ); + + class memory_space *get_global_memory() { return m_global_mem; } + class memory_space *get_tex_memory() { return m_tex_mem; } + class memory_space *get_surf_memory() { return m_surf_mem; } + class memory_space *get_param_memory() { return m_param_mem; } - class memory_space *get_global_memory() { return g_global_mem; } - class memory_space *get_tex_memory() { return g_tex_mem; } - class memory_space *get_surf_memory() { return g_surf_mem; } - class memory_space *get_param_memory() { return g_param_mem; } + void gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, const struct cudaArray* array); + void gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref); + const char* gpgpu_ptx_sim_findNamefromTexture(const struct textureReference* texref); + unsigned ptx_set_tex_cache_linesize(unsigned linesize); + const struct cudaArray* gpgpu_ptx_sim_accessArrayofTexture(struct textureReference* texref); + const struct textureReference* gpgpu_ptx_sim_accessTextureofName(const char* name); + int gpgpu_ptx_sim_sizeofTexture(const char* name); + + const struct textureReference* get_texref(const std::string &texname) const + { + std::map<std::string, const struct textureReference*>::const_iterator t=NameToTextureMap.find(texname); + assert( t != NameToTextureMap.end() ); + return t->second; + } + const struct cudaArray* get_texarray( const struct textureReference *texref ) const + { + std::map<const struct textureReference*,const struct cudaArray*>::const_iterator t=TextureToArrayMap.find(texref); + assert(t != TextureToArrayMap.end()); + return t->second; + } + const struct textureInfo* get_texinfo( const struct textureReference *texref ) const + { + std::map<const struct textureReference*, const struct textureInfo*>::const_iterator t=TextureToInfoMap.find(texref); + assert(t != TextureToInfoMap.end()); + return t->second; + } protected: - // functional simulation state - class memory_space *g_global_mem; - class memory_space *g_tex_mem; - class memory_space *g_surf_mem; - class memory_space *g_param_mem; + class memory_space *m_global_mem; + class memory_space *m_tex_mem; + class memory_space *m_surf_mem; + class memory_space *m_param_mem; + + unsigned long long m_dev_malloc; - unsigned long long g_dev_malloc; + std::map<const struct textureReference*,const struct cudaArray*> TextureToArrayMap; // texture bindings + std::map<const struct textureReference*, const struct textureInfo*> TextureToInfoMap; + std::map<std::string, const struct textureReference*> NameToTextureMap; + unsigned int g_texcache_linesize; }; struct gpgpu_ptx_sim_kernel_info diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 5293d7d..5d60010 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -95,24 +95,21 @@ addr_t g_debug_pc = 0xBEEF1518; FILE* ptx_inst_debug_file; unsigned g_ptx_sim_num_insn = 0; -std::map<const struct textureReference*,const struct cudaArray*> TextureToArrayMap; // texture bindings -std::map<const struct textureReference*, const struct textureInfo*> TextureToInfoMap; -std::map<std::string, const struct textureReference*> NameToTextureMap; -unsigned int g_texcache_linesize; unsigned gpgpu_param_num_shaders = 0; -void gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref) +void gpgpu_t::gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref) { std::string texname(name); NameToTextureMap[texname] = texref; } -const struct textureReference* gpgpu_ptx_sim_accessTextureofName(const char* name) { +const struct textureReference* gpgpu_t::gpgpu_ptx_sim_accessTextureofName(const char* name) +{ std::string texname(name); return NameToTextureMap[texname]; } -const char* gpgpu_ptx_sim_findNamefromTexture(const struct textureReference* texref) +const char* gpgpu_t::gpgpu_ptx_sim_findNamefromTexture(const struct textureReference* texref) { std::map<std::string, const struct textureReference*>::iterator itr = NameToTextureMap.begin(); while (itr != NameToTextureMap.end()) { @@ -140,7 +137,7 @@ unsigned int intLOGB2( unsigned int v ) { return r; } -void gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, const struct cudaArray* array) +void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, const struct cudaArray* array) { TextureToArrayMap[texref] = array; unsigned int texel_size_bits = array->desc.w + array->desc.x + array->desc.y + array->desc.z; @@ -195,11 +192,12 @@ void gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, con TextureToInfoMap[texref] = texInfo; } -const struct cudaArray* gpgpu_ptx_sim_accessArrayofTexture(struct textureReference* texref) { +const struct cudaArray* gpgpu_t::gpgpu_ptx_sim_accessArrayofTexture(struct textureReference* texref) +{ return TextureToArrayMap[texref]; } -int gpgpu_ptx_sim_sizeofTexture(const char* name) +int gpgpu_t::gpgpu_ptx_sim_sizeofTexture(const char* name) { std::string texname(name); const struct textureReference* texref = NameToTextureMap[texname]; @@ -378,71 +376,71 @@ addr_t generic_to_global( addr_t addr ) } -void* gpgpu_t::gpgpu_ptx_sim_malloc( size_t size ) +void* gpgpu_t::gpu_malloc( size_t size ) { - unsigned long long result = g_dev_malloc; - printf("GPGPU-Sim PTX: allocating %zu bytes on GPU starting at address 0x%Lx\n", size, g_dev_malloc ); + unsigned long long result = m_dev_malloc; + printf("GPGPU-Sim PTX: allocating %zu bytes on GPU starting at address 0x%Lx\n", size, m_dev_malloc ); fflush(stdout); - g_dev_malloc += size; - if (size%64) g_dev_malloc += (64 - size%64); //align to 64 byte boundaries + m_dev_malloc += size; + if (size%64) m_dev_malloc += (64 - size%64); //align to 64 byte boundaries return(void*) result; } -void* gpgpu_t::gpgpu_ptx_sim_mallocarray( size_t size ) +void* gpgpu_t::gpu_mallocarray( size_t size ) { - unsigned long long result = g_dev_malloc; - printf("GPGPU-Sim PTX: allocating %zu bytes on GPU starting at address 0x%Lx\n", size, g_dev_malloc ); + unsigned long long result = m_dev_malloc; + printf("GPGPU-Sim PTX: allocating %zu bytes on GPU starting at address 0x%Lx\n", size, m_dev_malloc ); fflush(stdout); - g_dev_malloc += size; - if (size%64) g_dev_malloc += (64 - size%64); //align to 64 byte boundaries + m_dev_malloc += size; + if (size%64) m_dev_malloc += (64 - size%64); //align to 64 byte boundaries return(void*) result; } -void gpgpu_t::gpgpu_ptx_sim_memcpy_to_gpu( size_t dst_start_addr, const void *src, size_t count ) +void gpgpu_t::memcpy_to_gpu( size_t dst_start_addr, const void *src, size_t count ) { printf("GPGPU-Sim PTX: copying %zu bytes from CPU[0x%Lx] to GPU[0x%Lx] ... ", count, (unsigned long long) src, (unsigned long long) dst_start_addr ); fflush(stdout); char *src_data = (char*)src; for (unsigned n=0; n < count; n ++ ) - g_global_mem->write(dst_start_addr+n,1, src_data+n,NULL,NULL); + m_global_mem->write(dst_start_addr+n,1, src_data+n,NULL,NULL); printf( " done.\n"); fflush(stdout); } -void gpgpu_t::gpgpu_ptx_sim_memcpy_from_gpu( void *dst, size_t src_start_addr, size_t count ) +void gpgpu_t::memcpy_from_gpu( void *dst, size_t src_start_addr, size_t count ) { printf("GPGPU-Sim PTX: copying %zu bytes from GPU[0x%Lx] to CPU[0x%Lx] ...", count, (unsigned long long) src_start_addr, (unsigned long long) dst ); fflush(stdout); unsigned char *dst_data = (unsigned char*)dst; for (unsigned n=0; n < count; n ++ ) - g_global_mem->read(src_start_addr+n,1,dst_data+n); + m_global_mem->read(src_start_addr+n,1,dst_data+n); printf( " done.\n"); fflush(stdout); } -void gpgpu_t::gpgpu_ptx_sim_memcpy_gpu_to_gpu( size_t dst, size_t src, size_t count ) +void gpgpu_t::memcpy_gpu_to_gpu( size_t dst, size_t src, size_t count ) { printf("GPGPU-Sim PTX: copying %zu bytes from GPU[0x%Lx] to GPU[0x%Lx] ...", count, (unsigned long long) src, (unsigned long long) dst ); fflush(stdout); for (unsigned n=0; n < count; n ++ ) { unsigned char tmp; - g_global_mem->read(src+n,1,&tmp); - g_global_mem->write(dst+n,1, &tmp,NULL,NULL); + m_global_mem->read(src+n,1,&tmp); + m_global_mem->write(dst+n,1, &tmp,NULL,NULL); } printf( " done.\n"); fflush(stdout); } -void gpgpu_t::gpgpu_ptx_sim_memset( size_t dst_start_addr, int c, size_t count ) +void gpgpu_t::gpu_memset( size_t dst_start_addr, int c, size_t count ) { printf("GPGPU-Sim PTX: setting %zu bytes of memory to 0x%x starting at 0x%Lx... ", count, (unsigned char) c, (unsigned long long) dst_start_addr ); fflush(stdout); unsigned char c_value = (unsigned char)c; for (unsigned n=0; n < count; n ++ ) - g_global_mem->write(dst_start_addr+n,1,&c_value,NULL,NULL); + m_global_mem->write(dst_start_addr+n,1,&c_value,NULL,NULL); printf( " done.\n"); fflush(stdout); } @@ -1474,7 +1472,7 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t kernel, dim3 gridDim, dim3 bloc fflush(stdout); } -unsigned ptx_set_tex_cache_linesize(unsigned linesize) +unsigned gpgpu_t::ptx_set_tex_cache_linesize(unsigned linesize) { g_texcache_linesize = linesize; return 0; diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index b72a157..b095afb 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -32,11 +32,6 @@ extern void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceNam extern void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ); extern void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ); -extern void gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, const struct cudaArray* array); -extern void gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref); -extern int gpgpu_ptx_sim_sizeofTexture(const char* name); -extern const char* gpgpu_ptx_sim_findNamefromTexture(const struct textureReference* texref); -extern const struct textureReference* gpgpu_ptx_sim_accessTextureofName(const char* name); extern void read_sim_environment_variables(); extern void ptxinfo_opencl_addinfo( std::map<std::string,function_info*> &kernels ); unsigned ptx_sim_init_thread( kernel_info_t &kernel, @@ -52,7 +47,6 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, const warp_inst_t *ptx_fetch_inst( address_type pc ); const struct gpgpu_ptx_sim_kernel_info* ptx_sim_kernel_info(class function_info *kernel); void ptx_print_insn( address_type pc, FILE *fp ); -unsigned int ptx_set_tex_cache_linesize( unsigned linesize); void set_param_gpgpu_num_shaders(int num_shaders); unsigned int get_converge_point(unsigned int pc, void *thd); const char *get_ptxinfo_kname(); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 335ea59..e67c779 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -81,8 +81,7 @@ #include <stdarg.h> -unsigned g_num_ptx_inst_uid=0; -unsigned cudasim_n_tex_insn=0; +unsigned ptx_instruction::g_num_ptx_inst_uid=0; const char *g_opcode_string[NUM_OPCODES] = { #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) STR, @@ -90,13 +89,8 @@ const char *g_opcode_string[NUM_OPCODES] = { #undef OP_DEF }; -extern std::map<struct textureReference*,struct cudaArray*> TextureToArrayMap; // texture bindings -extern std::map<struct textureReference*,struct textureInfo*> TextureToInfoMap; // texture bindings -extern std::map<std::string, struct textureReference*> NameToTextureMap; - void inst_not_implemented( const ptx_instruction * pI ) ; ptx_reg_t srcOperandModifiers(ptx_reg_t opData, operand_info opInfo, operand_info dstInfo, unsigned type, ptx_thread_info *thread); -unsigned unfound_register_warned = 0; void sign_extend( ptx_reg_t &data, unsigned src_size, const operand_info &dst ); @@ -114,6 +108,7 @@ void ptx_thread_info::set_reg( const symbol *reg, const ptx_reg_t &value ) ptx_reg_t ptx_thread_info::get_reg( const symbol *reg ) { + static bool unfound_register_warned = false; assert( reg != NULL ); assert( !m_regs.empty() ); reg_map_t::iterator regs_iter = m_regs.back().find(reg); @@ -128,7 +123,7 @@ ptx_reg_t ptx_thread_info::get_reg( const symbol *reg ) if( !unfound_register_warned ) { printf("GPGPU-Sim PTX: WARNING (%s) ** reading undefined register \'%s\' (cuid:%u). Setting to 0XDEADBEEF.\n", file_loc.c_str(), name.c_str(), call_uid ); - unfound_register_warned = 1; + unfound_register_warned = true; } regs_iter = m_regs.back().find(reg); } @@ -3532,7 +3527,6 @@ float tex_linf_sampling(memory_space* mem, unsigned tex_array_base, void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { - cudasim_n_tex_insn++; unsigned dimension = pI->dimension(); const operand_info &dst = pI->dst(); //the registers to which fetched texel will be placed const operand_info &src1 = pI->src1(); //the name of the texture @@ -3546,12 +3540,10 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread ) unsigned nelem = src2.get_vect_nelem(); thread->get_vector_operand_values(src2, ptx_tex_regs, nelem); //ptx_reg should be 4 entry vector type...coordinates into texture - assert(NameToTextureMap.find(texname) != NameToTextureMap.end());//use map to find texturerefence, then use map to find pointer to array - struct textureReference* texref = NameToTextureMap[texname]; - assert(TextureToArrayMap.find(texref) != TextureToArrayMap.end()); - struct cudaArray* cuArray = TextureToArrayMap[texref]; - assert(TextureToInfoMap.find(texref) != TextureToInfoMap.end()); - struct textureInfo* texInfo = TextureToInfoMap[texref]; + gpgpu_t *gpu = thread->get_gpu(); + const struct textureReference* texref = gpu->get_texref(texname); + const struct cudaArray* cuArray = gpu->get_texarray(texref); + const struct textureInfo* texInfo = gpu->get_texinfo(texref); //assume always 2D f32 input //access array with src2 coordinates diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 612326c..07f2e36 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1009,6 +1009,7 @@ private: virtual void pre_decode(); friend class function_info; + static unsigned g_num_ptx_inst_uid; }; class param_info { @@ -1343,56 +1344,6 @@ void copy_args_into_buffer_list( const ptx_instruction * pI, void copy_buffer_list_into_frame(ptx_thread_info * thread, arg_buffer_list_t &arg_values); void copy_buffer_to_frame(ptx_thread_info * thread, const arg_buffer_t &a); -#if !defined(__CUDA_RUNTIME_API_H__) -/*******************************/ -// These declarations should be identical to those in ./../../cuda-sim-dev/libcuda/texture_types.h -enum cudaChannelFormatKind { - cudaChannelFormatKindSigned, - cudaChannelFormatKindUnsigned, - cudaChannelFormatKindFloat -}; - -struct cudaChannelFormatDesc { - int x; - int y; - int z; - int w; - enum cudaChannelFormatKind f; -}; - -struct cudaArray { - void *devPtr; - int devPtr32; - struct cudaChannelFormatDesc desc; - int width; - int height; - int size; //in bytes - unsigned dimensions; -}; - -enum cudaTextureAddressMode { - cudaAddressModeWrap, - cudaAddressModeClamp -}; - -enum cudaTextureFilterMode { - cudaFilterModePoint, - cudaFilterModeLinear -}; - -enum cudaTextureReadMode { - cudaReadModeElementType, - cudaReadModeNormalizedFloat -}; - -struct textureReference { - int normalized; - enum cudaTextureFilterMode filterMode; - enum cudaTextureAddressMode addressMode[2]; - struct cudaChannelFormatDesc channelDesc; -}; -/**********************************/ -#endif struct textureInfo { unsigned int texel_size; //size in bytes, e.g. (channelDesc.x+y+z+w)/8 diff --git a/src/debug.cc b/src/debug.cc index 928c057..236870c 100644 --- a/src/debug.cc +++ b/src/debug.cc @@ -62,7 +62,7 @@ void gpgpu_sim::gpgpu_debug() if( b.is_watchpoint() ) { unsigned addr = b.get_addr(); unsigned new_value; - g_global_mem->read(addr,4,&new_value); + m_global_mem->read(addr,4,&new_value); if( new_value != b.get_value() || g_watchpoint_hits.find(num) != g_watchpoint_hits.end() ) { printf( "GPGPU-Sim PTX DBG: watch point %u triggered (old value=%x, new value=%x)\n", num,b.get_value(),new_value ); @@ -156,8 +156,8 @@ void gpgpu_sim::gpgpu_debug() unsigned addr; sscanf(tok,"%x",&addr); unsigned value; - g_global_mem->read(addr,4,&value); - g_global_mem->set_watch(addr,next_brkpt); + m_global_mem->read(addr,4,&value); + m_global_mem->set_watch(addr,next_brkpt); breakpoints[next_brkpt++] = brk_pt(addr,value); } else if( !strcmp(tok,"l") ) { if( brk_thd == NULL ) { diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 9ffb3e7..28ec00d 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -387,7 +387,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, m_L1T = new cache_t(L1T_name,m_config->gpgpu_cache_texl1_opt, 0,no_writes, m_sid,get_shader_texture_cache_id()); m_L1C = new cache_t(L1C_name,m_config->gpgpu_cache_constl1_opt,0,no_writes, m_sid,get_shader_constant_cache_id()); m_L1I = new cache_t(L1I_name,m_config->gpgpu_cache_il1_opt, 0,no_writes, m_sid,get_shader_instruction_cache_id()); - ptx_set_tex_cache_linesize(m_L1T->get_line_sz()); + m_gpu->ptx_set_tex_cache_linesize(m_L1T->get_line_sz()); m_mshr_unit = new mshr_shader_unit(m_config); m_pdom_warp = new pdom_warp_ctx_t*[config->max_warps_per_shader]; |