diff options
Diffstat (limited to 'libcuda/cuda_runtime_api.cc')
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 361 |
1 files changed, 283 insertions, 78 deletions
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index eed4017..edd4fb9 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -109,6 +109,7 @@ #include <stdarg.h> #include <iostream> #include <string> +#include <regex> #include <sstream> #include <fstream> #ifdef OPENGL_SUPPORT @@ -125,7 +126,9 @@ #include "host_defines.h" #include "builtin_types.h" #include "driver_types.h" +#if (CUDART_VERSION < 8000) #include "__cudaFatFormat.h" +#endif #include "../src/gpgpu-sim/gpu-sim.h" #include "../src/cuda-sim/ptx_loader.h" #include "../src/cuda-sim/cuda-sim.h" @@ -133,6 +136,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> @@ -178,6 +182,8 @@ cudaError_t g_last_cudaError = cudaSuccess; extern stream_manager *g_stream_manager; + + void register_ptx_function( const char *name, function_info *impl ) { // no longer need this @@ -260,10 +266,19 @@ struct CUctx_st { { 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; + if(s != NULL) { + function_info *f = s->get_pc(); + assert( f != NULL ); + m_kernel_lookup[hostFun] = f; + } + else { + printf("Warning: cannot find deviceFun %s\n", deviceFun); + m_kernel_lookup[hostFun] = NULL; + } + // assert( s != NULL ); + // function_info *f = s->get_pc(); + // assert( f != NULL ); + // m_kernel_lookup[hostFun] = f; } else { m_kernel_lookup[hostFun] = NULL; } @@ -311,7 +326,7 @@ private: gpgpu_ptx_sim_arg_list_t m_args; }; -class _cuda_device_id *GPGPUSim_Init() +struct _cuda_device_id *GPGPUSim_Init() { static _cuda_device_id *the_device = NULL; if( !the_device ) { @@ -319,9 +334,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; @@ -438,6 +453,10 @@ extern "C" { * * * * *******************************************************************************/ +cudaError_t cudaPeekAtLastError(void) +{ + return g_last_cudaError; +} __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) { @@ -532,6 +551,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(); @@ -966,6 +1001,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) { + return cudaStreamCreate(stream); +} + __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) { #if (CUDART_VERSION >= 3000) @@ -978,7 +1017,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__); @@ -1303,6 +1343,29 @@ std::string get_app_binary(){ return self_exe_path; } +static int get_app_cuda_version() { + int app_cuda_version = 0; + char fname[1024]; + snprintf(fname,1024,"_app_cuda_version_XXXXXX"); + int fd=mkstemp(fname); + close(fd); + std::string app_cuda_version_command = "ldd " + get_app_binary() + " | grep libcudart.so | sed 's/.*libcudart.so.\\(.*\\) =>.*/\\1/' > " + fname; + system(app_cuda_version_command.c_str()); + FILE * cmd = fopen(fname, "r"); + char buf[256]; + while (fgets(buf, sizeof(buf), cmd) != 0) { + std::cout << buf; + app_cuda_version = atoi(buf); + } + fclose(cmd); + if ( app_cuda_version == 0 ) { + printf( "Error - Cannot detect the app's CUDA version.\n" ); + exit(1); + } + return app_cuda_version; +} + + //! Call cuobjdump to extract everything (-elf -sass -ptx) /*! * This Function extract the whole PTX (for all the files) using cuobjdump @@ -1312,34 +1375,39 @@ std::string get_app_binary(){ * enabled * */ void extract_code_using_cuobjdump(){ - CUctx_st *context = GPGPUSim_Context(); - char command[1000]; + CUctx_st *context = GPGPUSim_Context(); + std::string command; - std::string app_binary = get_app_binary(); + std::string app_binary = get_app_binary(); char fname[1024]; snprintf(fname,1024,"_cuobjdump_complete_output_XXXXXX"); int fd=mkstemp(fname); close(fd); // Running cuobjdump using dynamic link to current process - snprintf(command,1000,"md5sum %s ", app_binary.c_str()); - printf("Running md5sum using \"%s\"\n", command); - system(command); + command = "md5sum " + app_binary; + printf("Running md5sum using \"%s\"\n", command.c_str()); + system(command.c_str()); // Running cuobjdump using dynamic link to current process - snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass %s > %s", app_binary.c_str(), fname); + // Needs the option '-all' to extract PTX from CDP-enabled binary + extern bool g_cdp_enabled; + if(!g_cdp_enabled) + command = "$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass " + app_binary + " > " + fname; + else + command = "$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass -all " + app_binary + " > " + fname; bool parse_output = true; - int result = system(command); + int result = system(command.c_str()); if(result) { if (context->get_device()->get_gpgpu()->get_config().experimental_lib_support() && (result == 65280)) { // Some CUDA application may exclusively use kernels provided by CUDA // libraries (e.g. CUBLAS). Skipping cuobjdump extraction from the // executable for this case. // 65280 is the return code from cuobjdump denoting the specific error (tested on CUDA 4.0/4.1/4.2) - printf("WARNING: Failed to execute: %s\n", command); + printf("WARNING: Failed to execute: %s\n", command.c_str()); printf(" Executable binary does not contain any GPU kernel.\n"); parse_output = false; } else { - printf("ERROR: Failed to execute: %s\n", command); + printf("ERROR: Failed to execute: %s\n", command.c_str()); exit(1); } } @@ -1363,7 +1431,7 @@ void extract_code_using_cuobjdump(){ cmd << "ldd " << app_binary << " | grep $CUDA_INSTALL_PATH | awk \'{print $3}\' > _tempfile_.txt"; int result = system(cmd.str().c_str()); if(result){ - std::cout << "Failed to execute: " << cmd << std::endl; + std::cout << "Failed to execute: " << cmd.str() << std::endl; exit(1); } std::ifstream libsf; @@ -1392,10 +1460,10 @@ void extract_code_using_cuobjdump(){ std::cout << "Running cuobjdump on " << line << std::endl; std::cout << "Using command: " << cmd.str() << std::endl; result = system(cmd.str().c_str()); - if(result) {printf("ERROR: Failed to execute: %s\n", command); exit(1);} + if(result) {printf("ERROR: Failed to execute: %s\n", command.c_str()); exit(1);} std::cout << "Done" << std::endl; - std::cout << "Trying to parse " << libcodfn << std::endl; + std::cout << "Trying to parse " << libcodfn.str() << std::endl; cuobjdump_in = fopen(libcodfn.str().c_str(), "r"); cuobjdump_parse(); fclose(cuobjdump_in); @@ -1457,18 +1525,22 @@ std::list<cuobjdumpSection*> pruneSectionList(std::list<cuobjdumpSection*> cuobj std::list<cuobjdumpSection*> prunedList; - //Find the highest capability (that is lower than the forces maximum) for each cubin file + //Find the highest capability (that is lower than the forced maximum) for each cubin file //and set it in cuobjdumpSectionMap. Do this only for ptx sections std::map<std::string, unsigned> cuobjdumpSectionMap; + int min_ptx_capability_found=0; for ( std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin(); iter != cuobjdumpSectionList.end(); iter++){ unsigned capability = (*iter)->getArch(); - if(dynamic_cast<cuobjdumpPTXSection*>(*iter) != NULL && - (capability <= forced_max_capability || - forced_max_capability==0)) { - if(cuobjdumpSectionMap[(*iter)->getIdentifier()] < capability) - cuobjdumpSectionMap[(*iter)->getIdentifier()] = capability; + if(dynamic_cast<cuobjdumpPTXSection*>(*iter) != NULL){ + if(capability<min_ptx_capability_found || min_ptx_capability_found==0) + min_ptx_capability_found=capability; + if (capability <= forced_max_capability || forced_max_capability==0) { + if((cuobjdumpSectionMap.find((*iter)->getIdentifier())==cuobjdumpSectionMap.end()) + || (cuobjdumpSectionMap[(*iter)->getIdentifier()] < capability)) + cuobjdumpSectionMap[(*iter)->getIdentifier()] = capability; + } } } @@ -1484,9 +1556,85 @@ std::list<cuobjdumpSection*> pruneSectionList(std::list<cuobjdumpSection*> cuobj delete *iter; } } + if(prunedList.empty()){ + printf("Error: No PTX sections found with sm capability that is lower than current forced maximum capability \n minimum ptx capability found = %u, maximum forced ptx capability = %u \n User might want to change either the forced maximum capability from gpgpusim configuration or update the compilation to generate the required PTX version\n",min_ptx_capability_found,forced_max_capability); + abort(); + } return prunedList; } +//! Merge all PTX sections that have a specific identifier into one file +std::list<cuobjdumpSection*> mergeMatchingSections(std::list<cuobjdumpSection*> cuobjdumpSectionList, std::string identifier){ + const char *ptxcode = ""; + std::list<cuobjdumpSection*>::iterator old_iter; + cuobjdumpPTXSection* old_ptxsection = NULL; + cuobjdumpPTXSection* ptxsection; + std::list<cuobjdumpSection*> mergedList; + + for ( std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin(); + iter != cuobjdumpSectionList.end(); + iter++){ + if((ptxsection=dynamic_cast<cuobjdumpPTXSection*>(*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()); + // 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 and PTX sections with non-matching identifiers + else { + mergedList.push_back(*iter); + } + } + + // Store the final PTX section + mergedList.push_back(*old_iter); + + return mergedList; +} + +//! Merge any PTX sections with matching identifiers +std::list<cuobjdumpSection*> mergeSections(std::list<cuobjdumpSection*> cuobjdumpSectionList){ + std::vector<std::string> identifier; + cuobjdumpPTXSection* ptxsection; + + // Add all identifiers present in PTX sections to a vector + for ( std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin(); + iter != cuobjdumpSectionList.end(); + iter++){ + if((ptxsection=dynamic_cast<cuobjdumpPTXSection*>(*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<std::string>::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<cuobjdumpSection*> sectionlist, const std::string identifier){ @@ -1511,7 +1659,7 @@ cuobjdumpELFSection* findELFSection(const std::string identifier){ if (sec!=NULL)return sec; sec = findELFSectionInList(libSectionList, identifier); if (sec!=NULL)return sec; - std::cout << "Cound not find " << identifier << std::endl; + std::cout << "Could not find " << identifier << std::endl; assert(0 && "Could not find the required ELF section"); return NULL; } @@ -1527,6 +1675,14 @@ cuobjdumpPTXSection* findPTXSectionInList(std::list<cuobjdumpSection*> sectionli if((ptxsection=dynamic_cast<cuobjdumpPTXSection*>(*iter)) != NULL){ if(ptxsection->getIdentifier() == identifier) return ptxsection; + else { + extern bool g_cdp_enabled; + if(g_cdp_enabled) { + printf("Warning: __cudaRegisterFatBinary needs %s, but find PTX section with %s\n", + identifier.c_str(), ptxsection->getIdentifier().c_str()); + return ptxsection; + } + } } } return NULL; @@ -1538,7 +1694,7 @@ cuobjdumpPTXSection* findPTXSection(const std::string identifier){ if (sec!=NULL)return sec; sec = findPTXSectionInList(libSectionList, identifier); if (sec!=NULL)return sec; - std::cout << "Cound not find " << identifier << std::endl; + std::cout << "Could not find " << identifier << std::endl; assert(0 && "Could not find the required PTX section"); return NULL; } @@ -1550,59 +1706,78 @@ void cuobjdumpInit(){ CUctx_st *context = GPGPUSim_Context(); extract_code_using_cuobjdump(); //extract all the output of cuobjdump to _cuobjdump_*.* cuobjdumpSectionList = pruneSectionList(cuobjdumpSectionList, context); + cuobjdumpSectionList = mergeSections(cuobjdumpSectionList); } std::map<int, std::string> fatbinmap; std::map<int, bool>fatbin_registered; +std::map<std::string, symbol_table*> name_symtab; //! Keep track of the association between filename and cubin handle -void cuobjdumpRegisterFatBinary(unsigned int handle, char* filename){ +void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename){ fatbinmap[handle] = filename; } //! Either submit PTX for simulation or convert SASS to PTXPlus and submit it void cuobjdumpParseBinary(unsigned int handle){ - if(fatbin_registered[handle]) return; - fatbin_registered[handle] = true; - CUctx_st *context = GPGPUSim_Context(); + if(fatbin_registered[handle]) return; + fatbin_registered[handle] = true; + CUctx_st *context = GPGPUSim_Context(); + std::string fname = fatbinmap[handle]; - std::string fname = fatbinmap[handle]; - cuobjdumpPTXSection* ptx = findPTXSection(fname); + if (name_symtab.find(fname) != name_symtab.end()) { + symbol_table *symtab = name_symtab[fname]; + context->add_binary(symtab, handle); + return; + } - 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) { - 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); - ptxcode = readfile(override_ptx_name); - } - if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) { - cuobjdumpELFSection* elfsection = findELFSection(ptx->getIdentifier()); - assert (elfsection!= NULL); - char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus( - ptx->getPTXfilename(), - elfsection->getELFfilename(), - elfsection->getSASSfilename()); - symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxplus_str, handle); - printf("Adding %s with cubin handle %u\n", ptx->getPTXfilename().c_str(), handle); - context->add_binary(symtab, handle); - gpgpu_ptxinfo_load_from_string( ptxcode, handle); - delete[] ptxplus_str; - } else { - symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxcode, handle); - printf("Adding %s with cubin handle %u\n", ptx->getPTXfilename().c_str(), handle); - context->add_binary(symtab, handle); - gpgpu_ptxinfo_load_from_string( ptxcode, 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()); + unsigned max_capability = 0; + for ( std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin(); + iter != cuobjdumpSectionList.end(); + iter++ ){ + unsigned capability = (*iter)->getArch(); + if (capability > max_capability) max_capability = capability; + } + printf("Using PTX version = %u\n", max_capability); + if (max_capability > 20) printf("WARNING: No guarantee that PTX will be parsed for SM version %u\n", max_capability); + + 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) { + 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); + ptxcode = readfile(override_ptx_name); + } + if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) { + cuobjdumpELFSection* elfsection = findELFSection(ptx->getIdentifier()); + assert (elfsection!= NULL); + char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus( + ptx->getPTXfilename(), + elfsection->getELFfilename(), + elfsection->getSASSfilename()); + symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxplus_str, handle); + printf("Adding %s with cubin handle %u\n", ptx->getPTXfilename().c_str(), handle); + context->add_binary(symtab, handle); + gpgpu_ptxinfo_load_from_string( ptxcode, handle ); + delete[] ptxplus_str; + } else { + symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxcode, handle); + printf("Adding %s with cubin handle %u\n", ptx->getPTXfilename().c_str(), handle); + context->add_binary(symtab, handle); + gpgpu_ptxinfo_load_from_string( ptxcode, 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 + //TODO: Remove temporarily files as per configurations } + void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) { #if (CUDART_VERSION < 2010) @@ -1616,17 +1791,29 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) if (sizeof(void*) == 4) printf("GPGPU-Sim PTX: FatBin file name extraction has not been tested on 32-bit system.\n"); - // FatBin handle from the .fatbin.c file (one of the intermediate files generated by NVCC) - typedef struct {int m; int v; const unsigned long long* d; char* f;} __fatDeviceText __attribute__ ((aligned (8))); - __fatDeviceText * fatDeviceText = (__fatDeviceText *) fatCubin; + // This code will get the CUDA version the app was compiled with. + // We need this to determine how to handle the parsing of the binary. + // Making this a runtime variable based on the app, enables GPGPU-Sim compiled + // with a newer version of CUDA to run apps compiled with older versions of + // CUDA. This is especially useful for PTXPLUS execution. + int app_cuda_version = get_app_cuda_version(); + assert( app_cuda_version == CUDART_VERSION / 1000 && "The app must be compiled with same major version as the simulator." ); + const char* filename; +#if CUDART_VERSION < 6000 + // FatBin handle from the .fatbin.c file (one of the intermediate files generated by NVCC) + typedef struct {int m; int v; const unsigned long long* d; char* f;} __fatDeviceText __attribute__ ((aligned (8))); + __fatDeviceText * fatDeviceText = (__fatDeviceText *) fatCubin; - // Extract the source code file name that generate the given FatBin. - // - Obtains the pointer to the actual fatbin structure from the FatBin handle (fatCubin). - // - An integer inside the fatbin structure contains the relative offset to the source code file name. - // - This offset differs among different CUDA and GCC versions. - char * pfatbin = (char*) fatDeviceText->d; - int offset = *((int*)(pfatbin+48)); - char * filename = (pfatbin+16+offset); + // Extract the source code file name that generate the given FatBin. + // - Obtains the pointer to the actual fatbin structure from the FatBin handle (fatCubin). + // - An integer inside the fatbin structure contains the relative offset to the source code file name. + // - This offset differs among different CUDA and GCC versions. + char * pfatbin = (char*) fatDeviceText->d; + int offset = *((int*)(pfatbin+48)); + filename = (pfatbin+16+offset); +#else + filename = "default"; +#endif // The extracted file name is associated with a fat_cubin_handle passed // into cudaLaunch(). Inside cudaLaunch(), the associated file name is @@ -1646,7 +1833,9 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) cuobjdumpRegisterFatBinary(fat_cubin_handle, filename); return (void**)fat_cubin_handle; - } else { + } +#if (CUDART_VERSION < 8000) + else { static unsigned source_num=1; unsigned long long fat_cubin_handle = next_fat_bin_handle++; __cudaFatCudaBinary *info = (__cudaFatCudaBinary *)fatCubin; @@ -1703,6 +1892,12 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) } return (void**)fat_cubin_handle; } +#else + else { + printf("ERROR ** __cudaRegisterFatBinary() needs to be updated\n"); + abort(); + } +#endif } void __cudaUnregisterFatBinary(void **fatCubinHandle) @@ -1795,10 +1990,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); @@ -1987,6 +2187,11 @@ __host__ cudaError_t CUDARTAPI cudaFuncSetCacheConfig(const char *func, enum cud context->get_device()->get_gpgpu()->set_cache_config(context->get_kernel(func)->get_name(), (FuncCache)cacheConfig); return g_last_cudaError = cudaSuccess; } + +//Jin: hack for cdp +__host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t value) { + return g_last_cudaError = cudaSuccess; +} #endif #endif |