aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--Jenkinsfile120
-rw-r--r--libcuda/cuda_runtime_api.cc369
-rw-r--r--setup_environment2
-rw-r--r--src/abstract_hardware_model.h3
-rw-r--r--src/cuda-sim/cuda-sim.cc16
-rw-r--r--src/cuda-sim/ptx.l2
-rw-r--r--src/cuda-sim/ptx_ir.cc32
-rw-r--r--src/cuda-sim/ptx_ir.h5
-rw-r--r--src/cuda-sim/ptx_loader.cc194
-rw-r--r--src/cuda-sim/ptx_loader.h1
10 files changed, 597 insertions, 147 deletions
diff --git a/Jenkinsfile b/Jenkinsfile
new file mode 100644
index 0000000..806231e
--- /dev/null
+++ b/Jenkinsfile
@@ -0,0 +1,120 @@
+pipeline {
+ agent {
+ label "purdue-cluster"
+ }
+
+ options {
+ disableConcurrentBuilds()
+ }
+
+ stages {
+ stage('simulator-build') {
+ steps {
+ parallel "4.2": {
+ sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/4.2_env_setup.sh &&\
+ source `pwd`/setup_environment &&\
+ make -j'
+ }, "9.1" : {
+ sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
+ source `pwd`/setup_environment &&\
+ make -j'
+ }
+ }
+ }
+ stage('simulations-build'){
+ steps{
+ sh 'rm -rf gpgpu-sim_simulations'
+ sh 'git clone [email protected]:TimRogersGroup/gpgpu-sim_simulations.git && \
+ cd gpgpu-sim_simulations && \
+ git checkout purdue-cluster && \
+ git pull && \
+ rm -r ./benchmarks/data_dirs && ln -s /home/tgrogers-raid/a/common/data_dirs benchmarks/'
+ sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/4.2_env_setup.sh &&\
+ source `pwd`/setup_environment &&\
+ cd gpgpu-sim_simulations && \
+ source ./benchmarks/src/setup_environment && \
+ make -j -C ./benchmarks/src rodinia_2.0-ft sdk-4.2 && \
+ make -C ./benchmarks/src data'
+ sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
+ source `pwd`/setup_environment &&\
+ cd gpgpu-sim_simulations && \
+ source ./benchmarks/src/setup_environment && \
+ make -j -C ./benchmarks/src/ rodinia_2.0-ft sdk-4.2 && \
+ make -C ./benchmarks/src data'
+ }
+ }
+ stage('regress'){
+ steps {
+
+ parallel "4.2-rodinia": {
+ sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/4.2_env_setup.sh &&\
+ source `pwd`/setup_environment &&\
+ ./gpgpu-sim_simulations/util/job_launching/run_simulations.py -B rodinia_2.0-ft -C GTX480,GTX480-PTXPLUS -N regress-$$ && \
+ PLOTDIR="jenkins/${JOB_NAME}/${BUILD_NUMBER}/4.2-rodinia" && ssh [email protected] mkdir -p /home/dynamo/a/tgrogers/website/gpgpu-sim-plots/$PLOTDIR && \
+ ./gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -v -N regress-$$ -s stats-$$.csv && \
+ ./gpgpu-sim_simulations/util/plotting/plot-get-stats.py -c stats-$$.csv -p [email protected]:~/website/gpgpu-sim-plots/$PLOTDIR -w https://engineering.purdue.edu/tgrogers/gpgpu-sim-plots/$PLOTDIR -n $PLOTDIR'
+ }, "9.1-rodinia": {
+ sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
+ source `pwd`/setup_environment &&\
+ ./gpgpu-sim_simulations/util/job_launching/run_simulations.py -B rodinia_2.0-ft -C GTX1080Ti -N regress-$$ && \
+ PLOTDIR="jenkins/${JOB_NAME}/${BUILD_NUMBER}/9.1-rodinia" && ssh [email protected] mkdir -p /home/dynamo/a/tgrogers/website/gpgpu-sim-plots/$PLOTDIR && \
+ ./gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -v -s stats-$$.csv -N regress-$$ && \
+ ./gpgpu-sim_simulations/util/plotting/plot-get-stats.py -c stats-$$.csv -p [email protected]:~/website/gpgpu-sim-plots/$PLOTDIR -w https://engineering.purdue.edu/tgrogers/gpgpu-sim-plots/$PLOTDIR -n $PLOTDIR'
+ }, "4.2-sdk-4.2": {
+ sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/4.2_env_setup.sh &&\
+ source `pwd`/setup_environment &&\
+ ./gpgpu-sim_simulations/util/job_launching/run_simulations.py -B sdk-4.2 -C GTX480 -N regress-$$ && \
+ PLOTDIR="jenkins/${JOB_NAME}/${BUILD_NUMBER}/4.2-sdk-4.2" && ssh [email protected] mkdir -p /home/dynamo/a/tgrogers/website/gpgpu-sim-plots/$PLOTDIR && \
+ ./gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -v -N regress-$$ -s stats-$$.csv && \
+ ./gpgpu-sim_simulations/util/plotting/plot-get-stats.py -c stats-$$.csv -p [email protected]:~/website/gpgpu-sim-plots/$PLOTDIR -w https://engineering.purdue.edu/tgrogers/gpgpu-sim-plots/$PLOTDIR -n $PLOTDIR'
+ }, "9.1-sdk-4.2": {
+ sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
+ source `pwd`/setup_environment &&\
+ ./gpgpu-sim_simulations/util/job_launching/run_simulations.py -B sdk-4.2 -C GTX1080Ti -N regress-$$ && \
+ PLOTDIR="jenkins/${JOB_NAME}/${BUILD_NUMBER}/9.1-sdk-4.2" && ssh [email protected] mkdir -p /home/dynamo/a/tgrogers/website/gpgpu-sim-plots/$PLOTDIR && \
+ ./gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -v -N regress-$$ -s stats-$$.csv && \
+ ./gpgpu-sim_simulations/util/plotting/plot-get-stats.py -c stats-$$.csv -p [email protected]:~/website/gpgpu-sim-plots/$PLOTDIR -w https://engineering.purdue.edu/tgrogers/gpgpu-sim-plots/$PLOTDIR -n $PLOTDIR'
+ }
+ }
+ }
+ stage('4.2-correlate'){
+ steps {
+ sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/4.2_env_setup.sh &&\
+ source `pwd`/setup_environment &&\
+ ./gpgpu-sim_simulations/util/job_launching/get_stats.py -R -K -k -B rodinia_2.0-ft -C GTX480,GTX480-PTXPLUS > stats-4.2.csv && \
+ PLOTDIR="jenkins/${JOB_NAME}/${BUILD_NUMBER}/correlate-4.2" && ssh [email protected] mkdir -p /home/dynamo/a/tgrogers/website/gpgpu-sim-plots/$PLOTDIR && \
+ sh ./gpgpu-sim_simulations/run_hw/get_hw_data.sh && rm -rf ./gpgpu-sim_simulations/util/plotting/correl-html &&\
+ ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c stats-4.2.csv -H ./gpgpu-sim_simulations/run_hw/ &&\
+ scp ./gpgpu-sim_simulations/util/plotting/correl-html/* [email protected]:/home/dynamo/a/tgrogers/website/gpgpu-sim-plots/$PLOTDIR'
+ }
+ }
+ stage('9.1-correlate'){
+ steps {
+ sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
+ source `pwd`/setup_environment &&\
+ ./gpgpu-sim_simulations/util/job_launching/get_stats.py -R -K -k -B sdk-4.2,rodinia_2.0-ft -C GTX1080Ti > stats-9.1.csv && \
+ PLOTDIR="jenkins/${JOB_NAME}/${BUILD_NUMBER}/correlate-9.1" && ssh [email protected] mkdir -p /home/dynamo/a/tgrogers/website/gpgpu-sim-plots/$PLOTDIR && \
+ sh ./gpgpu-sim_simulations/run_hw/get_hw_data.sh && rm -rf ./gpgpu-sim_simulations/util/plotting/correl-html &&\
+ ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c stats-9.1.csv -H ./gpgpu-sim_simulations/run_hw/ &&\
+ scp ./gpgpu-sim_simulations/util/plotting/correl-html/* [email protected]:/home/dynamo/a/tgrogers/website/gpgpu-sim-plots/$PLOTDIR'
+ }
+ }
+ }
+ post {
+ success {
+ emailext body: "See ${BUILD_URL}",
+ recipientProviders: [[$class: 'CulpritsRecipientProvider'],
+ [$class: 'RequesterRecipientProvider']],
+ subject: "[AALP Jenkins] Build #${BUILD_NUMBER} - Success!",
+ to: '[email protected]'
+ }
+ failure {
+ emailext body: "See ${BUILD_URL}",
+ recipientProviders: [[$class: 'CulpritsRecipientProvider'],
+ [$class: 'RequesterRecipientProvider']],
+ subject: "[AALP Jenkins] Build #${BUILD_NUMBER} - ${currentBuild.result}",
+ to: '[email protected]'
+ }
+ }
+}
+
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index cbe8a11..eda8d8e 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -145,6 +145,10 @@
#include <mach-o/dyld.h>
#endif
+std::map<void *,void **> pinned_memory; //support for pinned memories added
+std::map<void *, size_t> pinned_memory_size;
+int no_of_ptx=0;
+
extern void synchronize();
extern void exit_simulation();
@@ -474,6 +478,8 @@ __host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size)
GPGPUSim_Context();
*ptr = malloc(size);
if ( *ptr ) {
+ //track pinned memory size allocated in the host so that same amount of memory is also allocated in GPU.
+ pinned_memory_size[*ptr]=size;
return g_last_cudaError = cudaSuccess;
} else {
return g_last_cudaError = cudaErrorMemoryAllocation;
@@ -764,6 +770,16 @@ __host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count)
return g_last_cudaError = cudaSuccess;
}
+//memset operation is done but i think its not async?
+__host__ cudaError_t CUDARTAPI cudaMemsetAsync(void *mem, int c, size_t count, cudaStream_t stream=0)
+{
+ printf("GPGPU-Sim PTX: WARNING: Asynchronous memset not supported (%s)\n", __my_func__);
+ CUctx_st *context = GPGPUSim_Context();
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ gpu->gpu_memset((size_t)mem, c, count);
+ return g_last_cudaError = cudaSuccess;
+}
+
__host__ cudaError_t CUDARTAPI cudaMemset2D(void *mem, size_t pitch, int c, size_t width, size_t height)
{
cuda_not_implemented(__my_func__,__LINE__);
@@ -816,6 +832,58 @@ __host__ cudaError_t CUDARTAPI cudaGetDeviceProperties(struct cudaDeviceProp *pr
}
}
+#if (CUDART_VERSION > 5000)
+__host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device)
+{
+ const struct cudaDeviceProp *prop;
+ _cuda_device_id *dev = GPGPUSim_Init();
+ if (device <= dev->num_devices() ) {
+ prop = dev->get_prop();
+ switch (attr) {
+ case 5:
+ *value= prop->maxGridSize[0];
+ break;
+ case 6:
+ *value= prop->maxGridSize[1];
+ break;
+ case 7:
+ *value= prop->maxGridSize[2];
+ break;
+ case 10:
+ *value= prop->warpSize;
+ break;
+ case 12:
+ *value= prop->regsPerBlock;
+ break;
+ case 14:
+ *value= prop->textureAlignment ;
+ break;
+ case 16:
+ *value= prop->multiProcessorCount ;
+ break;
+ case 39:
+ *value= dev->get_gpgpu()->threads_per_core();
+ break;
+ case 75:
+ *value= 8 ;
+ break;
+ case 76:
+ *value= 3 ;
+ break;
+ case 78:
+ *value= 0 ; //TODO: as of now, we dont support stream priorities.
+ break;
+ default:
+ printf("ERROR: implement the attribute numbered %d \n",attr);
+ abort();
+ }
+ return g_last_cudaError = cudaSuccess;
+ } else {
+ return g_last_cudaError = cudaErrorInvalidDevice;
+ }
+}
+#endif
+
__host__ cudaError_t CUDARTAPI cudaChooseDevice(int *device, const struct cudaDeviceProp *prop)
{
_cuda_device_id *dev = GPGPUSim_Init();
@@ -969,7 +1037,16 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun )
printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun,
g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 );
kernel_info_t *grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context);
+ //do dynamic PDOM analysis for performance simulation scenario
std::string kname = grid->name();
+ function_info *kernel_func_info = grid->entry();
+ if (kernel_func_info->is_pdom_set()) {
+ printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", kname.c_str() );
+ } else {
+ printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", kname.c_str() );
+ kernel_func_info->do_pdom();
+ kernel_func_info->set_pdom();
+ }
dim3 gridDim = config.grid_dim();
dim3 blockDim = config.block_dim();
printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n",
@@ -999,8 +1076,17 @@ __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);
+//TODO: introduce priorities
+__host__ cudaError_t CUDARTAPI cudaStreamCreateWithPriority(cudaStream_t *stream, unsigned int flags, int priority) {
+ return cudaStreamCreate(stream);
+}
+
+__host__ cudaError_t CUDARTAPI cudaDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPriority) {
+ return cudaSuccess;
+}
+
+__host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *pStream, unsigned int flags) {
+ return cudaStreamCreate(pStream);
}
__host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
@@ -1341,6 +1427,26 @@ std::string get_app_binary(){
return self_exe_path;
}
+//above func gives abs path whereas this give just the name of application.
+char* get_app_binary_name(std::string abs_path){
+ char *self_exe_path;
+#ifdef __APPLE__
+ //TODO: get apple device and check the result.
+ printf("WARNING: not tested for Apple-mac devices \n");
+ abort();
+#else
+ char* buf = strdup(abs_path.c_str());
+ char *token = strtok(buf, "/");
+ while(token !=NULL){
+ self_exe_path = token;
+ token = strtok(NULL,"/");
+ }
+#endif
+ self_exe_path = strtok(self_exe_path, ".");
+ printf("self exe links to: %s\n", self_exe_path);
+ return self_exe_path;
+}
+
//! Call cuobjdump to extract everything (-elf -sass -ptx)
/*!
* This Function extract the whole PTX (for all the files) using cuobjdump
@@ -1350,105 +1456,154 @@ std::string get_app_binary(){
* enabled
* */
void extract_code_using_cuobjdump(){
- CUctx_st *context = GPGPUSim_Context();
- char command[1000];
+ CUctx_st *context = GPGPUSim_Context();
+ unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
- std::string app_binary = get_app_binary();
+ //prevent the dumping by cuobjdump everytime we execute the code!
+ const char *override_cuobjdump = getenv("CUOBJDUMP_SIM_FILE");
+ char command[1000], ptx_file[1000];
+ std::string app_binary = get_app_binary();
+ //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);
+ // Running cuobjdump using dynamic link to current process
+ // Needs the option '-all' to extract PTX from CDP-enabled binary
+ extern bool g_cdp_enabled;
- char fname[1024];
+ //dump ptx for all individial ptx files into sepearte files which is later used by ptxas.
+ char fname2[1024];
+ snprintf(fname2,1024,"_cuobjdump_list_ptx_XXXXXX");
+ int fd2=mkstemp(fname2);
+ close(fd2);
+ snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -lptx -arch=sm_%u %s > %s", forced_max_capability, app_binary.c_str(), fname2);
+ int result = system(command);
+ if( result != 0 ) {
+ printf("WARNING: Failed to execute cuobjdump to get list of ptx files \n");
+ exit(0);
+ } else {
+ /*
+ as we got list of ptx files, we need to extract one by one into seperate files so that ptxas can understand it.
+ In this way, the duplicate definitions in a single embedded file can be prevented.
+ No of lines in the file is equal to no of ptx fileis available.
+ */
+ FILE *fp = fopen(fname2,"r");
+ if (fp==NULL) {
+ printf("WARNING: cuobjdump file error! Could not open file %s \n", fname2);
+ exit(0);
+ } else {
+ for (char c = getc(fp); c != EOF; c = getc(fp))
+ if (c == '\n')
+ no_of_ptx = no_of_ptx + 1;
+ fclose(fp);
+ }
+ }
+ if(!g_cdp_enabled) {
+ //based on the list above, dump ptx files individually. Format of dumped ptx file is prog_name.unique_no.sm_<>.ptx
+ for (int index=1; index<= no_of_ptx; index++){
+ snprintf(ptx_file, 1000, "%s.%d.sm_%u.ptx", get_app_binary_name(app_binary), index, forced_max_capability);
+ printf("Extracting specific PTX file named %s \n",ptx_file);
+ snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -arch=sm_%u -xptx %s %s", forced_max_capability, ptx_file, app_binary.c_str());
+ if (system(command)!=0) {
+ printf("ERROR: command: %s failed \n",command);
+ exit(0);
+ }
+ }
+ }
+ //TODO: redundant to dump twice. how can it be prevented?
+ //dump only for specific arch
+ char fname[1024];
+ if ((override_cuobjdump == NULL) || (strlen(override_cuobjdump)==0)) {
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);
- // Running cuobjdump using dynamic link to current process
- // 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);
+ snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass -arch=sm_%u %s > %s", forced_max_capability, 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);
+ 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);
+ result = system(command);
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(" Executable binary does not contain any GPU kernel.\n");
- parse_output = false;
- } else {
- printf("ERROR: Failed to execute: %s\n", command);
- exit(1);
- }
- }
+ 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(" Executable binary does not contain any GPU kernel.\n");
+ parse_output = false;
+ } else {
+ printf("ERROR: Failed to execute: %s\n", command);
+ exit(1);
+ }
+ }
- if (parse_output) {
- printf("Parsing file %s\n", fname);
- cuobjdump_in = fopen(fname, "r");
+ if (parse_output) {
+ printf("Parsing file %s\n", fname);
+ cuobjdump_in = fopen(fname, "r");
- cuobjdump_parse();
- fclose(cuobjdump_in);
- printf("Done parsing!!!\n");
- } else {
- printf("Parsing skipped for %s\n", fname);
- }
+ cuobjdump_parse();
+ fclose(cuobjdump_in);
+ printf("Done parsing!!!\n");
+ } else {
+ printf("Parsing skipped for %s\n", fname);
+ }
- if (context->get_device()->get_gpgpu()->get_config().experimental_lib_support()){
- //Experimental library support
- //Currently only for cufft
+ if (context->get_device()->get_gpgpu()->get_config().experimental_lib_support()){
+ //Experimental library support
+ //Currently only for cufft
- std::stringstream cmd;
- 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;
- exit(1);
- }
- std::ifstream libsf;
- libsf.open("_tempfile_.txt");
- if(!libsf.is_open()) {
- std::cout << "Failed to open: _tempfile_.txt" << std::endl;
- exit(1);
- }
+ std::stringstream cmd;
+ 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;
+ exit(1);
+ }
+ std::ifstream libsf;
+ libsf.open("_tempfile_.txt");
+ if(!libsf.is_open()) {
+ std::cout << "Failed to open: _tempfile_.txt" << std::endl;
+ exit(1);
+ }
- //Save the original section list
- std::list<cuobjdumpSection*> tmpsl = cuobjdumpSectionList;
- cuobjdumpSectionList.clear();
+ //Save the original section list
+ std::list<cuobjdumpSection*> tmpsl = cuobjdumpSectionList;
+ cuobjdumpSectionList.clear();
- std::string line;
- std::getline(libsf, line);
- std::cout << "DOING: " << line << std::endl;
- int cnt=1;
- while(libsf.good()){
- std::stringstream libcodfn;
- libcodfn << "_cuobjdump_complete_lib_" << cnt << "_";
- cmd.str(""); //resetting
- cmd << "$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass ";
- cmd << line;
- cmd << " > ";
- cmd << libcodfn.str();
- 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);}
- std::cout << "Done" << std::endl;
+ std::string line;
+ std::getline(libsf, line);
+ std::cout << "DOING: " << line << std::endl;
+ int cnt=1;
+ while(libsf.good()){
+ std::stringstream libcodfn;
+ libcodfn << "_cuobjdump_complete_lib_" << cnt << "_";
+ cmd.str(""); //resetting
+ cmd << "$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass ";
+ cmd << line;
+ cmd << " > ";
+ cmd << libcodfn.str();
+ 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);}
+ std::cout << "Done" << std::endl;
- std::cout << "Trying to parse " << libcodfn << std::endl;
- cuobjdump_in = fopen(libcodfn.str().c_str(), "r");
- cuobjdump_parse();
- fclose(cuobjdump_in);
- std::getline(libsf, line);
- }
- libSectionList = cuobjdumpSectionList;
+ std::cout << "Trying to parse " << libcodfn << std::endl;
+ cuobjdump_in = fopen(libcodfn.str().c_str(), "r");
+ cuobjdump_parse();
+ fclose(cuobjdump_in);
+ std::getline(libsf, line);
+ }
+ libSectionList = cuobjdumpSectionList;
- //Restore the original section list
- cuobjdumpSectionList = tmpsl;
- }
+ //Restore the original section list
+ cuobjdumpSectionList = tmpsl;
+ }
+ } else {
+ printf("GPGPU-Sim PTX: overriding cuobjdump with '%s' (CUOBJDUMP_SIM_FILE is set)\n", override_cuobjdump);
+ snprintf(fname,1024, "%s",override_cuobjdump);
+ }
}
//! Read file into char*
@@ -1680,8 +1835,11 @@ cuobjdumpPTXSection* findPTXSection(const std::string identifier){
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);
+ const char* pre_load = getenv("CUOBJDUMP_SIM_FILE");
+ if (pre_load ==NULL || strlen(pre_load)==0){
+ cuobjdumpSectionList = pruneSectionList(cuobjdumpSectionList, context);
+ cuobjdumpSectionList = mergeSections(cuobjdumpSectionList);
+ }
}
std::map<int, std::string> fatbinmap;
@@ -1707,6 +1865,8 @@ void cuobjdumpParseBinary(unsigned int handle){
return;
}
+ //Why to search for capability value if we can directly find it in device info?
+ #if 0
unsigned max_capability = 0;
for ( std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin();
iter != cuobjdumpSectionList.end();
@@ -1715,12 +1875,17 @@ void cuobjdumpParseBinary(unsigned int handle){
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);
+ #endif
+ unsigned max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
- cuobjdumpPTXSection* ptx = findPTXSection(fname);
+ cuobjdumpPTXSection* ptx = NULL;
+ const char* pre_load = getenv("CUOBJDUMP_SIM_FILE");
+ if(pre_load==NULL || strlen(pre_load)==0)
+ 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 or strlen(getenv("PTX_SIM_USE_PTX_FILE"))==0) {
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);
@@ -1740,7 +1905,8 @@ void cuobjdumpParseBinary(unsigned int 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);
+ //if CUOBJDUMP_SIM_FILE is not set, ptx is NULL. So comment below.
+ //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, max_capability );
}
@@ -1857,6 +2023,7 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
return (void**)fat_cubin_handle;
}
#endif
+ return 0;
}
void __cudaUnregisterFatBinary(void **fatCubinHandle)
@@ -2072,6 +2239,9 @@ cudaError_t cudaGLUnregisterBufferObject(GLuint bufferObj)
cudaError_t CUDARTAPI cudaHostAlloc(void **pHost, size_t bytes, unsigned int flags)
{
*pHost = malloc(bytes);
+ //need to track the size allocated so that cudaHostGetDevicePointer() can function properly.
+ //TODO: vary this function behavior based on flags value (following nvidia documentation)
+ pinned_memory_size[*pHost]=bytes;
if( *pHost )
return g_last_cudaError = cudaSuccess;
else
@@ -2080,8 +2250,25 @@ cudaError_t CUDARTAPI cudaHostAlloc(void **pHost, size_t bytes, unsigned int fl
cudaError_t CUDARTAPI cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ //only cpu memory allocation happens in cudaHostAlloc. Linking with device pointer to pinned memory happens here.
+ //TODO: once kernel is executed, the contents in global pointer of GPU must be copied back to CPU host pointer!
+ flags=0;
+ CUctx_st* context = GPGPUSim_Context();
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ std::map<void *, size_t>::const_iterator i = pinned_memory_size.find(pHost);
+ assert(i != pinned_memory_size.end());
+ size_t size = i->second;
+ *pDevice = gpu->gpu_malloc(size);
+ if(g_debug_execution >= 3)
+ printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *pDevice);
+ if ( *pDevice ) {
+ pinned_memory[pHost]=pDevice;
+ //Copy contents in cpu to gpu
+ gpu->memcpy_to_gpu((size_t)*pDevice,pHost,size);
+ return g_last_cudaError = cudaSuccess;
+ } else {
+ return g_last_cudaError = cudaErrorMemoryAllocation;
+ }
}
cudaError_t CUDARTAPI cudaSetValidDevices(int *device_arr, int len)
diff --git a/setup_environment b/setup_environment
index 96cc362..f6a16c5 100644
--- a/setup_environment
+++ b/setup_environment
@@ -1,6 +1,6 @@
# see README before running this
-ps -p $$ | awk '/bash/ || / sh/ || /zsh/ {exit 1;}' && echo "ERROR ** source setup_environment must be run in a bash, zsh or sh shell; see README" && exit
+ps -p $$ | awk '/bash/ || / sh/ || /zsh/ {exit 1;}' && echo "WARNING ** source setup_environment must be run in a bash, zsh or sh shell; see README"
export GPGPUSIM_SETUP_ENVIRONMENT_WAS_RUN=
export GPGPUSIM_ROOT="$( cd "$( dirname "$BASH_SOURCE" )" && pwd )"
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index aaa4b00..67b36c7 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -182,6 +182,9 @@ void increment_x_then_y_then_z( dim3 &i, const dim3 &bound);
class stream_manager;
struct CUstream_st;
extern stream_manager * g_stream_manager;
+//support for pinned memories added
+extern std::map<void *,void **> pinned_memory;
+extern std::map<void *, size_t> pinned_memory_size;
class kernel_info_t {
public:
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index d4ace76..39a04dd 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -212,7 +212,9 @@ void function_info::ptx_assemble()
m_start_PC = PC;
addr_t n=0; // offset in m_instr_mem
- s_g_pc_to_insn.reserve(s_g_pc_to_insn.size() + MAX_INST_SIZE*m_instructions.size());
+ //Why s_g_pc_to_insn.size() is needed to reserve additional memory for insts? reserve is cumulative.
+ //s_g_pc_to_insn.reserve(s_g_pc_to_insn.size() + MAX_INST_SIZE*m_instructions.size());
+ s_g_pc_to_insn.reserve(MAX_INST_SIZE*m_instructions.size());
for ( i=m_instructions.begin(); i != m_instructions.end(); i++ ) {
ptx_instruction *pI = *i;
if ( pI->is_label() ) {
@@ -255,6 +257,8 @@ void function_info::ptx_assemble()
fflush(stdout);
printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", m_name.c_str() );
+ //disable pdom analysis here and do it at runtime
+#if 0
create_basic_blocks();
connect_basic_blocks();
bool modified = false;
@@ -278,6 +282,7 @@ void function_info::ptx_assemble()
print_postdominators();
print_ipostdominators();
}
+#endif
printf("GPGPU-Sim PTX: pre-decoding instructions for \'%s\'...\n", m_name.c_str() );
for ( unsigned ii=0; ii < n; ii += m_instr_mem[ii]->inst_size() ) { // handle branch instructions
@@ -1799,6 +1804,15 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL )
//using a shader core object for book keeping, it is not needed but as most function built for performance simulation need it we use it here
extern gpgpu_sim *g_the_gpu;
+ //before we execute, we should do PDOM analysis for functional simulation scenario.
+ function_info *kernel_func_info = kernel.entry();
+ if (kernel_func_info->is_pdom_set()) {
+ printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", kernel.name().c_str() );
+ } else {
+ printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", kernel.name().c_str() );
+ kernel_func_info->do_pdom();
+ kernel_func_info->set_pdom();
+ }
//we excute the kernel one CTA (Block) at the time, as synchronization functions work block wise
while(!kernel.no_more_ctas_to_run()){
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index 5471d6f..1b5d7f6 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -162,6 +162,7 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
\.file TC; BEGIN(INITIAL); return FILE_DIRECTIVE;
\.func TC; BEGIN(IN_FUNC_DECL); return FUNC_DIRECTIVE; // blocking opcode parsing in case the function has the same name as an opcode (e.g. sin(), cos())
\.global TC; return GLOBAL_DIRECTIVE;
+\.global.volatile TC; return GLOBAL_DIRECTIVE; //TODO: fix this!
\.local TC; return LOCAL_DIRECTIVE;
\.loc TC; return LOC_DIRECTIVE;
\.maxnctapersm TC; return MAXNCTAPERSM_DIRECTIVE;
@@ -233,6 +234,7 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
\.u32 TC; return U32_TYPE;
\.u64 TC; return U64_TYPE;
\.f16 TC; return F16_TYPE;
+\.f16x2 TC; return F16_TYPE; /* TODO: figure out what this should really be */
\.f32 TC; return F32_TYPE;
\.f64 TC; return F64_TYPE;
\.ff64 TC; return FF64_TYPE;
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index 8ebdcf8..17e91df 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -280,8 +280,10 @@ type_info *symbol_table::get_array_type( type_info *base_type, unsigned array_di
{
type_info_key t = base_type->get_key();
t.set_array_dim(array_dim);
- type_info *pt;
- pt = m_types[t] = new type_info(this,t);
+ type_info *pt = new type_info(this,t);
+ //Where else is m_types being used? As of now, I dont find any use of it and causing seg fault. So disabling m_types.
+ //TODO: find where m_types can be used in future and solve the seg fault.
+ //pt = m_types[t] = new type_info(this,t);
return pt;
}
@@ -575,6 +577,31 @@ bool function_info::connect_break_targets() //connecting break instructions with
return modified;
}
+void function_info::do_pdom() {
+ create_basic_blocks();
+ connect_basic_blocks();
+ bool modified = false;
+ do {
+ find_dominators();
+ find_idominators();
+ modified = connect_break_targets();
+ } while (modified == true);
+
+ if ( g_debug_execution>=50 ) {
+ print_basic_blocks();
+ print_basic_block_links();
+ print_basic_block_dot();
+ }
+ if ( g_debug_execution>=2 ) {
+ print_dominators();
+ }
+ find_postdominators();
+ find_ipostdominators();
+ if ( g_debug_execution>=50 ) {
+ print_postdominators();
+ print_ipostdominators();
+ }
+}
void intersect( std::set<int> &A, const std::set<int> &B )
{
// return intersection of A and B in A
@@ -1305,6 +1332,7 @@ function_info::function_info(int entry_point )
m_kernel_info.smem = 0;
m_local_mem_framesize = 0;
m_args_aligned_size = -1;
+ pdom_done = false; //initialize it to false
}
unsigned function_info::print_insn( unsigned pc, FILE * fp ) const
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index 9ad1571..26a2839 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -1178,7 +1178,7 @@ public:
//Muchnick's Adv. Compiler Design & Implemmntation Fig 7.15
void find_ipostdominators( );
void print_ipostdominators();
-
+ void do_pdom(); //function to call pdom analysis
unsigned get_num_reconvergence_pairs();
@@ -1274,6 +1274,8 @@ public:
m_local_mem_framesize = sz;
}
bool is_entry_point() const { return m_entry_point; }
+ bool is_pdom_set() const { return pdom_done; } //return pdom flag
+ void set_pdom() { pdom_done = true; } //set pdom flag
private:
unsigned m_uid;
@@ -1281,6 +1283,7 @@ private:
bool m_entry_point;
bool m_extern;
bool m_assembled;
+ bool pdom_done; //flag to check whether pdom is completed or not
std::string m_name;
ptx_instruction **m_instr_mem;
unsigned m_start_PC;
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 6c1b595..9ff0859 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -32,6 +32,7 @@
#include <unistd.h>
#include <dirent.h>
#include <fstream>
+#include <sstream>
/// globals
@@ -287,56 +288,86 @@ void fix_duplicate_errors(char fname2[1024]) {
}
}
+//we need the application name here too.
+char* get_app_binary_name(){
+ char exe_path[1025];
+ char *self_exe_path;
+#ifdef __APPLE__
+ //AMRUTH: get apple device and check the result.
+ printf("WARNING: not tested for Apple-mac devices \n");
+ abort();
+#else
+ std::stringstream exec_link;
+ exec_link << "/proc/self/exe";
+ ssize_t path_length = readlink(exec_link.str().c_str(), exe_path, 1024);
+ assert(path_length != -1);
+ exe_path[path_length] = '\0';
+
+ char *token = strtok(exe_path, "/");
+ while(token !=NULL){
+ self_exe_path = token;
+ token = strtok(NULL,"/");
+ }
+#endif
+ self_exe_path = strtok(self_exe_path, ".");
+ printf("self exe links to: %s\n", self_exe_path);
+ return self_exe_path;
+}
+
void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version )
{
- char fname[1024];
- snprintf(fname,1024,"_ptx_XXXXXX");
- int fd=mkstemp(fname);
- close(fd);
-
- printf("GPGPU-Sim PTX: extracting embedded .ptx to temporary file \"%s\"\n", fname);
- FILE *ptxfile = fopen(fname,"w");
- fprintf(ptxfile,"%s", p_for_info);
- fclose(ptxfile);
+ //do ptxas for individual files instead of one big embedded ptx. This prevents the duplicate defs and declarations.
+ char ptx_file[1000];
+ char *name=get_app_binary_name();
+ char commandline[4096], fname[1024], fname2[1024], final_tempfile_ptxinfo[1024], tempfile_ptxinfo[1024];
+ for (int index=1; index <= no_of_ptx; index++){
+ snprintf(ptx_file, 1000, "%s.%d.sm_%u.ptx", name, index, sm_version);
+ snprintf(fname,1024,"_ptx_XXXXXX");
+ int fd=mkstemp(fname);
+ close(fd);
- char fname2[1024];
- snprintf(fname2,1024,"_ptx2_XXXXXX");
- fd=mkstemp(fname2);
- close(fd);
- char commandline2[4096];
- snprintf(commandline2,4096,"cat %s | sed 's/.version 1.5/.version 1.4/' | sed 's/, texmode_independent//' | sed 's/\\(\\.extern \\.const\\[1\\] .b8 \\w\\+\\)\\[\\]/\\1\\[1\\]/' | sed 's/const\\[.\\]/const\\[0\\]/g' > %s", fname, fname2);
- printf("Running: %s\n", commandline2);
- int result = system(commandline2);
- if( result != 0 ) {
- printf("GPGPU-Sim PTX: ERROR ** while loading PTX (a) %d\n", result);
- printf(" Ensure you have write access to simulation directory\n");
- printf(" and have \'cat\' and \'sed\' in your path.\n");
- exit(1);
- }
+ printf("GPGPU-Sim PTX: extracting embedded .ptx to temporary file \"%s\"\n", fname);
+ snprintf(commandline,4096,"cat %s > %s",ptx_file, fname);
+ if (system(commandline) !=0) {
+ printf("ERROR: %s command failed\n", commandline);
+ exit(0);
+ }
+
+ snprintf(fname2,1024,"_ptx2_XXXXXX");
+ fd=mkstemp(fname2);
+ close(fd);
+ char commandline2[4096];
+ snprintf(commandline2,4096,"cat %s | sed 's/.version 1.5/.version 1.4/' | sed 's/, texmode_independent//' | sed 's/\\(\\.extern \\.const\\[1\\] .b8 \\w\\+\\)\\[\\]/\\1\\[1\\]/' | sed 's/const\\[.\\]/const\\[0\\]/g' > %s", fname, fname2);
+ printf("Running: %s\n", commandline2);
+ int result = system(commandline2);
+ if( result != 0 ) {
+ printf("GPGPU-Sim PTX: ERROR ** while loading PTX (a) %d\n", result);
+ printf(" Ensure you have write access to simulation directory\n");
+ printf(" and have \'cat\' and \'sed\' in your path.\n");
+ exit(1);
+ }
- char tempfile_ptxinfo[1024];
- snprintf(tempfile_ptxinfo,1024,"%sinfo",fname);
- char commandline[1024];
- char extra_flags[1024];
- extra_flags[0]=0;
+ snprintf(tempfile_ptxinfo,1024,"%sinfo",fname);
+ char extra_flags[1024];
+ extra_flags[0]=0;
-#if CUDART_VERSION >= 3000
- if (sm_version == 0) sm_version = 20;
- extern bool g_cdp_enabled;
- if(!g_cdp_enabled)
- snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version);
- else
- snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",sm_version);
-#endif
+ #if CUDART_VERSION >= 3000
+ if (sm_version == 0) sm_version = 20;
+ extern bool g_cdp_enabled;
+ if(!g_cdp_enabled)
+ snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version);
+ else
+ snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",sm_version);
+ #endif
- snprintf(commandline,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file /dev/null 2> %s",
+ snprintf(commandline,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file /dev/null 2> %s",
extra_flags, fname2, tempfile_ptxinfo);
- printf("GPGPU-Sim PTX: generating ptxinfo using \"%s\"\n", commandline);
- result = system(commandline);
- if( result != 0 ) {
- // 65280 = duplicate errors
- if (result == 65280) {
- ptxinfo_in = fopen(tempfile_ptxinfo,"r");
+ printf("GPGPU-Sim PTX: generating ptxinfo using \"%s\"\n", commandline);
+ result = system(commandline);
+ if( result != 0 ) {
+ // 65280 = duplicate errors
+ if (result == 65280) {
+ ptxinfo_in = fopen(tempfile_ptxinfo,"r");
g_ptxinfo_filename = tempfile_ptxinfo;
ptxinfo_parse();
@@ -345,26 +376,87 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
extra_flags, fname2, tempfile_ptxinfo);
printf("GPGPU-Sim PTX: regenerating ptxinfo using \"%s\"\n", commandline);
result = system(commandline);
+ }
+ if (result != 0) {
+ printf("GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result);
+ printf(" Ensure ptxas is in your path.\n");
+ exit(1);
+ }
+ }
+ }
+
+ //TODO: duplicate code! move it into a function so that it can be reused!
+ if(no_of_ptx==0) {
+ //For CDP, we dump everything. So no_of_ptx will be 0.
+ snprintf(fname,1024,"_ptx_XXXXXX");
+ int fd=mkstemp(fname);
+ close(fd);
+
+ printf("GPGPU-Sim PTX: extracting embedded .ptx to temporary file \"%s\"\n", fname);
+ FILE *ptxfile = fopen(fname,"w");
+ fprintf(ptxfile,"%s", p_for_info);
+ fclose(ptxfile);
+
+ snprintf(fname2,1024,"_ptx2_XXXXXX");
+ fd=mkstemp(fname2);
+ close(fd);
+ char commandline2[4096];
+ snprintf(commandline2,4096,"cat %s | sed 's/.version 1.5/.version 1.4/' | sed 's/, texmode_independent//' | sed 's/\\(\\.extern \\.const\\[1\\] .b8 \\w\\+\\)\\[\\]/\\1\\[1\\]/' | sed 's/const\\[.\\]/const\\[0\\]/g' > %s", fname, fname2);
+ printf("Running: %s\n", commandline2);
+ int result = system(commandline2);
+ if( result != 0 ) {
+ printf("GPGPU-Sim PTX: ERROR ** while loading PTX (a) %d\n", result);
+ printf(" Ensure you have write access to simulation directory\n");
+ printf(" and have \'cat\' and \'sed\' in your path.\n");
+ exit(1);
}
- if (result != 0) {
+ char tempfile_ptxinfo[1024];
+ snprintf(tempfile_ptxinfo,1024,"%sinfo",fname);
+ char extra_flags[1024];
+ extra_flags[0]=0;
+#if CUDART_VERSION >= 3000
+ snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version);
+#endif
+
+ snprintf(commandline,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file /dev/null 2> %s",
+ extra_flags, fname2, tempfile_ptxinfo);
+ printf("GPGPU-Sim PTX: generating ptxinfo using \"%s\"\n", commandline);
+ result = system(commandline);
+ if( result != 0 ) {
printf("GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result);
printf(" Ensure ptxas is in your path.\n");
exit(1);
}
}
- ptxinfo_in = fopen(tempfile_ptxinfo,"r");
- g_ptxinfo_filename = tempfile_ptxinfo;
+ //Now that we got resource usage per kernel in a ptx file, we dump all into one file and pass it to rest of the code as usual.
+ if(no_of_ptx>0){
+ char commandline3[4096];
+ snprintf(final_tempfile_ptxinfo,1024,"f_tempfile_ptx");
+ snprintf(commandline3,4096, "cat *info > %s", final_tempfile_ptxinfo);
+ if (system(commandline3)!=0) {
+ printf("ERROR: Either we dont have info files or cat is not working \n");
+ printf("ERROR: %s command failed\n",commandline3);
+ exit(1);
+ }
+ }
+
+ ptxinfo_in = fopen(final_tempfile_ptxinfo,"r");
+ if(no_of_ptx>0)
+ g_ptxinfo_filename = final_tempfile_ptxinfo;
+ else
+ g_ptxinfo_filename = tempfile_ptxinfo;
ptxinfo_parse();
if( ! g_save_embedded_ptx ) {
- snprintf(commandline,1024,"rm -f %s %s %s", fname, fname2, tempfile_ptxinfo);
+ if(no_of_ptx>0)
+ snprintf(commandline,1024,"rm -f %s %s %s *info", fname, fname2, final_tempfile_ptxinfo);
+ else
+ snprintf(commandline,1024,"rm -f %s %s %s *info", fname, fname2, tempfile_ptxinfo);
printf("GPGPU-Sim PTX: removing ptxinfo using \"%s\"\n", commandline);
- result = system(commandline);
- if( result != 0 ) {
- printf("GPGPU-Sim PTX: ERROR ** while loading PTX (c) %d\n", result);
+ if( system(commandline) != 0 ) {
+ printf("GPGPU-Sim PTX: ERROR ** while removing temporary files\n");
exit(1);
}
}
}
-
diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h
index d3d0c92..a8ecda3 100644
--- a/src/cuda-sim/ptx_loader.h
+++ b/src/cuda-sim/ptx_loader.h
@@ -30,6 +30,7 @@
#include <string>
extern bool g_override_embedded_ptx;
+extern int no_of_ptx; //counter to track number of ptx files to be extracted in an application.
class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num );
void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20 );
generated by cgit v1.3 (git 2.53.0) at 2026-07-20 15:36:41 +0000