From 584ebaa74a838680e6ed1fa13ac266e88c30c071 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Tue, 26 Jun 2018 13:20:39 -0700 Subject: exports and imports param data in new debug tool: WatchYourStep --- debug_tools/WatchYourStep/ptxjitplus/Makefile | 333 ++++++++++++++++++++++++++ 1 file changed, 333 insertions(+) create mode 100644 debug_tools/WatchYourStep/ptxjitplus/Makefile (limited to 'debug_tools/WatchYourStep/ptxjitplus/Makefile') diff --git a/debug_tools/WatchYourStep/ptxjitplus/Makefile b/debug_tools/WatchYourStep/ptxjitplus/Makefile new file mode 100644 index 0000000..d571567 --- /dev/null +++ b/debug_tools/WatchYourStep/ptxjitplus/Makefile @@ -0,0 +1,333 @@ +################################################################################ +# +# Copyright 1993-2015 NVIDIA Corporation. All rights reserved. +# +# NOTICE TO USER: +# +# This source code is subject to NVIDIA ownership rights under U.S. and +# international Copyright laws. +# +# NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE +# CODE FOR ANY PURPOSE. IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR +# IMPLIED WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH +# REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF +# MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE. +# IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL, +# OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS +# OF USE, DATA OR PROFITS, WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE +# OR OTHER TORTIOUS ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE +# OR PERFORMANCE OF THIS SOURCE CODE. +# +# U.S. Government End Users. This source code is a "commercial item" as +# that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting of +# "commercial computer software" and "commercial computer software +# documentation" as such terms are used in 48 C.F.R. 12.212 (SEPT 1995) +# and is provided to the U.S. Government only as a commercial end item. +# Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through +# 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the +# source code with only those rights set forth herein. +# +################################################################################ +# +# Makefile project only supported on Mac OS X and Linux Platforms) +# +################################################################################ + +# Location of the CUDA Toolkit +CUDA_PATH ?= "/home/jonathan/cuda8" + +############################## +# start deprecated interface # +############################## +ifeq ($(x86_64),1) + $(info WARNING - x86_64 variable has been deprecated) + $(info WARNING - please use TARGET_ARCH=x86_64 instead) + TARGET_ARCH ?= x86_64 +endif +ifeq ($(ARMv7),1) + $(info WARNING - ARMv7 variable has been deprecated) + $(info WARNING - please use TARGET_ARCH=armv7l instead) + TARGET_ARCH ?= armv7l +endif +ifeq ($(aarch64),1) + $(info WARNING - aarch64 variable has been deprecated) + $(info WARNING - please use TARGET_ARCH=aarch64 instead) + TARGET_ARCH ?= aarch64 +endif +ifeq ($(ppc64le),1) + $(info WARNING - ppc64le variable has been deprecated) + $(info WARNING - please use TARGET_ARCH=ppc64le instead) + TARGET_ARCH ?= ppc64le +endif +ifneq ($(GCC),) + $(info WARNING - GCC variable has been deprecated) + $(info WARNING - please use HOST_COMPILER=$(GCC) instead) + HOST_COMPILER ?= $(GCC) +endif +ifneq ($(abi),) + $(error ERROR - abi variable has been removed) +endif +############################ +# end deprecated interface # +############################ + +# architecture +HOST_ARCH := $(shell uname -m) +TARGET_ARCH ?= $(HOST_ARCH) +ifneq (,$(filter $(TARGET_ARCH),x86_64 aarch64 ppc64le armv7l)) + ifneq ($(TARGET_ARCH),$(HOST_ARCH)) + ifneq (,$(filter $(TARGET_ARCH),x86_64 aarch64 ppc64le)) + TARGET_SIZE := 64 + else ifneq (,$(filter $(TARGET_ARCH),armv7l)) + TARGET_SIZE := 32 + endif + else + TARGET_SIZE := $(shell getconf LONG_BIT) + endif +else + $(error ERROR - unsupported value $(TARGET_ARCH) for TARGET_ARCH!) +endif +ifneq ($(TARGET_ARCH),$(HOST_ARCH)) + ifeq (,$(filter $(HOST_ARCH)-$(TARGET_ARCH),aarch64-armv7l x86_64-armv7l x86_64-aarch64 x86_64-ppc64le)) + $(error ERROR - cross compiling from $(HOST_ARCH) to $(TARGET_ARCH) is not supported!) + endif +endif + +# When on native aarch64 system with userspace of 32-bit, change TARGET_ARCH to armv7l +ifeq ($(HOST_ARCH)-$(TARGET_ARCH)-$(TARGET_SIZE),aarch64-aarch64-32) + TARGET_ARCH = armv7l +endif + +# operating system +HOST_OS := $(shell uname -s 2>/dev/null | tr "[:upper:]" "[:lower:]") +TARGET_OS ?= $(HOST_OS) +ifeq (,$(filter $(TARGET_OS),linux darwin qnx android)) + $(error ERROR - unsupported value $(TARGET_OS) for TARGET_OS!) +endif + +# host compiler +ifeq ($(TARGET_OS),darwin) + ifeq ($(shell expr `xcodebuild -version | grep -i xcode | awk '{print $$2}' | cut -d'.' -f1` \>= 5),1) + HOST_COMPILER ?= clang++ + endif +else ifneq ($(TARGET_ARCH),$(HOST_ARCH)) + ifeq ($(HOST_ARCH)-$(TARGET_ARCH),x86_64-armv7l) + ifeq ($(TARGET_OS),linux) + HOST_COMPILER ?= arm-linux-gnueabihf-g++ + else ifeq ($(TARGET_OS),qnx) + ifeq ($(QNX_HOST),) + $(error ERROR - QNX_HOST must be passed to the QNX host toolchain) + endif + ifeq ($(QNX_TARGET),) + $(error ERROR - QNX_TARGET must be passed to the QNX target toolchain) + endif + export QNX_HOST + export QNX_TARGET + HOST_COMPILER ?= $(QNX_HOST)/usr/bin/arm-unknown-nto-qnx6.6.0eabi-g++ + else ifeq ($(TARGET_OS),android) + HOST_COMPILER ?= arm-linux-androideabi-g++ + endif + else ifeq ($(TARGET_ARCH),aarch64) + ifeq ($(TARGET_OS), linux) + HOST_COMPILER ?= aarch64-linux-gnu-g++ + else ifeq ($(TARGET_OS),qnx) + ifeq ($(QNX_HOST),) + $(error ERROR - QNX_HOST must be passed to the QNX host toolchain) + endif + ifeq ($(QNX_TARGET),) + $(error ERROR - QNX_TARGET must be passed to the QNX target toolchain) + endif + export QNX_HOST + export QNX_TARGET + HOST_COMPILER ?= $(QNX_HOST)/usr/bin/aarch64-unknown-nto-qnx7.0.0-g++ + else ifeq ($(TARGET_OS), android) + HOST_COMPILER ?= aarch64-linux-android-g++ + endif + else ifeq ($(TARGET_ARCH),ppc64le) + HOST_COMPILER ?= powerpc64le-linux-gnu-g++ + endif +endif +HOST_COMPILER ?= g++ +NVCC := $(CUDA_PATH)/bin/nvcc -ccbin $(HOST_COMPILER) + +# internal flags +NVCCFLAGS := -m${TARGET_SIZE} +CCFLAGS := +LDFLAGS := + +# build flags +ifeq ($(TARGET_OS),darwin) + LDFLAGS += -rpath $(CUDA_PATH)/lib + CCFLAGS += -arch $(HOST_ARCH) +else ifeq ($(HOST_ARCH)-$(TARGET_ARCH)-$(TARGET_OS),x86_64-armv7l-linux) + LDFLAGS += --dynamic-linker=/lib/ld-linux-armhf.so.3 + CCFLAGS += -mfloat-abi=hard +else ifeq ($(TARGET_OS),android) + LDFLAGS += -pie + CCFLAGS += -fpie -fpic -fexceptions +endif + +ifneq ($(TARGET_ARCH),$(HOST_ARCH)) + ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-linux) + ifneq ($(TARGET_FS),) + GCCVERSIONLTEQ46 := $(shell expr `$(HOST_COMPILER) -dumpversion` \<= 4.6) + ifeq ($(GCCVERSIONLTEQ46),1) + CCFLAGS += --sysroot=$(TARGET_FS) + endif + LDFLAGS += --sysroot=$(TARGET_FS) + LDFLAGS += -rpath-link=$(TARGET_FS)/lib + LDFLAGS += -rpath-link=$(TARGET_FS)/usr/lib + LDFLAGS += -rpath-link=$(TARGET_FS)/usr/lib/arm-linux-gnueabihf + endif + endif +endif + +# Debug build flags +ifeq ($(dbg),1) + NVCCFLAGS += -g -G + BUILD_TYPE := debug +else + BUILD_TYPE := release +endif + +ALL_CCFLAGS := +ALL_CCFLAGS += $(NVCCFLAGS) +ALL_CCFLAGS += $(EXTRA_NVCCFLAGS) +ALL_CCFLAGS += $(addprefix -Xcompiler ,$(CCFLAGS)) +ALL_CCFLAGS += $(addprefix -Xcompiler ,$(EXTRA_CCFLAGS)) + +UBUNTU = $(shell lsb_release -i -s 2>/dev/null | grep -i ubuntu) + +SAMPLE_ENABLED := 1 + +ALL_LDFLAGS := +ALL_LDFLAGS += $(ALL_CCFLAGS) +ALL_LDFLAGS += $(addprefix -Xlinker ,$(LDFLAGS)) +ALL_LDFLAGS += $(addprefix -Xlinker ,$(EXTRA_LDFLAGS)) + +# Common includes and paths for CUDA +INCLUDES := -I./inc +LIBRARIES := -ldl -lcudart + +################################################################################ + +# Gencode arguments +SMS ?= 20 30 35 37 50 52 60 61 + +ifeq ($(GENCODE_FLAGS),) +# Generate SASS code for each SM architecture listed in $(SMS) +$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) + +ifeq ($(SMS),) +# Generate PTX code from SM 20 +GENCODE_FLAGS += -gencode arch=compute_20,code=compute_20 +endif + +# Generate PTX code from the highest SM architecture in $(SMS) to guarantee forward-compatibility +HIGHEST_SM := $(lastword $(sort $(SMS))) +ifneq ($(HIGHEST_SM),) +GENCODE_FLAGS += -gencode arch=compute_$(HIGHEST_SM),code=compute_$(HIGHEST_SM) +endif +endif + +ifeq ($(TARGET_OS),darwin) + ALL_LDFLAGS += -Xcompiler -F/Library/Frameworks -Xlinker -framework -Xlinker CUDA +else + CUDA_SEARCH_PATH ?= + ifneq ($(TARGET_ARCH),$(HOST_ARCH)) + ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-linux) + ifneq ($(TARGET_FS),) + LIBRARIES += -L$(TARGET_FS)/usr/lib + CUDA_SEARCH_PATH += $(TARGET_FS)/usr/lib + CUDA_SEARCH_PATH += $(TARGET_FS)/usr/lib/arm-linux-gnueabihf + endif + CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/armv7-linux-gnueabihf/lib/stubs + CUDA_SEARCH_PATH += /usr/arm-linux-gnueabihf/lib + else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-linux) + CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/aarch64-linux/lib/stubs + else ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-android) + CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/armv7-linux-androideabi/lib/stubs + else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-android) + CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/aarch64-linux-androideabi/lib/stubs + else ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-qnx) + CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/ARMv7-linux-QNX/lib/stubs + else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-qnx) + CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/aarch64-qnx/lib/stubs + else ifeq ($(TARGET_ARCH)-$(TARGET_OS),ppc64le-linux) + CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/ppc64le-linux/lib/stubs + endif + else + ifneq ($(UBUNTU),) + CUDA_SEARCH_PATH += /usr/lib + else + CUDA_SEARCH_PATH += /usr/lib64 + endif + + ifeq ($(TARGET_ARCH),x86_64) + CUDA_SEARCH_PATH += $(CUDA_PATH)/lib64/stubs + endif + + ifeq ($(TARGET_ARCH),armv7l) + CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/armv7-linux-gnueabihf/lib/stubs + CUDA_SEARCH_PATH += /usr/lib/arm-linux-gnueabihf + endif + + ifeq ($(TARGET_ARCH),aarch64) + CUDA_SEARCH_PATH += /usr/lib + CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/aarch64-linux/lib/stubs + endif + + ifeq ($(TARGET_ARCH),ppc64le) + CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/ppc64le-linux/lib/stubs + CUDA_SEARCH_PATH += /usr/lib/powerpc64le-linux-gnu + endif + endif + + CUDALIB ?= $(shell find -L $(CUDA_SEARCH_PATH) -maxdepth 1 -name libcuda.so) + ifeq ("$(CUDALIB)","") + $(info >>> WARNING - libcuda.so not found, CUDA Driver is not installed. Please re-install the driver. <<<) + SAMPLE_ENABLED := 0 + endif + + LIBRARIES += -lcuda +endif + +LIBRARIES += -lcudart_static + +ifeq ($(SAMPLE_ENABLED),0) +EXEC ?= @echo "[@]" +endif + +################################################################################ + +# Target rules +all: build + +build: ptxjitplus + +check.deps: +ifeq ($(SAMPLE_ENABLED),0) + @echo "Sample will be waived due to the above missing dependencies" +else + @echo "Sample is ready - all dependencies have been met" +endif + +ptxjitplus.o:ptxjitplus.cpp + $(EXEC) $(NVCC) $(INCLUDES) --cudart shared $(ALL_CCFLAGS) $(GENCODE_FLAGS) -g -o $@ -c $< + +ptxjitplus: ptxjitplus.o ptxinst.o + $(EXEC) $(NVCC) $(ALL_LDFLAGS) --cudart shared $(GENCODE_FLAGS) -g -o $@ $+ $(LIBRARIES) + $(EXEC) mkdir -p ./bin/$(TARGET_ARCH)/$(TARGET_OS)/$(BUILD_TYPE) + $(EXEC) cp $@ ./bin/$(TARGET_ARCH)/$(TARGET_OS)/$(BUILD_TYPE) + +ptxinst.o:ptxinst.cpp + $(EXEC) $(NVCC) $(INCLUDES) --cudart shared $(ALL_CCFLAGS) $(GENCODE_FLAGS) -g -o $@ -c $< + +run: build + $(EXEC) ./ptxjitplus + +clean: + rm -f ptxjitplus ptxjitplus.o + rm -rf ./bin/$(TARGET_ARCH)/$(TARGET_OS)/$(BUILD_TYPE)/ptxjitplus + +clobber: clean -- cgit v1.3 From d072fa72c6ffce7f34520d24c8cb285e0d5b92b2 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Fri, 13 Jul 2018 10:46:47 -0700 Subject: Add more print statements to debug api calls Fixed bug with 32bit rem Fixed bug with launching kernels in ptxjitplus Removed extraneous files Cleaned up more pointers Added documentation for ptxjitplus --- debug_tools/WatchYourStep/ptxjitplus/Makefile | 5 +- debug_tools/WatchYourStep/ptxjitplus/ptxinst.cpp | 14 --- debug_tools/WatchYourStep/ptxjitplus/ptxinst.h | 14 --- .../WatchYourStep/ptxjitplus/ptxjitplus.cpp | 103 ++++++++++++--------- libcuda/cuda_runtime_api.cc | 36 +++++++ src/cuda-sim/cuda-sim.cc | 4 +- src/cuda-sim/instructions.cc | 16 +++- 7 files changed, 113 insertions(+), 79 deletions(-) delete mode 100644 debug_tools/WatchYourStep/ptxjitplus/ptxinst.cpp delete mode 100644 debug_tools/WatchYourStep/ptxjitplus/ptxinst.h (limited to 'debug_tools/WatchYourStep/ptxjitplus/Makefile') diff --git a/debug_tools/WatchYourStep/ptxjitplus/Makefile b/debug_tools/WatchYourStep/ptxjitplus/Makefile index d571567..b273ac2 100644 --- a/debug_tools/WatchYourStep/ptxjitplus/Makefile +++ b/debug_tools/WatchYourStep/ptxjitplus/Makefile @@ -315,14 +315,11 @@ endif ptxjitplus.o:ptxjitplus.cpp $(EXEC) $(NVCC) $(INCLUDES) --cudart shared $(ALL_CCFLAGS) $(GENCODE_FLAGS) -g -o $@ -c $< -ptxjitplus: ptxjitplus.o ptxinst.o +ptxjitplus: ptxjitplus.o $(EXEC) $(NVCC) $(ALL_LDFLAGS) --cudart shared $(GENCODE_FLAGS) -g -o $@ $+ $(LIBRARIES) $(EXEC) mkdir -p ./bin/$(TARGET_ARCH)/$(TARGET_OS)/$(BUILD_TYPE) $(EXEC) cp $@ ./bin/$(TARGET_ARCH)/$(TARGET_OS)/$(BUILD_TYPE) -ptxinst.o:ptxinst.cpp - $(EXEC) $(NVCC) $(INCLUDES) --cudart shared $(ALL_CCFLAGS) $(GENCODE_FLAGS) -g -o $@ -c $< - run: build $(EXEC) ./ptxjitplus diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxinst.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxinst.cpp deleted file mode 100644 index 6b39eb1..0000000 --- a/debug_tools/WatchYourStep/ptxjitplus/ptxinst.cpp +++ /dev/null @@ -1,14 +0,0 @@ -/* ptxinst.cpp - * Jonathan Lew - * University of British Columbia - */ -#include "ptxinst.h" - -void* instrument_ptx_from_function(std::string function, std::string path) -{ - return NULL; -} -void* instrument_ptx_from_string(std::string ptxcode) -{ - return NULL; -} diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxinst.h b/debug_tools/WatchYourStep/ptxjitplus/ptxinst.h deleted file mode 100644 index de2595e..0000000 --- a/debug_tools/WatchYourStep/ptxjitplus/ptxinst.h +++ /dev/null @@ -1,14 +0,0 @@ -/* ptxinst.h - * Jonathan Lew - * University of British Columbia - */ - -#ifndef _PTXINST_H_ -#define _PTXINST_H_ - -#include - -void* instrument_ptx_from_function(std::string function, std::string path); -void* instrument_ptx_from_string(std::string ptxcode); - -#endif diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp index 68964d8..b7f9f2d 100644 --- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp +++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp @@ -1,5 +1,4 @@ -/** - * Copyright 1993-2015 NVIDIA Corporation. All rights reserved. +/** Copyright 1993-2015 NVIDIA Corporation. All rights reserved. * * Please refer to the NVIDIA end user license agreement (EULA) associated * with this source code for terms and conditions that govern your use of @@ -7,9 +6,6 @@ * this software and related documentation outside the terms of the EULA * is strictly prohibited. * - */ - -/* * This sample uses the Driver API to just-in-time compile (JIT) a Kernel from PTX code. * Additionally, this sample demonstrates the seamless interoperability capability of CUDA runtime * Runtime and CUDA Driver API calls. @@ -17,6 +13,32 @@ * */ +/** + * Modified by: Jonathan Lew + * PTX JIT PLUS + * + * ********** + * User Guide + * ********** + * + * Welcome to WatchYourStep, a debugging tool that allows you launch individual + * kernels using parameters captured from cudaLaunch and outputs the values in + * the arrays from the kernel. It allows you to watch each step you program takes, + * kernel by kernel. + * + * 1. Set environment variables to create params.config* and ptx.config* files. + * a)export PTX_SIM_DEBUG=4 + * b)export PTX_JIT_PATH=[path to this file] + * c)export WYS_EXEC_PATH=[path to executable (program to debug)] + * d)export WYS_EXEC_NAME=[name of executable (program to debug)] + * e)Make sure all GPGPU-Sim path variables are set (see GPGPU-Sim documentation) + * 2. Run executable (program to debug) using GPGPU-Sim + * 3. export PTX_SIM_DEBUG=[less than 4 to not dump config files again] + * 4-1. Run one kernel at a time: export WYS_LAUNCH_NUM=[kernel to launch] and compile ptxjitplus and run ptxjitplus + * 4-2. Run all kernels: compile and run ". launchkernels 0 [max number of kernels]" in terminal + * 5. Find output in ../data/wys.out* where * is the launch number + */ + // System includes #include #include @@ -34,12 +56,12 @@ // sample include #include "ptxjitplus.h" -#include "ptxinst.h" -const char *sSDKname = "PTX Just In Time (JIT) Compilation (no-qatest)"; +const char *sSDKname = "PTX Just In Time (JIT) Compilation Plus"; char *wys_exec_path; char *wys_exec_name; char *wys_launch_num; +bool gpgpusim = false; dim3 gridDim, blockDim; std::string kernelName; @@ -116,6 +138,10 @@ void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUl void initializeData(std::vector& param_data, std::vector< std::pair >& param_info) { + char *gpgpusim_env = getenv("GPGPUSIM_SETUP_ENVIRONMENT_WAS_RUN"); + if (gpgpusim_env!=NULL&&gpgpusim_env[0] == '1'){ + gpgpusim=true; + } wys_exec_path = getenv("WYS_EXEC_PATH"); assert(wys_exec_path!=NULL); wys_exec_name = getenv("WYS_EXEC_NAME"); @@ -150,32 +176,19 @@ void initializeData(std::vector& param_data, std::vector< std::p err = fscanf(fin, "%lu : ", &len); info.first = len; assert( err==1 ); - //printf("%lu : ", len); unsigned char* params = (unsigned char*) malloc(len*sizeof(unsigned char)); for (size_t i=0; isecond[j]); - //} - //fprintf(fout, "\n"); - //fflush(fout); - //fclose(fout); } int main(int argc, char **argv) @@ -249,20 +262,28 @@ int main(int argc, char **argv) checkCudaErrors(cudaFree(d_tmp)); //maps param number to pointer to device data - std::map< size_t, unsigned char* > m_device_data; + std::map< size_t, void* > m_device_data; + std::map< size_t, void* > m_cleanup; + void * paramKernels[param_data.size()]; //Initialize param_data for kernel int paramOffset = 0; for( size_t i = 0; i m_output_data; - for(std::map< size_t, unsigned char* >::iterator i = m_device_data.begin(); i!=m_device_data.end(); i++){ + for(std::map< size_t, void* >::iterator i = m_device_data.begin(); i!=m_device_data.end(); i++){ unsigned char *h_data = 0; if ((h_data = (unsigned char *)malloc(param_info[i->first].first)) == NULL) { @@ -294,40 +315,32 @@ int main(int argc, char **argv) for(std::map< size_t, unsigned char* >::iterator i = m_output_data.begin(); i!=m_output_data.end(); i++){ fprintf(fout, "param %zu: size = %zu, data = ", i->first,param_info[i->first].first); for (size_t j = 0; jfirst].first; j++){ - fprintf(fout, " %u", i->second[j]); - if (j&&(!(j%20))){ + if (!(j%24)){ fprintf(fout, "\n"); } + fprintf(fout, " %u", i->second[j]); } fprintf(fout, "\n"); } fflush(fout); fclose(fout); -// int* h_data = (int*) m_output_data[0]; -// // Check the result -// bool dataGood = true; -// -// for (unsigned int i = 0 ; dataGood && i < nBlocks * nThreads ; i++) -// { -// if (h_data[i] != (int)i) -// { -// std::cerr << "Error at " << i << std::endl; -// dataGood = false; -// } -// } -// if(dataGood){ -// std::cout<<"OK!"<::iterator i = m_device_data.begin(); i!=m_device_data.end(); i++){ + for(std::map< size_t, void* >::iterator i = m_device_data.begin(); i!=m_device_data.end(); i++){ if (i->second){ checkCudaErrors(cudaFree(i->second)); i->second = 0; } } + for(std::map< size_t, void* >::iterator i = m_cleanup.begin(); i!=m_cleanup.end(); i++){ + if (i->second){ + free(i->second); + i->second = 0; + } + } + for(std::map< size_t, unsigned char* >::iterator i = m_output_data.begin(); i!=m_output_data.begin(); i++){ if (i->second) { @@ -342,5 +355,5 @@ int main(int argc, char **argv) hModule = 0; } - //return dataGood ? EXIT_SUCCESS : EXIT_FAILURE; + return 0; } diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 68b4017..866fa3b 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -2347,11 +2347,17 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) void __cudaUnregisterFatBinary(void **fatCubinHandle) { + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } ; } cudaError_t cudaDeviceReset ( void ) { // Should reset the simulated GPU + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } return g_last_cudaError = cudaSuccess; } cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ @@ -2397,6 +2403,9 @@ extern void __cudaRegisterVar( int constant, int global ) { + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } printf("GPGPU-Sim PTX: __cudaRegisterVar: hostVar = %p; deviceAddress = %s; deviceName = %s\n", hostVar, deviceAddress, deviceName); printf("GPGPU-Sim PTX: __cudaRegisterVar: Registering const memory space of %d bytes\n", size); if(GPGPUSim_Context()->get_device()->get_gpgpu()->get_config().use_cuobjdump()) @@ -2415,6 +2424,9 @@ void __cudaRegisterShared( void **devicePtr ) { + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } // we don't do anything here printf("GPGPU-Sim PTX: __cudaRegisterShared\n" ); } @@ -2444,6 +2456,9 @@ void __cudaRegisterTexture( int ext ) //passes in a newly created textureReference { + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } std::string devStr (deviceName); #if (CUDART_VERSION > 4020) if (devStr.size() > 2 && devStr.data()[0] == ':' && devStr.data()[1] == ':') @@ -2465,6 +2480,9 @@ typedef unsigned long GLuint; cudaError_t cudaGLRegisterBufferObject(GLuint bufferObj) { + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ ); return g_last_cudaError = cudaSuccess; } @@ -2481,6 +2499,9 @@ glbmap_entry_t* g_glbmap = NULL; cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj) { + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } #ifdef OPENGL_SUPPORT GLint buffer_size=0; CUctx_st* ctx = GPGPUSim_Context(); @@ -2534,6 +2555,9 @@ cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj) cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj) { + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } #ifdef OPENGL_SUPPORT glbmap_entry_t *p = g_glbmap; while ( p && p->m_bufferObj != bufferObj ) @@ -2558,6 +2582,9 @@ cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj) cudaError_t cudaGLUnregisterBufferObject(GLuint bufferObj) { + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ ); return g_last_cudaError = cudaSuccess; } @@ -2783,6 +2810,9 @@ extern FILE *ptxinfo_in; static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu ) { + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } printf( "GPGPU-Sim PTX: loading globals with explicit initializers... \n" ); fflush(stdout); int ng_bytes=0; @@ -2819,6 +2849,9 @@ static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsign static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu ) { + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } printf( "GPGPU-Sim PTX: loading constants with explicit initializers... " ); fflush(stdout); int nc_bytes = 0; @@ -2866,6 +2899,9 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, struct dim3 blockDim, CUctx_st* context ) { + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } function_info *entry = context->get_kernel(hostFun); kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry); if( entry == NULL ) { diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index df8d806..d2f096f 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1309,7 +1309,9 @@ void function_info::ptx_jit_config(std::map mallocPt do{ fprintf(fout, "%s", buff); fgets(buff, 1024, fin); - assert(!feof(fin)); + if(feof(fin)){ + break; + } } while(strstr(buff, "entry")==NULL); fclose(fin); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 034a7b9..9e18772 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -3190,7 +3190,21 @@ void rem_impl( const ptx_instruction *pI, ptx_thread_info *thread ) src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1); - data.u64 = src1_data.u64 % src2_data.u64; + switch ( i_type ) { + case S32_TYPE: + data.s32 = src1_data.s32 % src2_data.s32; + break; + case S64_TYPE: + data.s64 = src1_data.s64 % src2_data.s64; + break; + case U32_TYPE: + data.u32 = src1_data.u32 % src2_data.u32; + break; + case U64_TYPE: + data.u64 = src1_data.u64 % src2_data.u64; + break; + default: assert(0); break; + } thread->set_operand_value(dst,data, i_type, thread, pI); } -- cgit v1.3 From dd7b243078561235a1be3a7d4a8bf116b12684b7 Mon Sep 17 00:00:00 2001 From: J Date: Fri, 10 Aug 2018 16:36:02 -0700 Subject: Update cuda path in Makefile --- debug_tools/WatchYourStep/ptxjitplus/Makefile | 27 +++++++++++++-------------- 1 file changed, 13 insertions(+), 14 deletions(-) (limited to 'debug_tools/WatchYourStep/ptxjitplus/Makefile') diff --git a/debug_tools/WatchYourStep/ptxjitplus/Makefile b/debug_tools/WatchYourStep/ptxjitplus/Makefile index b273ac2..3cca514 100644 --- a/debug_tools/WatchYourStep/ptxjitplus/Makefile +++ b/debug_tools/WatchYourStep/ptxjitplus/Makefile @@ -34,7 +34,6 @@ ################################################################################ # Location of the CUDA Toolkit -CUDA_PATH ?= "/home/jonathan/cuda8" ############################## # start deprecated interface # @@ -148,7 +147,7 @@ else ifneq ($(TARGET_ARCH),$(HOST_ARCH)) endif endif HOST_COMPILER ?= g++ -NVCC := $(CUDA_PATH)/bin/nvcc -ccbin $(HOST_COMPILER) +NVCC := $(CUDA_INSTALL_PATH)/bin/nvcc -ccbin $(HOST_COMPILER) # internal flags NVCCFLAGS := -m${TARGET_SIZE} @@ -157,7 +156,7 @@ LDFLAGS := # build flags ifeq ($(TARGET_OS),darwin) - LDFLAGS += -rpath $(CUDA_PATH)/lib + LDFLAGS += -rpath $(CUDA_INSTALL_PATH)/lib CCFLAGS += -arch $(HOST_ARCH) else ifeq ($(HOST_ARCH)-$(TARGET_ARCH)-$(TARGET_OS),x86_64-armv7l-linux) LDFLAGS += --dynamic-linker=/lib/ld-linux-armhf.so.3 @@ -241,20 +240,20 @@ else CUDA_SEARCH_PATH += $(TARGET_FS)/usr/lib CUDA_SEARCH_PATH += $(TARGET_FS)/usr/lib/arm-linux-gnueabihf endif - CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/armv7-linux-gnueabihf/lib/stubs + CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/armv7-linux-gnueabihf/lib/stubs CUDA_SEARCH_PATH += /usr/arm-linux-gnueabihf/lib else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-linux) - CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/aarch64-linux/lib/stubs + CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/aarch64-linux/lib/stubs else ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-android) - CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/armv7-linux-androideabi/lib/stubs + CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/armv7-linux-androideabi/lib/stubs else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-android) - CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/aarch64-linux-androideabi/lib/stubs + CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/aarch64-linux-androideabi/lib/stubs else ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-qnx) - CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/ARMv7-linux-QNX/lib/stubs + CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/ARMv7-linux-QNX/lib/stubs else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-qnx) - CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/aarch64-qnx/lib/stubs + CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/aarch64-qnx/lib/stubs else ifeq ($(TARGET_ARCH)-$(TARGET_OS),ppc64le-linux) - CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/ppc64le-linux/lib/stubs + CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/ppc64le-linux/lib/stubs endif else ifneq ($(UBUNTU),) @@ -264,21 +263,21 @@ else endif ifeq ($(TARGET_ARCH),x86_64) - CUDA_SEARCH_PATH += $(CUDA_PATH)/lib64/stubs + CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/lib64/stubs endif ifeq ($(TARGET_ARCH),armv7l) - CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/armv7-linux-gnueabihf/lib/stubs + CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/armv7-linux-gnueabihf/lib/stubs CUDA_SEARCH_PATH += /usr/lib/arm-linux-gnueabihf endif ifeq ($(TARGET_ARCH),aarch64) CUDA_SEARCH_PATH += /usr/lib - CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/aarch64-linux/lib/stubs + CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/aarch64-linux/lib/stubs endif ifeq ($(TARGET_ARCH),ppc64le) - CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/ppc64le-linux/lib/stubs + CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/ppc64le-linux/lib/stubs CUDA_SEARCH_PATH += /usr/lib/powerpc64le-linux-gnu endif endif -- cgit v1.3