Adding support for printing out individual weft thread files, resolves …

…#6

README.md

@@ -174,6 +174,8 @@ Below is a summary of the command line flags that Weft supports.
  * `-n`: set the number of threads per CTA. This is required
                 if the CUDA kernel did not have a 
                 `__launch_bounds__` annotation
+ * `-p`: print out individual files for each thread of all Weft modeled 
+                instructions, this will generate one file per thread
  * `-s`: assume warp-synchronous execution when checking for races
  * `-t`: set the size of the thread pool for Weft to use; in
                 general, Weft is memory bound, so one or two threads per socket

src/instruction.cc

@@ -2416,11 +2416,21 @@ BarrierSync::BarrierSync(int n, int c, PTXBarrier *bar, Thread *thread)
 { 
 }
 
+void BarrierSync::print_instruction(FILE *target)
+{
+  fprintf(target,"bar.sync %d, %d;\n", name, count);
+}
+
 BarrierArrive::BarrierArrive(int n, int c, PTXBarrier *bar, Thread *thread)
   : WeftBarrier(n, c, bar, thread)
 {
 }
 
+void BarrierArrive::print_instruction(FILE *target)
+{
+  fprintf(target,"bar.arrive %d, %d;\n", name, count);
+}
+
 WeftAccess::WeftAccess(int addr, PTXSharedAccess *acc, 
                        Thread *thread, int acc_id)
   : WeftInstruction(acc, thread), address(addr), access(acc), access_id(acc_id)
