summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rwxr-xr-xdebug_tools/WatchYourStep/ptxjitplus/ptxjitplusbin139317 -> 0 bytes
-rw-r--r--debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp112
-rw-r--r--libcuda/cuda_runtime_api.cc12
-rw-r--r--src/cuda-sim/cuda-sim.cc57
-rw-r--r--src/cuda-sim/ptx_ir.h2
5 files changed, 156 insertions, 27 deletions
diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus
deleted file mode 100755
index ddc3435..0000000
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus
+++ /dev/null
Binary files differ
diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
index 9954e31..8645114 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
@@ -111,6 +111,87 @@ void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUl
checkCudaErrors(cuLinkDestroy(*lState));
}
+void function_info::debug_param( ) const
+{
+ char filename[] = "params.txt";
+ char buff[1024];
+ snprintf(buff,1024,"c++filt %s > %s", get_name().c_str(), filename);
+ system(buff);
+ FILE *fp = fopen(filename, "r");
+ fgets(buff, 1024, fp);
+ fclose(fp);
+
+ std::string fn(buff);
+ size_t pos1, pos2;
+ pos1 = fn.find("(");
+ pos2 = fn.find(")");
+ assert(pos2>pos1&&pos1>0);
+ strcpy(buff, fn.substr(pos1 + 1, pos2 - pos1 - 1).c_str());
+ printf("params: %s\n", buff);
+ char *tok;
+ std::vector<std::string> params;
+ tok = strtok(buff, ",");
+ while(tok!=NULL){
+ std::string param(tok);
+ param.erase(0, param.find_first_not_of(" "));
+ param.erase(param.find_last_not_of(" ")+1);
+ params.push_back(param);
+ tok = strtok(NULL, ",");
+ }
+ for (auto const& it : params){
+ std::cout<<it<<std::endl;
+ }
+
+ FILE *fout = fopen (filename, "w");
+ fprintf(fout, "Name of function:%s\n", fn.c_str());
+
+ for( std::map<unsigned,param_info>::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) {
+ const param_info &p = i->second;
+ std::string name = p.get_name();
+ param_t param_value = p.get_value();
+ if(params[i->first].find("const")!=std::string::npos){
+ fprintf(fout, "Input: ");
+ } else {
+ fprintf(fout, "Input/output: ");
+ }
+
+ symbol *param = m_symtab->lookup(name.c_str());
+ addr_t param_addr = param->get_address();
+ fprintf(fout, "%s: %#08x, ", name.c_str(), param_addr);
+
+ if(params[i->first].find("int")!=std::string::npos){
+ size_t len = param_value.size/sizeof(int);
+ int val[len];
+ memcpy((void*) val, param_value.pdata+param_value.offset, param_value.size);
+ fprintf(fout, "val (int) = ");
+ for (unsigned i = 0; i<len; i++){
+ fprintf(fout, "%d ", val[i]);
+ }
+ fprintf(fout, "\n");
+ } else if(params[i->first].find("float")!=std::string::npos){
+ size_t len = param_value.size/sizeof(float);
+ float val[len];
+ memcpy((void*) val, param_value.pdata+param_value.offset, param_value.size);
+ fprintf(fout, "val (float) = ");
+ for (unsigned i = 0; i<len; i++){
+ fprintf(fout, "%f ", val[i]);
+ }
+ fprintf(fout, "\n");
+ }else{
+ size_t len = param_value.size/sizeof(char);
+ char val[len];
+ memcpy((void*) val, param_value.pdata+param_value.offset, param_value.size);
+ fprintf(fout, "val (char) = ");
+ for (unsigned i = 0; i<len; i++){
+ fprintf(fout, "%c ", val[i]);
+ }
+ fprintf(fout, "\n");
+ }
+ }
+ fflush(fout);
+ fclose(fout);
+}
+
void* initializeData(std::vector< std::pair<size_t, unsigned char*> >& param_data)
{
char *wys_exec_path = getenv("WYS_EXEC_PATH");
@@ -160,19 +241,17 @@ int main(int argc, char **argv)
{
const unsigned int nThreads = 256;
const unsigned int nBlocks = 64;
- const size_t memSize = nThreads * nBlocks * sizeof(int);
CUmodule hModule = 0;
CUfunction hKernel = 0;
CUlinkState lState;
- int *d_data = 0;
- int *h_data = 0;
int cuda_device = 0;
cudaDeviceProp deviceProp;
printf("[%s] - Starting...\n", sSDKname);
std::vector< std::pair<size_t, unsigned char*> > param_data;
+ std::vector< std::pair<size_t, unsigned char*> > device_data;
void* storedReg = initializeData(param_data);
if (checkCmdLineFlag(argc, (const char **)argv, "device"))
@@ -213,15 +292,7 @@ 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
- if ((h_data = (int *)malloc(memSize)) == NULL)
- {
- std::cerr << "Could not allocate host memory" << std::endl;
- exit(EXIT_FAILURE);
- }
- checkCudaErrors(cudaMalloc(&d_data, memSize));
// JIT Compile the Kernel from PTX and get the Handles (Driver API)
ptxJIT(argc, argv, &hModule, &hKernel, &lState);
@@ -229,16 +300,29 @@ int main(int argc, char **argv)
// Set the kernel parameters (Driver API)
checkCudaErrors(cuFuncSetBlockShape(hKernel, nThreads, 1, 1));
- //param_data
+
+ //Initialize param_data for kernel
int paramOffset = 0;
- checkCudaErrors(cuParamSetv(hKernel, paramOffset, &d_data, sizeof(d_data)));
- paramOffset += sizeof(d_data);
+ for( std::vector< std::pair<size_t,unsigned char*> >::const_iterator i=param_data.begin(); i!=param_data.end(); i++ ) {
+ size_t memSize = nThreads * nBlocks * i->first;
+ unsigned char *d_data = 0;
+ checkCudaErrors(cudaMalloc((void**)&d_data, memSize));
+ checkCudaErrors(cudaMemcpy(d_data,i->first,cudaMemcpyHostToDevice));
+ checkCudaErrors(cuParamSetv(hKernel, paramOffset, &d_data, memSize));
+ paramOffset += i->first;
+ }
checkCudaErrors(cuParamSetSize(hKernel, paramOffset));
// Launch the kernel (Driver API_)
checkCudaErrors(cuLaunchGrid(hKernel, nBlocks, 1));
std::cout << "CUDA kernel launched" << std::endl;
+ int *h_data = 0;
+ if ((h_data = (int *)malloc(memSize)) == NULL)
+ {
+ std::cerr << "Could not allocate host memory" << std::endl;
+ exit(EXIT_FAILURE);
+ }
// Copy the result back to the host
checkCudaErrors(cudaMemcpy(h_data, d_data, memSize, cudaMemcpyDeviceToHost));
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index d2b855c..1e4f1df 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -150,6 +150,7 @@
std::map<void *,void **> pinned_memory; //support for pinned memories added
std::map<void *, size_t> pinned_memory_size;
+std::map<void *, size_t> g_devPtr_Size;
int no_of_ptx=0;
std::map<int, std::set<std::string> > version_filename;
@@ -487,8 +488,10 @@ __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size)
{
CUctx_st* context = GPGPUSim_Context();
*devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size);
- if(g_debug_execution >= 3)
+ if(g_debug_execution >= 3){
printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr);
+ g_devPtr_Size[*devPtr] = size;
+ }
if ( *devPtr ) {
return g_last_cudaError = cudaSuccess;
} else {
@@ -2374,8 +2377,10 @@ cudaError_t CUDARTAPI cudaHostGetDevicePointer(void **pDevice, void *pHost, unsi
assert(i != pinned_memory_size.end());
size_t size = i->second;
*pDevice = gpu->gpu_malloc(size);
- if(g_debug_execution >= 3)
+ if(g_debug_execution >= 3){
printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *pDevice);
+ g_devPtr_Size[*pDevice] = size;
+ }
if ( *pDevice ) {
pinned_memory[pHost]=pDevice;
//Copy contents in cpu to gpu
@@ -2617,8 +2622,9 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
g_ptx_kernel_count++;
fflush(stdout);
+
if(g_debug_execution >= 3){
- entry->debug_param();
+ entry->debug_param(g_devPtr_Size, result->get_param_memory());
}
return result;
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index bdb1e8d..8795e1c 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1226,33 +1226,72 @@ void function_info::list_param( FILE *fout ) const
fflush(fout);
}
-void function_info::debug_param() const
+void function_info::debug_param(std::map<void *, size_t> devPtr_Size, memory_space *param_mem) const
{
static unsigned long counter = 0;
- std::string gpgpusim_path(getenv("GPGPUSIM_ROOT"));
- assert(!gpgpusim_path.empty());
- std::string command = "mkdir " + gpgpusim_path + "/debug_tools/WatchYourStep/data";
+ std::vector< std::pair<size_t, unsigned char*> > param_data;
+ std::vector<bool> paramIsPointer;
+
+ char * gpgpusim_path = getenv("GPGPUSIM_ROOT");
+ assert(gpgpusim_path!=NULL);
+ std::string command = std::string("mkdir ") + gpgpusim_path + "/debug_tools/WatchYourStep/data";
system(command.c_str());
+ std::string filename(std::string(gpgpusim_path) + "/debug_tools/WatchYourStep/data/params.config" + std::to_string(counter++));
- std::string filename(gpgpusim_path + "/debug_tools/WatchYourStep/data/params.config" + std::to_string(counter++));
- std::vector< std::pair<size_t, unsigned char*> > param_data;
+ //initialize paramList
+ char buff[1024];
+ std::string filename_c(filename+"_c");
+ snprintf(buff,1024,"c++filt %s > %s", get_name().c_str(), filename_c.c_str());
+ system(buff);
+ FILE *fp = fopen(filename_c.c_str(), "r");
+ fgets(buff, 1024, fp);
+ fclose(fp);
+ std::string fn(buff);
+ size_t pos1, pos2;
+ pos1 = fn.find("(");
+ pos2 = fn.find(")");
+ assert(pos2>pos1&&pos1>0);
+ strcpy(buff, fn.substr(pos1 + 1, pos2 - pos1 - 1).c_str());
+ char *tok;
+ tok = strtok(buff, ",");
+ while(tok!=NULL){
+ std::string param(tok);
+ if(param.find("*")!=std::string::npos){
+ paramIsPointer.push_back(true);
+ }else{
+ paramIsPointer.push_back(false);
+ }
+ tok = strtok(NULL, ",");
+ }
for( std::map<unsigned,param_info>::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) {
const param_info &p = i->second;
+ std::string name = p.get_name();
+ symbol *param = m_symtab->lookup(name.c_str());
+ addr_t param_addr = param->get_address();
param_t param_value = p.get_value();
- unsigned char val[param_value.size];
- memcpy((void*) val, (void*)((char*)param_value.pdata+param_value.offset), param_value.size);
- param_data.push_back(std::pair<size_t, unsigned char*>(param_value.size,val));
+// if (paramIsPointer[i->first]){
+// assert(param_value.size==8);
+// }else{
+ unsigned char val[param_value.size];
+ param_mem->read(param_addr,param_value.size,(void*)val);
+ param_data.push_back(std::pair<size_t, unsigned char*>(param_value.size,val));
+ //}
}
FILE *fout = fopen (filename.c_str(), "w");
fprintf(fout, "%s\n", get_name().c_str());
+ size_t index = 0;
for( std::vector< std::pair<size_t,unsigned char*> >::const_iterator i=param_data.begin(); i!=param_data.end(); i++ ) {
+ if (paramIsPointer[index]){
+ fprintf(fout, "*");
+ }
fprintf(fout, "%lu :", i->first);
for (size_t j = 0; j<i->first; j++){
fprintf(fout, " %u", i->second[j]);
}
fprintf(fout, "\n");
+ index++;
}
fflush(fout);
fclose(fout);
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index 341f9b7..43907d3 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -1257,7 +1257,7 @@ public:
void finalize( memory_space *param_mem );
void param_to_shared( memory_space *shared_mem, symbol_table *symtab );
void list_param( FILE *fout ) const;
- void debug_param() const;
+ void debug_param(std::map<void *, size_t> devPtr_Size, memory_space *param_mem) const;
const struct gpgpu_ptx_sim_info* get_kernel_info () const
{
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 22:26:54 +0000