summaryrefslogtreecommitdiff
path: root/debug_tools
diff options
context:
space:
mode:
Diffstat (limited to 'debug_tools')
-rw-r--r--debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp69
-rw-r--r--debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h7
2 files changed, 41 insertions, 35 deletions
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<unsigned char*>& param_data, std::vector< std::pair<size_t, bool> >& param_info)
+void initializeData(std::vector<param>& 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<unsigned char*>& 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<size_t, bool> 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<len; i++)
{
@@ -183,9 +183,12 @@ void initializeData(std::vector<unsigned char*>& 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<unsigned char*> param_data;
- //parameter data size and isPointer
- std::vector< std::pair<size_t, bool> > param_info;
- initializeData(param_data,param_info);
-
-
+ std::vector<param> 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<param_data.size(); i++){
- if(param_info[i].second){
+
+ int index = 0;
+ for(std::vector<param>::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; j<param_info[i->first].first; j++){
+ fprintf(fout, "param %zu: size = %zu, data = ", i->first, v_params[i->first].size);
+ for (size_t j = 0; j<v_params[i->first].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:
*
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 22:25:20 +0000