diff options
| author | speverel <[email protected]> | 2016-06-07 14:24:24 -0700 |
|---|---|---|
| committer | speverel <[email protected]> | 2016-06-07 14:24:24 -0700 |
| commit | 7aeadc95cc50d266f93cdb3ada1c192d9b5a1046 (patch) | |
| tree | 8445ae73b991555e4c5f1ed5b97594d597dee9a6 | |
| parent | 8ed1522aede92bcafc60f69d7e03e6d48c44a86c (diff) | |
Added support for cudaMemcpyDefault flag in cudaMemcpy. Also increased the maximum allowable memory to 2GB and the compute version to 5.2.
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 23 | ||||
| -rw-r--r-- | src/abstract_hardware_model.h | 2 |
2 files changed, 21 insertions, 4 deletions
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 8049f43..018b387 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -134,6 +134,7 @@ #include "../src/cuda-sim/ptx_parser.h" #include "../src/gpgpusim_entrypoint.h" #include "../src/stream_manager.h" +#include "../src/abstract_hardware_model.h" #include <pthread.h> #include <semaphore.h> @@ -320,9 +321,9 @@ class _cuda_device_id *GPGPUSim_Init() 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->major = 5; + prop->minor = 2; + prop->totalGlobalMem = 0x80000000 /* 2 GB */; prop->memPitch = 0; prop->maxThreadsPerBlock = 512; prop->maxThreadsDim[0] = 512; @@ -533,6 +534,22 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t cou 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 if ( kind == cudaMemcpyDefault ) { + if ((size_t)src >= GLOBAL_HEAP_START) { + if ((size_t)dst >= GLOBAL_HEAP_START) + g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) ); // device to device + else + g_stream_manager->push( stream_operation((size_t)src,dst,count,0) ); // device to host + } + else { + if ((size_t)dst >= GLOBAL_HEAP_START) + g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) ); + else { + printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported transfer: host to host\n"); + abort(); + } + } + } else { printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported cudaMemcpyKind\n"); abort(); diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index ba4ea29..b29f918 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -334,7 +334,7 @@ protected: std::deque<simt_stack_entry> m_stack; }; -#define GLOBAL_HEAP_START 0x80000000 +#define GLOBAL_HEAP_START 0x703E20000 // start allocating from this address (lower values used for allocating globals in .ptx file) #define SHARED_MEM_SIZE_MAX (64*1024) #define LOCAL_MEM_SIZE_MAX (8*1024) |