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/ptxjitplus.h | 77 +++++++++++++++++++++++ 1 file changed, 77 insertions(+) create mode 100644 debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h (limited to 'debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h') diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h new file mode 100644 index 0000000..8fb5611 --- /dev/null +++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h @@ -0,0 +1,77 @@ +/* + * 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 + * this software. Any use, reproduction, disclosure, or distribution of + * this software and related documentation outside the terms of the EULA + * is strictly prohibited. + * + */ + +#ifndef _PTXJIT_H_ +#define _PTXJIT_H_ + +/* + * PTX is equivalent to the following kernel: + * + * __global__ void myKernel(int *data) + * { + * int tid = blockIdx.x * blockDim.x + threadIdx.x; + * data[tid] = tid; + * } + * + */ + +char myPtx64[] = "\n\ +.version 3.2\n\ +.target sm_20\n\ +.address_size 64\n\ +.visible .entry _Z8myKernelPi(\n\ + .param .u64 _Z8myKernelPi_param_0\n\ +)\n\ +{\n\ + .reg .s32 %r<5>;\n\ + .reg .s64 %rd<5>;\n\ + ld.param.u64 %rd1, [_Z8myKernelPi_param_0];\n\ + cvta.to.global.u64 %rd2, %rd1;\n\ + .loc 1 3 1\n\ + mov.u32 %r1, %ntid.x;\n\ + mov.u32 %r2, %ctaid.x;\n\ + mov.u32 %r3, %tid.x;\n\ + mad.lo.s32 %r4, %r1, %r2, %r3;\n\ + mul.wide.s32 %rd3, %r4, 4;\n\ + add.s64 %rd4, %rd2, %rd3;\n\ + .loc 1 4 1\n\ + st.global.u32 [%rd4], %r4;\n\ + .loc 1 5 2\n\ + ret;\n\ +}\n\ +"; + +char myPtx32[] = "\n\ +.version 3.2\n\ +.target sm_20\n\ +.address_size 32\n\ +.visible .entry _Z8myKernelPi(\n\ + .param .u32 _Z8myKernelPi_param_0\n\ +)\n\ +{\n\ + .reg .s32 %r<9>;\n\ + ld.param.u32 %r1, [_Z8myKernelPi_param_0];\n\ + cvta.to.global.u32 %r2, %r1;\n\ + .loc 1 3 1\n\ + mov.u32 %r3, %ntid.x;\n\ + mov.u32 %r4, %ctaid.x;\n\ + mov.u32 %r5, %tid.x;\n\ + mad.lo.s32 %r6, %r3, %r4, %r5;\n\ + .loc 1 4 1\n\ + shl.b32 %r7, %r6, 2;\n\ + add.s32 %r8, %r2, %r7;\n\ + st.global.u32 [%r8], %r6;\n\ + .loc 1 5 2\n\ + ret;\n\ +}\n\ +"; + +#endif -- cgit v1.3 From e22b3179ccdee8bfbd1e74f86d93e0d7fc9ae6dc Mon Sep 17 00:00:00 2001 From: J Date: Fri, 20 Jul 2018 11:40:44 -0700 Subject: load param offsets with refactored code and look in middle of malloc'd area for pointer --- .../WatchYourStep/ptxjitplus/ptxjitplus.cpp | 69 +++++++++++----------- debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h | 7 +++ src/cuda-sim/cuda-sim.cc | 23 +++++--- 3 files changed, 55 insertions(+), 44 deletions(-) (limited to 'debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h') diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp index 26a6d29..1065870 100644 --- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp +++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp @@ -136,7 +136,7 @@ void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUl checkCudaErrors(cuLinkDestroy(*lState)); } -void initializeData(std::vector& param_data, std::vector< std::pair >& param_info) +void initializeData(std::vector& v_params) { char *gpgpusim_env = getenv("GPGPUSIM_SETUP_ENVIRONMENT_WAS_RUN"); if (gpgpusim_env!=NULL&&gpgpusim_env[0] == '1'){ @@ -161,21 +161,21 @@ void initializeData(std::vector& param_data, std::vector< std::p fscanf(fin, "%u,%u,%u %u,%u,%u\n", &gridDim.x, &gridDim.y, &gridDim.z, &blockDim.x, &blockDim.y, &blockDim.z); //fill data structure to pass in params later while (!feof(fin)){ - std::pair info; + param p; int err; size_t len; unsigned val; int start = fgetc(fin); if (start == '*'){ - info.second=true; + p.isPointer = true; }else{ - info.second=false; + p.isPointer = false; int c = ungetc(start,fin); assert(c==start&&"Couldn't ungetc\n"); } err = fscanf(fin, "%lu : ", &len); - info.first = len; assert( err==1 ); + p.size = len; unsigned char* params = (unsigned char*) malloc(len*sizeof(unsigned char)); for (size_t i=0; i& param_data, std::vector< std::p assert( err==1 ); params[i] = (unsigned char) val; } - //TODO: parse param offset - param_info.push_back(info); - param_data.push_back(params); + p.data = params; + unsigned offset; + err = fscanf(fin, " : %u", &offset); + assert( err==1 ); + p.offset = offset; + v_params.push_back(p); err = fscanf(fin, "\n"); assert(err==0); } @@ -206,12 +209,8 @@ int main(int argc, char **argv) printf("[%s] - Starting...\n", sSDKname); //parameter data - std::vector param_data; - //parameter data size and isPointer - std::vector< std::pair > param_info; - initializeData(param_data,param_info); - - + std::vector v_params; + initializeData(v_params); if (checkCmdLineFlag(argc, (const char **)argv, "device")) { @@ -251,8 +250,6 @@ int main(int argc, char **argv) exit(EXIT_WAIVED); } - - // Allocate memory on host and device (Runtime API) // NOTE: The runtime API will create the GPU Context implicitly here int *d_tmp = 0; @@ -265,28 +262,31 @@ int main(int argc, char **argv) //maps param number to pointer to 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 + void * paramKernels[v_params.size()]; + //Initialize param data for kernel int paramOffset = 0; - for( size_t i = 0; i::iterator p = v_params.begin(); p!=v_params.end(); p++){ + if(p->isPointer){ unsigned char **d_data = (unsigned char **) malloc(sizeof(unsigned char **)); - checkCudaErrors(cudaMalloc((void**)d_data, param_info[i].first)); - checkCudaErrors(cudaMemcpy((void*)*d_data,(void*)param_data[i],param_info[i].first,cudaMemcpyHostToDevice)); + checkCudaErrors(cudaMalloc((void**)d_data, p->size)); + checkCudaErrors(cudaMemcpy((void*)*d_data,(void*)p->data,p->size,cudaMemcpyHostToDevice)); if (gpgpusim){ - checkCudaErrors(cuParamSetv(hKernel, paramOffset, d_data, sizeof(*d_data))); + checkCudaErrors(cuParamSetv(hKernel, p->offset, d_data, sizeof(*d_data))); } - paramKernels[i] = (void*)d_data; - m_device_data[i]=*d_data; - m_cleanup[i]=d_data; - paramOffset += 8; + paramKernels[index] = (void*)d_data; + m_device_data[index]=*d_data; + m_cleanup[index]=d_data; + paramOffset = p->offset + 8; }else{ if (gpgpusim){ - checkCudaErrors(cuParamSetv(hKernel, paramOffset, param_data[i], param_info[i].first)); + checkCudaErrors(cuParamSetv(hKernel, p->offset, p->data, p->size)); } - paramKernels[i] = (void*)param_data[i]; - paramOffset += param_info[i].first; + paramKernels[index] = (void*)p->data; + paramOffset = p->offset + p->size; } + index ++; } checkCudaErrors(cuParamSetSize(hKernel, paramOffset)); @@ -299,13 +299,13 @@ int main(int argc, char **argv) std::map< size_t, unsigned char* > m_output_data; 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) + if ((h_data = (unsigned char *)malloc(v_params[i->first].size)) == NULL) { std::cerr << "Could not allocate host memory" << std::endl; exit(EXIT_FAILURE); } // Copy the result back to the host - checkCudaErrors(cudaMemcpy(h_data, i->second, param_info[i->first].first, cudaMemcpyDeviceToHost)); + checkCudaErrors(cudaMemcpy(h_data, i->second, v_params[i->first].size, cudaMemcpyDeviceToHost)); m_output_data[i->first] = h_data; } @@ -313,8 +313,8 @@ int main(int argc, char **argv) FILE *fout = fopen(filename.c_str(), "w"); assert(fout); 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, "param %zu: size = %zu, data = ", i->first, v_params[i->first].size); + for (size_t j = 0; jfirst].size; j++){ if (!(j%24)){ fprintf(fout, "\n"); } @@ -325,7 +325,6 @@ int main(int argc, char **argv) fflush(fout); fclose(fout); - //Cleanup for(std::map< size_t, void* >::iterator i = m_device_data.begin(); i!=m_device_data.end(); i++){ if (i->second){ diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h index 8fb5611..3c6e181 100644 --- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h +++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h @@ -12,6 +12,13 @@ #ifndef _PTXJIT_H_ #define _PTXJIT_H_ +struct param{ + bool isPointer; + size_t size; + unsigned char *data; + unsigned offset; +}; + /* * PTX is equivalent to the following kernel: * diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index a499ce9..e0cd3f9 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1238,11 +1238,6 @@ void function_info::ptx_jit_config(std::map mallocPt char * wys_exec_path = getenv("WYS_EXEC_PATH"); assert(wys_exec_path!=NULL); std::string command = std::string("mkdir ") + gpgpusim_path + "/debug_tools/WatchYourStep/data"; - if(system(command.c_str())!=0){ - printf("WARNING: Failed to execute mkdir \n"); - printf("Problematic call: %s", command.c_str()); - abort(); - } std::string filename(std::string(gpgpusim_path) + "/debug_tools/WatchYourStep/data/params.config" + std::to_string(counter)); //initialize paramList @@ -1295,11 +1290,21 @@ void function_info::ptx_jit_config(std::map mallocPt if (paramIsPointer[i->first]){ //is pointer - assert(param_value.size==8); - assert(param_value.size==sizeof(void*) && mallocPtr_Size.find(*(unsigned long long*)param_value.pdata)!=mallocPtr_Size.end()); - //TODO: check in middle of malloc'd memory for pointer + assert(param_value.size==sizeof(void*)&&"MisID'd this param as pointer"); + size_t array_size = 0; + unsigned long long param_pointer = *(unsigned long long*)param_value.pdata; + if(mallocPtr_Size.find(param_pointer)!=mallocPtr_Size.end()){ + array_size = mallocPtr_Size[param_pointer]; + }else{ + for( std::map::iterator j=mallocPtr_Size.begin(); j!=mallocPtr_Size.end(); j++ ) { + if(param_pointer>j->first&¶m_pointerfirst + j->second){ + array_size = j->first + j->second - param_pointer; + break; + } + } + assert(array_size>0&&"pointer was not previously malloc'd"); + } - size_t array_size = mallocPtr_Size[*(unsigned long long*)param_value.pdata]; unsigned char* val = (unsigned char*) malloc(param_value.size); param_mem->read(param_addr,param_value.size,(void*)val); unsigned char* array_val = (unsigned char*) malloc(array_size); -- cgit v1.3