summaryrefslogtreecommitdiff
path: root/libcuda/cuda_runtime_api.cc
diff options
context:
space:
mode:
Diffstat (limited to 'libcuda/cuda_runtime_api.cc')
-rw-r--r--libcuda/cuda_runtime_api.cc170
1 files changed, 154 insertions, 16 deletions
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 75a2f31..cbe8a11 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
@@ -135,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>
@@ -262,10 +264,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;
}
@@ -321,9 +332,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;
@@ -440,6 +451,10 @@ extern "C" {
* *
* *
*******************************************************************************/
+cudaError_t cudaPeekAtLastError(void)
+{
+ return g_last_cudaError;
+}
__host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size)
{
@@ -534,6 +549,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();
@@ -968,6 +999,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)
@@ -980,7 +1015,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__);
@@ -1328,7 +1364,12 @@ void extract_code_using_cuobjdump(){
printf("Running md5sum using \"%s\"\n", command);
system(command);
// 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)
+ snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass %s > %s", app_binary.c_str(), fname);
+ else
+ snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass -all %s > %s", app_binary.c_str(), fname);
bool parse_output = true;
int result = system(command);
if(result) {
@@ -1497,6 +1538,78 @@ std::list<cuobjdumpSection*> pruneSectionList(std::list<cuobjdumpSection*> cuobj
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){
+ 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){
@@ -1521,7 +1634,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;
}
@@ -1537,6 +1650,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;
@@ -1548,7 +1669,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;
}
@@ -1560,10 +1681,12 @@ 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){
@@ -1576,6 +1699,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<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin();
@@ -1584,16 +1714,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);
@@ -1619,6 +1746,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
}
@@ -1821,10 +1949,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);
@@ -2013,6 +2146,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
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 23:24:53 +0000