Add more debug info

[ROCm/hip commit: ce92a53f25]

projects/hip/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.
+

projects/hip/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)

projects/hip/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;
+}

projects/hip/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

projects/hip/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:
+
+

projects/hip/samples/1_Utils/hipCommander/c.cmd

@@ -0,0 +1,3 @@
+loop,1000; H2D; NullKernel; D2H; endloop;
+streamsync;
+printTiming, 1000

projects/hip/samples/1_Utils/hipCommander/classic.cmd

@@ -0,0 +1 @@
+H2D; NullKernel, D2H, streamsync

projects/hip/samples/1_Utils/hipCommander/l2.hcm

@@ -0,0 +1,3 @@
+setstream,1;
+NullKernel; streamsync;
+loop,10000; H2D;  NullKernel; streamsync; endloop,1;

projects/hip/samples/1_Utils/hipCommander/loop.hcm

@@ -0,0 +1,3 @@
+loop,1000; H2D; NullKernel; D2H; endloop;
+streamsync;
+printTiming, 1000

projects/hip/samples/1_Utils/hipCommander/loop2.hcm

@@ -0,0 +1,2 @@
+setstream,1;
+loop,1000; NullKernel; syncstream; endloop,1,

projects/hip/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;
+    }
+}

projects/hip/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);

projects/hip/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;
+

projects/hip/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;

projects/hip/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;

projects/hip/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);

projects/hip/samples/1_Utils/hipCommander/setstream.hcm

@@ -0,0 +1,3 @@
+setstream,1; 
+setstream,2;  H2D; NullKernel; D2H; 
+streamsync

projects/hip/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);
+}

projects/hip/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);