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 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