From 46aad91327a265c2fea2cfe629cc38eadb629200 Mon Sep 17 00:00:00 2001 From: speverel Date: Thu, 2 Jun 2016 11:28:15 -0700 Subject: Added handling of .cc option for arithmetic instructions. NOTE: Only made changes to parse instructions. Carry functionality NOT fully implemented; .cc instructions function like their unmodified ueqivelents. Also modified GTX750Ti config to model L1 data cache as simply not being used for global loads (instead of not existing at all). Changed ptxinfo parsing to avoid crashing when info includes texture information. --- libcuda/cuda_runtime_api.cc | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 910bebd..e2626d2 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -109,6 +109,7 @@ #include #include #include +#include #include #include #ifdef OPENGL_SUPPORT @@ -1816,10 +1817,15 @@ void __cudaRegisterTexture( int ext ) //passes in a newly created textureReference { + std::string devStr (deviceName); + #if (CUDART_VERSION > 4020) + if (devStr.size() > 2 && devStr.data()[0] == ':' && devStr.data()[1] == ':') + devStr = devStr.replace(0, 2, ""); + #endif 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, dim, norm, ext); + gpu->gpgpu_ptx_sim_bindNameToTexture(devStr.data(), hostVar, dim, norm, ext); 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); -- cgit v1.3 From a2c48c472fc2ab159b826f3af1d1503c23eaed39 Mon Sep 17 00:00:00 2001 From: sspenst Date: Fri, 3 Jun 2016 14:10:20 -0700 Subject: Added mergeSectionList to combine any PTX files that remain after pruning into a single file. --- libcuda/cuda_runtime_api.cc | 44 ++++++++++++++++++++++++++++++++++++++++++++ libcuda/cuobjdump.y | 2 +- 2 files changed, 45 insertions(+), 1 deletion(-) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index e2626d2..8049f43 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1496,6 +1496,49 @@ std::list pruneSectionList(std::list cuobj return prunedList; } +//! Merge all remaining PTX sections in the section list into one PTX file +// NOTE: this function needs some tweaking to deal with repeated fucntion/variable declarations/definitions +std::list mergeSectionList(std::list cuobjdumpSectionList){ + char *ptxcode = ""; + std::list::iterator old_iter; + cuobjdumpPTXSection* old_ptxsection = NULL; + cuobjdumpPTXSection* ptxsection; + std::list mergedList; + + for ( std::list::iterator iter = cuobjdumpSectionList.begin(); + iter != cuobjdumpSectionList.end(); + iter++){ + if((ptxsection=dynamic_cast(*iter)) != NULL){ + // Read and remove the last PTX section + if (old_ptxsection != NULL) { + ptxcode = readfile(old_ptxsection->getPTXfilename()); + // remove ptx file? + delete *old_iter; + } + + // Append all the PTX from the last PTX section into the current PTX section + // Add 50 to ptxcode to ignore the information regarding version/target/address_size + if (strlen(ptxcode) >= 50) { + FILE *ptxfile = fopen((ptxsection->getPTXfilename()).c_str(), "a"); + fprintf(ptxfile, "%s", ptxcode + 50); + fclose(ptxfile); + } + + old_iter = iter; + old_ptxsection = ptxsection; + } + // Store all non-PTX sections + else { + mergedList.push_back(*iter); + } + } + + // Store the final PTX section + mergedList.push_front(*old_iter); + + return mergedList; +} + //! Within the section list, find the ELF section corresponding to a given identifier cuobjdumpELFSection* findELFSectionInList(std::list sectionlist, const std::string identifier){ @@ -1559,6 +1602,7 @@ void cuobjdumpInit(){ CUctx_st *context = GPGPUSim_Context(); extract_code_using_cuobjdump(); //extract all the output of cuobjdump to _cuobjdump_*.* cuobjdumpSectionList = pruneSectionList(cuobjdumpSectionList, context); + cuobjdumpSectionList = mergeSectionList(cuobjdumpSectionList); } std::map fatbinmap; diff --git a/libcuda/cuobjdump.y b/libcuda/cuobjdump.y index dce7e3d..9d61f25 100644 --- a/libcuda/cuobjdump.y +++ b/libcuda/cuobjdump.y @@ -114,7 +114,7 @@ compressedkeyword : H_COMPRESSED emptylines ptxcode : ptxcode PTXLINE {fprintf(ptxfile, "%s", $2);} | ; -elfcode : elfcode ELFLINE {printf(elffile, "%s", $2);} +elfcode : elfcode ELFLINE {fprintf(elffile, "%s", $2);} | ; sasscode : sasscode SASSLINE {fprintf(sassfile, "%s", $2);} -- cgit v1.3 From 7aeadc95cc50d266f93cdb3ada1c192d9b5a1046 Mon Sep 17 00:00:00 2001 From: speverel Date: Tue, 7 Jun 2016 14:24:24 -0700 Subject: Added support for cudaMemcpyDefault flag in cudaMemcpy. Also increased the maximum allowable memory to 2GB and the compute version to 5.2. --- libcuda/cuda_runtime_api.cc | 23 ++++++++++++++++++++--- src/abstract_hardware_model.h | 2 +- 2 files changed, 21 insertions(+), 4 deletions(-) (limited to 'libcuda') 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 #include @@ -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 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) -- cgit v1.3 From 587853a81f6fa6088b7f3d93fc8862a8b2610da7 Mon Sep 17 00:00:00 2001 From: sspenst Date: Tue, 7 Jun 2016 16:22:45 -0700 Subject: The ptx parser now recognizes the NC option for ld.global, however this option is not actually implemented --- libcuda/cuobjdump.y | 3 +-- src/cuda-sim/ptx.l | 2 ++ src/cuda-sim/ptx.y | 2 ++ src/cuda-sim/ptx_ir.cc | 2 ++ 4 files changed, 7 insertions(+), 2 deletions(-) (limited to 'libcuda') diff --git a/libcuda/cuobjdump.y b/libcuda/cuobjdump.y index 9d61f25..31760f7 100644 --- a/libcuda/cuobjdump.y +++ b/libcuda/cuobjdump.y @@ -82,8 +82,7 @@ section : PTXHEADER { snprintf(filename, 1024, "_cuobjdump_%d.elf", elfserial); elffile = fopen(filename, "w"); setCuobjdumpelffilename(filename); - } headerinfo identifier{ - } elfcode { + } headerinfo compressedkeyword identifier elfcode { fclose(elffile); snprintf(filename, 1024, "_cuobjdump_%d.sass", elfserial++); sassfile = fopen(filename, "w"); diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index 66ff48f..a44177b 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -326,6 +326,8 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; \.wb TC; return WB_OPTION; \.wt TC; return WT_OPTION; +\.nc TC; return NC_OPTION; + \.popc TC; return ATOMIC_POPC; \.and TC; return ATOMIC_AND; \.or TC; return ATOMIC_OR; diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index fca94db..4de39d1 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -193,6 +193,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. %token CV_OPTION; %token WB_OPTION; %token WT_OPTION; +%token NC_OPTION; %type function_decl_header %type function_decl @@ -449,6 +450,7 @@ option: type_spec | CV_OPTION { add_option(CV_OPTION); } | WB_OPTION { add_option(WB_OPTION); } | WT_OPTION { add_option(WT_OPTION); } + | NC_OPTION { add_option(NC_OPTION); } ; atomic_operation_spec: ATOMIC_AND { add_option(ATOMIC_AND); } diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 8f9c3d2..2eccabc 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1169,6 +1169,8 @@ ptx_instruction::ptx_instruction( int opcode, break; case EXTP_OPTION: break; + case NC_OPTION: + break; default: assert(0); break; -- cgit v1.3 From 41dbdd842f315645cc4cd05e3b6263b982d1be2e Mon Sep 17 00:00:00 2001 From: sspenst Date: Tue, 7 Jun 2016 16:43:09 -0700 Subject: Updated the PTX section merging to be able to deal with multiple identifiers --- libcuda/cuda_runtime_api.cc | 43 ++++++++++++++++++++++++++++++++++++------- 1 file changed, 36 insertions(+), 7 deletions(-) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 018b387..862e2b9 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1513,9 +1513,8 @@ std::list pruneSectionList(std::list cuobj return prunedList; } -//! Merge all remaining PTX sections in the section list into one PTX file -// NOTE: this function needs some tweaking to deal with repeated fucntion/variable declarations/definitions -std::list mergeSectionList(std::list cuobjdumpSectionList){ +//! Merge all PTX sections that have a specific identifier into one file +std::list mergeMatchingSections(std::list cuobjdumpSectionList, std::string identifier){ char *ptxcode = ""; std::list::iterator old_iter; cuobjdumpPTXSection* old_ptxsection = NULL; @@ -1525,7 +1524,8 @@ std::list mergeSectionList(std::list cuobj for ( std::list::iterator iter = cuobjdumpSectionList.begin(); iter != cuobjdumpSectionList.end(); iter++){ - if((ptxsection=dynamic_cast(*iter)) != NULL){ + if((ptxsection=dynamic_cast(*iter)) != NULL && + strcmp(ptxsection->getIdentifier().c_str(), identifier.c_str()) == 0){ // Read and remove the last PTX section if (old_ptxsection != NULL) { ptxcode = readfile(old_ptxsection->getPTXfilename()); @@ -1544,18 +1544,47 @@ std::list mergeSectionList(std::list cuobj old_iter = iter; old_ptxsection = ptxsection; } - // Store all non-PTX sections + // Store all non-PTX sections and PTX sections with non-matching identifiers else { mergedList.push_back(*iter); } } // Store the final PTX section - mergedList.push_front(*old_iter); + mergedList.push_back(*old_iter); return mergedList; } +//! Merge any PTX sections with matching identifiers +std::list mergeSections(std::list cuobjdumpSectionList){ + std::vector identifier; + cuobjdumpPTXSection* ptxsection; + + // Add all identifiers present in PTX sections to a vector + for ( std::list::iterator iter = cuobjdumpSectionList.begin(); + iter != cuobjdumpSectionList.end(); + iter++){ + if((ptxsection=dynamic_cast(*iter)) != NULL){ + std::string current_id = ptxsection->getIdentifier(); + + // If we haven't yet seen a given identifier, add it to the vector + if (std::find(identifier.begin(), identifier.end(), current_id) == identifier.end()) { + identifier.push_back(current_id); + } + } + } + + // Call mergeMatchingSections on all identifiers in the vector + for ( std::vector::iterator iter = identifier.begin(); + iter != identifier.end(); + iter++) { + cuobjdumpSectionList = mergeMatchingSections(cuobjdumpSectionList, *iter); + } + + return cuobjdumpSectionList; +} + //! Within the section list, find the ELF section corresponding to a given identifier cuobjdumpELFSection* findELFSectionInList(std::list sectionlist, const std::string identifier){ @@ -1619,7 +1648,7 @@ void cuobjdumpInit(){ CUctx_st *context = GPGPUSim_Context(); extract_code_using_cuobjdump(); //extract all the output of cuobjdump to _cuobjdump_*.* cuobjdumpSectionList = pruneSectionList(cuobjdumpSectionList, context); - cuobjdumpSectionList = mergeSectionList(cuobjdumpSectionList); + cuobjdumpSectionList = mergeSections(cuobjdumpSectionList); } std::map fatbinmap; -- cgit v1.3 From 5e9b6521e625d30f9f3efbb292fd22bdab5c884a Mon Sep 17 00:00:00 2001 From: sspenst Date: Thu, 9 Jun 2016 13:38:13 -0700 Subject: Added a basic implementation of cudaStreamCreateWithFlags --- libcuda/cuda_runtime_api.cc | 4 ++++ 1 file changed, 4 insertions(+) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 862e2b9..77456d4 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -984,6 +984,10 @@ __host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream) return g_last_cudaError = cudaSuccess; } +__host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *stream, unsigned int flags) { + cudaStreamCreate(stream); +} + __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) { #if (CUDART_VERSION >= 3000) -- cgit v1.3 From adabbc7070704d1be79465786d59e170030c7b0d Mon Sep 17 00:00:00 2001 From: sspenst Date: Thu, 9 Jun 2016 14:23:28 -0700 Subject: Forgot to add a return value --- libcuda/cuda_runtime_api.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 77456d4..4cc002d 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -985,7 +985,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream) } __host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *stream, unsigned int flags) { - cudaStreamCreate(stream); + return cudaStreamCreate(stream); } __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) -- cgit v1.3 From ac11fd57b93e22ee5a50ebf8f6d6b4d6dadbe3cb Mon Sep 17 00:00:00 2001 From: speverel Date: Thu, 9 Jun 2016 14:58:38 -0700 Subject: Modified runtime API to support synchronization with the null stream. --- libcuda/cuda_runtime_api.cc | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 77456d4..350ba0b 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -440,6 +440,10 @@ extern "C" { * * * * *******************************************************************************/ +cudaError_t cudaPeekAtLastError(void) +{ + return g_last_cudaError; +} __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) { @@ -1000,7 +1004,8 @@ __host__ cudaError_t CUDARTAPI cudaStreamSynchronize(cudaStream_t stream) { #if (CUDART_VERSION >= 3000) if( stream == NULL ) - return g_last_cudaError = cudaErrorInvalidResourceHandle; + synchronize(); + return g_last_cudaError = cudaSuccess; stream->synchronize(); #else printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); -- cgit v1.3 From 7356669b66425ee73ba61466f37fe25b886a3641 Mon Sep 17 00:00:00 2001 From: sspenst Date: Fri, 10 Jun 2016 14:32:00 -0700 Subject: Vastly improved runtime by relating symbol tables to file names --- libcuda/cuda_runtime_api.cc | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 5558c4b..e8a0e91 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1662,6 +1662,7 @@ void cuobjdumpInit(){ std::map fatbinmap; std::mapfatbin_registered; +std::map name_symtab; //! Keep track of the association between filename and cubin handle void cuobjdumpRegisterFatBinary(unsigned int handle, char* filename){ @@ -1674,6 +1675,13 @@ void cuobjdumpParseBinary(unsigned int handle){ if(fatbin_registered[handle]) return; fatbin_registered[handle] = true; CUctx_st *context = GPGPUSim_Context(); + std::string fname = fatbinmap[handle]; + + if (name_symtab.find(fname) != name_symtab.end()) { + symbol_table *symtab = name_symtab[fname]; + context->add_binary(symtab, handle); + return; + } unsigned max_capability = 0; for ( std::list::iterator iter = cuobjdumpSectionList.begin(); @@ -1682,16 +1690,13 @@ void cuobjdumpParseBinary(unsigned int handle){ unsigned capability = (*iter)->getArch(); if (capability > max_capability) max_capability = capability; } - if (max_capability > 20) printf("WARNING: No guarantee that PTX will be parsed for SM version %u\n", max_capability); - std::string fname = fatbinmap[handle]; cuobjdumpPTXSection* ptx = findPTXSection(fname); - symbol_table *symtab; char *ptxcode; const char *override_ptx_name = getenv("PTX_SIM_KERNELFILE"); - if (override_ptx_name == NULL or getenv("PTX_SIM_USE_PTX_FILE") == NULL) { + if (override_ptx_name == NULL or getenv("PTX_SIM_USE_PTX_FILE") == NULL) { ptxcode = readfile(ptx->getPTXfilename()); } else { printf("GPGPU-Sim PTX: overriding embedded ptx with '%s' (PTX_SIM_USE_PTX_FILE is set)\n", override_ptx_name); @@ -1717,6 +1722,7 @@ void cuobjdumpParseBinary(unsigned int handle){ } load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); + name_symtab[fname] = symtab; //TODO: Remove temporarily files as per configurations } -- cgit v1.3