summaryrefslogtreecommitdiff
path: root/debug_tools/WatchYourStep/ptxjitplus
diff options
context:
space:
mode:
Diffstat (limited to 'debug_tools/WatchYourStep/ptxjitplus')
-rw-r--r--debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp180
1 files changed, 63 insertions, 117 deletions
diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
index 8645114..c23d7f6 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
@@ -22,6 +22,7 @@
#include <math.h>
#include <string.h>
#include <stdio.h>
+#include <map>
// CUDA driver & runtime
#include <cuda.h>
@@ -111,88 +112,7 @@ 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)
+void initializeData(std::vector<unsigned char*>& param_data, std::vector< std::pair<size_t, bool> >& param_info)
{
char *wys_exec_path = getenv("WYS_EXEC_PATH");
assert(wys_exec_path!=NULL);
@@ -208,33 +128,41 @@ void* initializeData(std::vector< std::pair<size_t, unsigned char*> >& param_dat
assert(fin);
char buff[1024];
fscanf(fin, "%s\n", buff);
- //void *retval = instrument_ptx_from_function(std::string(buff), path_to_search);
- void *retval = NULL;
- printf("%s\n", buff);
+ printf("Processing :%s ...\n", buff);
+ fflush(stdout);
//fill data structure to pass in params later
while (!feof(fin)){
+ std::pair<size_t, bool> info;
int err;
size_t len;
unsigned val;
+ int start = fgetc(fin);
+ if (start == '*'){
+ info.second=true;
+ }else{
+ info.second=false;
+ int c = ungetc(start,fin);
+ assert(c==start&&"Couldn't ungetc\n");
+ }
err = fscanf(fin, "%lu : ", &len);
+ info.first = len;
assert( err==1 );
- printf("%lu : ", len);
+ //printf("%lu : ", len);
unsigned char params[len];
for (size_t i=0; i<len; i++)
{
err = fscanf(fin, "%u ", &val);
assert( err==1 );
- printf("%u ", val);
+ //printf("%u ", val);
params[i] = (unsigned char) val;
}
- param_data.push_back(std::pair<size_t, unsigned char*>(len, params));
+ param_info.push_back(info);
+ param_data.push_back(params);
err = fscanf(fin, "\n");
assert(err==0);
- printf("\n");
+ //printf("\n");
}
-
fclose(fin);
- return retval;
}
int main(int argc, char **argv)
@@ -250,9 +178,9 @@ int main(int argc, char **argv)
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);
+ std::vector<unsigned char*> param_data;
+ std::vector< std::pair<size_t, bool> > param_info;
+ initializeData(param_data,param_info);
if (checkCmdLineFlag(argc, (const char **)argv, "device"))
{
@@ -302,14 +230,20 @@ int main(int argc, char **argv)
checkCudaErrors(cuFuncSetBlockShape(hKernel, nThreads, 1, 1));
//Initialize param_data for kernel
+ std::map< size_t, unsigned char* > m_device_data;
int paramOffset = 0;
- 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;
+ for( size_t i = 0; i<param_data.size(); i++){
+ if(param_info[i].second){
+ unsigned char *d_data = 0;
+ checkCudaErrors(cudaMalloc((void**)&d_data, param_info[i].first));
+ checkCudaErrors(cudaMemcpy((void*)d_data,(void*)param_data[i],param_info[i].first,cudaMemcpyHostToDevice));
+ checkCudaErrors(cuParamSetv(hKernel, paramOffset, &d_data, sizeof(d_data)));
+ m_device_data[i]=d_data;
+ paramOffset += 8;
+ }else{
+ checkCudaErrors(cuParamSetv(hKernel, paramOffset, param_data[i], param_info[i].first));
+ paramOffset += param_info[i].first;
+ }
}
checkCudaErrors(cuParamSetSize(hKernel, paramOffset));
@@ -317,15 +251,24 @@ int main(int argc, char **argv)
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);
+ std::map< size_t, unsigned char* > m_output_data;
+ for(std::map< size_t, unsigned char* >::iterator i = m_device_data.begin(); i!=m_device_data.begin(); i++){
+ unsigned char *h_data = 0;
+ if ((h_data = (unsigned char *)malloc(param_info[i->first].first)) == 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));
+ m_output_data[i->first] = h_data;
+ }
+
+ for(std::map< size_t, unsigned char* >::iterator i = m_output_data.begin(); i!=m_output_data.begin(); i++){
+ //print data out to file
}
- // Copy the result back to the host
- checkCudaErrors(cudaMemcpy(h_data, d_data, memSize, cudaMemcpyDeviceToHost));
+ int* h_data;
// Check the result
bool dataGood = true;
@@ -338,17 +281,20 @@ int main(int argc, char **argv)
}
}
- // Cleanup
- if (d_data)
- {
- checkCudaErrors(cudaFree(d_data));
- d_data = 0;
+ //Cleanup
+ for(std::map< size_t, unsigned char* >::iterator i = m_device_data.begin(); i!=m_device_data.begin(); i++){
+ if (i->second){
+ checkCudaErrors(cudaFree(i->second));
+ i->second = 0;
+ }
}
- if (h_data)
- {
- free(h_data);
- h_data = 0;
+ for(std::map< size_t, unsigned char* >::iterator i = m_output_data.begin(); i!=m_output_data.begin(); i++){
+ if (i->second)
+ {
+ free(i->second);
+ i->second = 0;
+ }
}
if (hModule)
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 23:27:58 +0000