From 26b9853cc1f6a3b17de0e319b40f28a5703ad6bf Mon Sep 17 00:00:00 2001 From: Amruth Date: Sun, 25 Mar 2018 16:01:22 -0700 Subject: code for removing duplicates in embedded ptx --- src/cuda-sim/ptx_loader.cc | 155 ++++++++++++++++++++++++++++----------------- 1 file changed, 98 insertions(+), 57 deletions(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 6c1b595..d98ca07 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -32,6 +32,7 @@ #include #include #include +#include /// globals @@ -287,56 +288,87 @@ 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); - - 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); - } - - char tempfile_ptxinfo[1024]; - snprintf(tempfile_ptxinfo,1024,"%sinfo",fname); - char commandline[1024]; - 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 + //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]; + 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); - snprintf(commandline,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file /dev/null 2> %s", + 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 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 + + 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 +377,35 @@ 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) { + } + 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. + char commandline3[4096]; + char final_tempfile_ptxinfo[1024]; + 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"); + g_ptxinfo_filename = final_tempfile_ptxinfo; ptxinfo_parse(); if( ! g_save_embedded_ptx ) { - snprintf(commandline,1024,"rm -f %s %s %s", fname, fname2, tempfile_ptxinfo); + snprintf(commandline,1024,"rm -f %s %s %s *info", fname, fname2, final_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); } } } - -- cgit v1.3 From 0d07b73b58224724cff0576d91878817ec3c01c1 Mon Sep 17 00:00:00 2001 From: Amruth Date: Tue, 27 Mar 2018 15:52:52 -0700 Subject: considered CDP scenario --- libcuda/cuda_runtime_api.cc | 24 +++++++------- src/cuda-sim/ptx_loader.cc | 77 +++++++++++++++++++++++++++++++++++++-------- 2 files changed, 77 insertions(+), 24 deletions(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 97d702c..57c5ea1 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1083,8 +1083,8 @@ __host__ __device__ cudaError_t CUDARTAPI cudaDeviceGetStreamPriorityRange(int* return cudaSuccess; } -__host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *stream, unsigned int flags) { - return cudaStreamCreate(stream); +__host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *pStream, unsigned int flags) { + return cudaStreamCreate(pStream); } __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) @@ -1496,14 +1496,16 @@ void extract_code_using_cuobjdump(){ fclose(fp); } } - //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); + 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? @@ -1516,7 +1518,7 @@ void extract_code_using_cuobjdump(){ if(!g_cdp_enabled) 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 -arch=sm_%u -all %s > %s", forced_max_capability, 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; result = system(command); if(result) { diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index d98ca07..9ff0859 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -319,7 +319,7 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num //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]; + 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"); @@ -347,7 +347,6 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num exit(1); } - char tempfile_ptxinfo[1024]; snprintf(tempfile_ptxinfo,1024,"%sinfo",fname); char extra_flags[1024]; extra_flags[0]=0; @@ -385,23 +384,75 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num } } } + + //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); + } + 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); + } + } + //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. - char commandline3[4096]; - char final_tempfile_ptxinfo[1024]; - 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); - } + 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"); - g_ptxinfo_filename = final_tempfile_ptxinfo; + 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 *info", fname, fname2, final_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); if( system(commandline) != 0 ) { printf("GPGPU-Sim PTX: ERROR ** while removing temporary files\n"); -- cgit v1.3 From 6d36a6a20f6e0c7082243c0dabe1a6b734c5002b Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Sat, 31 Mar 2018 08:48:47 -0700 Subject: masked declaration leading to using unitialized buffer as file name, leading to crash this still fails regressions --- src/cuda-sim/ptx_loader.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 9ff0859..34870c4 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -410,7 +410,7 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num printf(" and have \'cat\' and \'sed\' in your path.\n"); exit(1); } - char tempfile_ptxinfo[1024]; + //char tempfile_ptxinfo[1024]; snprintf(tempfile_ptxinfo,1024,"%sinfo",fname); char extra_flags[1024]; extra_flags[0]=0; -- cgit v1.3 From 31ad7674de6dc4b25ba862bcd00b660fdb1a5cff Mon Sep 17 00:00:00 2001 From: Amruth Date: Wed, 4 Apr 2018 13:44:10 -0700 Subject: adding missing ptxas flags for cdp support --- README | 66 +++++++++++++++++++++++++++++++++++++++++++++ libcuda/cuda_runtime_api.cc | 4 +-- src/cuda-sim/ptx_loader.cc | 27 ++++++++++++------- 3 files changed, 85 insertions(+), 12 deletions(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/README b/README index 6e2d734..4426cdd 100644 --- a/README +++ b/README @@ -346,3 +346,69 @@ the applications you care about (implying these applications worked for you before you did the merge). You want to do this before making further changes to identify any compile time or runtime errors that occur due to the code merging process. + + +** Debugging failing GPGPU-Sim Regressions ** + +To debug failing GPGPU-Sim regression tests you need to run them locally. The fastest way to do this, assuming you are working with GPGPU-Sim versions more recent than the GPGPU-Sim dev branch circa March 28, 2018 (commit hash 2221d208a745a098a60b0d24c05007e92aaba092), is to install Docker. The instructions below were tested with Docker CE version 18.03 on Ubuntu and Mac OS. Docker will enable you to run the same set of regressions used by GPGPU-Sim when submitting a pull request to https://github.com/gpgpu-sim/gpgpu-sim_distribution and also allow you to log in and launch GPGPU-Sim in gdb so you can inspect failures. + +1. Install Docker. On Ubuntu 14.04 and 16.04 the following instructions work: https://docs.docker.com/install/linux/docker-ce/ubuntu/#uninstall-old-versions + +2. Clone GPGPU-Sim from your fork of GPGPU-Sim. For example: + + git clone https://github.com//gpgpu-sim_distribution.git + + +3. Run the following command (this is all one line) to run the regressions in docker: + + docker run --privileged -v `pwd`:/home/runner/gpgpu-sim_distribution:rw aamodt/gpgpu-sim_regress:latest /bin/bash -c "./start_torque.sh; chown -R runner /home/runner/gpgpu-sim_distribution; su - runner -c 'source /home/runner/gpgpu-sim_distribution/setup_environment && make -j -C /home/runner/gpgpu-sim_distribution && cd /home/runner/gpgpu-sim_simulations/ && git pull && /home/runner/gpgpu-sim_simulations/util/job_launching/run_simulations.py -c /home/runner/gpgpu-sim_simulations/util/job_launching/regression_recipies/rodinia_2.0-ft/configs.gtx1080ti.yml -N regress && /home/runner/gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -v -N regress’; tail -f /dev/null" + +Explanation: The last part of this command, "tail -f /dev/null" will keep the docker container running after the regressions finish. This enables you to log into the container to run the same tests inside gdb so you can debug. The "--privileged" part enables you to use breakpoints inside gdb in a container. The "-v" part maps the current directory (with the GPGPU-Sim source code you want to test) into the container. The string "aamodt/gpgpu-sim_regress:latest" is a tag for a container setup to run regressions which will be downloaded from docker hub. The portion starting with /bin/bash is a set of commands run inside a bash shell inside the container. E.g., the command start_torque.sh starts up a queue manager inside the container. + +If the above command stops with the message "fatal: unable to access 'https://github.com/tgrogers/gpgpu-sim_simulations.git/': Could not resolve host: github.com" this likely means your computer sits behind a firewall which is blocking access to Google's name servers (e.g., 8.8.8.8). To get around this you will need to modify th above command to point to your local DNS server. Lookup your DNS server IP address which we will call below. On Ubuntu run "ifconfig" to lookup the network interface connecting your computer to the network. Then run "nmcli device show " to find the IP address of your DNS server. Modify the above command to include "--dns " after "run", E.g., + + docker run --dns --privileged -v `pwd`:/home/runner/gpgpu-sim_distribution:rw aamodt/gpgpu-sim_regress:latest /bin/bash -c "./start_torque.sh; chown -R runner /home/runner/gpgpu-sim_distribution; su - runner -c 'source /home/runner/gpgpu-sim_distribution/setup_environment && make -j -C /home/runner/gpgpu-sim_distribution && cd /home/runner/gpgpu-sim_simulations/ && git pull && /home/runner/gpgpu-sim_simulations/util/job_launching/run_simulations.py -c /home/runner/gpgpu-sim_simulations/util/job_launching/regression_recipies/rodinia_2.0-ft/configs.gtx1080ti.yml -N regress && /home/runner/gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -v -N regress’; tail -f /dev/null" + +4. Find the CONTAINER ID associated with your docker container by running "docker ps". + +5. Log into the container by running the command: + + docker exec -it /bin/bash -c "su -l runner" + +The container is running Ubuntu 16.04 and has screen, cscope and vim installed (if you find a favorite Linux tool missing, it is fairly easy to create derived containers that have additional tools). + +6. Lookup the directory of the regression test you want to debug by going to the regression log file directory: + + cd /home/runner/gpgpu-sim_simulations/util/job_launching/logfiles + +7. The file "failed_job_log_sim_log.regress..txt" includes information about the failed test including its simulation directory. For the following example, I'll assume the first failing test was "hotspot-rodinia-2.0-ft-30_6_40___data_result_30_6_40_txt--GTX1080Ti" for which the simulation directory is /home/runner/gpgpu-sim_simulations/util/job_launching/../../sim_run_4.2/hotspot-rodinia-2.0-ft/30_6_40___data_result_30_6_40_txt/GTX1080Ti/ + +8. Change to the simulation directory using: + + cd + +E.g., "cd /home/runner/gpgpu-sim_simulations/util/job_launching/../../sim_run_4.2/hotspot-rodinia-2.0-ft/30_6_40___data_result_30_6_40_txt/GTX1080Ti/" + +This directory should contain a file called "torque.sim" that contains commands used to launch the simulation during regression tests. We will modify this file to enable us to re-run the regression test in gdb. This directory should also contain a file containing the standard output during the regression test. This file will end in .o where is the torque queue manager job number. For the running example for me this file is called "hotspot-rodinia-2.0-ft-30_6_40___data_result_30_6_40_txt.o2". Open this file to determine the LD_LIBRARY_PATH settings used when launching the simulation. Look for a line that starts "doing: export LD_LIBRARY_PATH" and copy the entire line starting with "export LD_LIBRARY_PATH ..." + +9. Paste the "export LD_LIBRARY_PATH ..." line into the bash shell to set LD_LIBRARY_PATH. E.g., + + export LD_LIBRARY_PATH=/home/runner/gpgpu-sim_simulations/util/job_launching/../../sim_run_4.2/gpgpu-sim-builds/libcudart_gpgpu-sim_git-commit-177d02254ae38b6331b17dd6cd139b570a03c589_modified_0.so:/gpgpu-sim/usr/local/gcc-4.5.4/lib64:/gpgpu-sim/usr/local/gcc-4.5.4/lib:/gpgpu-sim/usr/local/gcc-4.5.4/lib/gcc/x86_64-unknown-linux-gnu/lib64/:/gpgpu-sim/usr/local/gcc-4.5.4/lib/gcc/x86_64-unknown-linux-gnu/4.5.4/:/usr/lib/x86_64-linux-gnu:/home/runner/gpgpu-sim_distribution/lib/gcc-4.5.4/cuda-4020/release:/gpgpu-sim/usr/local/gcc-4.5.4/lib64:/gpgpu-sim/usr/local/gcc-4.5.4/lib:/gpgpu-sim/usr/local/gcc-4.5.4/lib/gcc/x86_64-unknown-linux-gnu/lib64/:/gpgpu-sim/usr/local/gcc-4.5.4/lib/gcc/x86_64-unknown-linux-gnu/4.5.4/:/usr/lib/x86_64-linux-gnu: + +10. In the same shell, build the debug version of GPGPU-Sim then return to the directory above: + + pushd ~/gpgpu-sim_distribution/ + source setup_environment debug + make + popd + +11. Open and edit torque.sim and preface the very last line with "gdb --args ". After editing the last line in torque.sim should look something like: + + gdb --args /home/runner/gpgpu-sim_simulations/util/job_launching/../../benchmarks/bin/4.2/release/hotspot-rodinia-2.0-ft 30 6 40 ./data/result_30_6_40.txt + +12. Re-run the regression test in gdb by sourcing the torque.sim file: + + . torque.sim + +This will put you in at the (gdb) prompt. Setup any breakpoints needed and run. + diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index ded1aee..c103244 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1085,8 +1085,8 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetStreamPriorityRange(int* leastPriori return cudaSuccess; } -__host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *pStream, unsigned int flags) { - return cudaStreamCreate(pStream); +__host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *stream, unsigned int flags) { + return cudaStreamCreate(stream); } __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 34870c4..7863be4 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -353,11 +353,7 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num #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); + 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", @@ -414,9 +410,15 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num 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 + + #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", extra_flags, fname2, tempfile_ptxinfo); @@ -448,11 +450,16 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num g_ptxinfo_filename = tempfile_ptxinfo; ptxinfo_parse(); + snprintf(commandline,1024,"rm -f *info"); + if( system(commandline) != 0 ) { + printf("GPGPU-Sim PTX: ERROR ** while removing temporary info files\n"); + exit(1); + } if( ! g_save_embedded_ptx ) { if(no_of_ptx>0) - snprintf(commandline,1024,"rm -f %s %s %s *info", fname, fname2, final_tempfile_ptxinfo); + snprintf(commandline,1024,"rm -f %s %s %s", fname, fname2, final_tempfile_ptxinfo); else - snprintf(commandline,1024,"rm -f %s %s %s *info", fname, fname2, tempfile_ptxinfo); + snprintf(commandline,1024,"rm -f %s %s %s", fname, fname2, tempfile_ptxinfo); printf("GPGPU-Sim PTX: removing ptxinfo using \"%s\"\n", commandline); if( system(commandline) != 0 ) { printf("GPGPU-Sim PTX: ERROR ** while removing temporary files\n"); -- cgit v1.3 From 335bb4b38974b1843e1786c5d86b0fcf2a7943c8 Mon Sep 17 00:00:00 2001 From: Amruth Date: Wed, 4 Apr 2018 16:01:32 -0700 Subject: Prog gets stuck because it doesnt recieve EOF --- src/cuda-sim/ptx_loader.cc | 3 +++ 1 file changed, 3 insertions(+) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 7863be4..03ea31a 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -448,6 +448,9 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num g_ptxinfo_filename = final_tempfile_ptxinfo; else g_ptxinfo_filename = tempfile_ptxinfo; + + //The program might get stuck because the parser didnt receive a EOF. + printf("NOTE: If the program is stuck, please press ctrl+d for ubuntu and ctrl+z for windows users \n"); ptxinfo_parse(); snprintf(commandline,1024,"rm -f *info"); -- cgit v1.3 From 69c57b077d6799f46ad43b99cc0e0049b7c43775 Mon Sep 17 00:00:00 2001 From: Amruth Date: Wed, 4 Apr 2018 16:06:42 -0700 Subject: updating EOF condition for Mac users --- src/cuda-sim/ptx_loader.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 03ea31a..4ddc6bf 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -450,7 +450,7 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num g_ptxinfo_filename = tempfile_ptxinfo; //The program might get stuck because the parser didnt receive a EOF. - printf("NOTE: If the program is stuck, please press ctrl+d for ubuntu and ctrl+z for windows users \n"); + printf("NOTE: If the program is stuck, please press ctrl+d for Ubuntu/Mac and ctrl+z for Windows users \n"); ptxinfo_parse(); snprintf(commandline,1024,"rm -f *info"); -- cgit v1.3 From e1dc8113aa2a51885541f96943bd8d90eaccd968 Mon Sep 17 00:00:00 2001 From: Amruth Date: Thu, 5 Apr 2018 16:31:44 -0700 Subject: fixing file pointer and attributes issues --- README | 1 + libcuda/cuda_runtime_api.cc | 34 ++++++++++++++++++++++++++++++++-- src/cuda-sim/cuda-sim.cc | 1 + src/cuda-sim/ptx_loader.cc | 5 ++--- 4 files changed, 36 insertions(+), 5 deletions(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/README b/README index 4426cdd..543177c 100644 --- a/README +++ b/README @@ -349,6 +349,7 @@ process. ** Debugging failing GPGPU-Sim Regressions ** +Credits: Tor M Aamodt To debug failing GPGPU-Sim regression tests you need to run them locally. The fastest way to do this, assuming you are working with GPGPU-Sim versions more recent than the GPGPU-Sim dev branch circa March 28, 2018 (commit hash 2221d208a745a098a60b0d24c05007e92aaba092), is to install Docker. The instructions below were tested with Docker CE version 18.03 on Ubuntu and Mac OS. Docker will enable you to run the same set of regressions used by GPGPU-Sim when submitting a pull request to https://github.com/gpgpu-sim/gpgpu-sim_distribution and also allow you to log in and launch GPGPU-Sim in gdb so you can inspect failures. diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index c103244..3fd88dc 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -349,7 +349,15 @@ class _cuda_device_id *GPGPUSim_Init() prop->maxGridSize[2] = 0x40000000; prop->totalConstMem = 0x40000000; prop->textureAlignment = 0; - prop->sharedMemPerBlock = the_gpu->shared_mem_size(); + /* + * TODO: Update the .config and xml files of all GPU config files with new value of sharedMemPerBlock. + * Previously, this was thought as sharedMemPerMultiprocessor and is being used in many places. + * Check whether all the instances of shared_mem_size(), gpgpu_shmem_size or sharedMemPerBlock are meant to use sharedMemPerBlock or sharedMemPerMultiprocessor. + */ + prop->sharedMemPerBlock = 49152; +#if (CUDART_VERSION > 5000) + prop->sharedMemPerMultiprocessor = the_gpu->shared_mem_size(); +#endif prop->regsPerBlock = the_gpu->num_registers_per_core(); prop->warpSize = the_gpu->wrp_size(); prop->clockRate = the_gpu->shader_clock(); @@ -840,6 +848,15 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic if (device <= dev->num_devices() ) { prop = dev->get_prop(); switch (attr) { + case 2: + *value= prop->maxThreadsDim[0]; + break; + case 3: + *value= prop->maxThreadsDim[1]; + break; + case 4: + *value= prop->maxThreadsDim[2]; + break; case 5: *value= prop->maxGridSize[0]; break; @@ -849,6 +866,12 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic case 7: *value= prop->maxGridSize[2]; break; + case 8: + *value= prop->sharedMemPerBlock; + break; + case 9: + *value= prop->totalConstMem; + break; case 10: *value= prop->warpSize; break; @@ -861,11 +884,14 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic case 16: *value= prop->multiProcessorCount ; break; + case 34: + *value= 0; + break; case 39: *value= dev->get_gpgpu()->threads_per_core(); break; case 75: - *value= 8 ; + *value= 9 ; break; case 76: *value= 3 ; @@ -873,6 +899,9 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic case 78: *value= 0 ; //TODO: as of now, we dont support stream priorities. break; + case 81: + *value= prop->sharedMemPerMultiprocessor; + break; default: printf("ERROR: implement the attribute numbered %d \n",attr); abort(); @@ -1882,6 +1911,7 @@ 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); + max_capability=context->get_device()->get_gpgpu()->get_config().get_forced_max_capability(); cuobjdumpPTXSection* ptx = NULL; const char* pre_load = getenv("CUOBJDUMP_SIM_FILE"); diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index dce35ca..b1eaf01 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -2155,3 +2155,4 @@ void functionalCoreSim::warp_exit( unsigned warp_id ) } } } + diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 4ddc6bf..33a4260 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -423,6 +423,7 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num 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); + fflush(stdout); result = system(commandline); if( result != 0 ) { printf("GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result); @@ -443,14 +444,12 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num } } - 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_in = fopen(g_ptxinfo_filename,"r"); - //The program might get stuck because the parser didnt receive a EOF. - printf("NOTE: If the program is stuck, please press ctrl+d for Ubuntu/Mac and ctrl+z for Windows users \n"); ptxinfo_parse(); snprintf(commandline,1024,"rm -f *info"); -- cgit v1.3 From 2278609c8e279e7dfba49211256e818e3152318b Mon Sep 17 00:00:00 2001 From: Jonathan Date: Fri, 1 Jun 2018 13:04:30 -0700 Subject: WIP attempting to parse ptx files in order --- libcuda/cuda_runtime_api.cc | 257 +++++++++++++++++++++++++------------------- src/cuda-sim/ptx_loader.cc | 20 ++++ src/cuda-sim/ptx_loader.h | 1 + 3 files changed, 168 insertions(+), 110 deletions(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index dc92522..ece958e 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -112,6 +112,7 @@ #include #include #include +#include #ifdef OPENGL_SUPPORT #define GL_GLEXT_PROTOTYPES #ifdef __APPLE__ @@ -148,6 +149,10 @@ std::map pinned_memory; //support for pinned memories added std::map pinned_memory_size; int no_of_ptx=0; +char ptx_list_file_name[1024]; +std::map fatbinmap; +std::mapfatbin_registered; +std::map name_symtab; extern void synchronize(); extern void exit_simulation(); @@ -1500,6 +1505,145 @@ char* get_app_binary_name(std::string abs_path){ return self_exe_path; } +//extracts all ptx files from binary and dumps into prog_name.unique_no.sm_<>.ptx files +void extract_ptx_files_using_cuobjdump(bool g_cdp_enabled){ + char command[1000]; + std::string app_binary = get_app_binary(); + + snprintf(ptx_list_file_name,1024,"_cuobjdump_list_ptx_XXXXXX"); + int fd2=mkstemp(ptx_list_file_name); + close(fd2); + //only want file names + snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -lptx %s | cut -d \":\" -f 2 | awk '{$1=$1}1' > %s", app_binary.c_str(), ptx_list_file_name); + if( system(command) != 0 ) { + printf("WARNING: Failed to execute cuobjdump to get list of ptx files \n"); + exit(0); + } + 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 + + std::ifstream infile(ptx_list_file_name); + std::string line; + while (std::getline(infile, line)) + { + //int pos = line.find(std::string(get_app_binary_name(app_binary))); + const char *ptx_file = line.c_str(); + printf("Extracting specific PTX file named %s \n",ptx_file); + snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -xptx %s %s", ptx_file, app_binary.c_str()); + if (system(command)!=0) { + printf("ERROR: command: %s failed \n",command); + exit(0); + } + no_of_ptx++; + } + } + + if(!no_of_ptx){ + printf("WARNING: Number of ptx in the executable file are 0. One of the reasons might be\n"); + printf("\t1. CDP is enabled\n"); + } +} + +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; + } + + std::map > version_filename; + + std::ifstream infile(ptx_list_file_name); + std::string line; + while (std::getline(infile, line)) + { + //int pos = line.find(std::string(get_app_binary_name(app_binary))); + const char *ptx_file = line.c_str(); + int pos1 = line.find("sm_"); + int pos2 = line.find_last_of("."); + if (pos1==std::string::npos&&pos2==std::string::npos){ + printf("ERROR: PTX list is not in correct format"); + exit(0); + } + std::string vstr = line.substr(pos1+3,pos2-pos1-3); + int version = atoi(vstr.c_str()); + if (version_filename.find(version)==version_filename.end()){ + version_filename[version] = std::set(); + } + version_filename[version].insert(line); + } + + symbol_table *symtab; + //loops through all ptx files from smallest sm version to largest + std::map >::iterator itr_m; + for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){ + std::set::iterator itr_s; + for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){ + std::string ptx_filename = *itr_s; + printf("GPGPU-Sim PTX: Parsing %s\n",ptx_filename.c_str()); + symtab = gpgpu_ptx_sim_load_ptx_from_filename( ptx_filename.c_str() ); + } + } + name_symtab[fname] = symtab; + context->add_binary(symtab, handle); + + +// unsigned max_capability = 0; +// for ( std::list::iterator iter = cuobjdumpSectionList.begin(); +// iter != cuobjdumpSectionList.end(); +// iter++){ +// 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); +// if (max_capability == 0) max_capability=context->get_device()->get_gpgpu()->get_config().get_forced_max_capability(); +// +// 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 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); +// 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, max_capability ); +// delete[] ptxplus_str; +// } else { +// symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxcode, 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 ); +// } + 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 +} + //! Call cuobjdump to extract everything (-elf -sass -ptx) /*! * This Function extract the whole PTX (for all the files) using cuobjdump @@ -1514,7 +1658,7 @@ void extract_code_using_cuobjdump(){ //prevent the dumping by cuobjdump everytime we execute the code! const char *override_cuobjdump = getenv("CUOBJDUMP_SIM_FILE"); - char command[1000], ptx_file[1000]; + char command[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()); @@ -1527,50 +1671,8 @@ void extract_code_using_cuobjdump(){ //dump ptx for all individial ptx files into sepearte files which is later used by ptxas. int result=0; #if (CUDART_VERSION >= 6000) - 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); - 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(no_of_ptx==0){ - printf("WARNING: Number of ptx in the executable file are 0. One of the reasons might be\n"); - printf("\t1. CDP is enabled\n"); - printf("\t2. cuobjdump -lptx doesnt recognize sm_%u\n",forced_max_capability); - printf("\t3. the application was not compiled with nvcc flag sm_%u\n",forced_max_capability); - } - } - 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); - } - } - } + extract_ptx_files_using_cuobjdump(g_cdp_enabled); + cuobjdumpParseBinary(1); #endif //TODO: redundant to dump twice. how can it be prevented? //dump only for specific arch @@ -1904,77 +2006,12 @@ 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, 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(); - 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(); - iter != cuobjdumpSectionList.end(); - iter++){ - 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); - if (max_capability == 0) max_capability=context->get_device()->get_gpgpu()->get_config().get_forced_max_capability(); - - 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 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); - 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, max_capability ); - delete[] ptxplus_str; - } else { - symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxcode, 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 ); - } - 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 -} void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) { diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 33a4260..8deafc6 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -185,6 +185,26 @@ symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source return symtab; } +symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ) +{ + symbol_table *symtab=init_parser(filename); + int errors = ptx_parse (); + if ( errors ) { + char fname[1024]; + snprintf(fname,1024,"_ptx_errors_XXXXXX"); + int fd=mkstemp(fname); + close(fd); + printf("GPGPU-Sim PTX: parser error detected, exiting... but first extracting .ptx to \"%s\"\n", fname); + FILE *ptxfile = fopen(fname,"w"); + fclose(ptxfile); + abort(); + exit(40); + } + + printf("GPGPU-Sim PTX: finished parsing EMBEDDED .ptx file %s\n",filename); + return symtab; +} + void fix_duplicate_errors(char fname2[1024]) { char tempfile[1024] = "_temp_ptx"; char commandline[1024]; diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index a8ecda3..c5c9dd8 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -33,6 +33,7 @@ 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 ); +class symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ); void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20 ); char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str); bool keep_intermediate_files(); -- cgit v1.3 From 73fea7152926dbc41cf008a4ed3402cc1feeefa7 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Mon, 4 Jun 2018 16:53:23 -0700 Subject: parses through all ptx files, TODO: need to impl dp4a --- libcuda/cuda_runtime_api.cc | 1 - src/cuda-sim/instructions.cc | 7 +++++++ src/cuda-sim/opcodes.def | 1 + src/cuda-sim/ptx.l | 1 + src/cuda-sim/ptx_loader.cc | 13 ------------- src/cuda-sim/ptx_parser.cc | 6 ++++++ 6 files changed, 15 insertions(+), 14 deletions(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index ece958e..d03caf7 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1546,7 +1546,6 @@ void extract_ptx_files_using_cuobjdump(bool g_cdp_enabled){ void cuobjdumpParseBinary(unsigned int handle){ - if(fatbin_registered[handle]) return; fatbin_registered[handle] = true; CUctx_st *context = GPGPUSim_Context(); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 35d1782..37438fa 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -2289,6 +2289,13 @@ void div_impl( const ptx_instruction *pI, ptx_thread_info *thread ) thread->set_operand_value(dst,data, i_type, thread,pI); } +void dp4a_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + printf("instruction not implemented yet"); + assert(0); + +} + void ex2_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { ptx_reg_t src1_data, src2_data, data; diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def index e1b1422..2dbfb78 100644 --- a/src/cuda-sim/opcodes.def +++ b/src/cuda-sim/opcodes.def @@ -60,6 +60,7 @@ OP_DEF(COS_OP,cos_impl,"cos",1,4) OP_DEF(CVT_OP,cvt_impl,"cvt",1,1) OP_DEF(CVTA_OP,cvta_impl,"cvta",1,1) OP_DEF(DIV_OP,div_impl,"div",1,1) +OP_DEF(DP4A_OP,dp4a_impl,"dp4a",1,1) OP_DEF(EX2_OP,ex2_impl,"ex2",1,4) OP_DEF(EXIT_OP,exit_impl,"exit",1,3) OP_DEF(FMA_OP,fma_impl,"fma",1,2) diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index 908c5be..3bb0d17 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -77,6 +77,7 @@ cos TC; ptx_lval.int_value = COS_OP; return OPCODE; cvt TC; ptx_lval.int_value = CVT_OP; return OPCODE; cvta TC; ptx_lval.int_value = CVTA_OP; return OPCODE; div TC; ptx_lval.int_value = DIV_OP; return OPCODE; +dp4a TC; ptx_lval.int_value = DP4A_OP; return OPCODE; ex2 TC; ptx_lval.int_value = EX2_OP; return OPCODE; exit TC; ptx_lval.int_value = EXIT_OP; return OPCODE; fma TC; ptx_lval.int_value = FMA_OP; return OPCODE; diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 8deafc6..0348af0 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -188,19 +188,6 @@ symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ) { symbol_table *symtab=init_parser(filename); - int errors = ptx_parse (); - if ( errors ) { - char fname[1024]; - snprintf(fname,1024,"_ptx_errors_XXXXXX"); - int fd=mkstemp(fname); - close(fd); - printf("GPGPU-Sim PTX: parser error detected, exiting... but first extracting .ptx to \"%s\"\n", fname); - FILE *ptxfile = fopen(fname,"w"); - fclose(ptxfile); - abort(); - exit(40); - } - printf("GPGPU-Sim PTX: finished parsing EMBEDDED .ptx file %s\n",filename); return symtab; } diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index a180da9..a51799a 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -32,6 +32,8 @@ extern int ptx_error( const char *s ); extern int ptx_lineno; +extern int ptx_parse(); +extern FILE *ptx_in; static const struct core_config *g_shader_core_config; void set_ptx_warp_size(const struct core_config * warp_size) @@ -135,6 +137,10 @@ symbol_table *init_parser( const char *ptx_filename ) g_ptx_token_decode[generic_space] = "generic_space"; g_ptx_token_decode[instruction_space] = "instruction_space"; + + ptx_in = fopen(ptx_filename, "r"); + ptx_parse(); + fclose(ptx_in); return g_global_symbol_table; } -- cgit v1.3 From b79708980b80b6e50b1612f5436655c724377779 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Tue, 5 Jun 2018 17:22:35 -0700 Subject: ptx parse only --- libcuda/cuda_runtime_api.cc | 48 ++++++++++++++++++++++----------------------- src/cuda-sim/ptx_loader.cc | 37 ++++------------------------------ src/cuda-sim/ptx_parser.cc | 5 ++--- 3 files changed, 29 insertions(+), 61 deletions(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index d03caf7..4efbd68 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -112,7 +112,6 @@ #include #include #include -#include #ifdef OPENGL_SUPPORT #define GL_GLEXT_PROTOTYPES #ifdef __APPLE__ @@ -149,7 +148,7 @@ std::map pinned_memory; //support for pinned memories added std::map pinned_memory_size; int no_of_ptx=0; -char ptx_list_file_name[1024]; +std::map > version_filename; std::map fatbinmap; std::mapfatbin_registered; std::map name_symtab; @@ -1510,6 +1509,7 @@ void extract_ptx_files_using_cuobjdump(bool g_cdp_enabled){ char command[1000]; std::string app_binary = get_app_binary(); + char ptx_list_file_name[1024]; snprintf(ptx_list_file_name,1024,"_cuobjdump_list_ptx_XXXXXX"); int fd2=mkstemp(ptx_list_file_name); close(fd2); @@ -1542,6 +1542,27 @@ void extract_ptx_files_using_cuobjdump(bool g_cdp_enabled){ printf("WARNING: Number of ptx in the executable file are 0. One of the reasons might be\n"); printf("\t1. CDP is enabled\n"); } + + std::ifstream infile(ptx_list_file_name); + std::string line; + while (std::getline(infile, line)) + { + //int pos = line.find(std::string(get_app_binary_name(app_binary))); + const char *ptx_file = line.c_str(); + int pos1 = line.find("sm_"); + int pos2 = line.find_last_of("."); + if (pos1==std::string::npos&&pos2==std::string::npos){ + printf("ERROR: PTX list is not in correct format"); + exit(0); + } + std::string vstr = line.substr(pos1+3,pos2-pos1-3); + int version = atoi(vstr.c_str()); + if (version_filename.find(version)==version_filename.end()){ + version_filename[version] = std::set(); + } + version_filename[version].insert(line); + } + } void cuobjdumpParseBinary(unsigned int handle){ @@ -1557,28 +1578,6 @@ void cuobjdumpParseBinary(unsigned int handle){ return; } - std::map > version_filename; - - std::ifstream infile(ptx_list_file_name); - std::string line; - while (std::getline(infile, line)) - { - //int pos = line.find(std::string(get_app_binary_name(app_binary))); - const char *ptx_file = line.c_str(); - int pos1 = line.find("sm_"); - int pos2 = line.find_last_of("."); - if (pos1==std::string::npos&&pos2==std::string::npos){ - printf("ERROR: PTX list is not in correct format"); - exit(0); - } - std::string vstr = line.substr(pos1+3,pos2-pos1-3); - int version = atoi(vstr.c_str()); - if (version_filename.find(version)==version_filename.end()){ - version_filename[version] = std::set(); - } - version_filename[version].insert(line); - } - symbol_table *symtab; //loops through all ptx files from smallest sm version to largest std::map >::iterator itr_m; @@ -1671,7 +1670,6 @@ void extract_code_using_cuobjdump(){ int result=0; #if (CUDART_VERSION >= 6000) extract_ptx_files_using_cuobjdump(g_cdp_enabled); - cuobjdumpParseBinary(1); #endif //TODO: redundant to dump twice. how can it be prevented? //dump only for specific arch diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 0348af0..97d5773 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -152,39 +152,6 @@ char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptxfilenam return ptxplus_str; } - -symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ) -{ - char buf[1024]; - snprintf(buf,1024,"_%u.ptx", source_num ); - if( g_save_embedded_ptx ) { - FILE *fp = fopen(buf,"w"); - fprintf(fp,"%s",p); - fclose(fp); - } - symbol_table *symtab=init_parser(buf); - ptx__scan_string(p); - int errors = ptx_parse (); - if ( errors ) { - char fname[1024]; - snprintf(fname,1024,"_ptx_errors_XXXXXX"); - int fd=mkstemp(fname); - close(fd); - printf("GPGPU-Sim PTX: parser error detected, exiting... but first extracting .ptx to \"%s\"\n", fname); - FILE *ptxfile = fopen(fname,"w"); - fprintf(ptxfile,"%s", p ); - fclose(ptxfile); - abort(); - exit(40); - } - - if ( g_debug_execution >= 100 ) - print_ptx_file(p,source_num,buf); - - printf("GPGPU-Sim PTX: finished parsing EMBEDDED .ptx file %s\n",buf); - return symtab; -} - symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ) { symbol_table *symtab=init_parser(filename); @@ -321,6 +288,10 @@ char* get_app_binary_name(){ return self_exe_path; } +void gpgpu_ptx_info_load_from_filename( const char *filename ) +{ +} + void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version ) { //do ptxas for individual files instead of one big embedded ptx. This prevents the duplicate defs and declarations. diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 06ca870..c418fac 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -142,9 +142,7 @@ symbol_table *init_parser( const char *ptx_filename ) g_global_allfiles_symbol_table = new symbol_table("global_allfiles", 0, NULL); g_global_symbol_table = g_current_symbol_table = g_global_allfiles_symbol_table; } -// else { -// g_global_symbol_table = g_current_symbol_table = new symbol_table("global",0,g_global_allfiles_symbol_table); -// } + ptx_lineno = 1; #define DEF(X,Y) g_ptx_token_decode[X] = Y; @@ -164,6 +162,7 @@ symbol_table *init_parser( const char *ptx_filename ) g_ptx_token_decode[global_space] = "global_space"; g_ptx_token_decode[generic_space] = "generic_space"; g_ptx_token_decode[instruction_space] = "instruction_space"; + init_directive_state(); init_instruction_state(); -- cgit v1.3 From 40b275992791f34ecbb856e358af78cd3097da7e Mon Sep 17 00:00:00 2001 From: Jonathan Date: Wed, 6 Jun 2018 14:30:41 -0700 Subject: reverted some changed to pass regression --- libcuda/cuda_runtime_api.cc | 165 ++++++++++++++++++++++---------------------- src/cuda-sim/ptx_loader.cc | 32 +++++++++ 2 files changed, 116 insertions(+), 81 deletions(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 4efbd68..08ee413 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -149,9 +149,6 @@ std::map pinned_memory; //support for pinned memories added std::map pinned_memory_size; int no_of_ptx=0; std::map > version_filename; -std::map fatbinmap; -std::mapfatbin_registered; -std::map name_symtab; extern void synchronize(); extern void exit_simulation(); @@ -1565,83 +1562,6 @@ void extract_ptx_files_using_cuobjdump(bool g_cdp_enabled){ } -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; - } - - symbol_table *symtab; - //loops through all ptx files from smallest sm version to largest - std::map >::iterator itr_m; - for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){ - std::set::iterator itr_s; - for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){ - std::string ptx_filename = *itr_s; - printf("GPGPU-Sim PTX: Parsing %s\n",ptx_filename.c_str()); - symtab = gpgpu_ptx_sim_load_ptx_from_filename( ptx_filename.c_str() ); - } - } - name_symtab[fname] = symtab; - context->add_binary(symtab, handle); - - -// unsigned max_capability = 0; -// for ( std::list::iterator iter = cuobjdumpSectionList.begin(); -// iter != cuobjdumpSectionList.end(); -// iter++){ -// 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); -// if (max_capability == 0) max_capability=context->get_device()->get_gpgpu()->get_config().get_forced_max_capability(); -// -// 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 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); -// 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, max_capability ); -// delete[] ptxplus_str; -// } else { -// symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxcode, 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 ); -// } - 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 -} - //! Call cuobjdump to extract everything (-elf -sass -ptx) /*! * This Function extract the whole PTX (for all the files) using cuobjdump @@ -1656,7 +1576,7 @@ void extract_code_using_cuobjdump(){ //prevent the dumping by cuobjdump everytime we execute the code! const char *override_cuobjdump = getenv("CUOBJDUMP_SIM_FILE"); - char command[1000]; + 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()); @@ -2003,12 +1923,95 @@ 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, 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(); + 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; + } + symbol_table *symtab; + +#if (CUDART_VERSION >= 6000) + //loops through all ptx files from smallest sm version to largest + std::map >::iterator itr_m; + for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){ + std::set::iterator itr_s; + for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){ + std::string ptx_filename = *itr_s; + printf("GPGPU-Sim PTX: Parsing %s\n",ptx_filename.c_str()); + symtab = gpgpu_ptx_sim_load_ptx_from_filename( ptx_filename.c_str() ); + } + } + name_symtab[fname] = symtab; + context->add_binary(symtab, 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()); + return; +#endif + + unsigned max_capability = 0; + for ( std::list::iterator iter = cuobjdumpSectionList.begin(); + iter != cuobjdumpSectionList.end(); + iter++){ + 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); + if (max_capability == 0) max_capability=context->get_device()->get_gpgpu()->get_config().get_forced_max_capability(); + + cuobjdumpPTXSection* ptx = NULL; + const char* pre_load = getenv("CUOBJDUMP_SIM_FILE"); + if(pre_load==NULL || strlen(pre_load)==0) + ptx = findPTXSection(fname); + char *ptxcode; + const char *override_ptx_name = getenv("PTX_SIM_KERNELFILE"); + 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); + 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, max_capability ); + delete[] ptxplus_str; + } else { + symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxcode, 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 ); + } + 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 +} void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) { diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 97d5773..2064b13 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -152,6 +152,38 @@ char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptxfilenam return ptxplus_str; } +symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ) +{ + char buf[1024]; + snprintf(buf,1024,"_%u.ptx", source_num ); + if( g_save_embedded_ptx ) { + FILE *fp = fopen(buf,"w"); + fprintf(fp,"%s",p); + fclose(fp); + } + symbol_table *symtab=init_parser(buf); + ptx__scan_string(p); + int errors = ptx_parse (); + if ( errors ) { + char fname[1024]; + snprintf(fname,1024,"_ptx_errors_XXXXXX"); + int fd=mkstemp(fname); + close(fd); + printf("GPGPU-Sim PTX: parser error detected, exiting... but first extracting .ptx to \"%s\"\n", fname); + FILE *ptxfile = fopen(fname,"w"); + fprintf(ptxfile,"%s", p ); + fclose(ptxfile); + abort(); + exit(40); + } + + if ( g_debug_execution >= 100 ) + print_ptx_file(p,source_num,buf); + + printf("GPGPU-Sim PTX: finished parsing EMBEDDED .ptx file %s\n",buf); + return symtab; +} + symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ) { symbol_table *symtab=init_parser(filename); -- cgit v1.3 From bd9a49ef576bfbbaf2be1d163eb99b7316f7bd20 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Mon, 23 Jul 2018 14:20:14 -0700 Subject: brought back ptxinfo --- libcuda/cuda_runtime_api.cc | 19 +++++++++++++++---- src/cuda-sim/ptx_loader.cc | 23 ++++++++++++++++++++++- src/cuda-sim/ptx_loader.h | 1 + 3 files changed, 38 insertions(+), 5 deletions(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index db4f58e..c87c6c3 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -152,7 +152,8 @@ std::map pinned_memory; //support for pinned memories added std::map pinned_memory_size; std::map g_mallocPtr_Size; int no_of_ptx=0; -std::map > version_filename; +//maps sm version number to set of filenames +std::map > version_filename; extern void synchronize(); extern void exit_simulation(); @@ -1718,7 +1719,8 @@ char* get_app_binary_name(std::string abs_path){ } //extracts all ptx files from binary and dumps into prog_name.unique_no.sm_<>.ptx files -void extract_ptx_files_using_cuobjdump(bool g_cdp_enabled){ +void extract_ptx_files_using_cuobjdump(){ + extern bool g_cdp_enabled; char command[1000]; std::string app_binary = get_app_binary(); @@ -1808,7 +1810,8 @@ void extract_code_using_cuobjdump(){ //dump ptx for all individial ptx files into sepearte files which is later used by ptxas. int result=0; #if (CUDART_VERSION >= 6000) - extract_ptx_files_using_cuobjdump(g_cdp_enabled); + extract_ptx_files_using_cuobjdump(); + return; #endif //TODO: redundant to dump twice. how can it be prevented? //dump only for specific arch @@ -2168,7 +2171,7 @@ void cuobjdumpParseBinary(unsigned int handle){ #if (CUDART_VERSION >= 6000) //loops through all ptx files from smallest sm version to largest - std::map >::iterator itr_m; + std::map >::iterator itr_m; for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){ std::set::iterator itr_s; for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){ @@ -2181,6 +2184,14 @@ void cuobjdumpParseBinary(unsigned int handle){ context->add_binary(symtab, 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()); + for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){ + std::set::iterator itr_s; + for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){ + std::string ptx_filename = *itr_s; + printf("GPGPU-Sim PTX: Loading PTXInfo from %s\n",ptx_filename.c_str()); + gpgpu_ptx_info_load_from_filename( ptx_filename.c_str(), itr_m->first ); + } + } return; #endif diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 2064b13..f5a1c77 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -320,8 +320,29 @@ char* get_app_binary_name(){ return self_exe_path; } -void gpgpu_ptx_info_load_from_filename( const char *filename ) +void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version) { + std::string ptxas_filename(std::string(filename) + "as"); + char buff[1024], extra_flags[256]; + extra_flags[0]=0; + 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); + snprintf(buff,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file /dev/null 2> %s", + extra_flags, filename, ptxas_filename.c_str()); + int result = system(buff); + 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); + } + + g_ptxinfo_filename = strdup(ptxas_filename.c_str()); + ptxinfo_in = fopen(g_ptxinfo_filename,"r"); + ptxinfo_parse(); + fclose(ptxinfo_in); } void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version ) diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index c5c9dd8..e5df6a9 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -35,6 +35,7 @@ extern int no_of_ptx; //counter to track number of ptx files to be extracted in class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); class symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ); void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20 ); +void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version ); char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str); bool keep_intermediate_files(); -- cgit v1.3