diff --git a/hipamd/samples/1_Utils/hipCommander/LICENSE.txt b/hipamd/samples/1_Utils/hipCommander/LICENSE.txt
new file mode 100644
index 0000000000..5d0d603232
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/LICENSE.txt
@@ -0,0 +1,27 @@
+
+Copyright (c) 2011, UT-Battelle, LLC
+All rights reserved.
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions are met:
+
+* Redistributions of source code must retain the above copyright
+  notice, this list of conditions and the following disclaimer.
+* Redistributions in binary form must reproduce the above copyright
+  notice, this list of conditions and the following disclaimer in the
+  documentation and/or other materials provided with the distribution.
+* Neither the name of Oak Ridge National Laboratory, nor UT-Battelle, LLC, nor
+  the names of its contributors may be used to endorse or promote products
+  derived from this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
diff --git a/hipamd/samples/1_Utils/hipCommander/Makefile b/hipamd/samples/1_Utils/hipCommander/Makefile
new file mode 100644
index 0000000000..e770c636a4
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/Makefile
@@ -0,0 +1,33 @@
+HIP_PATH?= $(wildcard /opt/rocm/hip)
+ifeq (,$(HIP_PATH))
+	HIP_PATH=../../..
+endif
+HIPCC=$(HIP_PATH)/bin/hipcc
+
+EXE=hipCommander
+OPT=-O3
+#CXXFLAGS = -O3 -g
+CXXFLAGS =  $(OPT) --std=c++11
+
+HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform)
+ifeq (${HIP_PLATFORM}, hcc)
+	CXXFLAGS += " -stdlib=libc++"
+endif
+
+CODE_OBJECTS=nullkernel.hsaco
+
+all: ${EXE} ${CODE_OBJECTS}
+
+$(EXE): hipCommander.cpp 
+	$(HIPCC) $(CXXFLAGS) $^ -o $@
+
+nullkernel.hsaco : nullkernel.hip.cpp
+	$(HIPCC) --genco nullkernel.hip -o nullkernel.hsaco
+
+
+install: $(EXE)
+	cp $(EXE) $(HIP_PATH)/bin
+
+
+clean:
+	rm -f *.o *.co $(EXE)
diff --git a/hipamd/samples/1_Utils/hipCommander/ResultDatabase.cpp b/hipamd/samples/1_Utils/hipCommander/ResultDatabase.cpp
new file mode 100644
index 0000000000..2ec686f260
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/ResultDatabase.cpp
@@ -0,0 +1,527 @@
+#include "ResultDatabase.h"
+
+#include <cfloat>
+#include <algorithm>
+#include <cmath>
+#include <iomanip>
+
+using namespace std;
+
+bool ResultDatabase::Result::operator<(const Result &rhs) const
+{
+    if (test < rhs.test)
+        return true;
+    if (test > rhs.test)
+        return false;
+    if (atts < rhs.atts)
+        return true;
+    if (atts > rhs.atts)
+        return false;
+    return false; // less-operator returns false on equal
+}
+
+double ResultDatabase::Result::GetMin() const
+{
+    double r = FLT_MAX;
+    for (int i=0; i<value.size(); i++)
+    {
+        r = min(r, value[i]);
+    }
+    return r;
+}
+
+double ResultDatabase::Result::GetMax() const
+{
+    double r = -FLT_MAX;
+    for (int i=0; i<value.size(); i++)
+    {
+        r = max(r, value[i]);
+    }
+    return r;
+}
+
+double ResultDatabase::Result::GetMedian() const
+{
+    return GetPercentile(50);
+}
+
+double ResultDatabase::Result::GetPercentile(double q) const
+{
+    int n = value.size();
+    if (n == 0)
+        return FLT_MAX;
+    if (n == 1)
+        return value[0];
+
+    if (q <= 0)
+        return value[0];
+    if (q >= 100)
+        return value[n-1];
+
+    double index = ((n + 1.) * q / 100.) - 1;
+
+    vector<double> sorted = value;
+    sort(sorted.begin(), sorted.end());
+
+    if (n == 2)
+        return (sorted[0] * (1 - q/100.)  +  sorted[1] * (q/100.));
+
+    int index_lo = int(index);
+    double frac = index - index_lo;
+    if (frac == 0)
+        return sorted[index_lo];
+
+    double lo = sorted[index_lo];
+    double hi = sorted[index_lo + 1];
+    return lo + (hi-lo)*frac;
+}
+
+double ResultDatabase::Result::GetMean() const
+{
+    double r = 0;
+    for (int i=0; i<value.size(); i++)
+    {
+        r += value[i];
+    }
+    return r / double(value.size());
+}
+
+double ResultDatabase::Result::GetStdDev() const
+{
+    double r = 0;
+    double u = GetMean();
+    if (u == FLT_MAX)
+        return FLT_MAX;
+    for (int i=0; i<value.size(); i++)
+    {
+        r += (value[i] - u) * (value[i] - u);
+    }
+    r = sqrt(r / value.size());
+    return r;
+}
+
+
+void ResultDatabase::AddResults(const string &test,
+                                const string &atts,
+                                const string &unit,
+                                const vector<double> &values)
+{
+    for (int i=0; i<values.size(); i++)
+    {
+        AddResult(test, atts, unit, values[i]);
+    }
+}
+
+static string RemoveAllButLeadingSpaces(const string &a)
+{
+    string b;
+    int n = a.length();
+    int i = 0;
+    while (i<n && a[i] == ' ')
+    {
+        b += a[i];
+        ++i;
+    }
+    for (; i<n; i++)
+    {
+        if (a[i] != ' ' && a[i] != '\t')
+            b += a[i];
+    }
+    return b;
+}
+
+void ResultDatabase::AddResult(const string &test_orig,
+                               const string &atts_orig,
+                               const string &unit_orig,
+                               double value)
+{
+    string test = RemoveAllButLeadingSpaces(test_orig);
+    string atts = RemoveAllButLeadingSpaces(atts_orig);
+    string unit = RemoveAllButLeadingSpaces(unit_orig);
+    int index;
+    for (index = 0; index < results.size(); index++)
+    {
+        if (results[index].test == test &&
+            results[index].atts == atts)
+        {
+            if (results[index].unit != unit)
+                throw "Internal error: mixed units";
+
+            break;
+        }
+    }
+
+    if (index >= results.size())
+    {
+        Result r;
+        r.test = test;
+        r.atts = atts;
+        r.unit = unit;
+        results.push_back(r);
+    }
+
+    results[index].value.push_back(value);
+}
+
+// ****************************************************************************
+//  Method:  ResultDatabase::DumpDetailed
+//
+//  Purpose:
+//    Writes the full results, including all trials.
+//
+//  Arguments:
+//    out        where to print
+//
+//  Programmer:  Jeremy Meredith
+//  Creation:    August 14, 2009
+//
+//  Modifications:
+//    Jeremy Meredith, Wed Nov 10 14:25:17 EST 2010
+//    Renamed to DumpDetailed to make room for a DumpSummary.
+//
+//    Jeremy Meredith, Thu Nov 11 11:39:57 EST 2010
+//    Added note about (*) missing value tag.
+//
+//    Jeremy Meredith, Tue Nov 23 13:57:02 EST 2010
+//    Changed note about missing values to be worded a little better.
+//
+// ****************************************************************************
+void ResultDatabase::DumpDetailed(ostream &out)
+{
+    vector<Result> sorted(results);
+    sort(sorted.begin(), sorted.end());
+
+    const int testNameW = 24 ;
+    const int attW = 12;
+    const int fieldW = 11;
+    out << std::fixed << right << std::setprecision(4);
+
+    int maxtrials = 1;
+    for (int i=0; i<sorted.size(); i++)
+    {
+        if (sorted[i].value.size() > maxtrials)
+            maxtrials = sorted[i].value.size();
+    }
+
+    // TODO: in big parallel runs, the "trials" are the procs
+    // and we really don't want to print them all out....
+    out << setw(testNameW) << "test\t"  
+        << setw(attW) << "atts\t"
+        << setw(fieldW) 
+        << "median\t"
+        << "mean\t"
+        << "stddev\t"
+        << "min\t"
+        << "max\t";
+    for (int i=0; i<maxtrials; i++)
+        out << "trial"<<i<<"\t";
+    out << endl;
+
+    for (int i=0; i<sorted.size(); i++)
+    {
+        Result &r = sorted[i];
+        out << setw(testNameW) << r.test + "\t";
+        out << setw(attW) << r.atts + "\t";
+        out << setw(fieldW) << r.unit + "\t";
+        if (r.GetMedian() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetMedian() << "\t";
+        if (r.GetMean() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetMean()   << "\t";
+        if (r.GetStdDev() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetStdDev() << "\t";
+        if (r.GetMin() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetMin()    << "\t";
+        if (r.GetMax() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetMax()    << "\t";
+        for (int j=0; j<r.value.size(); j++)
+        {
+            if (r.value[j] == FLT_MAX)
+                out << "N/A\t";
+            else
+                out << r.value[j] << "\t";
+        }
+
+        out << endl;
+    }
+    out << endl
+        << "Note: Any results marked with (*) had missing values." << endl
+        << "      This can occur on systems with a mixture of" << endl
+        << "      device types or architectural capabilities." << endl;
+}
+
+
+// ****************************************************************************
+//  Method:  ResultDatabase::DumpDetailed
+//
+//  Purpose:
+//    Writes the summary results (min/max/stddev/med/mean), but not
+//    every individual trial.
+//
+//  Arguments:
+//    out        where to print
+//
+//  Programmer:  Jeremy Meredith
+//  Creation:    November 10, 2010
+//
+//  Modifications:
+//    Jeremy Meredith, Thu Nov 11 11:39:57 EST 2010
+//    Added note about (*) missing value tag.
+//
+// ****************************************************************************
+void ResultDatabase::DumpSummary(ostream &out)
+{
+    vector<Result> sorted(results);
+    sort(sorted.begin(), sorted.end());
+
+    const int testNameW = 24 ;
+    const int attW = 12;
+    const int fieldW = 9;
+    out << std::fixed << right << std::setprecision(4);
+
+    // TODO: in big parallel runs, the "trials" are the procs
+    // and we really don't want to print them all out....
+    out << setw(testNameW) << "test\t"  
+        << setw(attW) << "atts\t"
+        << setw(fieldW) 
+        << "units\t"
+        << "median\t"
+        << "mean\t"
+        << "stddev\t"
+        << "min\t"
+        << "max\t";
+    out << endl;
+
+    for (int i=0; i<sorted.size(); i++)
+    {
+        Result &r = sorted[i];
+        out << setw(testNameW) << r.test + "\t";
+        out << setw(attW) << r.atts + "\t";
+        out << setw(fieldW) << r.unit + "\t";
+        if (r.GetMedian() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetMedian() << "\t";
+        if (r.GetMean() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetMean()   << "\t";
+        if (r.GetStdDev() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetStdDev() << "\t";
+        if (r.GetMin() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetMin()    << "\t";
+        if (r.GetMax() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetMax()    << "\t";
+
+        out << endl;
+    }
+    out << endl
+        << "Note: results marked with (*) had missing values such as" << endl
+        << "might occur with a mixture of architectural capabilities." << endl;
+}
+
+// ****************************************************************************
+//  Method:  ResultDatabase::ClearAllResults
+//
+//  Purpose:
+//    Clears all existing results from the ResultDatabase; used for multiple passes
+//    of the same test or multiple tests.
+//
+//  Arguments:
+//
+//  Programmer:  Jeffrey Young
+//  Creation:    September 10th, 2014
+//
+//  Modifications:
+//
+//
+// ****************************************************************************
+void ResultDatabase::ClearAllResults()
+{
+	results.clear();	
+}
+
+// ****************************************************************************
+//  Method:  ResultDatabase::DumpCsv
+//
+//  Purpose:
+//    Writes either detailed or summary results (min/max/stddev/med/mean), but not
+//    every individual trial.
+//
+//  Arguments:
+//    out        file to print CSV results
+//
+//  Programmer:  Jeffrey Young
+//  Creation:    August 28th, 2014
+//
+//  Modifications:
+//
+// ****************************************************************************
+void ResultDatabase::DumpCsv(string fileName)
+{
+    bool emptyFile;
+    vector<Result> sorted(results);
+
+    sort(sorted.begin(), sorted.end());
+
+    //Check to see if the file is empty - if so, add the headers
+    emptyFile = this->IsFileEmpty(fileName);
+
+    //Open file and append by default
+    ofstream out;
+    out.open(fileName.c_str(), std::ofstream::out | std::ofstream::app); 
+
+    //Add headers only for empty files
+    if(emptyFile)
+    {
+    // TODO: in big parallel runs, the "trials" are the procs
+    // and we really don't want to print them all out....
+    out << "test, "
+        << "atts, "
+        << "units, "
+        << "median, "
+        << "mean, "
+        << "stddev, "
+        << "min, "
+        << "max, ";
+    out << endl;
+    }
+
+    for (int i=0; i<sorted.size(); i++)
+    {
+        Result &r = sorted[i];
+        out << r.test << ", ";
+        out << r.atts << ", ";
+        out << r.unit << ", ";
+        if (r.GetMedian() == FLT_MAX)
+            out << "N/A, ";
+        else
+            out << r.GetMedian() << ", ";
+        if (r.GetMean() == FLT_MAX)
+            out << "N/A, ";
+        else
+            out << r.GetMean()   << ", ";
+        if (r.GetStdDev() == FLT_MAX)
+            out << "N/A, ";
+        else
+            out << r.GetStdDev() << ", ";
+        if (r.GetMin() == FLT_MAX)
+            out << "N/A, ";
+        else
+            out << r.GetMin()    << ", ";
+        if (r.GetMax() == FLT_MAX)
+            out << "N/A, ";
+        else
+            out << r.GetMax()    << ", ";
+
+        out << endl;
+    }
+    out << endl;
+
+    out.close();
+}
+
+// ****************************************************************************
+//  Method:  ResultDatabase::IsFileEmpty
+//
+//  Purpose:
+//    Returns whether a file is empty - used as a helper for CSV printing
+//
+//  Arguments:
+//    file  The input file to check for emptiness
+//
+//  Programmer:  Jeffrey Young
+//  Creation:    August 28th, 2014
+//
+//  Modifications:
+//
+// ****************************************************************************
+
+bool ResultDatabase::IsFileEmpty(string fileName)
+{
+      bool fileEmpty;
+
+      ifstream file(fileName.c_str());
+
+      //If the file doesn't exist it is by definition empty
+      if(!file.good())
+      {
+        return true;
+      }
+      else
+      {
+        fileEmpty = (bool)(file.peek() == ifstream::traits_type::eof());
+        file.close();
+        
+	return fileEmpty;
+      }
+  
+      //Otherwise, return false  
+        return false;
+}
+
+
+
+// ****************************************************************************
+//  Method:  ResultDatabase::GetResultsForTest
+//
+//  Purpose:
+//    Returns a vector of results for just one test name.
+//
+//  Arguments:
+//    test       the name of the test results to search for
+//
+//  Programmer:  Jeremy Meredith
+//  Creation:    December  3, 2010
+//
+//  Modifications:
+//
+// ****************************************************************************
+vector<ResultDatabase::Result>
+ResultDatabase::GetResultsForTest(const string &test)
+{
+    // get only the given test results
+    vector<Result> retval;
+    for (int i=0; i<results.size(); i++)
+    {
+        Result &r = results[i];
+        if (r.test == test)
+            retval.push_back(r);
+    }
+    return retval;
+}
+
+// ****************************************************************************
+//  Method:  ResultDatabase::GetResults
+//
+//  Purpose:
+//    Returns all the results.
+//
+//  Arguments:
+//
+//  Programmer:  Jeremy Meredith
+//  Creation:    December  3, 2010
+//
+//  Modifications:
+//
+// ****************************************************************************
+const vector<ResultDatabase::Result> &
+ResultDatabase::GetResults() const
+{
+    return results;
+}
diff --git a/hipamd/samples/1_Utils/hipCommander/ResultDatabase.h b/hipamd/samples/1_Utils/hipCommander/ResultDatabase.h
new file mode 100644
index 0000000000..4b63a02a1f
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/ResultDatabase.h
@@ -0,0 +1,100 @@
+#ifndef RESULT_DATABASE_H
+#define RESULT_DATABASE_H
+
+#include <string>
+#include <vector>
+#include <iostream>
+#include <fstream>
+#include <cfloat>
+using std::string;
+using std::vector;
+using std::ostream;
+using std::ofstream;
+using std::ifstream;
+
+
+// ****************************************************************************
+// Class:  ResultDatabase
+//
+// Purpose:
+//   Track numerical results as they are generated.
+//   Print statistics of raw results.
+//
+// Programmer:  Jeremy Meredith
+// Creation:    June 12, 2009
+//
+// Modifications:
+//    Jeremy Meredith, Wed Nov 10 14:20:47 EST 2010
+//    Split timing reports into detailed and summary.  E.g. for serial code,
+//    we might report all trial values, but skip them in parallel.
+//
+//    Jeremy Meredith, Thu Nov 11 11:40:18 EST 2010
+//    Added check for missing value tag.
+//
+//    Jeremy Meredith, Mon Nov 22 13:37:10 EST 2010
+//    Added percentile statistic.
+//
+//    Jeremy Meredith, Fri Dec  3 16:30:31 EST 2010
+//    Added a method to extract a subset of results based on test name.  Also,
+//    the Result class is now public, so that clients can use them directly.
+//    Added a GetResults method as well, and made several functions const.
+//
+// ****************************************************************************
+class ResultDatabase
+{
+  public:
+    //
+    // A performance result for a single SHOC benchmark run.
+    //
+    struct Result
+    {
+        string test;  // e.g. "readback"
+        string atts;  // e.g. "pagelocked 4k^2"
+        string unit;  // e.g. "MB/sec"
+        vector<double> value; // e.g. "837.14"
+        double GetMin() const;
+        double GetMax() const;
+        double GetMedian() const;
+        double GetPercentile(double q) const;
+        double GetMean() const;
+        double GetStdDev() const;
+
+        bool operator<(const Result &rhs) const;
+
+        bool HadAnyFLTMAXValues() const
+        {
+            for (int i=0; i<value.size(); ++i)
+            {
+                if (value[i] >= FLT_MAX)
+                    return true;
+            }
+            return false;
+        }
+    };
+
+  protected:
+    vector<Result> results;
+
+  public:
+    void AddResult(const string &test,
+                   const string &atts,
+                   const string &unit,
+                   double value);
+    void AddResults(const string &test,
+                    const string &atts,
+                    const string &unit,
+                    const vector<double> &values);
+    vector<Result>        GetResultsForTest(const string &test);
+    const vector<Result> &GetResults() const;
+    void ClearAllResults();
+    void DumpDetailed(ostream&);
+    void DumpSummary(ostream&);
+    void DumpCsv(string fileName);
+
+  private:
+    bool IsFileEmpty(string fileName);
+
+};
+
+
+#endif
diff --git a/hipamd/samples/1_Utils/hipCommander/TODO b/hipamd/samples/1_Utils/hipCommander/TODO
new file mode 100644
index 0000000000..4c835cfced
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/TODO
@@ -0,0 +1,50 @@
+_ Add AQL kernel.
+_ Fix &*kernel command so the kernel name/type is an argument not a new command.
+
+_ Add command to parse only.
+_ Add regression to parse all the hcm files.
+
+_ Partition HCC, HIP, HSA, OpenCL commands into separate files.
+
+
+_ Show time for back-to-back copies.
+_ Add variables.
+  %loopcnt
+
+    ./hipCommander %loopcnt=4  
+
+_ Add datasize command.
+
+
+_ Add ( ) to parsing.
+_ Add argument parsing and checking.
+
+_ Add verbose option to print each step of setup.
+  - print deliniater between setup and run.  Add run start message.
+
+  - print sizes of all buffers.
+  - print each command before running.
+  - show start/stop of timer routine.
+
+_ 
+_ Clear documentation on what each oepration does.
+_ Add time instrumentation for each command.
+_ Add pcie atomic.
+
+
+_ Add tests for negative cases, ie endloop w/o opening loop.
+
+
+README tips
+---
+- HIP_API_TRACE combined with -v is useful to track the exact commands generates by hipCommander.
+
+
+Other ideas:
+---
+[ ] Perf guide : stream creation very slow on HCC and should be avoided.
+
+
+Scratch:
+
+
diff --git a/hipamd/samples/1_Utils/hipCommander/c.cmd b/hipamd/samples/1_Utils/hipCommander/c.cmd
new file mode 100644
index 0000000000..db11071203
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/c.cmd
@@ -0,0 +1,3 @@
+loop,1000; H2D; NullKernel; D2H; endloop;
+streamsync;
+printTiming, 1000
diff --git a/hipamd/samples/1_Utils/hipCommander/classic.cmd b/hipamd/samples/1_Utils/hipCommander/classic.cmd
new file mode 100644
index 0000000000..c149eec5f7
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/classic.cmd
@@ -0,0 +1 @@
+H2D; NullKernel, D2H, streamsync
diff --git a/hipamd/samples/1_Utils/hipCommander/hipCommander.cpp b/hipamd/samples/1_Utils/hipCommander/hipCommander.cpp
new file mode 100644
index 0000000000..9c07d066ba
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/hipCommander.cpp
@@ -0,0 +1,1096 @@
+#include <stdio.h>
+#include <iostream>
+#include <sstream>
+#include <iomanip>
+#include <algorithm>
+#include <string>
+#include <typeinfo>
+
+#include <hip/hip_runtime.h>
+#ifdef __HIP_PLATFORM_HCC__
+#include <elf.h>
+#include <hsa/hsa.h>
+#include <hc.hpp>
+#include <hip/hcc_detail/hcc_acc.h>
+#endif
+
+#include <sys/time.h>
+
+#include "ResultDatabase.h"
+#include "nullkernel.hip.cpp"
+
+bool g_printedTiming = false;
+
+// Cmdline parms:
+int           p_device  = 0;
+const char*   p_command = "H2D; NullKernel; D2H";
+const char*   p_file = nullptr;
+unsigned      p_verbose = 0x0;
+unsigned      p_db      = 0x0;
+unsigned      p_blockingSync = 0x0;
+
+//---
+int           p_iterations   = 1;
+
+#define KNRM  "\x1B[0m"
+#define KRED  "\x1B[31m"
+#define KGRN  "\x1B[32m"
+
+
+#define failed(...) \
+    printf ("error: ");\
+    printf (__VA_ARGS__);\
+    printf ("\n");\
+    abort();
+
+
+#define HIPCHECK(error) \
+{\
+    hipError_t localError = error; \
+    if (localError != hipSuccess) { \
+      printf("%serror: '%s'(%d) from %s at %s:%d%s\n", \
+          KRED,hipGetErrorString(localError), localError,\
+          #error,\
+      __FILE__, __LINE__,KNRM); \
+        failed("API returned error code.");\
+    }\
+}
+#define HIPASSERT(condition, msg) \
+    if (! (condition) ) { \
+        failed("%sassertion %s at %s:%d: %s%s\n", \
+               KRED, #condition,\
+              __FILE__, __LINE__,msg, KNRM); \
+    }
+
+
+
+
+
+
+
+int parseInt(const char *str, int *output)
+{
+    char *next;
+    *output = strtol(str, &next, 0);
+    return !strlen(next);
+}
+
+
+void printConfig() {
+    hipDeviceProp_t props;
+    HIPCHECK(hipGetDeviceProperties(&props, p_device));
+
+    printf ("Device:%s Mem=%.1fGB #CUs=%d Freq=%.0fMhz\n", props.name, props.totalGlobalMem/1024.0/1024.0/1024.0, props.multiProcessorCount, props.clockRate/1000.0);
+}
+
+
+
+
+void help() {
+    printf ("Usage: hipBusBandwidth [OPTIONS]\n");
+    printf ("  --file, -f               : Read string of commands from file\n");
+    printf ("  --command, -c            : String specifying commands to run.\n");
+    printf ("  --iterations, -i         : Number of copy iterations to run.\n");
+    printf ("  --device, -d             : Device ID to use (0..numDevices).\n");
+    printf ("  --verbose, -v            : Verbose printing of status.  Fore more info, combine with HIP_TRACE_API on ROCm\n");
+};
+
+
+
+int parseStandardArguments(int argc, char *argv[])
+{
+    for (int i = 1; i < argc; i++) {
+        const char *arg = argv[i];
+
+        if (!strcmp(arg, " ")) {
+            // skip NULL args.
+        } else if (!strcmp(arg, "--iterations") || (!strcmp(arg, "-i"))) {
+            if (++i >= argc || !parseInt(argv[i], &p_iterations)) {
+               failed("Bad --iterations argument"); 
+            }
+
+        } else if (!strcmp(arg, "--device") || (!strcmp(arg, "-d"))) {
+            if (++i >= argc || !parseInt(argv[i], &p_device)) {
+               failed("Bad --device argument"); 
+            }
+
+        } else if (!strcmp(arg, "--file") || (!strcmp(arg, "-f"))) {
+            if (++i >= argc) {
+               failed("Bad --file argument"); 
+            } else {
+                p_file = argv[i];
+            }
+
+        } else if (!strcmp(arg, "--commands") || (!strcmp(arg, "-c"))) {
+            if (++i >= argc) {
+               failed("Bad --commands argument"); 
+            } else {
+                p_command = argv[i];
+            }
+
+        } else if (!strcmp(arg, "--verbose") || (!strcmp(arg, "-v"))) {
+            p_verbose = 1;
+
+        } else if (!strcmp(arg, "--blockingSync") || (!strcmp(arg, "-B"))) {
+            p_blockingSync = 1;
+
+
+        } else if (!strcmp(arg, "--help")  || (!strcmp(arg, "-h"))) {
+            help();
+            exit(EXIT_SUCCESS);
+
+        } else {
+            failed("Bad argument '%s'", arg);
+        }
+    } 
+
+    return 0;
+};
+
+// Returns the current system time in microseconds
+inline long long get_time()
+{
+    struct timeval tv;
+    gettimeofday(&tv, 0);
+    return (tv.tv_sec * 1000000) + tv.tv_usec;
+}
+
+
+class Command;
+
+
+//=================================================================================================
+// A stream of commands , specified as a string.
+class CommandStream {
+public:
+    // State that is inherited by sub-blocks:
+    struct CommandStreamState {
+        hipStream_t                 _currentStream;
+        std::vector<hipStream_t>    _streams;
+        vector<CommandStream *>       _subBlocks;
+    };
+public:
+    CommandStream(std::string commandStreamString, int iterations);
+    ~CommandStream();
+
+    hipStream_t currentStream() const { return _state._currentStream; };
+
+    void print(const std::string &indent="") const;
+    void printBrief(std::ostream &s=std::cout) const ;
+    void run();
+    void recordTime();
+    void printTiming(int iterations=0);
+
+    CommandStream *currentCommandStream()  {
+        return _parseInSubBlock ? _state._subBlocks.back() : this;
+    };
+
+    void enterSubBlock(CommandStream *commandStream) {
+        _parseInSubBlock = true;
+        _state._subBlocks.push_back(commandStream);
+    };
+
+    void exitSubBlock() {
+        _parseInSubBlock = false;
+    };
+
+
+    void setParent(CommandStream *parentCmdStream) 
+    { 
+        _parentCommandStream = parentCmdStream; 
+        _state               = parentCmdStream->_state;
+    };
+    CommandStream * getParent() { return _parentCommandStream; };
+
+    void setStream(int streamIndex);
+
+    CommandStreamState &getState() { return _state; };
+
+private:
+    static void tokenize(const std::string &s, char delim, std::vector<std::string> &tokens);
+    void parse(const std::string fullCmd);
+
+protected:
+    CommandStreamState              _state;
+private:
+
+
+    // List of commands to run in this stream:
+    std::vector<Command*>        _commands;
+
+
+
+    // Number of iterations to run the command loop
+    int                         _iterations;
+
+
+
+
+    // Us to run the the command-stream.  Only valid after run is called.
+    long long                   _startTime;
+    double                      _elapsedUs;
+
+    // Track nested loop of command streams:
+    CommandStream                *_parentCommandStream;
+
+    // Track if we are parsing commands in the subblock.
+    bool                          _parseInSubBlock;
+
+};
+
+
+//=================================================================================================
+class Command {
+public:
+
+    // @p minArgs : Minimum arguments for command.  -1 = don't check.
+    // @p maxArgs : Minimum arguments for command.  0 means min=max, ie exact #arguments expected.  -1 = don't check max.
+    Command(CommandStream *cmdStream, const std::vector<std::string> &args, int minArgs=0, int maxArgs=0) 
+        : _commandStream(cmdStream),
+          _args(args)
+    {
+        int numArgs = args.size() - 1;
+
+        if ((minArgs != -1 ) && (numArgs < minArgs)) {
+            // TODO - print full command here.
+            failed ("Not enough arguments for command %s.  (Expected %d, got %d)", args[0].c_str(), minArgs, numArgs);
+        }
+
+        // Check for an exact number of arguments:
+        if (maxArgs == 0) {
+            maxArgs = minArgs;
+        }
+        if ((maxArgs != -1 ) && (numArgs > maxArgs)) {
+            failed ("Too many arguments for command %s.  (Expected %d, got %d)", args[0].c_str(), maxArgs, numArgs);
+        }
+    };
+
+    void printBrief(std::ostream &s=std::cout) const 
+    {
+        s << _args[0];
+    }
+
+    virtual ~Command() {};
+
+    virtual void print(const std::string &indent = "") const {
+        std::cout << indent << "[";
+        std::for_each(_args.begin(), _args.end(), [] (const std::string &s) {
+            std::cout << s;
+        });
+        std::cout << "]";
+    };
+
+
+    virtual void run() = 0;
+
+protected:
+    int readIntArg(int argIndex, const std::string &argName)
+    {
+        // TODO - catch references to non-existant arguments here.
+        int argVal;
+        try {
+            argVal = std::stoi(_args[argIndex]);
+        } catch (std::invalid_argument) {
+            failed ("Command %s has bad %s argument ('%s')", _args[0].c_str(), argName.c_str(), _args[argIndex].c_str());
+        }
+        return argVal;
+    }
+protected:
+    CommandStream               *_commandStream;
+    std::vector<std::string>     _args;
+};
+
+
+#define FILENAME "nullkernel.hsaco"
+#define KERNEL_NAME "NullKernel"
+
+
+#ifdef __HIP_PLATFORM_HCC__
+//=================================================================================================
+// Use Aql to launch the NULL kernel.
+class AqlKernelCommand : public Command 
+{
+public:
+    AqlKernelCommand(CommandStream *cmdStream, const std::vector<std::string> args) :
+        Command(cmdStream, args) 
+    {
+        hc::accelerator_view *av;
+        HIPCHECK(hipHccGetAcceleratorView(cmdStream->currentStream(), &av));
+
+        hc::accelerator acc = av->get_accelerator();
+
+        hsa_region_t systemRegion = *(hsa_region_t*)acc.get_hsa_am_system_region();
+
+        _hsaAgent = *(hsa_agent_t*) acc.get_hsa_agent();
+
+        std::ifstream file(FILENAME, std::ios::binary | std::ios::ate);
+        std::streamsize fsize = file.tellg();
+        file.seekg(0, std::ios::beg);
+
+        std::vector<char> buffer(fsize);
+        if (file.read(buffer.data(), fsize))
+        {
+            uint64_t elfSize = ElfSize(&buffer[0]);
+
+            assert(fsize == elfSize);
+
+            //TODO - replace module load code with explicit module load and unload.
+
+            hipModule_t module;
+            HIPCHECK(hipModuleLoadData(&module, &buffer[0]));
+            HIPCHECK(hipModuleGetFunction(&_function, module, KERNEL_NAME));
+
+        } else {
+            failed("could not open code object '%s'\n", FILENAME);
+        }
+    };
+
+    ~AqlKernelCommand() {};
+
+    void run() override {
+#define LEN 64
+		uint32_t len = LEN;
+		uint32_t one = 1;
+
+        float *Ad = NULL;
+
+        size_t argSize = 36;
+        char argBuffer[argSize];
+        *(uint32_t*) (&argBuffer[0]) = len;
+        *(uint32_t*) (&argBuffer[4]) = one;
+        *(uint32_t*) (&argBuffer[8]) = one;
+        *(uint32_t*) (&argBuffer[12]) = len;
+        *(uint32_t*) (&argBuffer[16]) = one;
+        *(uint32_t*) (&argBuffer[20]) = one;
+        *(float**)   (&argBuffer[24]) = Ad; // Ad pointer argument
+
+
+        void *config[] = {
+          HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0],
+          HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize,
+          HIP_LAUNCH_PARAM_END
+        };
+
+        hipModuleLaunchKernel(_function, len, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); 
+    };
+
+
+public:
+    hsa_queue_t      _hsaQueue;
+    hsa_agent_t      _hsaAgent;
+
+    hipFunction_t _function;
+
+private:
+    static uint64_t ElfSize(const void *emi){
+        const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi;
+        const Elf64_Shdr *shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff);
+
+        uint64_t max_offset = ehdr->e_shoff;
+        uint64_t total_size = max_offset + ehdr->e_shentsize * ehdr->e_shnum;
+
+        for(uint16_t i=0;i < ehdr->e_shnum;++i){
+            uint64_t cur_offset = static_cast<uint64_t>(shdr[i].sh_offset);
+            if(max_offset < cur_offset){
+                max_offset = cur_offset;
+                total_size = max_offset;
+                if(SHT_NOBITS != shdr[i].sh_type){
+                    total_size += static_cast<uint64_t>(shdr[i].sh_size);
+                }
+            }
+        }
+        return total_size;
+    }
+};
+#endif
+
+//=================================================================================================
+// HCC optimizes away fully NULL kernel calls, so run one that is nearly null:
+class ModuleKernelCommand : public Command 
+{
+public:
+    ModuleKernelCommand(CommandStream *cmdStream, const std::vector<std::string> args) :
+        Command(cmdStream, args),
+        _stream (cmdStream->currentStream()) {
+
+            hipModule_t module;
+            HIPCHECK(hipModuleLoad(&module, FILENAME));
+            HIPCHECK(hipModuleGetFunction(&_function, module, KERNEL_NAME));
+        };
+    ~ModuleKernelCommand() {};
+
+    void run() override {
+#define LEN 64
+		uint32_t len = LEN;
+		uint32_t one = 1;
+
+        float *Ad = NULL;
+
+        size_t argSize = 36;
+        char argBuffer[argSize];
+        *(uint32_t*) (&argBuffer[0]) = len;
+        *(uint32_t*) (&argBuffer[4]) = one;
+        *(uint32_t*) (&argBuffer[8]) = one;
+        *(uint32_t*) (&argBuffer[12]) = len;
+        *(uint32_t*) (&argBuffer[16]) = one;
+        *(uint32_t*) (&argBuffer[20]) = one;
+        *(float**)  (&argBuffer[24]) = Ad; // Ad pointer argument
+
+
+        void *config[] = {
+          HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0],
+          HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize,
+          HIP_LAUNCH_PARAM_END
+        };
+
+        hipModuleLaunchKernel(_function, len, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); 
+    };
+
+
+public:
+    hipFunction_t _function;
+    hipStream_t _stream;
+};
+
+
+class KernelCommand : public Command 
+{
+public:
+    enum Type  {Null, VectorAdd};
+    KernelCommand(CommandStream *cmdStream, const std::vector<std::string> args, Type kind) : 
+        Command(cmdStream, args),
+        _kind(kind),
+        _stream (cmdStream->currentStream()) {
+        };
+    ~KernelCommand() {};
+
+
+    void run() override {
+        static const int gridX  = 64;
+        static const int groupX = 64;
+
+        switch (_kind) {
+            case Null:
+                hipLaunchKernel(NullKernel, dim3(gridX/groupX), dim3(gridX), 0, _stream, nullptr);
+                break;
+            case VectorAdd:
+                assert(0); // TODO
+                break;
+        };
+    }
+private:
+    Type        _kind;
+    hipStream_t _stream;
+};
+
+
+#ifdef __HIP_PLATFORM_HCC__
+//=================================================================================================
+class PfeCommand : public Command 
+{
+public:
+    PfeCommand(CommandStream *cmdStream, const std::vector<std::string> args, hipStream_t stream = 0) : 
+        Command(cmdStream, args)
+    {
+        HIPCHECK(hipHccGetAcceleratorView(stream, &_av));
+    }
+
+    ~PfeCommand() { }
+
+
+    void run() override {
+        static const int gridX  = 64;
+        static const int groupX = 64;
+        auto cf = hc::parallel_for_each(*_av, hc::extent<1>(gridX).tile(groupX),
+        [=](hc::index<1>& idx) __HC__ {
+        });
+    }
+private:
+    hc::accelerator_view *_av;
+};
+#endif
+
+
+
+//=================================================================================================
+class CopyCommand : public Command 
+{
+enum MemType {PinnedHost, UnpinnedHost, Device} ;
+
+public:
+    CopyCommand(CommandStream *cmdStream, const std::vector<std::string> &args, hipMemcpyKind kind, bool isAsync, bool isPinnedHost) ;
+
+    ~CopyCommand() 
+    {
+        if (_dst) {
+            dealloc(_dst, _dstType);
+            _dst = NULL;
+        };
+
+        if (_src) {
+            dealloc(_src, _srcType);
+            _src = NULL;
+        }
+    }
+
+
+    void run() override {
+        if (_isAsync) {
+            HIPCHECK(hipMemcpyAsync(_dst, _src, _sizeBytes, _kind, _stream));
+        } else {
+            HIPCHECK(hipMemcpy(_dst, _src, _sizeBytes, _kind));
+        }
+    };
+
+private:
+    void * alloc(size_t size, MemType memType) {
+        void * p;
+        if (memType == Device) {
+            HIPCHECK(hipMalloc(&p, size));
+
+        } else if (memType == PinnedHost) {
+            HIPCHECK(hipHostMalloc(&p, size));
+
+        } else if (memType == UnpinnedHost) {
+            p = (char*)malloc(size);
+            HIPASSERT(p, "malloc failed");
+
+        } else {
+            HIPASSERT(0, "unsupported memType");
+        }
+
+        return p;
+    };
+
+
+    void dealloc(void *p, MemType memType) {
+        if (memType == Device) {
+            HIPCHECK(hipFree(p));
+        } else if (memType == PinnedHost) {
+            HIPCHECK(hipHostFree(p));
+        } else if (memType == UnpinnedHost) {
+            free(p);
+        } else {
+            HIPASSERT(0, "unsupported memType");
+        }
+    }
+
+
+private:
+    bool          _isAsync;
+    hipStream_t   _stream;
+    hipMemcpyKind _kind;
+
+    size_t  _sizeBytes;
+    void    *_dst;
+    MemType _dstType;
+
+    void    *_src;
+    MemType _srcType;
+};
+
+
+//=================================================================================================
+class DeviceSyncCommand : public Command 
+{
+public:
+    DeviceSyncCommand(CommandStream *cmdStream, const std::vector<std::string> &args) :
+        Command(cmdStream, args) {};
+
+    void run() override {
+        HIPCHECK(hipDeviceSynchronize());
+    };
+};
+
+
+//=================================================================================================
+class StreamSyncCommand : public Command 
+{
+public:
+    StreamSyncCommand(CommandStream *cmdStream, const std::vector<std::string> &args) :
+        Command(cmdStream, args),
+        _stream(cmdStream->currentStream())
+    {};
+
+    const char *help() {
+        return "synchronizes the current stream";
+    };
+
+
+
+    void run() override {
+        HIPCHECK(hipStreamSynchronize(_stream));
+    };
+
+private:
+    hipStream_t _stream;
+};
+
+
+//=================================================================================================
+
+//=================================================================================================
+class LoopCommand : public Command 
+{
+public:
+    LoopCommand(CommandStream *parentCmdStream, const std::vector<std::string> &args) :
+        Command(parentCmdStream, args, 1)
+    {
+        int loopCnt;
+        try {
+            loopCnt = std::stoi(args[1]);
+        } catch (std::invalid_argument) {
+            failed ("bad LOOP_CNT=%s", args[1].c_str());
+        }
+
+        _commandStream = new CommandStream("", loopCnt);
+        _commandStream->setParent(parentCmdStream);
+        parentCmdStream->enterSubBlock(_commandStream);
+        
+    };
+
+
+    void print(const std::string &indent = "") const override {
+        Command::print();
+        _commandStream->print (indent + "  ");
+    };
+
+    void run() override {
+        _commandStream->run();
+    };
+};
+
+
+//=================================================================================================
+class EndBlockCommand : public Command 
+{
+public:
+    EndBlockCommand(CommandStream *blockCmdStream, CommandStream *parentCmdStream, const std::vector<std::string> &args) :
+        Command(parentCmdStream, args, 0, 1),
+        _blockCmdStream(blockCmdStream),
+        _printTiming(0)
+    {
+        int argCnt = args.size()-1;
+        if (argCnt >= 1 ) {
+            _printTiming = readIntArg(1, "PRINT_TIMING");
+        }
+
+        if (parentCmdStream == nullptr) {
+            failed ("%s without corresponding command to start block", args[0].c_str());
+        }
+        parentCmdStream->exitSubBlock();
+    };
+
+    void run() override {
+        if (_printTiming) {
+            _blockCmdStream->printTiming();
+        }
+
+    };
+private:
+
+    CommandStream *_blockCmdStream;
+
+    // print the stream when loop exits.
+    int _printTiming;
+};
+
+
+//=================================================================================================
+class SetStreamCommand : public Command 
+{
+public:
+    SetStreamCommand(CommandStream *cmdStream, const std::vector<std::string> &args) :
+        Command(cmdStream, args, 1) 
+    {
+        int streamIndex = readIntArg(1, "STREAM_INDEX");
+
+        cmdStream->setStream(streamIndex);
+
+    };
+
+    void run() override {
+    };
+};
+
+
+//=================================================================================================
+class PrintTimingCommand : public Command 
+{
+public:
+    PrintTimingCommand(CommandStream *cmdStream, const std::vector<std::string> &args) 
+        : Command(cmdStream, args, 1) 
+    {
+        _iterations = readIntArg(1, "ITERATIONS");
+    };
+
+    void run() override {
+        _commandStream->printTiming(_iterations);
+    };
+
+private:
+    int             _iterations;
+};
+
+
+//=================================================================================================
+CopyCommand::CopyCommand(CommandStream *cmdStream, const std::vector<std::string> &args, 
+                         hipMemcpyKind kind, bool isAsync, bool isPinnedHost) :
+        Command(cmdStream, args) ,
+        _isAsync(isAsync),
+        _kind(kind),
+        _stream(cmdStream->currentStream())
+    {
+        switch (kind) {
+            case hipMemcpyDeviceToHost:
+                _srcType = Device;  
+                _dstType = isPinnedHost ? PinnedHost : UnpinnedHost;
+                break;
+            case hipMemcpyHostToDevice:
+                _srcType = isPinnedHost ? PinnedHost : UnpinnedHost;
+                _dstType = Device;  
+                break;
+            default:
+                HIPASSERT(0, "Unknown hipMemcpyKind");
+        };
+
+        _sizeBytes = 64;  //TODO, support reading from arg.
+
+        _dst = alloc(_sizeBytes, _dstType);
+        _src = alloc(_sizeBytes, _srcType);
+    };
+
+
+//=================================================================================================
+//=================================================================================================
+// Implementations:
+//=================================================================================================
+
+//=================================================================================================
+CommandStream::CommandStream(std::string commandStreamString, int iterations)
+    : _iterations(iterations),
+      _startTime(0),
+      _elapsedUs(0.0),
+      _parentCommandStream(nullptr),
+      _parseInSubBlock(false)
+{
+    std::vector<std::string> tokens;
+    tokenize(commandStreamString, ';', tokens);
+
+
+    std::for_each(tokens.begin(), tokens.end(), [&] (const std::string s) {
+        this->parse(s);
+    });
+
+    setStream(0);
+}
+
+
+CommandStream::~CommandStream()
+{
+    std::for_each(_state._streams.begin(), _state._streams.end(), [&] (hipStream_t s) {
+        if (s) {
+            HIPCHECK(hipStreamDestroy(s));
+        }
+    });
+
+    std::for_each(_commands.begin(), _commands.end(), [&] (Command *c) {
+        delete c;
+    });
+
+
+}
+
+
+void CommandStream::setStream(int streamIndex)
+{
+
+    if (streamIndex >= _state._streams.size()) {
+        _state._streams.resize(streamIndex+1);
+    }
+
+    if (streamIndex && (_state._streams[streamIndex] == nullptr)) {
+        // Create new stream:
+        hipStream_t stream;
+        HIPCHECK(hipStreamCreate(&stream));
+        _state._streams[streamIndex] = stream;
+        _state._currentStream = stream;
+    } else {
+        // Use existing stream:
+        
+        _state._currentStream = _state._streams[streamIndex];
+    }
+   
+}
+
+
+void CommandStream::tokenize(const std::string &s, char delim, std::vector<std::string> &tokens)
+{
+    std::stringstream ss;
+    ss.str(s);
+    std::string item;
+    while (getline(ss, item, delim)) {
+        item.erase (std::remove (item.begin(), item.end(), ' '), item.end()); // remove whitespace.
+        tokens.push_back(item);
+    }
+}
+
+void trim(std::string *s)
+{
+    // trim whitespace from begin and end:
+    const char *t =  "\t\n\r\f\v";
+    s->erase(0, s->find_first_not_of(t));
+    s->erase(s->find_last_not_of(t)+1);
+}
+
+void ltrim(std::string *s)
+{
+    // trim whitespace from begin and end:
+    const char *t =  "\t\n\r\f\v";
+    s->erase(0, s->find_first_not_of(t));
+}
+
+void CommandStream::parse(std::string fullCmd)
+{
+    //convert to lower-case:
+    std::transform(fullCmd.begin(), fullCmd.end(), fullCmd.begin(), ::tolower);
+    trim(&fullCmd);
+
+    if (p_db) {
+        printf ("parse: <%s>\n", fullCmd.c_str());
+    }
+
+
+
+    std::string c;
+    std::vector<std::string> args;
+    size_t leftParenZ = fullCmd.find_first_of('(');
+    if (leftParenZ == string::npos) {
+        c = fullCmd;
+        args.push_back(c);
+    } else {
+        c = fullCmd.substr(0, leftParenZ);
+        args.push_back(c);
+        size_t rightParenZ = fullCmd.find_first_of(')', leftParenZ);
+        std::string argStr = fullCmd.substr(leftParenZ+1, rightParenZ-leftParenZ-1);
+        //printf ("c=%s argstr='%s' leftParenZ=%zu rightParenZ=%zu\n", c.c_str(), argStr.c_str(), leftParenZ, rightParenZ);
+        tokenize(argStr, ',', args);
+
+    }
+
+    
+    
+
+
+    if ((args.size()==0) || (fullCmd.c_str()[0] == '#') ) {
+        if (p_db) {
+            printf ("  skip comment\n");
+        }
+        return;
+    } 
+
+
+
+    Command *cmd = NULL;
+    CommandStream *cmdStream = currentCommandStream();
+    
+    if (c == "h2d") {
+        cmd = new CopyCommand(cmdStream, args, hipMemcpyHostToDevice, true/*isAsync*/, true/*isPinned*/);
+        //= h2d
+        //= Performs an async host-to-device copy of array A_h to A_d.
+        //= The size of these arrays may be set with the datasize command.
+
+    } else if (c == "d2h") {
+        cmd = new CopyCommand(cmdStream, args, hipMemcpyDeviceToHost, true/*isAsync*/, true/*isPinned*/);
+        //= d2h
+        //= Performs an async device-to-host copy of array A_d to A_h.
+        //= The size of these arrays may be set with the datasize command.
+        
+    } else if (c == "modulekernel") {
+        cmd = new ModuleKernelCommand(cmdStream, args);
+
+    } else if (c == "nullkernel") {
+        cmd = new KernelCommand(cmdStream, args, KernelCommand::Null);
+        //= nullkernel
+        //= Dispatches a null kernel to the device.
+
+    } else if (c == "vectoraddkernel") {
+        cmd = new KernelCommand(cmdStream, args, KernelCommand::VectorAdd);
+
+#ifdef __HIP_PLATFORM_HCC__
+    } else if (c == "nullpfe") {
+        cmd = new PfeCommand(cmdStream, args);
+
+    } else if (c == "aqlkernel") {
+        cmd = new AqlKernelCommand(cmdStream, args);
+#endif
+
+    } else if (c == "devicesync") {
+        cmd = new DeviceSyncCommand(cmdStream, args);
+
+    } else if (c == "streamsync") {
+        //= streamsync
+        //= Execute hipStreamSynchronize.  
+        //= This will cause the host thread to wait until the current stream
+        //= completes all pending operations.
+        cmd = new StreamSyncCommand(cmdStream, args);
+
+    } else if (c == "setstream") {
+        //= setstream(STREAM_INDEX);
+        //= Set current stream used by subsequent commands.  
+        //= STREAM_INDEX is index starting from 0...N.    
+        //= This function will create new stream on first call to setstream or re-use previous 
+        //= stream if setstream has already been called with STREAM_INDEX.
+        //= STREAM_INDEX=0 will use the default "null" stream associated with the device, and will not create a new stream.  
+        //=  The default stream has special, conservative synchronization properties. 
+        
+        cmd = new SetStreamCommand(cmdStream, args);
+
+    } else if (c == "printtiming") {
+        cmd = new PrintTimingCommand(cmdStream, args);
+
+    } else if (c == "loop") {
+        //= loop(LOOP_CNT)
+        //= Loop over next set of commands (until 'endloop' command) for LOOP_CNT iterations.
+        //= Loops can be nested.
+        
+        cmd = new LoopCommand(cmdStream, args);
+
+    } else if (c == "endloop") {
+        //= endloop
+        //= End a looped sequence. Must be paired with a preceding loop command. 
+        //= Command between the `loop` and `endloop` must be executed 
+       
+        CommandStream * parentCmdStream = cmdStream->getParent() ;
+        cmd = new EndBlockCommand(cmdStream, parentCmdStream, args);
+        cmdStream = parentCmdStream;
+
+    } else {
+        std::cerr << "error: Bad command '" << fullCmd << "\n";
+        HIPASSERT(0, "bad command in command-stream");
+    }
+
+    if (cmd) {
+        cmdStream->_commands.push_back(cmd);
+    }
+}
+
+
+
+
+void CommandStream::print(const std::string &indent) const
+{
+    for (auto cmdI = _commands.begin(); cmdI != _commands.end(); cmdI++) {
+        (*cmdI)->print(indent);
+    };
+}
+
+
+void CommandStream::printBrief(std::ostream &s) const 
+{
+    for (auto cmdI = _commands.begin(); cmdI != _commands.end(); cmdI++) {
+        (*cmdI)->printBrief(s);
+        s << ";";
+    };
+}
+
+void CommandStream::run() 
+{
+    _startTime = get_time();
+    for (int i=0; i<_iterations; i++) {
+        for (auto cmdI = _commands.begin(); cmdI != _commands.end(); cmdI++) {
+            if (p_verbose) {
+                (*cmdI)->print();
+            }
+            (*cmdI)->run();
+        }
+    }
+
+    // Record time, if not already stored.  (an earlier printTime command will also store the time)
+    recordTime();
+};
+
+void CommandStream::recordTime()
+{
+    if (_elapsedUs == 0.0) {
+        auto stopTime = get_time();
+        _elapsedUs = stopTime - _startTime;
+    }
+}
+
+
+void CommandStream::printTiming(int iterations) 
+{
+
+    if ((_state._subBlocks.size() == 1) && (_commands.size()==1)) {
+        //printf ("print just the loop\n");
+        _state._subBlocks.front()->printTiming(iterations);
+    } else {
+        g_printedTiming = true;
+
+        recordTime();
+        if (iterations == 0) {
+            iterations = _iterations;
+        }
+        std::cout << "command<";  printBrief(std::cout);
+        std::cout << ">," ;
+        printf ("    iterations,%d,   total_time,%6.3f,  time/iteration,%6.3f\n", iterations, _elapsedUs, _elapsedUs/iterations);
+    }
+};
+
+
+
+
+
+//=================================================================================================
+int main(int argc, char *argv[])
+{
+    parseStandardArguments(argc, argv);
+
+    printConfig();
+
+    CommandStream *cs;
+
+    if (p_blockingSync) {
+#ifdef __HIP_PLATFORM_HCC__
+        printf ("setting BlockingSync for AMD\n");
+        setenv("HIP_BLOCKING_SYNC", "1", 1);
+
+#endif
+#ifdef __HIP_PLATFORM_NVCC__
+        printf ("setting cudaDeviceBlockingSync\n");
+        HIPCHECK(hipSetDeviceFlags(cudaDeviceBlockingSync));
+#endif
+    };
+
+
+    if (p_file) {
+        // TODO - catch exception on file IO here:
+        std::ifstream file(p_file);
+        std::string str;
+        std::string file_contents;
+        while (std::getline(file, str))
+        {
+            file_contents += str;
+        }  
+
+        cs = new CommandStream(file_contents, p_iterations);
+
+    } else {
+        cs = new CommandStream(p_command, p_iterations);
+    }
+
+    cs->print();
+    printf ("------\n");
+
+    cs->run();
+    if (!g_printedTiming) {
+        cs->printTiming();
+    }
+
+    delete cs;
+}
+
+
+
+// TODO - add error checking for arguments.
diff --git a/hipamd/samples/1_Utils/hipCommander/l2.hcm b/hipamd/samples/1_Utils/hipCommander/l2.hcm
new file mode 100644
index 0000000000..b541bd6a66
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/l2.hcm
@@ -0,0 +1,3 @@
+setstream,1;
+NullKernel; streamsync;
+loop,10000; H2D;  NullKernel; streamsync; endloop,1;
diff --git a/hipamd/samples/1_Utils/hipCommander/loop.hcm b/hipamd/samples/1_Utils/hipCommander/loop.hcm
new file mode 100644
index 0000000000..db11071203
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/loop.hcm
@@ -0,0 +1,3 @@
+loop,1000; H2D; NullKernel; D2H; endloop;
+streamsync;
+printTiming, 1000
diff --git a/hipamd/samples/1_Utils/hipCommander/loop2.hcm b/hipamd/samples/1_Utils/hipCommander/loop2.hcm
new file mode 100644
index 0000000000..b8a14aa156
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/loop2.hcm
@@ -0,0 +1,2 @@
+setstream,1;
+loop,1000; NullKernel; syncstream; endloop,1,
diff --git a/hipamd/samples/1_Utils/hipCommander/nullkernel.hip.cpp b/hipamd/samples/1_Utils/hipCommander/nullkernel.hip.cpp
new file mode 100644
index 0000000000..890e9bdc1e
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/nullkernel.hip.cpp
@@ -0,0 +1,7 @@
+#include "hip/hip_runtime.h"
+
+extern "C" __global__ void NullKernel(hipLaunchParm lp, float* Ad){
+    if (Ad) {
+        Ad[0] = 42;
+    }
+}
diff --git a/hipamd/samples/1_Utils/hipCommander/nullkernel.hsaco b/hipamd/samples/1_Utils/hipCommander/nullkernel.hsaco
new file mode 100755
index 0000000000..585b55cce5
Binary files /dev/null and b/hipamd/samples/1_Utils/hipCommander/nullkernel.hsaco differ
diff --git a/hipamd/samples/1_Utils/hipCommander/perf/latency2.hcm b/hipamd/samples/1_Utils/hipCommander/perf/latency2.hcm
new file mode 100644
index 0000000000..e43960dc5a
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/perf/latency2.hcm
@@ -0,0 +1,10 @@
+setstream(1);
+NullKernel; streamsync;
+loop(30000); NullKernel; streamsync; endloop(1);
+loop(30000); H2D; H2D;  NullKernel; streamsync; endloop(1);
+loop(30000); H2D; H2D; H2D;  NullKernel; streamsync; endloop(1);
+
+loop(30000); H2D; NullKernel; D2H; streamsync; endloop(1);
+loop(30000); NullKernel; D2H; streamsync; endloop(1);
+loop(30000); NullKernel; D2H; D2H; streamsync; endloop(1);
+loop(30000); NullKernel; D2H; D2H; D2H; streamsync; endloop(1);
diff --git a/hipamd/samples/1_Utils/hipCommander/perf/latency_hostsync.hcm b/hipamd/samples/1_Utils/hipCommander/perf/latency_hostsync.hcm
new file mode 100644
index 0000000000..511ab355d5
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/perf/latency_hostsync.hcm
@@ -0,0 +1,8 @@
+setstream,1;
+NullKernel; streamsync;
+loop,100000; NullKernel; streamsync; endloop,1;
+
+loop,100000; H2D;  streamsync; NullKernel; streamsync; endloop,1;
+
+loop,100000; H2D;  streamsync; NullKernel; streamsync; D2H;  streamsync; endloop,1;
+
diff --git a/hipamd/samples/1_Utils/hipCommander/perf/latency_nosync.hcm b/hipamd/samples/1_Utils/hipCommander/perf/latency_nosync.hcm
new file mode 100644
index 0000000000..c89d738be9
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/perf/latency_nosync.hcm
@@ -0,0 +1,5 @@
+setstream,1;
+NullKernel; streamsync;
+loop,100000; NullKernel; streamsync; endloop,1;
+loop,100000; H2D;  NullKernel; streamsync; endloop,1;
+loop,100000; H2D;  NullKernel; D2H;  streamsync; endloop,1;
diff --git a/hipamd/samples/1_Utils/hipCommander/perf/latency_nullstream.hcm b/hipamd/samples/1_Utils/hipCommander/perf/latency_nullstream.hcm
new file mode 100644
index 0000000000..69345b23b0
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/perf/latency_nullstream.hcm
@@ -0,0 +1,7 @@
+setstream,0;
+NullKernel; streamsync;
+loop,100000; NullKernel; streamsync; endloop,1;
+
+loop,100000; H2D;  NullKernel; streamsync; endloop,1;
+
+loop,100000; H2D;  NullKernel; D2H;  streamsync; endloop,1;
diff --git a/hipamd/samples/1_Utils/hipCommander/perf/modulelaunch_latency.hcm b/hipamd/samples/1_Utils/hipCommander/perf/modulelaunch_latency.hcm
new file mode 100644
index 0000000000..d1d4091fad
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/perf/modulelaunch_latency.hcm
@@ -0,0 +1,5 @@
+setstream(1);
+NullKernel; streamsync;
+loop(100);  ModuleKernel; streamsync; endloop(1);
+loop(100);  AqlKernel; streamsync; endloop(1);
+loop(3000); NullKernel; streamsync; endloop(1);
diff --git a/hipamd/samples/1_Utils/hipCommander/setstream.hcm b/hipamd/samples/1_Utils/hipCommander/setstream.hcm
new file mode 100644
index 0000000000..a7bdd093b8
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/setstream.hcm
@@ -0,0 +1,3 @@
+setstream,1; 
+setstream,2;  H2D; NullKernel; D2H; 
+streamsync
diff --git a/hipamd/samples/1_Utils/hipCommander/testcase.cpp b/hipamd/samples/1_Utils/hipCommander/testcase.cpp
new file mode 100644
index 0000000000..0eeae7c50f
--- /dev/null
+++ b/hipamd/samples/1_Utils/hipCommander/testcase.cpp
@@ -0,0 +1,21 @@
+#include <hip/hip_runtime.h>
+
+static const int BLOCKSIZEX=32;
+static const int BLOCKSIZEY=16;
+
+__global__ void fails(hipLaunchParm lp, float* pErrorI)
+{
+    if(pErrorI!=0)
+    {
+        pErrorI[0]=1;
+    }
+}
+
+int main()
+{
+    dim3 blocks(1,1);
+    dim3 threads(BLOCKSIZEX,BLOCKSIZEY);
+    float error;
+
+    hipLaunchKernel(HIP_KERNEL_NAME(fails), blocks, threads, 0, 0, &error);
+}
diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp
index 79d117a0c4..36790271de 100644
--- a/hipamd/src/hip_hcc.cpp
+++ b/hipamd/src/hip_hcc.cpp
@@ -1749,13 +1749,17 @@ void ihipStream_t::resolveHcMemcpyDirection(unsigned hipMemKind,
 
         if (HIP_FORCE_P2P_HOST & 0x1) {
             *forceUnpinnedCopy = true;
-            tprintf (DB_COPY, "P2P.  Copy engine (dev:%d) can see src and dst but HIP_FORCE_P2P_HOST=0, forcing copy through staging buffers.\n", (*copyDevice)->getDeviceNum());
+            tprintf (DB_COPY, "P2P.  Copy engine (dev:%d agent=0x%lx) can see src and dst but HIP_FORCE_P2P_HOST=0, forcing copy through staging buffers.\n", 
+                    (*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle);
+
         } else {
-            tprintf (DB_COPY, "P2P.  Copy engine (dev:%d) can see src and dst.\n",  (*copyDevice)->getDeviceNum());
+            tprintf (DB_COPY, "P2P.  Copy engine (dev:%d agent=0x%lx) can see src and dst.\n",  
+                    (*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle);
         }
     } else {
         *forceUnpinnedCopy = true;
-        tprintf (DB_COPY, "P2P: copy engine(dev:%d) cannot see both host and device pointers - forcing copy with unpinned engine.\n", (*copyDevice)->getDeviceNum());
+        tprintf (DB_COPY, "P2P: copy engine(dev:%d agent=0x%lx) cannot see both host and device pointers - forcing copy with unpinned engine.\n",
+                    (*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle);
     }
 }
 
@@ -1789,10 +1793,10 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes,
                  dst, dstPtrInfo._appId, dstPtrInfo._isInDeviceMem,
                  src, srcPtrInfo._appId, srcPtrInfo._isInDeviceMem,
                  sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy);
-        tprintf (DB_COPY, "  dst=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d, isDevMem=%d\n",
+        tprintf (DB_COPY, "  dst=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d isDevMem=%d\n",
                  dst, dstPtrInfo._hostPointer, dstPtrInfo._devicePointer, dstPtrInfo._sizeBytes,
                  dstPtrInfo._appId, dstTracked, dstPtrInfo._isInDeviceMem);
-        tprintf (DB_COPY, "  src=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d, isDevMem=%d\n",
+        tprintf (DB_COPY, "  src=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d isDevMem=%d\n",
                  src, srcPtrInfo._hostPointer, srcPtrInfo._devicePointer, srcPtrInfo._sizeBytes,
                  srcPtrInfo._appId, srcTracked, srcPtrInfo._isInDeviceMem);
 
@@ -1846,10 +1850,10 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
                  dst, dstPtrInfo._appId, dstPtrInfo._isInDeviceMem,
                  src, srcPtrInfo._appId, srcPtrInfo._isInDeviceMem,
                  sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy);
-        tprintf (DB_COPY, "  dst=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d, isDevMem=%d\n",
+        tprintf (DB_COPY, "  dst=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d isDevMem=%d\n",
                  dst, dstPtrInfo._hostPointer, dstPtrInfo._devicePointer, dstPtrInfo._sizeBytes,
                  dstPtrInfo._appId, dstTracked, dstPtrInfo._isInDeviceMem);
-        tprintf (DB_COPY, "  src=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d, isDevMem=%d\n",
+        tprintf (DB_COPY, "  src=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d isDevMem=%d\n",
                  src, srcPtrInfo._hostPointer, srcPtrInfo._devicePointer, srcPtrInfo._sizeBytes,
                  srcPtrInfo._appId, srcTracked, srcPtrInfo._isInDeviceMem);