aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--Jenkinsfile120
-rw-r--r--README92
-rw-r--r--bitbucket-pipelines.yml15
-rw-r--r--configs/GeForceGTX1080Ti/gpgpusim.config2
-rw-r--r--libcuda/cuda_runtime_api.cc413
-rw-r--r--setup_environment2
-rw-r--r--src/abstract_hardware_model.h3
-rw-r--r--src/cuda-sim/cuda-math.h3
-rw-r--r--src/cuda-sim/cuda-sim.cc27
-rw-r--r--src/cuda-sim/cuda_device_runtime.cc14
-rw-r--r--src/cuda-sim/instructions.cc18
-rw-r--r--src/cuda-sim/ptx.l7
-rw-r--r--src/cuda-sim/ptx_ir.cc41
-rw-r--r--src/cuda-sim/ptx_ir.h21
-rw-r--r--src/cuda-sim/ptx_loader.cc201
-rw-r--r--src/cuda-sim/ptx_loader.h1
-rw-r--r--src/gpgpu-sim/gpu-sim.cc8
-rw-r--r--src/gpgpu-sim/gpu-sim.h1
-rw-r--r--src/gpgpu-sim/shader.h2
19 files changed, 837 insertions, 154 deletions
diff --git a/Jenkinsfile b/Jenkinsfile
new file mode 100644
index 0000000..57fc6c3
--- /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 -fr ./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/README b/README
index 6e2d734..bf5aa62 100644
--- a/README
+++ b/README
@@ -235,6 +235,14 @@ The documentation resides at doc/doxygen/html.
Step 3: Run
============
+Before we run, we need to make sure the application's executable file is dynamically linked to CUDA runtime library. This can be done during compilation of your program by introducing the nvcc flag "--cudart shared" in makefile (quotes should be excluded).
+
+To confirm the same, type the follwoing command:
+
+ldd <your_application_name>
+
+You should see that your application is using libcudart.so file in GPGPUSim directory.
+
Copy the contents of configs/QuadroFX5800/ or configs/GTX480/ to your
application's working directory. These files configure the microarchitecture
models to resemble the respective GPGPU architectures.
@@ -346,3 +354,87 @@ 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.
+
+
+4. MISCELLANEOUS
+
+4.1 Speeding up the execution
+
+Some applications take several hours to execute on GPGPUSim. This is because the simulator has to dump the PTX, analyze them and get resource usage statistics. This can be avoided everytime we execute the program in the following way:
+
+Step 1: Execute the program by enabling “-save_embedded_ptx 1” in config file, execute the code and let cuobjdump command dump all necessary files. After this process, you will get 2 new files namely: _cuobjdump_complete_output_<some_random_name> and _1.ptx
+
+Step 2: Create new environment variables or include the below in your .bashrc file:
+ a. export PTX_SIM_USE_PTX_FILE=_1.ptx
+ b. export PTX_SIM_KERNELFILE=_1.ptx
+ c. export CUOBJDUMP_SIM_FILE=_cuobjdump_complete_output_<some_random_name>
+
+Step 3: Disable -save_embedded_ptx flag, execute the code again. This will skip the dumping by cuobjdump and directly goes to executing the program thus saving time.
+
+
+4.2 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.
+
+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/<YOUR GITHUB USERNAME>/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 <DNS_IP_ADDRESS> below. On Ubuntu run "ifconfig" to lookup the network interface connecting your computer to the network. Then run "nmcli device show <interface name>" to find the IP address of your DNS server. Modify the above command to include "--dns <DNS_IP_ADDRESS>" after "run", E.g.,
+
+ docker run --dns <DNS_IP_ADDRESS> --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 <CONTAINER_ID> /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.<DATE>.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 <simulation_directory>
+
+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<number> where <number> 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/bitbucket-pipelines.yml b/bitbucket-pipelines.yml
new file mode 100644
index 0000000..6e331d4
--- /dev/null
+++ b/bitbucket-pipelines.yml
@@ -0,0 +1,15 @@
+# This is a sample build configuration for C++ � Make.
+# Check our guides at https://confluence.atlassian.com/x/5Q4SMw for more examples.
+# Only use spaces to indent your .yml configuration.
+# -----
+# You can specify a custom docker image from Docker Hub as your build environment.
+image: tgrogers/gpgpu-sim_regress:latest
+
+pipelines:
+ default:
+ - step:
+ script: # Modify the commands below to build your repository.
+ - docker run -v `pwd`:/home/runner/gpgpu-sim_distribution:rw tgrogers/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'"
+ services:
+ - docker
+ \ No newline at end of file
diff --git a/configs/GeForceGTX1080Ti/gpgpusim.config b/configs/GeForceGTX1080Ti/gpgpusim.config
index 47c2b6a..a5ada34 100644
--- a/configs/GeForceGTX1080Ti/gpgpusim.config
+++ b/configs/GeForceGTX1080Ti/gpgpusim.config
@@ -3,7 +3,7 @@
# functional simulator specification
-gpgpu_ptx_instruction_classification 0
-gpgpu_ptx_sim_mode 0
--gpgpu_ptx_force_max_capability 20
+-gpgpu_ptx_force_max_capability 61
# SASS execution (only supported with CUDA >= 4.0)
-gpgpu_ptx_convert_to_ptxplus 0
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 9bdb993..ef46f00 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();
@@ -345,7 +349,24 @@ 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 = the_gpu->shared_mem_per_block();
+
+ #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();
@@ -474,6 +495,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 +787,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 +849,82 @@ __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 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;
+ case 6:
+ *value= prop->maxGridSize[1];
+ break;
+ 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;
+ case 12:
+ *value= prop->regsPerBlock;
+ break;
+ case 14:
+ *value= prop->textureAlignment ;
+ break;
+ case 16:
+ *value= prop->multiProcessorCount ;
+ break;
+ case 34:
+ *value= 0;
+ break;
+ case 39:
+ *value= dev->get_gpgpu()->threads_per_core();
+ break;
+ case 75:
+ *value= 9 ;
+ break;
+ case 76:
+ *value= 3 ;
+ break;
+ case 78:
+ *value= 0 ; //TODO: as of now, we dont support stream priorities.
+ break;
+ case 81:
+ *value= prop->sharedMemPerMultiprocessor;
+ break;
+ case 82:
+ *value= prop->regsPerBlock;
+ break;
+ default:
+ printf("ERROR: Attribute number %d unimplemented \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 +1078,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,6 +1117,15 @@ __host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream)
return g_last_cudaError = cudaSuccess;
}
+//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 *stream, unsigned int flags) {
return cudaStreamCreate(stream);
}
@@ -1341,6 +1468,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 +1497,163 @@ 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.
+ 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);
+ }
+ }
+ }
+#endif
+ //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 %s > %s", app_binary.c_str(), fname);
else
- snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass -all %s > %s", app_binary.c_str(), fname);
+ 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.str() << 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.str() << 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.str() << 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.str() << 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 +1885,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;
@@ -1715,12 +1923,16 @@ 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);
+ if (max_capability == 0) 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 +1952,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 );
}
@@ -2077,6 +2290,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
@@ -2085,8 +2301,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 3f2476e..a9813c2 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 7125b6b..f5708bc 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-math.h b/src/cuda-sim/cuda-math.h
index a3db0df..afac330 100644
--- a/src/cuda-sim/cuda-math.h
+++ b/src/cuda-sim/cuda-math.h
@@ -67,6 +67,8 @@
#ifndef CUDA_MATH
#define CUDA_MATH
+#include <cmath>
+
// cuda math implementations
#undef max
#undef min
@@ -148,6 +150,7 @@ float __ll2float_rd(long long int a) {
#include <device_types.h>
#include <fenv.h>
+
// 32-bit integer to float
float __int2float_rn(int a) {
int orig_rnd_mode = fegetround();
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index a34b99b..946043a 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() ) {
@@ -250,11 +252,13 @@ void function_info::ptx_assemble()
target.set_type(label_t);
}
}
-
+ m_n = n;
printf(" done.\n");
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
+ printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", m_name.c_str() );
create_basic_blocks();
connect_basic_blocks();
bool modified = false;
@@ -288,6 +292,7 @@ void function_info::ptx_assemble()
fflush(stdout);
m_assembled = true;
+#endif
}
addr_t shared_to_generic( unsigned smid, addr_t addr )
@@ -1134,8 +1139,13 @@ void function_info::finalize( memory_space *param_mem )
}
// copy the parameter over word-by-word so that parameter that crosses a memory page can be copied over
//Jin: copy parameter using aligned rules
+ const type_info *paramtype = param->type();
+ int align_amount = paramtype->get_key().get_alignment_spec();
+ align_amount = (align_amount == -1) ? size : align_amount;
+ param_address = (param_address + align_amount - 1) / align_amount * align_amount; //aligned
+
const size_t word_size = 4;
- param_address = (param_address + size - 1) / size * size; //aligned with size
+ //param_address = (param_address + size - 1) / size * size; //aligned with size
for (size_t idx = 0; idx < size; idx += word_size) {
const char *pdata = reinterpret_cast<const char*>(param_value.pdata) + idx; // cast to char * for ptr arithmetic
param_mem->write(param_address + idx, word_size, pdata,NULL,NULL);
@@ -1799,6 +1809,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/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc
index 4a8ffe5..b399133 100644
--- a/src/cuda-sim/cuda_device_runtime.cc
+++ b/src/cuda-sim/cuda_device_runtime.cc
@@ -177,6 +177,20 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info *
//device_grid = op.grid;
device_kernel_entry = config.entry;
DEV_RUNTIME_REPORT("find device kernel " << device_kernel_entry->get_name());
+
+ //PDOM analysis is done for Parent kernel but not for child kernel.
+ if (device_kernel_entry->is_pdom_set()) {
+ printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", device_kernel_entry->get_name().c_str() );
+ } else {
+ printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", device_kernel_entry->get_name().c_str() );
+ /*
+ * Some of the instructions like printf() gives the gpgpusim the wrong impression that it is a function call.
+ * As printf() doesnt have a body like functions do, doing pdom analysis for printf() causes a crash.
+ */
+ if (device_kernel_entry->get_function_size() >0)
+ device_kernel_entry->do_pdom();
+ device_kernel_entry->set_pdom();
+ }
//copy data in parameter_buffer to device kernel param memory
unsigned device_kernel_arg_size = device_kernel_entry->get_args_aligned_size();
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 71286c9..e53aaab 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -154,7 +154,9 @@ ptx_reg_t ptx_thread_info::get_operand_value( const operand_info &op, operand_in
} else if ( op.is_local() ) {
result.u64 = op.get_symbol()->get_address();
} else if ( op.is_function_address() ) {
- result.u64 = (size_t)op.get_symbol()->get_pc();
+ result.u64 = (size_t)op.get_symbol()->get_pc();
+ } else if ( op.is_param_kernel()) {
+ result.u64 = op.get_symbol()->get_address();
} else {
const char *name = op.name().c_str();
printf("GPGPU-Sim PTX: ERROR ** get_operand_value : unknown operand type for %s\n", name );
@@ -1466,7 +1468,19 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread )
const operand_info &target = pI->func_addr();
assert( target.is_function_address() );
const symbol *func_addr = target.get_symbol();
- const function_info *target_func = func_addr->get_pc();
+ function_info *target_func = func_addr->get_pc();
+ if (target_func->is_pdom_set()) {
+ printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", target_func->get_name().c_str() );
+ } else {
+ printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", target_func->get_name().c_str() );
+ /*
+ * Some of the instructions like printf() gives the gpgpusim the wrong impression that it is a function call.
+ * As printf() doesnt have a body like functions do, doing pdom analysis for printf() causes a crash.
+ */
+ if (target_func->get_function_size() >0)
+ target_func->do_pdom();
+ target_func->set_pdom();
+ }
// check that number of args and return match function requirements
if( pI->has_return() ^ target_func->has_return() ) {
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index 5471d6f..908c5be 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -36,7 +36,8 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#include "ptx.tab.h"
#include <string.h>
-char linebuf[1024];
+#define LINEBUF_SIZE (64*1024)
+char linebuf[LINEBUF_SIZE];
unsigned col = 0;
#define TC col+=strlen(ptx_text);
#define CHECK_UNSIGNED \
@@ -162,6 +163,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 +235,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;
@@ -382,7 +385,7 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
"//"[^\n]* TC; // eat single
-\n.* col=0; strncpy(linebuf, yytext + 1, 1024); yyless( 1 );
+\n.* col=0; strncpy(linebuf, yytext + 1, LINEBUF_SIZE); yyless( 1 );
" " TC;
"\t" TC;
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index 8ebdcf8..be25dbe 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,40 @@ 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();
+ }
+ printf("GPGPU-Sim PTX: pre-decoding instructions for \'%s\'...\n", m_name.c_str() );
+ for ( unsigned ii=0; ii < m_n; ii += m_instr_mem[ii]->inst_size() ) { // handle branch instructions
+ ptx_instruction *pI = m_instr_mem[ii];
+ pI->pre_decode();
+ }
+ printf("GPGPU-Sim PTX: ... done pre-decoding instructions for \'%s\'.\n", m_name.c_str() );
+ fflush(stdout);
+ m_assembled = true;
+}
void intersect( std::set<int> &A, const std::set<int> &B )
{
// return intersection of A and B in A
@@ -1305,6 +1341,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..58d5f49 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -91,6 +91,7 @@ public:
bool is_tex() const { return m_space_spec == tex_space;}
bool is_func_addr() const { return m_is_function?true:false; }
int scalar_type() const { return m_scalar_type_spec;}
+ int get_alignment_spec() const { return m_alignment_spec;}
unsigned type_decode( size_t &size, int &t ) const;
static unsigned type_decode( int type, size_t &size, int &t );
memory_space_t get_memory_space() const { return m_space_spec; }
@@ -163,6 +164,7 @@ public:
m_is_global = false;
m_is_local = false;
m_is_param_local = false;
+ m_is_param_kernel = false;
m_is_tex = false;
m_is_func_addr = false;
m_reg_num_valid = false;
@@ -176,6 +178,7 @@ public:
if ( type ) m_is_global = type->get_key().is_global();
if ( type ) m_is_local = type->get_key().is_local();
if ( type ) m_is_param_local = type->get_key().is_param_local();
+ if ( type ) m_is_param_kernel = type->get_key().is_param_kernel();
if ( type ) m_is_tex = type->get_key().is_tex();
if ( type ) m_is_func_addr = type->get_key().is_func_addr();
}
@@ -226,6 +229,7 @@ public:
bool is_global() const { return m_is_global;}
bool is_local() const { return m_is_local;}
bool is_param_local() const { return m_is_param_local; }
+ bool is_param_kernel() const { return m_is_param_kernel; }
bool is_tex() const { return m_is_tex;}
bool is_func_addr() const { return m_is_func_addr; }
bool is_reg() const
@@ -283,6 +287,7 @@ private:
bool m_is_global;
bool m_is_local;
bool m_is_param_local;
+ bool m_is_param_kernel;
bool m_is_tex;
bool m_is_func_addr;
unsigned m_reg_num;
@@ -399,6 +404,8 @@ public:
m_type = symbolic_t;
} else if ( addr->is_param_local() ) {
m_type = symbolic_t;
+ } else if ( addr->is_param_kernel() ) {
+ m_type = symbolic_t;
} else if ( addr->is_tex() ) {
m_type = symbolic_t;
} else if ( addr->is_func_addr() ) {
@@ -675,6 +682,13 @@ public:
return m_value.m_symbolic->type()->get_key().is_param_local();
}
+ bool is_param_kernel() const
+ {
+ if ( m_type != symbolic_t )
+ return false;
+ return m_value.m_symbolic->type()->get_key().is_param_kernel();
+ }
+
bool is_vector() const
{
if ( m_vector) return true;
@@ -1178,7 +1192,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 +1288,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 +1297,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;
@@ -1305,6 +1322,8 @@ private:
//parameter size for device kernels
int m_args_aligned_size;
+
+ addr_t m_n; // offset in m_instr_mem (used in do_pdom)
};
class arg_buffer_t {
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 6c1b595..33a4260 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,82 @@ 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)
+ #if CUDART_VERSION >= 3000
+ if (sm_version == 0) sm_version = 20;
snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version);
- else
- snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",sm_version);
-#endif
+ #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 +372,100 @@ 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
+ 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);
+ fflush(stdout);
+ 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);
+ }
+ }
+
+ if(no_of_ptx>0)
+ g_ptxinfo_filename = final_tempfile_ptxinfo;
+ else
+ g_ptxinfo_filename = tempfile_ptxinfo;
+ ptxinfo_in = fopen(g_ptxinfo_filename,"r");
+
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 ) {
- 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", fname, fname2, final_tempfile_ptxinfo);
+ else
+ snprintf(commandline,1024,"rm -f %s %s %s", 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 );
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index bd4c00c..892b245 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -274,6 +274,9 @@ void shader_core_config::reg_options(class OptionParser * opp)
option_parser_register(opp, "-gpgpu_n_ldst_response_buffer_size", OPT_UINT32, &ldst_unit_response_queue_size,
"number of response packets in ld/st unit ejection buffer",
"2");
+ option_parser_register(opp, "-gpgpu_shmem_per_block", OPT_UINT32, &gpgpu_shmem_per_block,
+ "Size of shared memory per thread block or CTA (default 48kB)",
+ "49152");
option_parser_register(opp, "-gpgpu_shmem_size", OPT_UINT32, &gpgpu_shmem_size,
"Size of shared memory per shader core (default 16kB)",
"16384");
@@ -670,6 +673,11 @@ int gpgpu_sim::shared_mem_size() const
return m_shader_config->gpgpu_shmem_size;
}
+int gpgpu_sim::shared_mem_per_block() const
+{
+ return m_shader_config->gpgpu_shmem_per_block;
+}
+
int gpgpu_sim::num_registers_per_core() const
{
return m_shader_config->gpgpu_shader_registers;
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index 7d92c66..8d1c4fc 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -389,6 +389,7 @@ public:
void get_pdom_stack_top_info( unsigned sid, unsigned tid, unsigned *pc, unsigned *rpc );
int shared_mem_size() const;
+ int shared_mem_per_block() const;
int num_registers_per_core() const;
int wrp_size() const;
int shader_clock() const;
diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h
index bdd8dbe..fbddd18 100644
--- a/src/gpgpu-sim/shader.h
+++ b/src/gpgpu-sim/shader.h
@@ -1285,7 +1285,7 @@ struct shader_core_config : public core_config
unsigned max_cta_per_core; //Limit on number of concurrent CTAs in shader core
unsigned max_barriers_per_cta;
char * gpgpu_scheduler_string;
-
+ unsigned gpgpu_shmem_per_block;
char* pipeline_widths_string;
int pipe_widths[N_PIPELINE_STAGES];
generated by cgit v1.3 (git 2.53.0) at 2026-07-21 12:00:23 +0000