@@ -2464,9 +2474,19 @@ SharedWrite::SharedWrite(int addr, PTXSharedAccess *acc,
 {
 }
 
+void SharedWrite::print_instruction(FILE *target)
+{
+  fprintf(target,"write shared[%d];\n", address);
+}
+
 SharedRead::SharedRead(int addr, PTXSharedAccess *acc, 
                        Thread *thread, int acc_id /*=-1*/)
   : WeftAccess(addr, acc, thread, acc_id)
 {
 }
 
+void SharedRead::print_instruction(FILE *target)
+{
+  fprintf(target,"read shared[%d];\n", address);
+}
+

src/instruction.h

@@ -725,6 +725,8 @@ class WeftInstruction {
 public:
   void initialize_happens(Happens *happens);
   inline Happens* get_happens(void) const { return happens_relationship; }
+public:
+  virtual void print_instruction(FILE *target) = 0;
 public:
   PTXInstruction *const instruction;
   Thread *const thread;
@@ -747,6 +749,8 @@ class WeftBarrier : public WeftInstruction {
 public:
   void set_instance(BarrierInstance *instance);
   inline BarrierInstance* get_instance(void) const { return instance; }
+public:
+  virtual void print_instruction(FILE *target) = 0;
 public:
   const int name;
   const int count;
@@ -766,6 +770,7 @@ class BarrierSync : public WeftBarrier {
 public:
   virtual bool is_sync(void) const { return true; }
   virtual BarrierSync* as_sync(void) { return this; }
+  virtual void print_instruction(FILE *target);
 };
 
 class BarrierArrive : public WeftBarrier {
@@ -779,6 +784,7 @@ class BarrierArrive : public WeftBarrier {
 public:
   virtual bool is_arrive(void) const { return true; }
   virtual BarrierArrive* as_arrive(void) { return this; }
+  virtual void print_instruction(FILE *target);
 };
 
 class WeftAccess : public WeftInstruction {
@@ -795,6 +801,8 @@ class WeftAccess : public WeftInstruction {
 public:
   bool has_happens_relationship(WeftAccess *other);
   bool is_warp_synchronous(WeftAccess *other);
+public:
+  virtual void print_instruction(FILE *target) = 0;
 public:
   const int address;
   PTXSharedAccess *const access;
@@ -813,6 +821,7 @@ class SharedWrite : public WeftAccess {
 public:
   virtual bool is_write(void) const { return true; }
   virtual SharedWrite* as_write(void) { return this; }
+  virtual void print_instruction(FILE *target);
 };
 
 class SharedRead : public WeftAccess {
@@ -827,6 +836,7 @@ class SharedRead : public WeftAccess {
 public:
   virtual bool is_read(void) const { return true; }
   virtual SharedRead* as_read(void) { return this; }
+  virtual void print_instruction(FILE *target);
 };
 
 #endif // __INSTRUCTION_H__

src/program.cc

@@ -537,6 +537,29 @@ int Thread::accumulate_instruction_counts(std::vector<int> &total_counts)
   return total;
 }
 
+void Thread::dump_weft_thread(void)
+{
+  // Open up a file for this thread and then 
+  // print out all of our weft instructions
+  char file_name[1024];
+  program->weft->get_file_prefix(file_name, 1024-32);
+  char buffer[32];
+  snprintf(buffer, 31, "_%d_%d_%d.weft", tid_x, tid_y, tid_z);
+  strncat(file_name, buffer, 31);
+  FILE *weft_file = fopen(file_name, "w");
+  if (weft_file == NULL)
+  {
+    fprintf(stderr, "WEFT WARNING: Failed to open file %s\n", file_name);
+    return ;
+  }
+  for (std::vector<WeftInstruction*>::const_iterator it = 
+        instructions.begin(); it != instructions.end(); it++)
+  {
+    (*it)->print_instruction(weft_file);
+  }
+  assert(fclose(weft_file) == 0);
+}
+
 void Thread::update_shared_memory(WeftAccess *access)
 {
   shared_memory->update_accesses(access);
@@ -700,3 +723,13 @@ void UpdateThreadTask::execute(void)
   thread->update_happens_relationships();
 }
 
+DumpThreadTask::DumpThreadTask(Thread *t)
+  : WeftTask(), thread(t)
+{
+}
+
+void DumpThreadTask::execute(void)
+{
+  thread->dump_weft_thread();
+}
+

src/program.h

@@ -117,6 +117,7 @@ class Thread {
 public:
   void profile_instruction(PTXInstruction *instruction);
   int accumulate_instruction_counts(std::vector<int> &total_counts);
+  void dump_weft_thread(void);
 public:
   void update_shared_memory(WeftAccess *access);
 public:

src/weft.cc

@@ -39,7 +39,7 @@ Weft::Weft(int argc, char **argv)
   : file_name(NULL), max_num_threads(-1), 
     thread_pool_size(1), max_num_barriers(1),
     verbose(false), detailed(false), instrument(false), 
-    warnings(false), warp_synchronous(false),
+    warnings(false), warp_synchronous(false), print_files(false),
     program(NULL), shared_memory(NULL), graph(NULL),
     worker_threads(NULL), pending_count(0)
 {
@@ -132,6 +132,11 @@ void Weft::parse_inputs(int argc, char **argv)
       instrument = true;
       continue;
     }
+    if (!strcmp(argv[i],"-p"))
+    {
+      print_files = true;
+      continue;
+    }
     if (!strcmp(argv[i],"-n"))
     {
       std::string threads(argv[++i]);
@@ -189,6 +194,7 @@ void Weft::parse_inputs(int argc, char **argv)
     fprintf(stdout,"  Instrument: %s\n", (instrument ? "yes" : "no"));
     fprintf(stdout,"  Report Warnings: %s\n", (warnings ? "yes" : "no"));
     fprintf(stdout,"  Warp-Synchronous Execution: %s\n", (warnings ? "yes" : "no"));
+    fprintf(stdout,"  Dump Weft thread files: %s\n", (print_files ? "yes" : "no"));
   }
 }
 
@@ -256,6 +262,7 @@ void Weft::report_usage(int error, const char *error_str)
   fprintf(stderr,"      can be an integer or an x-separated tuple e.g. 32x32x2 or 32x1\n");
   fprintf(stderr,"      Weft will still only simulate a single CTA specified by '-b'\n");
   fprintf(stderr,"  -i: instrument execution\n");
+  fprintf(stderr,"  -p: print individual Weft thread files (one file per thread!)\n");
   fprintf(stderr,"  -n: number of threads per CTA\n");
   fprintf(stderr,"      can be an integer or an x-separated tuple e.g. 64x2 or 32x8x1\n");
   fprintf(stderr,"  -s: assume warp-synchronous execution\n");
@@ -378,6 +385,20 @@ void Weft::emulate_threads(void)
 
   if (instrument)
     stop_instrumentation(1/*stage*/);
+
+  // If we want to dump thread-specific files, do that now
+  // Note that we don't include this in the timing
+  if (print_files)
+  {
+    initialize_count(max_num_threads);
+    for (std::vector<Thread*>::const_iterator it = threads.begin();
+          it != threads.end(); it++)
+    {
+      DumpThreadTask *dump_task = new DumpThreadTask(*it); 
+      enqueue_task(dump_task);
+    }
+    wait_until_done();
+  }
 }
 
 void Weft::construct_dependence_graph(void)
@@ -522,6 +543,14 @@ void Weft::fill_grid_dim(int *array)
     array[i] = grid_dim[i];
 }
 
+void Weft::get_file_prefix(char *buffer, size_t count)
+{
+  std::string full_name(file_name);
+  assert(full_name.find(".ptx") != std::string::npos);
+  std::string base = full_name.substr(0, full_name.find(".ptx"));
+  strncpy(buffer, base.c_str(), count);
+}
+
 void Weft::start_threadpool(void)
 {
   assert(thread_pool_size > 0);

src/weft.h

@@ -182,7 +182,21 @@ class RaceCheckTask : public WeftTask {
 public:
   virtual void execute(void);
 public:
-  Address *address;
+  Address *const address;
+};
+
+class DumpThreadTask : public WeftTask {
+public:
+  DumpThreadTask(Thread *thread);
+  DumpThreadTask(const DumpThreadTask &rhs) : thread(NULL) { assert(false); }
+  virtual ~DumpThreadTask(void) { }
+public:
+  DumpThreadTask& operator=(const DumpThreadTask &rhs)
+    { assert(false); return *this; }
+public:
+  virtual void execute(void);
+public:
+  Thread *const thread;
 };
 
 class Weft {
@@ -215,6 +229,7 @@ class Weft {
   void fill_block_dim(int *array);
   void fill_block_id(int *array);
   void fill_grid_dim(int *array);
+  void get_file_prefix(char *buffer, size_t count);
 protected:
   void start_threadpool(void);
   void stop_threadpool(void);
@@ -245,6 +260,7 @@ class Weft {
   bool instrument;
   bool warnings;
   bool warp_synchronous;
+  bool print_files;
 protected:
   Program *program;
   std::vector<Thread*> threads;