hsa-runtime integration

Change-Id: I48968966ffe164218ebff88d0e3a1268e96bf1dd

runtime/hsa-amd-aqlprofile/test/CMakeLists.txt

@@ -0,0 +1,48 @@
+#
+# Header files include path(s).
+#
+include_directories ( $ENV{ROCR_INC_DIR} )
+include_directories ( ${API_DIR} )
+include_directories ( ${TEST_DIR}/util )
+include_directories ( ${TEST_DIR}/ctrl )
+
+#
+# Specify the directory containing the libraries of HsaRt
+# to be linked against for building a Hsa Perf application
+#
+LINK_DIRECTORIES($ENV{ROCR_LIB_DIR})
+find_library ( ROCR_LIB NAMES hsa-runtime64 PATHS $ENV{ROCR_LIB_DIR} )
+
+#
+# Set Name for Common library and build it as a
+# static library to be linked with others
+#
+set ( UTIL_LIB "util${ONLY64STR}" )
+add_subdirectory ( ${TEST_DIR}/util "${PROJECT_BINARY_DIR}/util" )
+
+#
+# Build the test library
+#
+set ( TEST_NAME simple_convolution )
+include_directories ( ${TEST_DIR}/${TEST_NAME} )
+set ( LIB_NAME "${TEST_NAME}${ONLY64STR}" )
+add_library ( ${LIB_NAME} STATIC ${TEST_DIR}/${TEST_NAME}/${TEST_NAME}.cpp )
+target_link_libraries( ${LIB_NAME} c stdc++ )
+set ( TEST_LIBS ${LIB_NAME} )
+
+#
+# Build the test control
+#
+set ( SRC_LIST ${TEST_DIR}/ctrl/test.cpp )
+set ( SRC_LIST ${SRC_LIST} ${TEST_DIR}/ctrl/test_pmgr.cpp )
+set ( SRC_LIST ${SRC_LIST} ${TEST_DIR}/ctrl/test_hsa.cpp )
+set ( LIB_LIST ${TEST_LIBS} ${UTIL_LIB} ${CORE_UTILS_LIB} ${ROCR_LIB} )
+set ( EXE_NAME "ctrl" )
+add_executable ( ${EXE_NAME} ${SRC_LIST} )
+target_link_libraries( ${EXE_NAME} ${LIB_LIST} c stdc++ dl pthread rt atomic )
+
+#
+# Copy the test files
+#
+execute_process ( COMMAND sh -xc "cp ${TEST_DIR}/${TEST_NAME}/*.hsaco ${PROJECT_BINARY_DIR}" )
+execute_process ( COMMAND sh -xc "cp ${TEST_DIR}/run.sh ${PROJECT_BINARY_DIR}" )

runtime/hsa-amd-aqlprofile/test/binary_search/binary_search.cc

@@ -0,0 +1,876 @@
+/*
+ * =============================================================================
+ *   ROC Runtime Conformance Release License
+ * =============================================================================
+ * The University of Illinois/NCSA
+ * Open Source License (NCSA)
+ *
+ * Copyright (c) 2017, Advanced Micro Devices, Inc.
+ * All rights reserved.
+ *
+ * Developed by:
+ *
+ *                 AMD Research and AMD ROC Software Development
+ *
+ *                 Advanced Micro Devices, Inc.
+ *
+ *                 www.amd.com
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal with the Software without restriction, including without limitation
+ * the rights to use, copy, modify, merge, publish, distribute, sublicense,
+ * and/or sell copies of the Software, and to permit persons to whom the
+ * Software is furnished to do so, subject to the following conditions:
+ *
+ *  - Redistributions of source code must retain the above copyright notice,
+ *    this list of conditions and the following disclaimers.
+ *  - Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimers in
+ *    the documentation and/or other materials provided with the distribution.
+ *  - Neither the names of <Name of Development Group, Name of Institution>,
+ *    nor the names of its contributors may be used to endorse or promote
+ *    products derived from this Software without specific prior written
+ *    permission.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
+ * THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
+ * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
+ * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
+ * DEALINGS WITH THE SOFTWARE.
+ *
+ */
+
+#include <assert.h>
+#include <stdint.h>
+#include <string.h>
+#include <fcntl.h>
+#include <unistd.h>
+#include <string>
+#include <iostream>
+#include <climits>
+#include "hsa/hsa.h"
+#include "hsa/hsa_ext_amd.h"
+
+#define RET_IF_HSA_ERR(err) { \
+  if ((err) != HSA_STATUS_SUCCESS) { \
+    std::cout << "hsa api call failure at line " << __LINE__ << ", file: " << \
+              __FILE__ << ". Call returned " << err << std::endl; \
+    return (err); \
+  } \
+}
+
+static const uint32_t kBinarySearchLength = 512;
+static const uint32_t kBinarySearchFindMe = 108;
+static const uint32_t kWorkGroupSize = 256;
+
+// Hold all the info specific to binary search
+typedef struct BinarySearch {
+  // Binary Search parameters
+  uint32_t length;
+  uint32_t work_group_size;
+  uint32_t work_grid_size;
+  uint32_t num_sub_divisions;
+  uint32_t find_me;
+
+  // Buffers needed for this application
+  uint32_t* input;
+  uint32_t* input_arr;
+  uint32_t* input_arr_local;
+  uint32_t* output;
+  // Keneral argument buffers and addresses
+  void* kern_arg_buffer;  // Begin of allocated memory
+  //  this pointer to be deallocated
+  void* kern_arg_address;  // Properly aligned address to be used in aql
+  // packet (don't use for deallocation)
+
+  // Kernel code
+  std::string kernel_file_name;
+  std::string kernel_name;
+  uint32_t kernarg_size;
+  uint32_t kernarg_align;
+
+  // HSA/RocR objects needed for this application
+  hsa_agent_t gpu_dev;
+  hsa_agent_t cpu_dev;
+  hsa_signal_t signal;
+  hsa_queue_t* queue;
+  hsa_amd_memory_pool_t cpu_pool;
+  hsa_amd_memory_pool_t gpu_pool;
+  hsa_amd_memory_pool_t kern_arg_pool;
+
+  // Other items we need to populate AQL packet
+  uint64_t kernel_object;
+  uint32_t group_segment_size;   ///< Kernel group seg size
+  uint32_t private_segment_size;   ///< Kernel private seg size
+} BinarySearch;
+
+void InitializeBinarySearch(BinarySearch* bs) {
+  bs->kernel_file_name = "./binary_search_kernels.hsaco";
+  bs->kernel_name = "binarySearch";
+  bs->length = 512;
+  bs->find_me = 108;
+  bs->work_group_size = 256;
+  bs->num_sub_divisions = bs->length / bs->work_group_size;
+}
+
+// This function is called by the call-back functions used to find an agent of
+// the specified hsa_device_type_t. Note that it cannot be called directly from
+// hsa_iterate_agents() as it does not match the prototype of the call-back
+// function. It must be wrapped by a function with the correct prototype.
+//
+// Return values:
+//  HSA_STATUS_INFO_BREAK -- "agent" is of the specified type (dev_type)
+//  HSA_STATUS_SUCCESS -- "agent" is not of the specified type
+//  Other -- Some error occurred
+static hsa_status_t FindAgent(hsa_agent_t agent, void* data,
+                              hsa_device_type_t dev_type) {
+  if (data == nullptr) {
+    return HSA_STATUS_ERROR_INVALID_ARGUMENT;
+  }
+
+  // See if the provided agent matches the input type (dev_type)
+  hsa_device_type_t hsa_device_type;
+  hsa_status_t hsa_error_code = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE,
+                                &hsa_device_type);
+  RET_IF_HSA_ERR(hsa_error_code);
+
+  if (hsa_device_type == dev_type) {
+    *(reinterpret_cast<hsa_agent_t*>(data)) = agent;
+    return HSA_STATUS_INFO_BREAK;
+  }
+
+  return HSA_STATUS_SUCCESS;
+}
+
+// This is the call-back function used to find a GPU type agent. Note that the
+// prototype of this function is dictated by the HSA specification
+hsa_status_t FindGPUDevice(hsa_agent_t agent, void* data) {
+  return FindAgent(agent, data, HSA_DEVICE_TYPE_GPU);
+}
+
+// This is the call-back function used to find a CPU type agent. Note that the
+// prototype of this function is dictated by the HSA specification
+hsa_status_t FindCPUDevice(hsa_agent_t agent, void* data) {
+  return FindAgent(agent, data, HSA_DEVICE_TYPE_CPU);
+}
+
+// Find the CPU and GPU agents we need to run this sample, and save them in the
+// BinarySearch structure for later use.
+hsa_status_t FindDevices(BinarySearch* bs) {
+  hsa_status_t err;
+
+  // Note that hsa_iterate_agents iterate through all known agents until
+  // HSA_STATUS_SUCCESS is not returned. The call-backs are implemented such
+  // that HSA_STATUS_INFO_BREAK means we found an agent of the specified type.
+  // This value is returned by hsa_iterate_agents.
+  bs->gpu_dev.handle = 0;
+  err = hsa_iterate_agents(FindGPUDevice, &bs->gpu_dev);
+
+  if (err != HSA_STATUS_INFO_BREAK) {
+    return HSA_STATUS_ERROR;
+  }
+
+  bs->cpu_dev.handle = 0;
+  err = hsa_iterate_agents(FindCPUDevice, &bs->cpu_dev);
+
+  if (err != HSA_STATUS_INFO_BREAK) {
+    return HSA_STATUS_ERROR;
+  }
+
+  if (0 == bs->gpu_dev.handle) {
+    std::cout << "GPU Device is not Created properly!" << std::endl;
+    RET_IF_HSA_ERR(HSA_STATUS_ERROR);
+  }
+
+  if (0 == bs->cpu_dev.handle) {
+    std::cout << "CPU Device is not Created properly!" << std::endl;
+    RET_IF_HSA_ERR(HSA_STATUS_ERROR);
+  }
+
+  return HSA_STATUS_SUCCESS;
+}
+
+// This function checks to see if the provided
+// pool has the HSA_AMD_SEGMENT_GLOBAL property. If the kern_arg flag is true,
+// the function adds an additional requirement that the pool have the
+// HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT property. If kern_arg is false,
+// pools must NOT have this property.
+// Upon finding a pool that meets these conditions, HSA_STATUS_INFO_BREAK is
+// returned. HSA_STATUS_SUCCESS is returned if no errors were encountered, but
+// no pool was found meeting the requirements. If an error is encountered, we
+// return that error.
+
+// Note that this function does not match the required prototype for the
+// hsa_amd_agent_iterate_memory_pools call back function, and therefore must be
+// wrapped by a function with the correct prototype.
+static hsa_status_t
+FindGlobalPool(hsa_amd_memory_pool_t pool, void* data, bool kern_arg) {
+  hsa_status_t err;
+  hsa_amd_segment_t segment;
+  uint32_t flag;
+
+  if (nullptr == data) {
+    return HSA_STATUS_ERROR_INVALID_ARGUMENT;
+  }
+
+  err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT,
+                                     &segment);
+  RET_IF_HSA_ERR(err);
+
+  if (HSA_AMD_SEGMENT_GLOBAL != segment) {
+    return HSA_STATUS_SUCCESS;
+  }
+
+  err = hsa_amd_memory_pool_get_info(pool,
+                                HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag);
+  RET_IF_HSA_ERR(err);
+
+  uint32_t karg_st = flag & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT;
+
+  if ((karg_st == 0 && kern_arg) ||
+      (karg_st != 0 && !kern_arg)) {
+    return HSA_STATUS_SUCCESS;
+  }
+
+  *(reinterpret_cast<hsa_amd_memory_pool_t*>(data)) = pool;
+  return HSA_STATUS_INFO_BREAK;
+}
+
+// This is the call-back function for hsa_amd_agent_iterate_memory_pools() that
+// finds a pool with the properties of HSA_AMD_SEGMENT_GLOBAL and that is NOT
+// HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT
+hsa_status_t FindStandardPool(hsa_amd_memory_pool_t pool, void* data) {
+  return FindGlobalPool(pool, data, false);
+}
+
+// This is the call-back function for hsa_amd_agent_iterate_memory_pools() that
+// finds a pool with the properties of HSA_AMD_SEGMENT_GLOBAL and that IS
+// HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT
+hsa_status_t FindKernArgPool(hsa_amd_memory_pool_t pool, void* data) {
+  return FindGlobalPool(pool, data, true);
+}
+
+// Find memory pools that we will need to allocate from for this sample
+// application. We will need memory associated with the host CPU, the GPU
+// executing the kernels, and for kernel arguments. This function will
+// save the found pools to the BinarySearch structure for use elsewhere
+// in this program.
+hsa_status_t FindPools(BinarySearch* bs) {
+  hsa_status_t err;
+
+  err = hsa_amd_agent_iterate_memory_pools(bs->cpu_dev, FindStandardPool,
+        &bs->cpu_pool);
+
+  if (err != HSA_STATUS_INFO_BREAK) {
+    return HSA_STATUS_ERROR;
+  }
+
+  err = hsa_amd_agent_iterate_memory_pools(bs->gpu_dev, FindStandardPool,
+        &bs->gpu_pool);
+
+  if (err != HSA_STATUS_INFO_BREAK) {
+    return HSA_STATUS_ERROR;
+  }
+
+  err = hsa_amd_agent_iterate_memory_pools(bs->cpu_dev,
+        FindKernArgPool, &bs->kern_arg_pool);
+
+  if (err != HSA_STATUS_INFO_BREAK) {
+    return HSA_STATUS_ERROR;
+  }
+
+  return HSA_STATUS_SUCCESS;
+}
+
+// Once the needed memory pools have been found and the BinarySearch structure
+// has been updated with these handles, this function is then used to allocate
+// memory from those pools.
+// Devices with which a pool is associated already have access to the pool.
+// However, other devices may also need to read or write to that memory. Below,
+// we see how we can grant access to other devices to address this issue.
+hsa_status_t AllocateAndInitBuffers(BinarySearch* bs) {
+  hsa_status_t err;
+  uint32_t out_length = 4 * sizeof(uint32_t);
+  uint32_t in_length = bs->num_sub_divisions * 2 * sizeof(uint32_t);
+
+  // In all of these examples, we want both the cpu and gpu to have access to
+  // the buffer in question. We use the array of agents below in the susequent
+  // calls to hsa_amd_agents_allow_access() for this purpose.
+  hsa_agent_t ag_list[2] = {bs->gpu_dev, bs->cpu_dev};
+
+  err = hsa_amd_memory_pool_allocate(bs->cpu_pool, in_length, 0,
+                                     reinterpret_cast<void**>(&bs->input));
+  RET_IF_HSA_ERR(err);
+  err = hsa_amd_agents_allow_access(2, ag_list, NULL, bs->input);
+  RET_IF_HSA_ERR(err);
+  (void)memset(bs->input, 0, in_length);
+
+  err = hsa_amd_memory_pool_allocate(bs->cpu_pool, out_length, 0,
+                                     reinterpret_cast<void**>(&bs->output));
+  RET_IF_HSA_ERR(err);
+  err = hsa_amd_agents_allow_access(2, ag_list, NULL, bs->output);
+  RET_IF_HSA_ERR(err);
+  (void)memset(bs->input, 0, in_length);
+
+  err = hsa_amd_memory_pool_allocate(bs->cpu_pool, in_length, 0,
+                                     reinterpret_cast<void**>(&bs->input_arr));
+  RET_IF_HSA_ERR(err);
+  err = hsa_amd_agents_allow_access(2, ag_list, NULL, bs->input_arr);
+  RET_IF_HSA_ERR(err);
+  (void)memset(bs->input, 0, in_length);
+
+  err = hsa_amd_memory_pool_allocate(bs->cpu_pool, in_length, 0,
+                               reinterpret_cast<void**>(&bs->input_arr_local));
+  RET_IF_HSA_ERR(err);
+  err = hsa_amd_agents_allow_access(2, ag_list, NULL, bs->input_arr_local);
+  RET_IF_HSA_ERR(err);
+
+  // Binary-search application specific code...
+  // Initialize input buffer with random values in an increasing order
+  uint32_t max = bs->length * 20;
+  bs->input[0] = 0;
+
+  uint32_t seed = (unsigned int)time(NULL);
+  srand(seed);
+
+  for (uint32_t i = 1; i < bs->length; ++i) {
+    bs->input[i] = bs->input[i - 1] +
+     static_cast<uint32_t>(max * rand_r(&seed) / static_cast<float>(RAND_MAX));
+  }
+
+// #define VERBOSE 1
+#ifdef VERBOSE
+  std::cout << "Input array values:" << std::endl;
+
+  for (uint32_t i = 0; i < bs->length; ++i) {
+    std::cout << "input[" << i << "] = " << bs->input[i] << " ";
+
+    if (i % 4 == 0) {
+      std::cout << std::endl;
+    }
+  }
+
+  std::cout << std::endl;
+#endif
+
+  return err;
+}
+
+// The code in this function illustrates how to load a kernel from
+// pre-compiled code. The goal is to get a handle that can be later
+// used in an AQL packet and also to extract information about kernel
+// that we will need. All of the information hand kernel handle will
+// be saved to the BinarySearch structure. It will be used when we
+// populate the AQL packet.
+hsa_status_t LoadKernelFromObjFile(BinarySearch* bs) {
+  hsa_status_t err;
+  hsa_code_object_reader_t code_obj_rdr = {0};
+  hsa_executable_t executable = {0};
+
+  hsa_file_t file_handle = open(bs->kernel_file_name.c_str(), O_RDONLY);
+
+  if (file_handle == -1) {
+    std::cout << "failed to open " << bs->kernel_file_name.c_str() <<
+              " at line " << __LINE__ << ", errno: " << errno << std::endl;
+    return HSA_STATUS_ERROR;
+  }
+
+  err = hsa_code_object_reader_create_from_file(file_handle, &code_obj_rdr);
+  RET_IF_HSA_ERR(err);
+  close(file_handle);
+
+  err = hsa_executable_create_alt(HSA_PROFILE_FULL,
+                HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, NULL, &executable);
+  RET_IF_HSA_ERR(err);
+
+  err = hsa_executable_load_agent_code_object(executable, bs->gpu_dev,
+        code_obj_rdr, NULL, NULL);
+  RET_IF_HSA_ERR(err);
+
+  err = hsa_executable_freeze(executable, NULL);
+  RET_IF_HSA_ERR(err);
+
+  hsa_executable_symbol_t kern_sym;
+  err = hsa_executable_get_symbol(executable, NULL, bs->kernel_name.c_str(),
+                                  bs->gpu_dev, 0, &kern_sym);
+  RET_IF_HSA_ERR(err);
+
+  err = hsa_executable_symbol_get_info(kern_sym,
+                                    HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT,
+                                                          &bs->kernel_object);
+  RET_IF_HSA_ERR(err);
+
+  err = hsa_executable_symbol_get_info(kern_sym,
+                      HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE,
+                                                   &bs->private_segment_size);
+  RET_IF_HSA_ERR(err);
+
+  err = hsa_executable_symbol_get_info(kern_sym,
+                        HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE,
+                                                     &bs->group_segment_size);
+  RET_IF_HSA_ERR(err);
+
+  err = hsa_executable_symbol_get_info(kern_sym,
+                      HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE,
+                                                           &bs->kernarg_size);
+  RET_IF_HSA_ERR(err);
+
+  err = hsa_executable_symbol_get_info(kern_sym,
+                 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT,
+                                                          &bs->kernarg_align);
+  RET_IF_HSA_ERR(err);
+
+  return err;
+}
+
+// This function shows how to do an asynchronous copy. We have to create a
+// signal and use the signal to notify us when the copy has completed.
+hsa_status_t AgentMemcpy(void* dst, const void* src,
+                         size_t size, hsa_agent_t dst_ag, hsa_agent_t src_ag) {
+  hsa_signal_t s;
+  hsa_status_t err;
+
+  err = hsa_signal_create(1, 0, NULL, &s);
+  RET_IF_HSA_ERR(err);
+
+  err = hsa_amd_memory_async_copy(dst, dst_ag, src, src_ag, size, 0, NULL, s);
+  RET_IF_HSA_ERR(err);
+
+  if (hsa_signal_wait_scacquire(s, HSA_SIGNAL_CONDITION_LT, 1,
+                                UINT64_MAX, HSA_WAIT_STATE_BLOCKED) != 0) {
+    err = HSA_STATUS_ERROR;
+    std::cout << "Async copy signal error" << std::endl;
+
+    RET_IF_HSA_ERR(err);
+  }
+
+  err = hsa_signal_destroy(s);
+
+  RET_IF_HSA_ERR(err);
+
+  return err;
+}
+
+// AlignDown and AlignUp are 2 utility functions we use to find an aligned
+// boundary either below or above a given value (address). The function will
+// return a value that has the specified alignment.
+static intptr_t
+AlignDown(intptr_t value, size_t alignment) {
+  return (intptr_t) (value & ~(alignment - 1));
+}
+static void*
+AlignUp(void* value, size_t alignment) {
+  return reinterpret_cast<void*>(AlignDown((uintptr_t)
+           (reinterpret_cast<uintptr_t>(value) + alignment - 1), alignment));
+}
+
+// This function populates the AQL patch with the information
+// we have collected and stored in the BinarySearch structure thus far.
+void PopulateAQLPacket(BinarySearch const* bs,
+                       hsa_kernel_dispatch_packet_t* aql) {
+  aql->header = 0;  // Dummy val. for now. Set this right before doorbell ring
+  aql->setup = 1;
+  aql->workgroup_size_x = bs->work_group_size;
+  aql->workgroup_size_y = 1;
+  aql->workgroup_size_z = 1;
+  aql->grid_size_x = bs->work_grid_size;
+  aql->grid_size_y = 1;
+  aql->grid_size_z = 1;
+  aql->private_segment_size = bs->private_segment_size;
+  aql->group_segment_size = bs->group_segment_size;
+  aql->kernel_object = bs->kernel_object;
+  aql->kernarg_address = bs->kern_arg_address;
+  aql->completion_signal = bs->signal;
+
+  return;
+}
+/*
+ * Write everything in the provided AQL packet to the queue except the first 32
+ * bits which include the header and setup fields. That should be done
+ * last.
+ */
+void WriteAQLToQueue(hsa_kernel_dispatch_packet_t const* in_aql,
+                     hsa_queue_t* q) {
+  void* queue_base = q->base_address;
+  const uint32_t queue_mask = q->size - 1;
+  uint64_t que_idx = hsa_queue_add_write_index_relaxed(q, 1);
+
+  hsa_kernel_dispatch_packet_t* queue_aql_packet;
+
+  queue_aql_packet =
+    &(reinterpret_cast<hsa_kernel_dispatch_packet_t*>(queue_base))
+    [que_idx & queue_mask];
+
+  queue_aql_packet->workgroup_size_x = in_aql->workgroup_size_x;
+  queue_aql_packet->workgroup_size_y = in_aql->workgroup_size_y;
+  queue_aql_packet->workgroup_size_z = in_aql->workgroup_size_z;
+  queue_aql_packet->grid_size_x = in_aql->grid_size_x;
+  queue_aql_packet->grid_size_y = in_aql->grid_size_y;
+  queue_aql_packet->grid_size_z = in_aql->grid_size_z;
+  queue_aql_packet->private_segment_size = in_aql->private_segment_size;
+  queue_aql_packet->group_segment_size = in_aql->group_segment_size;
+  queue_aql_packet->kernel_object = in_aql->kernel_object;
+  queue_aql_packet->kernarg_address = in_aql->kernarg_address;
+  queue_aql_packet->completion_signal = in_aql->completion_signal;
+}
+
+// This function allocates memory from the kern_arg pool we already found, and
+// then sets the argument values needed by the kernel code.
+hsa_status_t AllocAndSetKernArgs(BinarySearch* bs, void* args,
+                                 size_t arg_size, void** aql_buf_ptr) {
+  void* kern_arg_buf = nullptr;
+  hsa_status_t err;
+  size_t buf_size;
+  size_t req_align;
+
+  // The kernel code must be written to memory at the correct alignment. We
+  // already queried the executable to get the correct alignment, which is
+  // stored in bs->kernarg_align. In case the memory returned from
+  // hsa_amd_memory_pool is not of the correct alignment, we request a little
+  // more than what we need in case we need to adjust.
+  req_align = bs->kernarg_align;
+  // Allocate enough extra space for alignment adjustments if ncessary
+  buf_size = arg_size + (req_align << 1);
+
+  err = hsa_amd_memory_pool_allocate(bs->kern_arg_pool, buf_size, 0,
+                                     reinterpret_cast<void**>(&kern_arg_buf));
+  RET_IF_HSA_ERR(err);
+
+  // Address of the allocated buffer
+  bs->kern_arg_buffer = kern_arg_buf;
+
+  // Addr. of kern arg start.
+  bs->kern_arg_address = AlignUp(kern_arg_buf, req_align);
+
+  assert(arg_size >= bs->kernarg_size);
+  assert(((uintptr_t)bs->kern_arg_address + arg_size) <
+         ((uintptr_t)bs->kern_arg_buffer + buf_size));
+
+  (void)memcpy(bs->kern_arg_address, args, arg_size);
+  RET_IF_HSA_ERR(err);
+
+  // Make sure both the CPU and GPU can access the kernel arguments
+  hsa_agent_t ag_list[2] = {bs->gpu_dev, bs->cpu_dev};
+  err = hsa_amd_agents_allow_access(2, ag_list, NULL, bs->kern_arg_buffer);
+  RET_IF_HSA_ERR(err);
+
+  // Save this info in our BinarySearch structure for later.
+  *aql_buf_ptr = bs->kern_arg_address;
+
+  return HSA_STATUS_SUCCESS;
+}
+
+// This wrapper atomically writes the provided header and setup to the
+// provided AQL packet. The provided AQL packet address should be in the
+// queue memory space.
+inline void AtomicSetPacketHeader(uint16_t header, uint16_t setup,
+                                  hsa_kernel_dispatch_packet_t* queue_packet) {
+  __atomic_store_n(reinterpret_cast<uint32_t*>(queue_packet),
+                   header | (setup << 16), __ATOMIC_RELEASE);
+}
+
+// Once all the required data for kernel execution is collected (in this
+// application it is stored in the BinarySearch structure) we can put it in
+// an AQL packet and ring the queue door bell to tell the command processor to
+// execute it.
+hsa_status_t Run(BinarySearch* bs) {
+  hsa_status_t err;
+
+  std::cout << "Executing kernel " << bs->kernel_name << std::endl;
+
+  // Adjust the size of workgroup
+  // This is mostly application specific.
+  if (bs->work_group_size > 64) {
+    bs->work_group_size = 64;
+    bs->num_sub_divisions = bs->length / bs->work_group_size;
+
+    if (bs->num_sub_divisions < bs->work_group_size) {
+      bs->num_sub_divisions = bs->work_group_size;
+    }
+
+    bs->work_grid_size = bs->num_sub_divisions;
+  }
+
+  // Explanation of BinarySearch algorithm.
+  /*
+   * Since a plain binary search on the GPU would not achieve much benefit
+   * over the GPU we are doing an N'ary search. We split the array into N
+   * segments every pass and therefore get log (base N) passes instead of log
+   * (base 2) passes.
+   *
+   * In every pass, only the thread that can potentially have the element we
+   * are looking for writes to the output array. For ex: if we are looking to
+   * find 4567 in the array and every thread is searching over a segment of
+   * 1000 values and the input array is 1, 2, 3, 4,... then the first thread
+   * is searching in 1 to 1000, the second one from 1001 to 2000, etc. The
+   * first one does not write to the output. The second one doesn't either.
+   * The fifth one however is from 4001 to 5000. So it can potentially have
+   * the element 4567 which lies between them.
+   *
+   * This particular thread writes to the output the lower bound, upper bound
+   * and whether the element equals the lower bound element. So, it would be
+   * 4001, 5000, 0
+   *
+   * The next pass would subdivide 4001 to 5000 into smaller segments and
+   * continue the same process from there.
+   *
+   * When a pass returns 1 in the third element, it means the element has been
+   * found and we can stop executing the kernel. If the element is not found,
+   * then the execution stops after looking at segment of size 1.
+   */
+
+  uint32_t global_lower_bound = 0;
+  uint32_t global_upper_bound = bs->length - 1;
+  uint32_t sub_div_size = (global_upper_bound - global_lower_bound + 1) /
+                          bs->num_sub_divisions;
+
+  if ((bs->input[0] > bs->find_me) ||
+      (bs->input[bs->length - 1] < bs->find_me)) {
+    bs->output[0] = 0;
+    bs->output[1] = bs->length - 1;
+    bs->output[2] = 0;
+    std::cout << "Returning too early" << std::endl;
+    return HSA_STATUS_SUCCESS;
+  }
+
+  bs->output[3] = 1;
+
+  // Setup the kernel args
+  // See the meta-data for the compiled OpenCL kernel code to ascertain
+  // the sizes, padding and alignment required for kernel arguments.
+  // This can be seen by executing
+  // $ amdgcn-amd-amdhsa-readelf -aw ./binary_search_kernels.hsaco
+  // The kernel code will expect the following arguments aligned as shown.
+  typedef uint32_t uint2[2];
+  typedef uint32_t uint4[4];
+  struct __attribute__((aligned(16))) local_args_t {
+    uint4* outputArray;
+    uint2*  sortedArray;
+    uint32_t findMe;
+    uint32_t pad;
+    uint64_t global_offset_x;
+    uint64_t global_offset_y;
+    uint64_t global_offset_z;
+  } local_args;
+
+  local_args.outputArray = reinterpret_cast<uint4*>(bs->output);
+  local_args.sortedArray = reinterpret_cast<uint2*>(bs->input_arr_local);
+  local_args.findMe = bs->find_me;
+  local_args.global_offset_x = 0;
+  local_args.global_offset_y = 0;
+  local_args.global_offset_z = 0;
+
+  // Copy the kernel args structure into kernel arg memory
+  err = AllocAndSetKernArgs(bs, &local_args, sizeof(local_args),
+                            &bs->kern_arg_address);
+  RET_IF_HSA_ERR(err);
+
+  // Populate an AQL packet with the info we've gathered
+  hsa_kernel_dispatch_packet_t aql;
+  PopulateAQLPacket(bs, &aql);
+
+  uint32_t in_length = bs->num_sub_divisions * 2 * sizeof(uint32_t);
+
+  while ((sub_div_size > 1) && (bs->output[3] != 0)) {
+    for (uint32_t i = 0 ; i < bs->num_sub_divisions; i++) {
+      int idx1 = i * sub_div_size;
+      int idx2 = ((i + 1) * sub_div_size) - 1;
+      bs->input_arr[2 * i] = bs->input[idx1];
+      bs->input_arr[2 * i + 1] = bs->input[idx2];
+    }
+
+    // Copy kernel parameter from system memory to local memory
+    err = AgentMemcpy(reinterpret_cast<uint8_t*>(bs->input_arr_local),
+                      reinterpret_cast<uint8_t*>(bs->input_arr),
+                                        in_length, bs->gpu_dev, bs->cpu_dev);
+
+    RET_IF_HSA_ERR(err);
+
+    // Reset output buffer to zero
+    bs->output[3] = 0;
+
+    // Dispatch kernel with global work size, work group size with ONE dimesion
+    // and wait for kernel to complete
+
+    // Compute the write index of queue and copy Aql packet into it
+    uint64_t que_idx = hsa_queue_load_write_index_relaxed(bs->queue);
+
+    const uint32_t mask = bs->queue->size - 1;
+
+    // This function simply copies the data we've collected so far into our
+    // local AQL packet, except the the setup and header fields.
+    WriteAQLToQueue(&aql, bs->queue);
+
+    uint32_t aql_header = HSA_PACKET_TYPE_KERNEL_DISPATCH;
+    aql_header |= HSA_FENCE_SCOPE_SYSTEM <<
+                  HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE;
+    aql_header |= HSA_FENCE_SCOPE_SYSTEM <<
+                  HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE;
+
+    // Set the packet's type, acquire and release fences. This should be done
+    // atomically after all the other fields have been set, using release
+    // memory ordering to ensure all the fields are set when the door bell
+    // signal is activated.
+    void* q_base = bs->queue->base_address;
+
+    AtomicSetPacketHeader(aql_header, aql.setup,
+                      &(reinterpret_cast<hsa_kernel_dispatch_packet_t*>
+                                                   (q_base))[que_idx & mask]);
+
+    // Increment the write index and ring the doorbell to dispatch kernel.
+    hsa_queue_store_write_index_relaxed(bs->queue, (que_idx + 1));
+    hsa_signal_store_relaxed(bs->queue->doorbell_signal, que_idx);
+
+    // Wait on the dispatch signal until the kernel is finished.
+    // Modify the wait condition to HSA_WAIT_STATE_ACTIVE (instead of
+    // HSA_WAIT_STATE_BLOCKED) if polling is needed instead of blocking, as we
+    // have below.
+    // The call below will block until the condition is met. Below we have said
+    // the condition is that the signal value (initiailzed to 1) associated with
+    // the queue is less than 1. When the kernel associated with the queued AQL
+    // packet has completed execution, the signal value is automatically
+    // decremented by the packet processor.
+    hsa_signal_value_t value = hsa_signal_wait_scacquire(bs->signal,
+                               HSA_SIGNAL_CONDITION_LT, 1,
+                               UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
+
+    // value should be 0, or we timed-out
+    if (value) {
+      std::cout << "Timed out waiting for kernel to complete?" << std::endl;
+      RET_IF_HSA_ERR(HSA_STATUS_ERROR);
+    }
+
+    // Reset the signal to its initial value for the next iteration
+    hsa_signal_store_screlease(bs->signal, 1);
+
+    // Binary search algorithm stuff...
+    global_lower_bound = bs->output[0] * sub_div_size;
+    global_upper_bound = global_lower_bound + sub_div_size - 1;
+    sub_div_size = (global_upper_bound - global_lower_bound + 1) /
+                   bs->num_sub_divisions;
+  }
+
+  uint32_t element_index = UINT_MAX;
+
+  for (uint32_t i = global_lower_bound; i <= global_upper_bound; i++) {
+    if (bs->input[i] == bs->find_me) {
+      element_index = i;
+      bs->output[0] = i;
+      bs->output[1] = i + 1;
+      bs->output[2] = 1;
+      break;
+    }
+
+    // Element is not found in region specified
+    // by global lower bound to global upper bound
+    bs->output[2] = 0;
+  }
+
+  uint32_t is_elem_found = bs->output[2];
+
+  std::cout << "Lower bound = " << global_lower_bound << std::endl;
+  std::cout << "Upper bound = " << global_upper_bound << std::endl;
+  std::cout << "Element search for = " << bs->find_me << std::endl;
+
+
+  if (is_elem_found == 1) {
+    std::cout << "Element found at index " << element_index << std::endl;
+  } else {
+    std::cout << "Element value " << bs->find_me << " not found" << std::endl;
+  }
+
+  return HSA_STATUS_SUCCESS;
+}
+
+// Release all the RocR resources we have acquired in this application.
+hsa_status_t CleanUp(BinarySearch* bs) {
+  hsa_status_t err;
+
+  err = hsa_amd_memory_pool_free(bs->input);
+  RET_IF_HSA_ERR(err);
+
+  err = hsa_amd_memory_pool_free(bs->output);
+  RET_IF_HSA_ERR(err);
+
+  err = hsa_amd_memory_pool_free(bs->input_arr);
+  RET_IF_HSA_ERR(err);
+
+  err = hsa_amd_memory_pool_free(bs->kern_arg_buffer);
+  RET_IF_HSA_ERR(err);
+
+  err = hsa_queue_destroy(bs->queue);
+  RET_IF_HSA_ERR(err);
+
+  err = hsa_signal_destroy(bs->signal);
+  RET_IF_HSA_ERR(err);
+
+  err = hsa_shut_down();
+  RET_IF_HSA_ERR(err);
+
+  return HSA_STATUS_SUCCESS;
+}
+
+int main(int argc, char* argv[]) {
+  // This BinarySearch structure (bs) below holds all of the appl. specific
+  // info we need to run the sample. This includes algorithm specific
+  // information as well as handles to RocR/HSA objects.
+
+  // The basic structure of this sample is to fill in this structure with the
+  // required RocR/HSA handles to RocR resources (e.g., agents, memory pools,
+  // queues, etc.) and then dispatch the packets to the queue, and examine the
+  // output.
+
+  BinarySearch bs;
+  hsa_status_t err;
+
+  // Set some working values specific to this application
+  InitializeBinarySearch(&bs);
+
+  // hsa_init() initializes internal data structures and causes devices
+  // (agents), memory pools and other resources to be discovered.
+  err = hsa_init();
+  RET_IF_HSA_ERR(err);
+
+  // Find the agents needed for the sample
+  err = FindDevices(&bs);
+  RET_IF_HSA_ERR(err);
+
+  // Create the completion signal used when dispatching a packet
+  err = hsa_signal_create(1, 0, NULL, &bs.signal);
+  RET_IF_HSA_ERR(err);
+
+  // Create a queue to submit our binary search AQL packets
+  err = hsa_queue_create(bs.gpu_dev, 128, HSA_QUEUE_TYPE_MULTI, NULL, NULL,
+                         UINT32_MAX, UINT32_MAX, &bs.queue);
+  RET_IF_HSA_ERR(err);
+
+  // Find the HSA memory pools we need to run this sample
+  err = FindPools(&bs);
+  RET_IF_HSA_ERR(err);
+
+  // Allocate memory from the correct memory pool, and initialize them as
+  // neeeded for the algorihm.
+  err = AllocateAndInitBuffers(&bs);
+  RET_IF_HSA_ERR(err);
+
+  // Create a kernel object from the pre-compiled kernel, and read some
+  // attributes associated with the kernel that we will need.
+  err = LoadKernelFromObjFile(&bs);
+  RET_IF_HSA_ERR(err);
+
+  // Fill in the AQL packet, assign the kernel arguments, enqueue the packet,
+  // "ring" the doorbell, and wait for completion.
+  err = Run(&bs);
+  RET_IF_HSA_ERR(err);
+
+  // Release all the RocR resources we've acquired and shutdown HSA.
+  err = CleanUp(&bs);
+
+  return 0;
+}
+
+
+#undef RET_IF_HSA_ERR

runtime/hsa-amd-aqlprofile/test/binary_search/binary_search_kernels.cl

@@ -0,0 +1,127 @@
+/*
+ * =============================================================================
+ *   ROC Runtime Conformance Release License
+ * =============================================================================
+ * The University of Illinois/NCSA
+ * Open Source License (NCSA)
+ *
+ * Copyright (c) 2017, Advanced Micro Devices, Inc.
+ * All rights reserved.
+ *
+ * Developed by:
+ *
+ *                 AMD Research and AMD ROC Software Development
+ *
+ *                 Advanced Micro Devices, Inc.
+ *
+ *                 www.amd.com
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal with the Software without restriction, including without limitation
+ * the rights to use, copy, modify, merge, publish, distribute, sublicense,
+ * and/or sell copies of the Software, and to permit persons to whom the
+ * Software is furnished to do so, subject to the following conditions:
+ *
+ *  - Redistributions of source code must retain the above copyright notice,
+ *    this list of conditions and the following disclaimers.
+ *  - Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimers in
+ *    the documentation and/or other materials provided with the distribution.
+ *  - Neither the names of <Name of Development Group, Name of Institution>,
+ *    nor the names of its contributors may be used to endorse or promote
+ *    products derived from this Software without specific prior written
+ *    permission.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
+ * THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
+ * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
+ * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
+ * DEALINGS WITH THE SOFTWARE.
+ *
+ */
+
+/**
+ * One instance of this kernel call is a thread.
+ * Each thread finds out the segment in which it should look for the element.
+ * After that, it checks if the element is between the lower bound and upper
+ * bound of its segment. If yes, then this segment becomes the total
+ * searchspace for the next pass.
+ *
+ * To achieve this, it writes the lower bound and upper bound to the output
+ * array. In case the element at the left end (lower bound) matches the element
+ * we are looking for, that is marked in the output and we no longer need to
+ * look any further.
+ */
+ 
+__kernel void
+binarySearch(__global uint4 * outputArray,
+             __const __global uint2  * sortedArray,
+             const   unsigned int findMe) {
+  unsigned int tid = get_global_id(0);
+
+  // Then we find the elements  for this thread
+  uint2 element = sortedArray[tid];
+
+
+  // If the element to be found does not lie between
+  // them, then nothing left to do in this thread
+  if((element.x > findMe) || (element.y < findMe)) {
+    return;
+  } else {
+    // However, if the element does lie between the lower
+    // and upper bounds of this thread's searchspace
+    // we need to narrow down the search further in this
+    // search space 
+    // The search space for this thread is marked in the
+    // output as being the total search space for the next pass
+    outputArray[0].x = tid;
+    outputArray[0].w = 1;
+  }
+}
+
+
+__kernel void
+binarySearch_mulkeys(__global int *keys,
+                     __global uint *input,
+                     const unsigned int numKeys,
+                     __global int *output) {
+
+  int gid = get_global_id(0);
+  int lBound = gid * 256;
+  int uBound = lBound + 255;
+
+  for(int i = 0; i < numKeys; i++) {
+    if(keys[i] >= input[lBound] && keys[i] <= input[uBound])
+      output[i]=lBound;
+  }
+
+}
+
+
+__kernel void
+binarySearch_mulkeysConcurrent(__global uint *keys,
+                               __global uint *input,
+                               const unsigned int inputSize, // num. of inputs
+                               const unsigned int numSubdivisions,
+                               __global int *output) {
+
+  int lBound = (get_global_id(0) % numSubdivisions) * (inputSize / numSubdivisions);
+  int uBound = lBound + inputSize / numSubdivisions;
+  int myKey = keys[get_global_id(0) / numSubdivisions];
+  int mid;
+
+  while(uBound >= lBound) {
+    mid = (lBound + uBound) / 2;
+    if(input[mid] == myKey) {
+      output[get_global_id(0) / numSubdivisions] = mid;
+      return;
+    } else if(input[mid] > myKey) {
+      uBound = mid - 1;
+    } else {
+      lBound = mid + 1;
+    }
+  }
+}

runtime/hsa-amd-aqlprofile/test/ctrl/test.cpp

@@ -0,0 +1,91 @@
+/******************************************************************************
+
+Copyright ©2013 Advanced Micro Devices, Inc. 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.
+
+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.
+
+*******************************************************************************/
+
+#include "test_assert.h"
+#include "simple_convolution.h"
+#include "test_hsa.h"
+#include "test_pgen_pmc.h"
+#include "test_pgen_sqtt.h"
+
+int main(int argc, char* argv[]) {
+#if defined(NDEBUG)
+  clog.rdbuf(NULL);
+#endif
+
+  bool ret_val = true;
+
+  // Create SimpleConvolution test object
+  TestKernel* test_kernel = new SimpleConvolution();
+  TestAql* test_aql = new TestHSA(test_kernel);
+
+  const bool pmc_enable = (getenv("ROCR_ENABLE_PMC") != NULL);
+  const bool sqtt_enable = (getenv("ROCR_ENABLE_SQTT") != NULL);
+  if (pmc_enable)
+    test_aql = new TestPGenPMC(test_aql);
+  else if (sqtt_enable)
+    test_aql = new TestPGenSQTT(test_aql);
+  test_assert(test_aql != NULL);
+  if (test_aql == NULL) return 1;
+
+  // Initialization of Hsa Runtime
+  ret_val = test_aql->initialize(argc, argv);
+  if (ret_val == false) {
+    std::cout << "Error in the test initialization" << std::endl;
+    test_assert(ret_val);
+    return 1;
+  }
+
+  // Setup Hsa resources needed for execution
+  ret_val = test_aql->setup();
+  if (ret_val == false) {
+    std::cout << "Error in creating hsa resources" << std::endl;
+    test_assert(ret_val);
+    return 1;
+  }
+
+  // Run SimpleConvolution kernel
+  ret_val = test_aql->run();
+  if (ret_val == false) {
+    std::cout << "Error in running the test kernel" << std::endl;
+    test_assert(ret_val);
+    return 1;
+  }
+
+  // Verify the results of the execution
+  ret_val = test_aql->verify_results();
+  if (ret_val) {
+    std::cout << "Test : Passed" << std::endl;
+  } else {
+    std::cout << "Test : Failed" << std::endl;
+  }
+
+  // Print time taken by sample
+  test_aql->print_time();
+  test_aql->cleanup();
+
+  return (ret_val) ? 0 : 1;
+}

runtime/hsa-amd-aqlprofile/test/ctrl/test_aql.h

@@ -0,0 +1,78 @@
+/******************************************************************************
+
+Copyright ©2013 Advanced Micro Devices, Inc. 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.
+
+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.
+
+*******************************************************************************/
+
+#ifndef _TEST_AQL_H_
+#define _TEST_AQL_H_
+
+#include "hsa.h"
+#include "hsa_rsrc_factory.h"
+#include "hsa_ven_amd_aqlprofile.h"
+
+// Test AQL interface
+class TestAql {
+  TestAql* const test_aql;
+
+ public:
+  explicit TestAql(TestAql* t = 0) : test_aql(t) {}
+  virtual ~TestAql() {}
+
+  TestAql* testAql() { return test_aql; }
+  virtual AgentInfo* getAgentInfo() { return (test_aql) ? test_aql->getAgentInfo() : 0; }
+  virtual hsa_queue_t* getQueue() { return (test_aql) ? test_aql->getQueue() : 0; }
+  virtual HsaRsrcFactory* getRsrcFactory() { return (test_aql) ? test_aql->getRsrcFactory() : 0; }
+
+  // Initialize application environment including setting
+  // up of various configuration parameters based on
+  // command line arguments
+  // @return bool true on success and false on failure
+  virtual bool initialize(int argc, char** argv) {
+    return (test_aql) ? test_aql->initialize(argc, argv) : true;
+  }
+
+  // Setup application parameters for exectuion
+  // @return bool true on success and false on failure
+  virtual bool setup() { return (test_aql) ? test_aql->setup() : true; }
+
+  // Run the kernel
+  // @return bool true on success and false on failure
+  virtual bool run() { return (test_aql) ? test_aql->run() : true; }
+
+  // Verify results
+  // @return bool true on success and false on failure
+  virtual bool verify_results() { return (test_aql) ? test_aql->verify_results() : true; }
+
+  // Print to console the time taken to execute kernel
+  virtual void print_time() {
+    if (test_aql) test_aql->print_time();
+  }
+
+  // Release resources e.g. memory allocations
+  // @return bool true on success and false on failure
+  virtual bool cleanup() { return (test_aql) ? test_aql->cleanup() : true; }
+};
+
+#endif  // _TEST_AQL_H_

runtime/hsa-amd-aqlprofile/test/ctrl/test_assert.h

@@ -0,0 +1,13 @@
+#ifndef _TEST_ASSERT_H_
+#define _TEST_ASSERT_H_
+
+#define test_assert(cond)                                                                          \
+  {                                                                                                \
+    if (!(cond)) {                                                                                 \
+      std::cout << "ASSERT FAILED(" << #cond << ") at \"" << __FILE__ << "\" line " << __LINE__    \
+                << std::endl;                                                                      \
+      exit(-1);                                                                                    \
+    }                                                                                              \
+  }
+
+#endif  // _TEST_ASSERT_H_

runtime/hsa-amd-aqlprofile/test/ctrl/test_hsa.cpp

@@ -0,0 +1,237 @@
+/******************************************************************************
+
+Copyright ©2013 Advanced Micro Devices, Inc. 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.
+
+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.
+
+*******************************************************************************/
+
+#include "test_assert.h"
+#include <atomic>
+
+//#include "os.h"
+#include "helper_funcs.h"
+#include "hsa_rsrc_factory.h"
+#include "test_hsa.h"
+
+bool TestHSA::initialize(int arg_cnt, char** arg_list) {
+  std::cout << "TestHSA::initialize :" << std::endl;
+  // Initialize command line arguments
+  hsa_cmdline_arg_cnt = arg_cnt;
+  hsa_cmdline_arg_list = arg_list;
+
+  // Instantiate a Timer object
+  setup_timer_idx_ = hsa_timer_.CreateTimer();
+  dispatch_timer_idx_ = hsa_timer_.CreateTimer();
+
+  // Instantiate an instance of Hsa Resources Factory
+  hsa_rsrc_ = new HsaRsrcFactory();
+
+  // Print properties of the agents
+  hsa_rsrc_->PrintGpuAgents("> GPU agents");
+
+  // Create an instance of Gpu agent
+  const char* p = getenv("ROCR_AGENT_IND");
+  const uint32_t agent_ind = (p == NULL) ? 0 : atol(p);
+  if (!hsa_rsrc_->GetGpuAgentInfo(agent_ind, &agent_info_)) {
+    std::cout << "> error: agent[" << agent_ind << "] is not found" << std::endl;
+    return false;
+  }
+  std::cout << "> Using agent[" << agent_ind << "] : " << agent_info_->name << std::endl;
+
+  // Create an instance of Aql Queue
+  uint32_t num_pkts = 128;
+  hsa_rsrc_->CreateQueue(agent_info_, num_pkts, &hsa_queue_);
+
+  // Obtain handle of signal
+  hsa_rsrc_->CreateSignal(1, &hsa_signal_);
+
+  // Obtain the code object file name
+  std::string agentName(agent_info_->name);
+  if (agentName.compare(0, 4, "gfx8") == 0) {
+    brig_path_obj_.append("gfx8");
+  } else if (agentName.compare(0, 4, "gfx9") == 0) {
+    brig_path_obj_.append("gfx9");
+  } else {
+    test_assert(false);
+    return false;
+  }
+  brig_path_obj_.append("_" + name_ + ".hsaco");
+
+  return true;
+}
+
+bool TestHSA::setup() {
+  std::cout << "TestHSA::setup :" << std::endl;
+
+  // Start the timer object
+  hsa_timer_.StartTimer(setup_timer_idx_);
+
+  mem_map_t& mem_map = test_->get_mem_map();
+  for (mem_it_t it = mem_map.begin(); it != mem_map.end(); ++it) {
+    mem_descr_t& des = it->second;
+    void* ptr = (des.local) ? hsa_rsrc_->AllocateLocalMemory(agent_info_, des.size)
+                            : hsa_rsrc_->AllocateSysMemory(agent_info_, des.size);
+    des.ptr = ptr;
+    test_assert(ptr != NULL);
+    if (ptr == NULL) return false;
+  }
+  test_->init();
+
+  // Load and Finalize Kernel Code Descriptor
+  char* brig_path = (char*)brig_path_obj_.c_str();
+  const bool ret_val =
+      hsa_rsrc_->LoadAndFinalize(agent_info_, brig_path, strdup(name_.c_str()), &kernel_code_desc_);
+  if (ret_val == false) {
+    std::cout << "Error in loading and finalizing Kernel" << std::endl;
+    return ret_val;
+  }
+
+  // Stop the timer object
+  hsa_timer_.StopTimer(setup_timer_idx_);
+  setup_time_taken_ = hsa_timer_.ReadTimer(setup_timer_idx_);
+  total_time_taken_ = setup_time_taken_;
+
+  return true;
+}
+
+bool TestHSA::run() {
+  std::cout << "TestHSA::run :" << std::endl;
+
+  const uint32_t work_group_size = 64;
+  const uint32_t work_grid_size = test_->get_elements_count();
+  uint32_t group_segment_size = 0;
+  uint32_t private_segment_size = 0;
+  const size_t kernarg_segment_size = test_->get_kernarg_size();
+  uint64_t code_handle = 0;
+
+  // Retrieve the amount of group memory needed
+  hsa_executable_symbol_get_info(
+      kernel_code_desc_, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, &group_segment_size);
+
+  // Retrieve the amount of private memory needed
+  hsa_executable_symbol_get_info(kernel_code_desc_,
+                                 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE,
+                                 &private_segment_size);
+
+  // Check the kernel args size
+  size_t size_info = 0;
+  hsa_executable_symbol_get_info(
+      kernel_code_desc_, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &size_info);
+  test_assert(kernarg_segment_size == size_info);
+  if (kernarg_segment_size != size_info) return false;
+
+  // Retrieve handle of the code block
+  hsa_executable_symbol_get_info(kernel_code_desc_, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT,
+                                 &code_handle);
+
+  // Initialize the dispatch packet.
+  hsa_kernel_dispatch_packet_t aql;
+  memset(&aql, 0, sizeof(aql));
+  // Set the packet's type, barrier bit, acquire and release fences
+  aql.header = HSA_PACKET_TYPE_KERNEL_DISPATCH;
+  aql.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE;
+  aql.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE;
+  // Populate Aql packet with default values
+  aql.setup = 1;
+  aql.grid_size_x = work_grid_size;
+  aql.grid_size_y = 1;
+  aql.grid_size_z = 1;
+  aql.workgroup_size_x = work_group_size;
+  aql.workgroup_size_y = 1;
+  aql.workgroup_size_z = 1;
+  // Bind the kernel code descriptor and arguments
+  aql.kernel_object = code_handle;
+  aql.kernarg_address = test_->get_kernarg_ptr();
+  aql.group_segment_size = group_segment_size;
+  aql.private_segment_size = private_segment_size;
+  // Initialize Aql packet with handle of signal
+  aql.completion_signal = hsa_signal_;
+
+  // Compute the write index of queue and copy Aql packet into it
+  const uint64_t que_idx = hsa_queue_load_write_index_relaxed(hsa_queue_);
+  const uint32_t mask = hsa_queue_->size - 1;
+
+  std::cout << "> Executing kernel: \"" << name_ << "\"" << std::endl;
+
+  // Start the timer object
+  hsa_timer_.StartTimer(dispatch_timer_idx_);
+
+  // Disable packet so that submission to HW is complete
+  const auto header = aql.header;
+  const uint8_t packet_type_mask = (1 << HSA_PACKET_HEADER_WIDTH_TYPE) - 1;
+  aql.header &= (~packet_type_mask) << HSA_PACKET_HEADER_TYPE;
+  aql.header |= HSA_PACKET_TYPE_INVALID << HSA_PACKET_HEADER_TYPE;
+
+  // Copy Aql packet into queue buffer
+  ((hsa_kernel_dispatch_packet_t*)(hsa_queue_->base_address))[que_idx & mask] = aql;
+
+  // After AQL packet is fully copied into queue buffer
+  // update packet header from invalid state to valid state
+  std::atomic_thread_fence(std::memory_order_release);
+  ((hsa_kernel_dispatch_packet_t*)(hsa_queue_->base_address))[que_idx & mask].header = header;
+
+  // Increment the write index and ring the doorbell to dispatch the kernel.
+  hsa_queue_store_write_index_relaxed(hsa_queue_, (que_idx + 1));
+  hsa_signal_store_relaxed(hsa_queue_->doorbell_signal, que_idx);
+
+  std::cout << "> Waiting on kernel dispatch signal" << std::endl;
+
+  // Wait on the dispatch signal until the kernel is finished.
+  // Update wait condition to HSA_WAIT_STATE_ACTIVE for Polling
+  hsa_signal_value_t value = hsa_signal_wait_acquire(hsa_signal_, HSA_SIGNAL_CONDITION_LT, 1,
+                                                     (uint64_t)-1, HSA_WAIT_STATE_BLOCKED);
+
+  // Stop the timer object
+  hsa_timer_.StopTimer(dispatch_timer_idx_);
+  dispatch_time_taken_ = hsa_timer_.ReadTimer(dispatch_timer_idx_);
+  total_time_taken_ += dispatch_time_taken_;
+
+  // Copy kernel buffers from local memory into system memory
+  hsa_rsrc_->TransferData((uint8_t*)test_->get_output_ptr(), (uint8_t*)test_->get_local_ptr(),
+                          test_->get_output_size(), false);
+  test_->print_output();
+
+  return true;
+}
+
+bool TestHSA::verify_results() {
+  // Compare the results and see if they match
+  const int32_t cmp_val =
+      memcmp(test_->get_output_ptr(), test_->get_refout_ptr(), test_->get_output_size());
+  return (cmp_val == 0);
+}
+
+void TestHSA::print_time() {
+  std::cout << "Time taken for Setup by " << this->name_ << " : " << this->setup_time_taken_
+            << std::endl;
+  std::cout << "Time taken for Dispatch by " << this->name_ << " : " << this->dispatch_time_taken_
+            << std::endl;
+  std::cout << "Time taken in Total by " << this->name_ << " : " << this->total_time_taken_
+            << std::endl;
+}
+
+bool TestHSA::cleanup() {
+  // shutdown Hsa Runtime system
+  hsa_status_t ret_val = hsa_shut_down();
+  return (HSA_STATUS_SUCCESS == ret_val);
+}

runtime/hsa-amd-aqlprofile/test/ctrl/test_hsa.h

@@ -0,0 +1,115 @@
+/******************************************************************************
+
+Copyright ©2013 Advanced Micro Devices, Inc. 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.
+
+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.
+
+*******************************************************************************/
+
+#ifndef _TEST_HSA_H_
+#define _TEST_HSA_H_
+
+#include "test_aql.h"
+#include "test_kernel.h"
+#include "hsa_rsrc_factory.h"
+
+// Class implements HSA test
+class TestHSA : public TestAql {
+ public:
+  // Constructor
+  explicit TestHSA(TestKernel* test) : test_(test), name_(test->Name()) {
+    total_time_taken_ = 0;
+    setup_time_taken_ = 0;
+    dispatch_time_taken_ = 0;
+  }
+
+  // Get methods for Agent Info, HAS queue, HSA Resourcse Manager
+  AgentInfo* getAgentInfo() { return agent_info_; }
+  hsa_queue_t* getQueue() { return hsa_queue_; }
+  HsaRsrcFactory* getRsrcFactory() { return hsa_rsrc_; }
+
+  // Initialize application environment including setting
+  // up of various configuration parameters based on
+  // command line arguments
+  // @return bool true on success and false on failure
+  bool initialize(int argc, char** argv);
+
+  // Setup application parameters for exectuion
+  // @return bool true on success and false on failure
+  bool setup();
+
+  // Run the BinarySearch kernel
+  // @return bool true on success and false on failure
+  bool run();
+
+  // Verify against reference implementation
+  // @return bool true on success and false on failure
+  bool verify_results();
+
+  // Print to console the time taken to execute kernel
+  void print_time();
+
+  // Release resources e.g. memory allocations
+  // @return bool true on success and false on failure
+  bool cleanup();
+
+ private:
+  typedef TestKernel::mem_descr_t mem_descr_t;
+  typedef TestKernel::mem_map_t mem_map_t;
+  typedef TestKernel::mem_it_t mem_it_t;
+
+  // Test object
+  TestKernel* test_;
+
+  // Path of Brig file
+  std::string brig_path_obj_;
+
+  // Used to track time taken to run the sample
+  double total_time_taken_;
+  double setup_time_taken_;
+  double dispatch_time_taken_;
+
+  // Handle to an Hsa Gpu Agent
+  AgentInfo* agent_info_;
+
+  // Handle to an Hsa Queue
+  hsa_queue_t* hsa_queue_;
+
+  // Handle of signal
+  hsa_signal_t hsa_signal_;
+
+  // Handle of Kernel Code Descriptor
+  hsa_executable_symbol_t kernel_code_desc_;
+
+  // Instance of timer object
+  uint32_t setup_timer_idx_;
+  uint32_t dispatch_timer_idx_;
+  PerfTimer hsa_timer_;
+
+  // Instance of Hsa Resources Factory
+  HsaRsrcFactory* hsa_rsrc_;
+
+  // Test kernel name
+  std::string name_;
+};
+
+#endif  // _TEST_HSA_H_

runtime/hsa-amd-aqlprofile/test/ctrl/test_kernel.h

@@ -0,0 +1,105 @@
+/******************************************************************************
+
+Copyright ©2013 Advanced Micro Devices, Inc. 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.
+
+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.
+
+*******************************************************************************/
+
+#ifndef _TEST_KERNEL_H_
+#define _TEST_KERNEL_H_
+
+#include <map>
+#include <stdint.h>
+
+// Class implements Kernel test
+class TestKernel {
+ public:
+  // Memory descriptors IDs
+  enum { INPUT_DES_ID, OUTPUT_DES_ID, LOCAL_DES_ID, MASK_DES_ID, KERNARG_DES_ID, REFOUT_DES_ID };
+
+  // Memory descriptors vector declaration
+  struct mem_descr_t {
+    void* ptr;
+    uint32_t size;
+    bool local;
+  };
+
+  // Memory map declaration
+  typedef std::map<uint32_t, mem_descr_t> mem_map_t;
+  typedef mem_map_t::iterator mem_it_t;
+  typedef mem_map_t::const_iterator mem_const_it_t;
+
+  // Initialize method
+  virtual void init() = 0;
+
+  // Return kernel memory map
+  mem_map_t& get_mem_map() { return mem_map_; }
+
+  // Return NULL descriptor
+  static mem_descr_t null_descriptor() { return {0, 0, 0}; }
+
+  // Methods to get the kernel attributes
+  void* get_kernarg_ptr() const { return get_descr(KERNARG_DES_ID).ptr; }
+  uint32_t get_kernarg_size() const { return get_descr(KERNARG_DES_ID).size; }
+  void* get_output_ptr() const { return get_descr(OUTPUT_DES_ID).ptr; }
+  uint32_t get_output_size() const { return get_descr(OUTPUT_DES_ID).size; }
+  void* get_local_ptr() const { return get_descr(LOCAL_DES_ID).ptr; }
+  void* get_refout_ptr() const { return get_descr(REFOUT_DES_ID).ptr; }
+  virtual uint32_t get_elements_count() const = 0;
+
+  // Print output
+  virtual void print_output() const = 0;
+
+  // Return name
+  virtual std::string Name() const = 0;
+
+ protected:
+  // Set system memory descriptor
+  bool set_sys_descr(const uint32_t& id, const uint32_t& size) {
+    return set_mem_descr(id, size, false);
+  }
+
+  // Set local memory descriptor
+  bool set_local_descr(const uint32_t& id, const uint32_t& size) {
+    return set_mem_descr(id, size, true);
+  }
+
+  // Get memory descriptor
+  mem_descr_t get_descr(const uint32_t& id) const {
+    mem_const_it_t it = mem_map_.find(id);
+    return (it != mem_map_.end()) ? it->second : null_descriptor();
+  }
+
+ private:
+  // Set memory descriptor
+  bool set_mem_descr(const uint32_t& id, const uint32_t& size, const bool& local) {
+    const mem_descr_t des = {NULL, size, local};
+    auto ret = mem_map_.insert(mem_map_t::value_type(id, des));
+    return ret.second;
+  }
+
+  // Kernel memory map object
+  mem_map_t mem_map_;
+};
+
+#endif  // _TEST_KERNEL_H_

runtime/hsa-amd-aqlprofile/test/ctrl/test_pgen.h

@@ -0,0 +1,45 @@
+/******************************************************************************
+
+Copyright ©2013 Advanced Micro Devices, Inc. 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.
+
+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.
+
+*******************************************************************************/
+
+#ifndef _TEST_PGEN_H_
+#define _TEST_PGEN_H_
+
+#include "test_pmgr.h"
+
+// SimpleConvolution: Class implements OpenCL SimpleConvolution sample
+class TestPGen : public TestPMgr {
+ protected:
+  typedef hsa_ext_amd_aql_pm4_packet_t packet_t;
+
+  packet_t* PrePacket() { return reinterpret_cast<packet_t*>(&prePacket); }
+  packet_t* PostPacket() { return reinterpret_cast<packet_t*>(&postPacket); }
+
+ public:
+  explicit TestPGen(TestAql* t) : TestPMgr(t) {}
+};
+
+#endif  // _TEST_PGEN_H_

runtime/hsa-amd-aqlprofile/test/ctrl/test_pgen_pmc.h

@@ -0,0 +1,163 @@
+/******************************************************************************
+
+Copyright ©2013 Advanced Micro Devices, Inc. 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.
+
+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.
+
+*******************************************************************************/
+
+#ifndef _TEST_PGEN_PMC_H_
+#define _TEST_PGEN_PMC_H_
+
+#include "test_assert.h"
+#include "test_pgen.h"
+
+#include <vector>
+
+hsa_status_t TestPGenPMC_Callback(hsa_ven_amd_aqlprofile_info_type_t info_type,
+                                  hsa_ven_amd_aqlprofile_info_data_t* info_data,
+                                  void* callback_data) {
+  hsa_status_t status = HSA_STATUS_SUCCESS;
+  typedef std::vector<hsa_ven_amd_aqlprofile_info_data_t> passed_data_t;
+  reinterpret_cast<passed_data_t*>(callback_data)->push_back(*info_data);
+  return status;
+}
+
+// SimpleConvolution: Class implements OpenCL SimpleConvolution sample
+class TestPGenPMC : public TestPGen {
+  const static uint32_t buffer_alignment = 0x1000;  // 4K
+
+  hsa_agent_t agent;
+  hsa_ven_amd_aqlprofile_profile_t profile;
+  hsa_ven_amd_aqlprofile_event_t* events;
+
+  bool buildPackets() { return true; }
+
+  bool dumpData() {
+    std::cout << "TestPGenPMC::dumpData :" << std::endl;
+
+    typedef std::vector<hsa_ven_amd_aqlprofile_info_data_t> callback_data_t;
+
+    callback_data_t data;
+    api.hsa_ven_amd_aqlprofile_iterate_data(&profile, TestPGenPMC_Callback, &data);
+    for (callback_data_t::iterator it = data.begin(); it != data.end(); ++it) {
+      std::cout << dec << "event( block(" << it->pmc_data.event.block_name << "_"
+                << it->pmc_data.event.block_index << "), id(" << it->pmc_data.event.counter_id
+                << ")), sample(" << it->sample_id << "), result(" << it->pmc_data.result << ")"
+                << std::endl;
+    }
+
+    return true;
+  }
+
+ public:
+  explicit TestPGenPMC(TestAql* t) : TestPGen(t) { std::cout << "Test: PGen PMC" << std::endl; }
+
+  bool initialize(int arg_cnt, char** arg_list) {
+    if (!TestPMgr::initialize(arg_cnt, arg_list)) return false;
+
+    hsa_status_t status;
+    hsa_agent_t agent;
+    uint32_t command_buffer_alignment;
+    uint32_t command_buffer_size;
+    uint32_t output_buffer_alignment;
+    uint32_t output_buffer_size;
+
+    // GPU identificator
+    agent = getAgentInfo()->dev_id;
+
+    // Instantiation of the profile object
+    // //////////////////////////////////////////////////////////////
+    // Set the event fields
+    const hsa_ven_amd_aqlprofile_event_t events_arr[] = {
+        {HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 4 /*WAVES*/},
+        {HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 14 /*ITEMS*/},
+        {HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 47 /*WAVE_READY*/},
+        {HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCC, 2, 1 /*CYCLE*/},
+        {HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCC, 2, 3 /*REQ*/},
+        {HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCC, 2, 22 /*WRITEBACK*/},
+        {HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_CPC, 0, 0 /*ALWAYS_COUNT*/},
+        {HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_CPC, 0, 8 /*ME1_STALL_WAIT_ON_RCIU_READ*/},
+    };
+    const size_t event_count = sizeof(events_arr) / sizeof(hsa_ven_amd_aqlprofile_event_t);
+    events = new hsa_ven_amd_aqlprofile_event_t[event_count];
+    memcpy(events, events_arr, sizeof(events_arr));
+
+    // Initialization the profile
+    memset(&profile, 0, sizeof(profile));
+    profile.agent = agent;
+    profile.type = HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_PMC;
+
+    // set enabled events list
+    profile.events = events;
+    profile.event_count = event_count;
+
+    // Profile buffers attributes
+    command_buffer_alignment = buffer_alignment;
+    status = api.hsa_ven_amd_aqlprofile_get_info(
+        &profile, HSA_VEN_AMD_AQLPROFILE_INFO_COMMAND_BUFFER_SIZE, &command_buffer_size);
+    if (status != HSA_STATUS_SUCCESS) {
+      const char* str = "";
+      api.hsa_ven_amd_aqlprofile_error_string(&str);
+      std::cout << "aqlprofile err: " << str << std::endl;
+    }
+    test_assert(status == HSA_STATUS_SUCCESS);
+
+    output_buffer_alignment = buffer_alignment;
+    status = api.hsa_ven_amd_aqlprofile_get_info(
+        &profile, HSA_VEN_AMD_AQLPROFILE_INFO_PMC_DATA_SIZE, &output_buffer_size);
+    test_assert(status == HSA_STATUS_SUCCESS);
+
+    // Application is allocating the command buffer
+    // Allocate(command_buffer_alignment, command_buffer_size,
+    //          MODE_HOST_ACC|MODE_DEV_ACC|MODE_EXEC_DATA)
+    profile.command_buffer.ptr =
+        getRsrcFactory()->AllocateSysMemory(getAgentInfo(), command_buffer_size);
+    profile.command_buffer.size = command_buffer_size;
+
+    // Application is allocating the output buffer
+    // Allocate(output_buffer_alignment, output_buffer_size,
+    //          MODE_HOST_ACC|MODE_DEV_ACC)
+    profile.output_buffer.ptr =
+        getRsrcFactory()->AllocateSysMemory(getAgentInfo(), output_buffer_size);
+    profile.output_buffer.size = output_buffer_size;
+    memset(profile.output_buffer.ptr, 0x77, output_buffer_size);
+
+    // Populating the AQL start packet
+    status = api.hsa_ven_amd_aqlprofile_start(&profile, PrePacket());
+    if (status != HSA_STATUS_SUCCESS) {
+      const char* str;
+      api.hsa_ven_amd_aqlprofile_error_string(&str);
+      std::cout << "aqlprofile err: " << str << std::endl;
+    }
+    test_assert(status == HSA_STATUS_SUCCESS);
+    if (status != HSA_STATUS_SUCCESS) return false;
+
+    // Populating the AQL stop packet
+    status = api.hsa_ven_amd_aqlprofile_stop(&profile, PostPacket());
+    test_assert(status == HSA_STATUS_SUCCESS);
+
+    return (status == HSA_STATUS_SUCCESS);
+  }
+};
+
+#endif  // _TEST_PGEN_PMC_H_

runtime/hsa-amd-aqlprofile/test/ctrl/test_pgen_sqtt.h

@@ -0,0 +1,162 @@
+/******************************************************************************
+
+Copyright ©2013 Advanced Micro Devices, Inc. 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.
+
+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.
+
+*******************************************************************************/
+
+#ifndef _TEST_PGEN_SQTT_H_
+#define _TEST_PGEN_SQTT_H_
+
+#include <iostream>
+#include <iomanip>
+#include <fstream>
+#include <vector>
+
+#include "test_assert.h"
+#include "test_pgen.h"
+
+hsa_status_t TestPGenSQTT_Callback(hsa_ven_amd_aqlprofile_info_type_t info_type,
+                                   hsa_ven_amd_aqlprofile_info_data_t* info_data,
+                                   void* callback_data) {
+  hsa_status_t status = HSA_STATUS_SUCCESS;
+  typedef std::vector<hsa_ven_amd_aqlprofile_info_data_t> passed_data_t;
+  reinterpret_cast<passed_data_t*>(callback_data)->push_back(*info_data);
+  return status;
+}
+
+// SimpleConvolution: Class implements OpenCL SimpleConvolution sample
+class TestPGenSQTT : public TestPGen {
+  const static uint32_t buffer_alignment = 0x1000;  // 4K
+  const static uint32_t buffer_size = 0x2000000;    // 32M
+
+  hsa_agent_t agent;
+  hsa_ven_amd_aqlprofile_profile_t profile;
+
+  bool buildPackets() { return true; }
+
+  bool dumpData() {
+    std::cout << "TestPGenSQTT::dumpData :" << std::endl;
+
+    typedef std::vector<hsa_ven_amd_aqlprofile_info_data_t> callback_data_t;
+
+    callback_data_t data;
+    api.hsa_ven_amd_aqlprofile_iterate_data(&profile, TestPGenSQTT_Callback, &data);
+    for (callback_data_t::iterator it = data.begin(); it != data.end(); ++it) {
+      std::cout << "> sample(" << dec << it->sample_id << ") ptr(" << hex << it->sqtt_data.ptr
+                << ") size(" << dec << it->sqtt_data.size << ")" << std::endl;
+
+      void* sys_buf = getRsrcFactory()->AllocateSysMemory(getAgentInfo(), it->sqtt_data.size);
+      test_assert(sys_buf != NULL);
+      if (sys_buf == NULL) return HSA_STATUS_ERROR;
+
+      hsa_status_t status = hsa_memory_copy(sys_buf, it->sqtt_data.ptr, it->sqtt_data.size);
+      test_assert(status == HSA_STATUS_SUCCESS);
+      if (status != HSA_STATUS_SUCCESS) return status;
+
+      std::string file_name;
+      file_name.append("sqtt_dump_");
+      file_name.append(std::to_string(it->sample_id));
+      file_name.append(".txt");
+      std::ofstream out_file;
+      out_file.open(file_name);
+
+      // Write the buffer in terms of shorts (16 bits)
+      short* sqtt_data = (short*)sys_buf;
+      for (int i = 0; i < (it->sqtt_data.size / sizeof(short)); ++i) {
+        out_file << std::setw(4) << std::setfill('0') << std::hex << sqtt_data[i] << "\n";
+      }
+
+      out_file.close();
+    }
+
+    return true;
+  }
+
+ public:
+  explicit TestPGenSQTT(TestAql* t) : TestPGen(t) { std::cout << "Test: PGen SQTT" << std::endl; }
+
+  bool initialize(int arg_cnt, char** arg_list) {
+    if (!TestPMgr::initialize(arg_cnt, arg_list)) return false;
+
+    hsa_status_t status;
+    hsa_agent_t agent;
+    uint32_t command_buffer_alignment;
+    uint32_t command_buffer_size;
+    uint32_t output_buffer_alignment;
+    uint32_t output_buffer_size;
+
+    // GPU identificator
+    agent = getAgentInfo()->dev_id;
+
+    // Instantiation of the profile object
+    // //////////////////////////////////////////////////////////////
+    // Set the parameters
+    // parameters = ....;
+
+    // Initialization the profile
+    memset(&profile, 0, sizeof(profile));
+    profile.agent = agent;
+    profile.type = HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_SQTT;
+
+    // set parameters
+    // profile.parameters = &event;
+    // profile.parameter_count = 1;
+
+    // Profile buffers attributes
+    command_buffer_alignment = buffer_alignment;
+    status = api.hsa_ven_amd_aqlprofile_get_info(
+        &profile, HSA_VEN_AMD_AQLPROFILE_INFO_COMMAND_BUFFER_SIZE, &command_buffer_size);
+    test_assert(status == HSA_STATUS_SUCCESS);
+
+    output_buffer_alignment = buffer_alignment;
+    output_buffer_size = buffer_size;
+
+    // Application is allocating the command buffer
+    // AllocateSystem(command_buffer_alignment, command_buffer_size,
+    //                MODE_HOST_ACC|MODE_DEV_ACC|MODE_EXEC_DATA)
+    profile.command_buffer.ptr =
+        getRsrcFactory()->AllocateSysMemory(getAgentInfo(), command_buffer_size);
+    profile.command_buffer.size = command_buffer_size;
+
+    // Application is allocating the output buffer
+    // AllocateLocal(output_buffer_alignment, output_buffer_size,
+    //               MODE_DEV_ACC)
+    profile.output_buffer.ptr =
+        getRsrcFactory()->AllocateLocalMemory(getAgentInfo(), output_buffer_size);
+    profile.output_buffer.size = output_buffer_size;
+
+    // Populating the AQL start packet
+    status = api.hsa_ven_amd_aqlprofile_start(&profile, PrePacket());
+    test_assert(status == HSA_STATUS_SUCCESS);
+    if (status != HSA_STATUS_SUCCESS) return false;
+
+    // Populating the AQL stop packet
+    status = api.hsa_ven_amd_aqlprofile_stop(&profile, PostPacket());
+    test_assert(status == HSA_STATUS_SUCCESS);
+
+    return (status == HSA_STATUS_SUCCESS);
+  }
+};
+
+#endif  // _TEST_PGEN_SQTT_H_

runtime/hsa-amd-aqlprofile/test/ctrl/test_pmgr.cpp

@@ -0,0 +1,130 @@
+/******************************************************************************
+
+Copyright ©2013 Advanced Micro Devices, Inc. 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.
+
+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.
+
+*******************************************************************************/
+
+#include <atomic>
+#include "test_assert.h"
+
+#include "test_pmgr.h"
+
+bool TestPMgr::addPacketGfx9(const packet_t* packet) {
+  packet_t aql_packet = *packet;
+
+  // Compute the write index of queue and copy Aql packet into it
+  uint64_t que_idx = hsa_queue_load_write_index_relaxed(getQueue());
+  const uint32_t mask = getQueue()->size - 1;
+
+  // Disable packet so that submission to HW is complete
+  const auto header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
+  aql_packet.header &= (~((1 << HSA_PACKET_HEADER_WIDTH_TYPE) - 1)) << HSA_PACKET_HEADER_TYPE;
+  aql_packet.header |= HSA_PACKET_TYPE_INVALID << HSA_PACKET_HEADER_TYPE;
+
+  // Copy Aql packet into queue buffer
+  ((packet_t*)(getQueue()->base_address))[que_idx & mask] = aql_packet;
+
+  // After AQL packet is fully copied into queue buffer
+  // update packet header from invalid state to valid state
+  std::atomic_thread_fence(std::memory_order_release);
+  ((packet_t*)(getQueue()->base_address))[que_idx & mask].header = header;
+
+  // Increment the write index and ring the doorbell to dispatch the kernel.
+  hsa_queue_store_write_index_relaxed(getQueue(), (que_idx + 1));
+  hsa_signal_store_relaxed(getQueue()->doorbell_signal, que_idx);
+
+  return true;
+}
+
+bool TestPMgr::addPacketGfx8(const packet_t* packet) {
+  // Create legacy devices PM4 data
+  const hsa_ext_amd_aql_pm4_packet_t* aql_packet = (const hsa_ext_amd_aql_pm4_packet_t*)packet;
+  slot_pm4_s data;
+  api.hsa_ven_amd_aqlprofile_legacy_get_pm4(aql_packet, reinterpret_cast<void*>(data.words));
+
+  // Compute the write index of queue and copy Aql packet into it
+  uint64_t que_idx = hsa_queue_load_write_index_relaxed(getQueue());
+  const uint32_t mask = getQueue()->size - 1;
+
+  // Copy Aql packet into queue buffer
+  packet_t* ptr = ((packet_t*)(getQueue()->base_address)) + (que_idx & mask);
+  slot_pm4_t* slot_pm4 = (slot_pm4_t*)ptr;
+  slot_pm4->store(data, std::memory_order_relaxed);
+
+  // Increment the write index and ring the doorbell to dispatch the kernel.
+  hsa_queue_store_write_index_relaxed(getQueue(), (que_idx + SLOT_PM4_SIZE_AQLP));
+  hsa_signal_store_relaxed(getQueue()->doorbell_signal, que_idx + SLOT_PM4_SIZE_AQLP - 1);
+
+  return true;
+}
+
+bool TestPMgr::addPacket(const packet_t* packet) {
+  const char* agent_name = getAgentInfo()->name;
+  return (strncmp(agent_name, "gfx8", 4) == 0) ? addPacketGfx8(packet) : addPacketGfx9(packet);
+}
+
+bool TestPMgr::run() {
+  // Build Aql Pkts
+  const bool active = buildPackets();
+  if (active) {
+    // Submit Pre-Dispatch Aql packet
+    addPacket(&prePacket);
+  }
+
+  testAql()->run();
+
+  if (active) {
+    // Set post packet completion signal
+    postPacket.completion_signal = postSignal;
+
+    // Submit Post-Dispatch Aql packet
+    addPacket(&postPacket);
+
+    // Wait for Post-Dispatch packet to complete
+    hsa_signal_wait_acquire(postSignal, HSA_SIGNAL_CONDITION_LT, 1, (uint64_t)-1,
+                            HSA_WAIT_STATE_BLOCKED);
+
+    // Dumping profiling data
+    dumpData();
+  }
+
+  return true;
+}
+
+bool TestPMgr::initialize(int argc, char** argv) {
+  TestAql::initialize(argc, argv);
+  hsa_status_t status = hsa_signal_create(1, 0, NULL, &postSignal);
+  test_assert(status == HSA_STATUS_SUCCESS);
+  return (status == HSA_STATUS_SUCCESS);
+}
+
+TestPMgr::TestPMgr(TestAql* t) : TestAql(t) {
+  dummySignal.handle = 0;
+  postSignal = dummySignal;
+
+  hsa_status_t status = hsa_init();
+  test_assert(status == HSA_STATUS_SUCCESS);
+  status = hsa_system_get_extension_table(HSA_EXTENSION_AMD_AQLPROFILE, 1, 0, &api);
+  test_assert(status == HSA_STATUS_SUCCESS);
+}

runtime/hsa-amd-aqlprofile/test/ctrl/test_pmgr.h

@@ -0,0 +1,71 @@
+/******************************************************************************
+
+Copyright ©2013 Advanced Micro Devices, Inc. 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.
+
+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.
+
+*******************************************************************************/
+
+#ifndef _TEST_PMGR_H_
+#define _TEST_PMGR_H_
+
+#include <atomic>
+
+#include "hsa.h"
+#include "test_aql.h"
+#include "hsa_ven_amd_aqlprofile.h"
+
+// SimpleConvolution: Class implements OpenCL SimpleConvolution sample
+class TestPMgr : public TestAql {
+ public:
+  typedef hsa_ext_amd_aql_pm4_packet_t packet_t;
+  explicit TestPMgr(TestAql* t);
+  bool run();
+
+ protected:
+  packet_t prePacket;
+  packet_t postPacket;
+  hsa_signal_t dummySignal;
+  hsa_signal_t postSignal;
+
+  hsa_ven_amd_aqlprofile_1_00_pfn_t api;
+
+  virtual bool buildPackets() { return false; }
+  virtual bool dumpData() { return false; }
+  virtual bool initialize(int argc, char** argv);
+
+ private:
+  enum {
+    SLOT_PM4_SIZE_DW = HSA_VEN_AMD_AQLPROFILE_LEGACY_PM4_PACKET_SIZE / sizeof(uint32_t),
+    SLOT_PM4_SIZE_AQLP = HSA_VEN_AMD_AQLPROFILE_LEGACY_PM4_PACKET_SIZE / sizeof(packet_t)
+  };
+  struct slot_pm4_s {
+    uint32_t words[SLOT_PM4_SIZE_DW];
+  };
+  typedef std::atomic<slot_pm4_s> slot_pm4_t;
+
+  bool addPacket(const packet_t* packet);
+  bool addPacketGfx8(const packet_t* packet);
+  bool addPacketGfx9(const packet_t* packet);
+};
+
+#endif  // _TEST_PMGR_H_

runtime/hsa-amd-aqlprofile/test/run.sh

@@ -0,0 +1,30 @@
+#/bin/sh
+set -x
+
+tbin=./test/ctrl
+
+CDIR=`pwd`
+export LD_LIBRARY_PATH=$CDIR
+
+export HSA_ENABLE_SDMA=0
+export HSA_EMULATE_AQL=1
+
+echo
+echo "Run simple convolution kernel"
+unset ROCR_ENABLE_PMC
+unset ROCR_ENABLE_SQTT
+eval $tbin
+
+echo
+echo "Run with PMC"
+export ROCR_ENABLE_PMC=1
+unset ROCR_ENABLE_SQTT
+eval $tbin
+
+echo
+echo "Run with SQTT"
+unset ROCR_ENABLE_PMC
+export ROCR_ENABLE_SQTT=1
+eval $tbin
+
+

runtime/hsa-amd-aqlprofile/test/simple_convolution/simple_convolution.cl

@@ -0,0 +1,81 @@
+/******************************************************************************
+
+Copyright ©2013 Advanced Micro Devices, Inc. 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.
+
+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.
+
+********************************************************************************/
+
+/**
+ * SimpleConvolution is where each pixel of the output image
+ * is the weighted sum of the neighborhood pixels of the input image
+ * The neighborhood is defined by the dimensions of the mask and 
+ * weight of each neighbor is defined by the mask itself.
+ * @param output Output matrix after performing convolution
+ * @param input  Input  matrix on which convolution is to be performed
+ * @param mask   mask matrix using which convolution was to be performed
+ * @param inputDimensions dimensions of the input matrix
+ * @param maskDimensions  dimensions of the mask matrix
+ */
+__kernel void simpleConvolution(__global  uint  * output,
+                                __global  uint  * input,
+                                __global  float  * mask,
+                                const     uint2  inputDimensions,
+                                const     uint2  maskDimensions) {
+
+  uint tid   = get_global_id(0);
+
+  uint width  = inputDimensions.x;
+  uint height = inputDimensions.y;
+
+  uint x      = tid%width;
+  uint y      = tid/width;
+
+  uint maskWidth  = maskDimensions.x;
+  uint maskHeight = maskDimensions.y;
+
+  uint vstep = (maskWidth  -1)/2;
+  uint hstep = (maskHeight -1)/2;
+
+  // find the left, right, top and bottom indices such that
+  // the indices do not go beyond image boundaires
+  uint left    = (x           <  vstep) ? 0         : (x - vstep);
+  uint right   = ((x + vstep) >= width) ? width - 1 : (x + vstep); 
+  uint top     = (y           <  hstep) ? 0         : (y - hstep);
+  uint bottom  = ((y + hstep) >= height)? height - 1: (y + hstep); 
+
+  // initializing wighted sum value
+  float sumFX = 0;
+
+  for(uint i = left; i <= right; ++i) {
+    for(uint j = top ; j <= bottom; ++j) {
+      // performing wighted sum within the mask boundaries
+      uint maskIndex = (j - (y - hstep)) * maskWidth  + (i - (x - vstep));
+      uint index     = j                 * width      + i;
+      sumFX += ((float)input[index] * mask[maskIndex]);
+    }
+  }
+
+  // To round to the nearest integer
+  sumFX += 0.5f;
+  output[tid] = (uint)sumFX;
+}

runtime/hsa-amd-aqlprofile/test/simple_convolution/simple_convolution.cpp

@@ -0,0 +1,160 @@
+/******************************************************************************
+
+Copyright ©2013 Advanced Micro Devices, Inc. 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.
+
+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.
+
+*******************************************************************************/
+
+#include <iostream>
+#include <string.h>
+
+#include "helper_funcs.h"
+#include "simple_convolution.h"
+
+SimpleConvolution::SimpleConvolution() {
+  width_ = 64;
+  height_ = 64;
+  mask_width_ = 3;
+  mask_height_ = mask_width_;
+
+  if (!isPowerOf2(width_)) {
+    width_ = roundToPowerOf2(width_);
+  }
+
+  if (!isPowerOf2(height_)) {
+    height_ = roundToPowerOf2(height_);
+  }
+
+  if (!(mask_width_ % 2)) {
+    mask_width_++;
+  }
+
+  if (!(mask_height_ % 2)) {
+    mask_height_++;
+  }
+
+  if (width_ * height_ < 256) {
+    width_ = 64;
+    height_ = 64;
+  }
+
+  const uint32_t input_size_bytes = width_ * height_ * sizeof(uint32_t);
+  const uint32_t mask_size_bytes = mask_width_ * mask_height_ * sizeof(float);
+
+  set_sys_descr(KERNARG_DES_ID, sizeof(kernel_args_t));
+  set_sys_descr(INPUT_DES_ID, input_size_bytes);
+  set_sys_descr(OUTPUT_DES_ID, input_size_bytes);
+  set_local_descr(LOCAL_DES_ID, input_size_bytes);
+  set_sys_descr(MASK_DES_ID, mask_size_bytes);
+  set_sys_descr(REFOUT_DES_ID, input_size_bytes);
+}
+
+void SimpleConvolution::init() {
+  std::cout << "SimpleConvolution::init :" << std::endl;
+
+  mem_descr_t input_des = get_descr(INPUT_DES_ID);
+  mem_descr_t local_des = get_descr(LOCAL_DES_ID);
+  mem_descr_t mask_des = get_descr(MASK_DES_ID);
+  mem_descr_t refout_des = get_descr(REFOUT_DES_ID);
+  mem_descr_t kernarg_des = get_descr(KERNARG_DES_ID);
+
+  uint32_t* input = (uint32_t*)input_des.ptr;
+  uint32_t* output_local = (uint32_t*)local_des.ptr;
+  float* mask = (float*)mask_des.ptr;
+  kernel_args_t* kernel_args = (kernel_args_t*)kernarg_des.ptr;
+
+  // random initialisation of input
+  fillRandom<uint32_t>(input, width_, height_, 0, 255);
+
+  // Fill a blurr filter or some other filter of your choice
+  const float val = 1.0f / (mask_width_ * 2.0f - 1.0f);
+  for (uint32_t i = 0; i < (mask_width_ * mask_height_); i++) {
+    mask[i] = 0;
+  }
+  for (uint32_t i = 0; i < mask_width_; i++) {
+    uint32_t y = mask_height_ / 2;
+    mask[y * mask_width_ + i] = val;
+  }
+  for (uint32_t i = 0; i < mask_height_; i++) {
+    uint32_t x = mask_width_ / 2;
+    mask[i * mask_width_ + x] = val;
+  }
+
+  // Print the INPUT array.
+  printArray<uint32_t>("> Input[0]", input, width_, 1);
+  printArray<float>("> Mask", mask, mask_width_, mask_height_);
+
+  // Fill the kernel args
+  kernel_args->arg1 = output_local;
+  kernel_args->arg2 = input;
+  kernel_args->arg3 = mask;
+  kernel_args->arg4 = width_;
+  kernel_args->arg41 = height_;
+  kernel_args->arg5 = mask_width_;
+  kernel_args->arg51 = mask_height_;
+
+  // Calculate the reference output
+  memset(refout_des.ptr, 0, refout_des.size);
+  reference_impl((uint32_t*)refout_des.ptr, input, mask, width_, height_, mask_width_,
+                 mask_height_);
+}
+
+void SimpleConvolution::print_output() const {
+  printArray<uint32_t>("> Output[0]", (uint32_t*)get_output_ptr(), width_, 1);
+}
+
+bool SimpleConvolution::reference_impl(uint32_t* output, const uint32_t* input, const float* mask,
+                                       const uint32_t width, const uint32_t height,
+                                       const uint32_t mask_width, const uint32_t mask_height) {
+  const uint32_t vstep = (mask_width - 1) / 2;
+  const uint32_t hstep = (mask_height - 1) / 2;
+
+  // for each pixel in the input
+  for (uint32_t x = 0; x < width; x++) {
+    for (uint32_t y = 0; y < height; y++) {
+      // find the left, right, top and bottom indices such that
+      // the indices do not go beyond image boundaires
+      const uint32_t left = (x < vstep) ? 0 : (x - vstep);
+      const uint32_t right = ((x + vstep) >= width) ? width - 1 : (x + vstep);
+      const uint32_t top = (y < hstep) ? 0 : (y - hstep);
+      const uint32_t bottom = ((y + hstep) >= height) ? height - 1 : (y + hstep);
+
+      // initializing wighted sum value
+      float sum_fx = 0;
+      for (uint32_t i = left; i <= right; ++i) {
+        for (uint32_t j = top; j <= bottom; ++j) {
+          // performing wighted sum within the mask boundaries
+          uint32_t mask_idx = (j - (y - hstep)) * mask_width + (i - (x - vstep));
+          uint32_t index = j * width + i;
+
+          // to round to the nearest integer
+          sum_fx += ((float)input[index] * mask[mask_idx]);
+        }
+      }
+      sum_fx += 0.5f;
+      output[y * width + x] = uint32_t(sum_fx);
+    }
+  }
+
+  return true;
+}

runtime/hsa-amd-aqlprofile/test/simple_convolution/simple_convolution.h

@@ -0,0 +1,90 @@
+/******************************************************************************
+
+Copyright ©2013 Advanced Micro Devices, Inc. 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.
+
+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.
+
+*******************************************************************************/
+
+#ifndef _SIMPLE_CONVOLUTION_H_
+#define _SIMPLE_CONVOLUTION_H_
+
+#include <vector>
+#include <map>
+
+#include "test_kernel.h"
+
+// SimpleConvolution: Class implements OpenCL SimpleConvolution sample
+class SimpleConvolution : public TestKernel {
+ public:
+  // Constructor
+  SimpleConvolution();
+
+  // Initialize method
+  void init();
+
+  // Return number of compute elements
+  uint32_t get_elements_count() const { return width_ * height_; }
+
+  // Print output
+  void print_output() const;
+
+  // Return name
+  std::string Name() const { return std::string("simpleConvolution"); }
+
+ private:
+  // Local kernel arguments declaration
+  struct kernel_args_t {
+    void* arg1;
+    void* arg2;
+    void* arg3;
+    uint32_t arg4;
+    uint32_t arg41;
+    uint32_t arg5;
+    uint32_t arg51;
+  };
+
+  // Width of the Input array
+  uint32_t width_;
+
+  // Height of the Input array
+  uint32_t height_;
+
+  // Mask dimensions
+  uint32_t mask_width_;
+
+  // Mask dimensions
+  uint32_t mask_height_;
+
+  // Reference CPU implementation of Simple Convolution
+  // @param output Output matrix after performing convolution
+  // @param input  Input  matrix on which convolution is to be performed
+  // @param mask   mask matrix using which convolution was to be performed
+  // @param input_dimensions dimensions of the input matrix
+  // @param mask_dimensions  dimensions of the mask matrix
+  // @return bool true on success and false on failure
+  bool reference_impl(uint32_t* output, const uint32_t* input, const float* mask,
+                      const uint32_t width, const uint32_t height, const uint32_t maskWidth,
+                      const uint32_t maskHeight);
+};
+
+#endif  // _SIMPLE_CONVOLUTION_H_

runtime/hsa-amd-aqlprofile/test/simple_convolution/simple_convolution.hsail

@@ -0,0 +1,154 @@
+module &m:1:0:$full:$large:$default;
+extension "amd:gcn";
+extension "IMAGE";
+
+decl prog function &abort()();
+
+prog kernel &__OpenCL_SimpleConvolution(kernarg_u64 %__global_offset_0,
+                                        kernarg_u64 %output,
+                                        kernarg_u64 %input,
+                                        kernarg_u64 %mask,
+                                        kernarg_u32 %inputDimensions[2],
+                                        kernarg_u32 %maskDimensions[2]) {
+
+	pragma  "AMD RTI", "ARGSTART:__OpenCL_SimpleConvolution";
+	pragma  "AMD RTI", "version:3:1:104";
+	pragma  "AMD RTI", "device:generic";
+	pragma  "AMD RTI", "uniqueid:1024";
+	pragma  "AMD RTI", "memory:private:0";
+	pragma  "AMD RTI", "memory:region:0";
+	pragma  "AMD RTI", "memory:local:0";
+	pragma  "AMD RTI", "value:__global_offset_0:u64:1:1:0";
+	pragma  "AMD RTI", "pointer:output:u32:1:1:96:uav:7:4:RW:0:0:0";
+	pragma  "AMD RTI", "pointer:input:u32:1:1:112:uav:7:4:RW:0:0:0";
+	pragma  "AMD RTI", "pointer:mask:float:1:1:128:uav:7:4:RW:0:0:0";
+	pragma  "AMD RTI", "value:inputDimensions:u32:2:1:144";
+	pragma  "AMD RTI", "constarg:4:inputDimensions";
+	pragma  "AMD RTI", "value:maskDimensions:u32:2:1:160";
+	pragma  "AMD RTI", "constarg:5:maskDimensions";
+	pragma  "AMD RTI", "function:1:0";
+	pragma  "AMD RTI", "memory:64bitABI";
+	pragma  "AMD RTI", "privateid:8";
+	pragma  "AMD RTI", "enqueue_kernel:0";
+	pragma  "AMD RTI", "kernel_index:0";
+	pragma  "AMD RTI", "reflection:0:size_t";
+	pragma  "AMD RTI", "reflection:1:uint*";
+	pragma  "AMD RTI", "reflection:2:uint*";
+	pragma  "AMD RTI", "reflection:3:float*";
+	pragma  "AMD RTI", "reflection:4:uint2";
+	pragma  "AMD RTI", "reflection:5:uint2";
+	pragma  "AMD RTI", "ARGEND:__OpenCL_SimpleConvolution";
+
+  @__OpenCL_SimpleConvolution_Entry:
+	
+  // BB#0:    // %entry
+
+	workitemabsid_u32	$s6, 0;
+	cvt_u64_u32	$d0, $s6;
+	ld_kernarg_align(8)_width(all)_u64	$d4, [%__global_offset_0];
+	add_u64	$d0, $d0, $d4;
+	cvt_u32_u64	$s5, $d0;
+	ld_v2_kernarg_align(4)_width(all)_u32	($s0, $s4), [%inputDimensions];
+	ld_v2_kernarg_align(4)_width(all)_u32	($s1, $s9), [%maskDimensions];
+	rem_u32	$s7, $s5, $s0;
+	add_u32	$s2, $s1, 4294967295;
+	shr_u32	$s8, $s2, 1;
+	add_u32	$s2, $s7, $s8;
+	add_u32	$s3, $s0, 4294967295;
+	cmp_ge_b1_u32	$c0, $s2, $s0;
+	cmov_b32	$s2, $c0, $s3, $s2;
+	sub_u32	$s3, $s7, $s8;
+	cmp_lt_b1_u32	$c0, $s7, $s8;
+	cmov_b32	$s3, $c0, 0, $s3;
+	ld_kernarg_align(8)_width(all)_u64	$d1, [%output];
+	cmp_le_b1_u32	$c0, $s3, $s2;
+	cbr_b1	$c0, @BB0_2;
+
+	// BB#1:
+	
+  mov_b32	$s6, 0;
+	br	@BB0_6;
+
+	// @BB0_2:    // %for.cond32.preheader.lr.ph
+  
+  @BB0_2:
+	
+  div_u32	$s5, $s5, $s0;
+	add_u32	$s9, $s9, 4294967295;
+	shr_u32	$s9, $s9, 1;
+	add_u32	$s10, $s5, $s9;
+	add_u32	$s11, $s4, 4294967295;
+	cmp_ge_b1_u32	$c0, $s10, $s4;
+	cmov_b32	$s4, $c0, $s11, $s10;
+	sub_u32	$s10, $s5, $s9;
+	cmp_lt_b1_u32	$c0, $s5, $s9;
+	cmov_b32	$s5, $c0, 0, $s10;
+	ld_kernarg_align(8)_width(all)_u64	$d2, [%mask];
+	ld_kernarg_align(8)_width(all)_u64	$d3, [%input];
+	cvt_u64_u32	$d5, $s6;
+	add_u64	$d4, $d4, $d5;
+	cvt_u32_u64	$s6, $d4;
+	div_u32	$s6, $s6, $s0;
+	max_u32	$s10, $s9, $s6;
+	sub_u32	$s12, $s10, $s6;
+	max_u32	$s11, $s7, $s8;
+	mov_b32	$s6, 0;
+	mad_u32	$s12, $s1, $s12, $s11;
+	sub_u32	$s7, $s12, $s7;
+	sub_u32	$s9, $s10, $s9;
+	mad_u32	$s9, $s0, $s9, $s11;
+	sub_u32	$s8, $s9, $s8;
+
+	// @BB0_3:    // %for.cond32.preheader
+
+  @BB0_3:
+	
+  cmp_gt_b1_u32	$c0, $s5, $s4;
+	mov_b32	$s9, $s7;
+	mov_b32	$s10, $s8;
+	mov_b32	$s11, $s5;
+	cbr_b1	$c0, @BB0_5;
+
+  // @BB0_4:    // %for.body35
+
+  @BB0_4:
+	
+	cvt_u64_u32	$d4, $s9;
+	shl_u64	$d4, $d4, 2;
+	add_u64	$d4, $d2, $d4;
+	ld_global_align(4)_f32	$s12, [$d4];
+	cvt_u64_u32	$d4, $s10;
+	shl_u64	$d4, $d4, 2;
+	add_u64	$d4, $d3, $d4;
+	ld_global_align(4)_u32	$s13, [$d4];
+	cvt_f32_u32	$s13, $s13;
+	mul_ftz_f32	$s12, $s13, $s12;
+	add_u32	$s9, $s9, $s1;
+	add_u32	$s10, $s10, $s0;
+	add_u32	$s11, $s11, 1;
+	add_ftz_f32	$s6, $s6, $s12;
+	cmp_le_b1_u32	$c0, $s11, $s4;
+	cbr_b1	$c0, @BB0_4;
+
+	// @BB0_5:    // %for.inc48
+  
+  @BB0_5:
+	
+  add_u32	$s7, $s7, 1;
+	add_u32	$s8, $s8, 1;
+	add_u32	$s3, $s3, 1;
+	cmp_le_b1_u32	$c0, $s3, $s2;
+	cbr_b1	$c0, @BB0_3;
+
+	// @BB0_6:    // %for.end50
+
+  @BB0_6:
+	
+  and_b64	$d0, $d0, 4294967295;
+	shl_u64	$d0, $d0, 2;
+	add_u64	$d0, $d1, $d0;
+	add_ftz_f32	$s0, $s6, 0F3f000000;
+	cvt_ftz_u32_f32	$s0, $s0;
+	st_global_align(4)_u32	$s0, [$d0];
+	ret;
+};

runtime/hsa-amd-aqlprofile/test/util/CMakeLists.txt

@@ -0,0 +1,15 @@
+#
+# Source files for Rocr Utils library
+#
+file( GLOB MODULE_SRC "*.cpp" )
+
+#
+# Header files include path(s).
+#
+include_directories ( $ENV{ROCR_INC_DIR} )
+
+#
+# Build Utils as a Static Library object
+#
+add_library( ${UTIL_LIB} STATIC ${MODULE_SRC} )
+target_link_libraries( ${UTIL_LIB} c stdc++ dl pthread rt )

runtime/hsa-amd-aqlprofile/test/util/helper_funcs.cpp

@@ -0,0 +1,230 @@
+/**********************************************************************
+Copyright ©2013 Advanced Micro Devices, Inc. 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.
+
+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.
+********************************************************************/
+
+#include <iostream>
+#include <sstream>
+#include <string>
+#include <cmath>
+#include <time.h>
+
+#include "helper_funcs.h"
+
+#ifndef _WIN32
+#include <unistd.h>
+#endif
+
+void error(std::string errorMsg) { std::cout << "Error: " << errorMsg << std::endl; }
+
+/*
+ * Prints no more than 256 elements of the given array.
+ * Prints full array if length is less than 256.
+ * Prints Array name followed by elements.
+ */
+template <typename T>
+void printArray(const std::string header, const T* data, const int width, const int height) {
+  std::cout << header << " :\n";
+  for (int i = 0; i < height; i++) {
+    std::cout << "> ";
+    for (int j = 0; j < width; j++) {
+      std::cout << data[i * width + j] << " ";
+    }
+    std::cout << "\n";
+  }
+}
+
+template <typename T>
+bool fillRandom(T* arrayPtr, const int width, const int height, const T rangeMin, const T rangeMax,
+                unsigned int seed) {
+  if (!arrayPtr) {
+    error("Cannot fill array. NULL pointer.");
+    return false;
+  }
+
+  if (!seed) seed = (unsigned int)time(NULL);
+
+  srand(seed);
+  double range = double(rangeMax - rangeMin) + 1.0;
+
+  /* random initialisation of input */
+  for (int i = 0; i < height; i++)
+    for (int j = 0; j < width; j++) {
+      int index = i * width + j;
+      arrayPtr[index] = rangeMin + T(range * rand() / (RAND_MAX + 1.0));
+    }
+
+  return true;
+}
+
+template <typename T> bool fillPos(T* arrayPtr, const int width, const int height) {
+  if (!arrayPtr) {
+    error("Cannot fill array. NULL pointer.");
+    return false;
+  }
+
+  /* initialisation of input with positions*/
+  for (T i = 0; i < height; i++)
+    for (T j = 0; j < width; j++) {
+      T index = i * width + j;
+      arrayPtr[index] = index;
+    }
+
+  return true;
+}
+
+template <typename T>
+bool fillConstant(T* arrayPtr, const int width, const int height, const T val) {
+  if (!arrayPtr) {
+    error("Cannot fill array. NULL pointer.");
+    return false;
+  }
+
+  /* initialisation of input with constant value*/
+  for (int i = 0; i < height; i++)
+    for (int j = 0; j < width; j++) {
+      int index = i * width + j;
+      arrayPtr[index] = val;
+    }
+
+  return true;
+}
+
+template <typename T> T roundToPowerOf2(T val) {
+  int bytes = sizeof(T);
+
+  val--;
+  for (int i = 0; i < bytes; i++) val |= val >> (1 << i);
+  val++;
+
+  return val;
+}
+
+template <typename T> bool isPowerOf2(T val) {
+  long long _val = val;
+  return (((_val & (-_val)) - _val == 0) && (_val != 0));
+}
+
+template <typename T> std::string toString(T t, std::ios_base& (*r)(std::ios_base&)) {
+  std::ostringstream output;
+  output << r << t;
+  return output.str();
+}
+
+bool compare(const float* refData, const float* data, const int length, const float epsilon) {
+  float error = 0.0f;
+  float ref = 0.0f;
+
+  for (int i = 1; i < length; ++i) {
+    float diff = refData[i] - data[i];
+    error += diff * diff;
+    ref += refData[i] * refData[i];
+  }
+
+  float normRef = ::sqrtf((float)ref);
+  if (::fabs((float)ref) < 1e-7f) {
+    return false;
+  }
+  float normError = ::sqrtf((float)error);
+  error = normError / normRef;
+
+  return error < epsilon;
+}
+
+bool compare(const double* refData, const double* data, const int length, const double epsilon) {
+  double error = 0.0;
+  double ref = 0.0;
+
+  for (int i = 1; i < length; ++i) {
+    double diff = refData[i] - data[i];
+    error += diff * diff;
+    ref += refData[i] * refData[i];
+  }
+
+  double normRef = ::sqrt((double)ref);
+  if (::fabs((double)ref) < 1e-7) {
+    return false;
+  }
+  double normError = ::sqrt((double)error);
+  error = normError / normRef;
+
+  return error < epsilon;
+}
+
+/////////////////////////////////////////////////////////////////
+// Template Instantiations
+/////////////////////////////////////////////////////////////////
+template void printArray<short>(const std::string, const short*, int, int);
+template void printArray<unsigned char>(const std::string, const unsigned char*, int, int);
+template void printArray<unsigned int>(const std::string, const unsigned int*, int, int);
+template void printArray<int>(const std::string, const int*, int, int);
+template void printArray<long>(const std::string, const long*, int, int);
+template void printArray<float>(const std::string, const float*, int, int);
+template void printArray<double>(const std::string, const double*, int, int);
+
+template bool fillRandom<unsigned char>(unsigned char* arrayPtr, const int width, const int height,
+                                        unsigned char rangeMin, unsigned char rangeMax,
+                                        unsigned int seed);
+template bool fillRandom<unsigned int>(unsigned int* arrayPtr, const int width, const int height,
+                                       unsigned int rangeMin, unsigned int rangeMax,
+                                       unsigned int seed);
+template bool fillRandom<int>(int* arrayPtr, const int width, const int height, int rangeMin,
+                              int rangeMax, unsigned int seed);
+template bool fillRandom<long>(long* arrayPtr, const int width, const int height, long rangeMin,
+                               long rangeMax, unsigned int seed);
+template bool fillRandom<float>(float* arrayPtr, const int width, const int height, float rangeMin,
+                                float rangeMax, unsigned int seed);
+template bool fillRandom<double>(double* arrayPtr, const int width, const int height,
+                                 double rangeMin, double rangeMax, unsigned int seed);
+
+template short roundToPowerOf2<short>(short val);
+template unsigned int roundToPowerOf2<unsigned int>(unsigned int val);
+template int roundToPowerOf2<int>(int val);
+template long roundToPowerOf2<long>(long val);
+
+template bool isPowerOf2<short>(short val);
+template bool isPowerOf2<unsigned int>(unsigned int val);
+template bool isPowerOf2<int>(int val);
+template bool isPowerOf2<long>(long val);
+
+template <> bool fillPos<short>(short* arrayPtr, const int width, const int height);
+template <> bool fillPos<unsigned int>(unsigned int* arrayPtr, const int width, const int height);
+template <> bool fillPos<int>(int* arrayPtr, const int width, const int height);
+template <> bool fillPos<long>(long* arrayPtr, const int width, const int height);
+
+template <>
+bool fillConstant<short>(short* arrayPtr, const int width, const int height, const short val);
+template <>
+bool fillConstant(unsigned int* arrayPtr, const int width, const int height,
+                  const unsigned int val);
+template <> bool fillConstant(int* arrayPtr, const int width, const int height, const int val);
+template <> bool fillConstant(long* arrayPtr, const int width, const int height, const long val);
+template <> bool fillConstant(long* arrayPtr, const int width, const int height, const long val);
+template <> bool fillConstant(long* arrayPtr, const int width, const int height, const long val);
+
+template std::string toString<char>(char t, std::ios_base& (*r)(std::ios_base&));
+template std::string toString<short>(short t, std::ios_base& (*r)(std::ios_base&));
+template std::string toString<unsigned int>(unsigned int t, std::ios_base& (*r)(std::ios_base&));
+template std::string toString<int>(int t, std::ios_base& (*r)(std::ios_base&));
+template std::string toString<long>(long t, std::ios_base& (*r)(std::ios_base&));
+template std::string toString<float>(float t, std::ios_base& (*r)(std::ios_base&));
+template std::string toString<double>(double t, std::ios_base& (*r)(std::ios_base&));

runtime/hsa-amd-aqlprofile/test/util/helper_funcs.h

@@ -0,0 +1,90 @@
+/**********************************************************************
+Copyright ©2013 Advanced Micro Devices, Inc. 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.
+
+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.
+********************************************************************/
+#ifndef _HELPER_FUNCS_H_
+#define _HELPER_FUNCS_H_
+
+#include <string>
+
+/**
+ * compare template version
+ * compare data to check error
+ * @param refData templated input
+ * @param data templated input
+ * @param length number of values to compare
+ * @param epsilon errorWindow
+ */
+bool compare(const float* refData, const float* data, const int length,
+             const float epsilon = 1e-6f);
+bool compare(const double* refData, const double* data, const int length,
+             const double epsilon = 1e-6);
+
+/**
+ * printArray
+ * displays a array on std::out
+ */
+template <typename T>
+void printArray(const std::string header, const T* data, const int width, const int height);
+
+
+/**
+ * fillRandom
+ * fill array with random values
+ */
+template <typename T>
+bool fillRandom(T* arrayPtr, const int width, const int height, const T rangeMin, const T rangeMax,
+                unsigned int seed = 123);
+
+/**
+ * fillPos
+ * fill the specified positions
+ */
+template <typename T> bool fillPos(T* arrayPtr, const int width, const int height);
+
+/**
+ * fillConstant
+ * fill the array with constant value
+ */
+template <typename T>
+bool fillConstant(T* arrayPtr, const int width, const int height, const T val);
+
+
+/**
+ * roundToPowerOf2
+ * rounds to a power of 2
+ */
+template <typename T> T roundToPowerOf2(T val);
+
+/**
+ * isPowerOf2
+ * checks if input is a power of 2
+ */
+template <typename T> bool isPowerOf2(T val);
+
+/**
+ * toString
+ * convert a T type to string
+ */
+template <typename T> std::string toString(T t, std::ios_base& (*r)(std::ios_base&));
+
+#endif  // _HELPER_FUNCS_H_

runtime/hsa-amd-aqlprofile/test/util/hsa_rsrc_factory.cpp

@@ -0,0 +1,473 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdint.h>
+#include <string.h>
+#include <cassert>
+
+#include <fstream>
+#include <iostream>
+#include <vector>
+#include <string>
+
+#include "hsa.h"
+#include "hsa_rsrc_factory.h"
+#include "hsa_ext_finalize.h"
+
+using namespace std;
+
+// Provide access to command line arguments passed in by user
+uint32_t hsa_cmdline_arg_cnt;
+char** hsa_cmdline_arg_list;
+
+// Callback function to find and bind kernarg region of an agent
+static hsa_status_t find_memregions(hsa_region_t region, void* data) {
+  hsa_region_global_flag_t flags;
+  hsa_region_segment_t segment_id;
+
+  hsa_region_get_info(region, HSA_REGION_INFO_SEGMENT, &segment_id);
+  if (segment_id != HSA_REGION_SEGMENT_GLOBAL) {
+    return HSA_STATUS_SUCCESS;
+  }
+
+  AgentInfo* agent_info = (AgentInfo*)data;
+  hsa_region_get_info(region, HSA_REGION_INFO_GLOBAL_FLAGS, &flags);
+  if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) {
+    agent_info->coarse_region = region;
+  }
+
+  if (flags & HSA_REGION_GLOBAL_FLAG_KERNARG) {
+    agent_info->kernarg_region = region;
+  }
+
+  return HSA_STATUS_SUCCESS;
+}
+
+// Callback function to get the number of agents
+static hsa_status_t get_hsa_agents(hsa_agent_t agent, void* data) {
+  // Copy handle of agent and increment number of agents reported
+  HsaRsrcFactory* rsrcFactory = reinterpret_cast<HsaRsrcFactory*>(data);
+
+  // Determine if device is a Gpu agent
+  hsa_status_t status;
+  hsa_device_type_t type;
+  status = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type);
+  if (type == HSA_DEVICE_TYPE_DSP) {
+    return HSA_STATUS_SUCCESS;
+  }
+
+  if (type == HSA_DEVICE_TYPE_CPU) {
+    AgentInfo* agent_info = reinterpret_cast<AgentInfo*>(malloc(sizeof(AgentInfo)));
+    agent_info->dev_id = agent;
+    agent_info->dev_type = HSA_DEVICE_TYPE_CPU;
+    rsrcFactory->AddAgentInfo(agent_info, false);
+    return HSA_STATUS_SUCCESS;
+  }
+
+  // Device is a Gpu agent, build an instance of AgentInfo
+  AgentInfo* agent_info = reinterpret_cast<AgentInfo*>(malloc(sizeof(AgentInfo)));
+  agent_info->dev_id = agent;
+  agent_info->dev_type = HSA_DEVICE_TYPE_GPU;
+  hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, agent_info->name);
+  agent_info->max_wave_size = 0;
+  hsa_agent_get_info(agent, HSA_AGENT_INFO_WAVEFRONT_SIZE, &agent_info->max_wave_size);
+  agent_info->max_queue_size = 0;
+  hsa_agent_get_info(agent, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &agent_info->max_queue_size);
+  agent_info->profile = hsa_profile_t(108);
+  hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &agent_info->profile);
+
+  // Initialize memory regions to zero
+  agent_info->kernarg_region.handle = 0;
+  agent_info->coarse_region.handle = 0;
+
+  // Find and Bind Memory regions of the Gpu agent
+  hsa_agent_iterate_regions(agent, find_memregions, agent_info);
+
+  // Save the instance of AgentInfo
+  rsrcFactory->AddAgentInfo(agent_info, true);
+  return HSA_STATUS_SUCCESS;
+}
+
+// Definitions for Static Data members of the class
+char* HsaRsrcFactory::brig_path_ = NULL;
+uint32_t HsaRsrcFactory::num_cus_ = 4;
+uint32_t HsaRsrcFactory::num_waves_;
+uint32_t HsaRsrcFactory::num_workitems_;
+uint32_t HsaRsrcFactory::kernel_loop_count_;
+bool HsaRsrcFactory::print_debug_info_ = false;
+
+char* HsaRsrcFactory::num_cus_key_ = "num_cus";
+char* HsaRsrcFactory::brig_path_key_ = "brig_path";
+char* HsaRsrcFactory::num_waves_key_ = "waves_per_cu";
+char* HsaRsrcFactory::num_workitems_key_ = "workitems_per_wave";
+char* HsaRsrcFactory::print_debug_key_ = "print_debug";
+char* HsaRsrcFactory::kernel_loop_count_key_ = "kernel_loop_count";
+
+// Constructor of the class
+HsaRsrcFactory::HsaRsrcFactory() {
+  // Initialize the Hsa Runtime
+  hsa_status_t status = hsa_init();
+  check("Error in hsa_init", status);
+
+  // Discover the set of Gpu devices available on the platform
+  status = hsa_iterate_agents(get_hsa_agents, this);
+  check("Error Calling hsa_iterate_agents", status);
+
+  // Process command line arguments
+  ProcessCmdline();
+}
+
+// Destructor of the class
+HsaRsrcFactory::~HsaRsrcFactory() {}
+
+// Get the count of Hsa Gpu Agents available on the platform
+//
+// @return uint32_t Number of Gpu agents on platform
+//
+uint32_t HsaRsrcFactory::GetCountOfGpuAgents() { return uint32_t(gpu_list_.size()); }
+
+// Get the count of Hsa Cpu Agents available on the platform
+//
+// @return uint32_t Number of Cpu agents on platform
+//
+uint32_t HsaRsrcFactory::GetCountOfCpuAgents() { return uint32_t(cpu_list_.size()); }
+
+// Get the AgentInfo handle of a Gpu device
+//
+// @param idx Gpu Agent at specified index
+//
+// @param agent_info Output parameter updated with AgentInfo
+//
+// @return bool true if successful, false otherwise
+//
+bool HsaRsrcFactory::GetGpuAgentInfo(uint32_t idx, AgentInfo** agent_info) {
+  // Determine if request is valid
+  uint32_t size = uint32_t(gpu_list_.size());
+  if (idx >= size) {
+    return false;
+  }
+
+  // Copy AgentInfo from specified index
+  *agent_info = gpu_list_[idx];
+  return true;
+}
+
+// Get the AgentInfo handle of a Cpu device
+//
+// @param idx Cpu Agent at specified index
+//
+// @param agent_info Output parameter updated with AgentInfo
+//
+// @return bool true if successful, false otherwise
+//
+bool HsaRsrcFactory::GetCpuAgentInfo(uint32_t idx, AgentInfo** agent_info) {
+  // Determine if request is valid
+  uint32_t size = uint32_t(cpu_list_.size());
+  if (idx >= size) {
+    return false;
+  }
+
+  // Copy AgentInfo from specified index
+  *agent_info = cpu_list_[idx];
+  return true;
+}
+
+// Create a Queue object and return its handle. The queue object is expected
+// to support user requested number of Aql dispatch packets.
+//
+// @param agent_info Gpu Agent on which to create a queue object
+//
+// @param num_Pkts Number of packets to be held by queue
+//
+// @param queue Output parameter updated with handle of queue object
+//
+// @return bool true if successful, false otherwise
+//
+bool HsaRsrcFactory::CreateQueue(AgentInfo* agent_info, uint32_t num_pkts, hsa_queue_t** queue) {
+  hsa_status_t status;
+  status = hsa_queue_create(agent_info->dev_id, num_pkts, HSA_QUEUE_TYPE_MULTI, NULL, NULL,
+                            UINT32_MAX, UINT32_MAX, queue);
+  return (status == HSA_STATUS_SUCCESS);
+}
+
+// Create a Signal object and return its handle.
+//
+// @param value Initial value of signal object
+//
+// @param signal Output parameter updated with handle of signal object
+//
+// @return bool true if successful, false otherwise
+//
+bool HsaRsrcFactory::CreateSignal(uint32_t value, hsa_signal_t* signal) {
+  hsa_status_t status;
+  status = hsa_signal_create(value, 0, NULL, signal);
+  return (status == HSA_STATUS_SUCCESS);
+}
+
+// Allocate memory for use by a kernel of specified size in specified
+// agent's memory region. Currently supports Global segment whose Kernarg
+// flag set.
+//
+// @param agent_info Agent from whose memory region to allocate
+//
+// @param size Size of memory in terms of bytes
+//
+// @return uint8_t* Pointer to buffer, null if allocation fails.
+//
+uint8_t* HsaRsrcFactory::AllocateLocalMemory(AgentInfo* agent_info, size_t size) {
+  hsa_status_t status;
+  uint8_t* buffer = NULL;
+
+  if (agent_info->coarse_region.handle != 0) {
+    // Allocate in local memory if it is available
+    status = hsa_memory_allocate(agent_info->coarse_region, size, (void**)&buffer);
+    if (status == HSA_STATUS_SUCCESS) {
+      status = hsa_memory_assign_agent(buffer, agent_info->dev_id, HSA_ACCESS_PERMISSION_RW);
+    }
+  } else {
+    // Allocate in system memory if local memory is not available
+    status = hsa_memory_allocate(agent_info->kernarg_region, size, (void**)&buffer);
+  }
+
+  return (status == HSA_STATUS_SUCCESS) ? buffer : NULL;
+}
+
+// Allocate memory tp pass kernel parameters.
+//
+// @param agent_info Agent from whose memory region to allocate
+//
+// @param size Size of memory in terms of bytes
+//
+// @return uint8_t* Pointer to buffer, null if allocation fails.
+//
+uint8_t* HsaRsrcFactory::AllocateSysMemory(AgentInfo* agent_info, size_t size) {
+  hsa_status_t status;
+  uint8_t* buffer = NULL;
+  status = hsa_memory_allocate(agent_info->kernarg_region, size, (void**)&buffer);
+  return (status == HSA_STATUS_SUCCESS) ? buffer : NULL;
+}
+
+bool HsaRsrcFactory::TransferData(uint8_t* dest_buff, uint8_t* src_buff, uint32_t length,
+                                  bool host_to_dev) {
+  hsa_status_t status;
+  status = hsa_memory_copy(dest_buff, src_buff, length);
+  return (status == HSA_STATUS_SUCCESS);
+}
+
+// Fake method for compilation steps only
+uint8_t* HsaRsrcFactory::AllocateMemory(AgentInfo* agent_info, size_t size) {
+  hsa_status_t status;
+  uint8_t* buffer = NULL;
+  status = hsa_memory_allocate(agent_info->kernarg_region, size, (void**)&buffer);
+  return (status == HSA_STATUS_SUCCESS) ? buffer : NULL;
+}
+
+// Loads an Assembled Brig file and Finalizes it into Device Isa
+//
+// @param agent_info Gpu device for which to finalize
+//
+// @param brig_path File path of the Assembled Brig file
+//
+// @param kernel_name Name of the kernel to finalize
+//
+// @param code_desc Handle of finalized Code Descriptor that could
+// be used to submit for execution
+//
+// @return bool true if successful, false otherwise
+//
+bool HsaRsrcFactory::LoadAndFinalize(AgentInfo* agent_info, const char* brig_path,
+                                     char* kernel_name, hsa_executable_symbol_t* code_desc) {
+  // Finalize the Hsail object into code object
+  hsa_status_t status;
+  hsa_code_object_t code_object;
+
+  // Build the code object filename
+  std::string filename(brig_path);
+  std::cout << "Code object filename: " << filename << std::endl;
+
+  // Open the file containing code object
+  std::ifstream codeStream(filename.c_str(), std::ios::binary | std::ios::ate);
+  if (!codeStream) {
+    std::cout << "Error: failed to load " << filename << std::endl;
+    assert(false);
+    return false;
+  }
+
+  // Allocate memory to read in code object from file
+  size_t size = std::string::size_type(codeStream.tellg());
+  char* codeBuff = (char*)AllocateSysMemory(agent_info, size);
+  if (!codeBuff) {
+    std::cout << "Error: failed to allocate memory for code object." << std::endl;
+    assert(false);
+    return false;
+  }
+
+  // Read the code object into allocated memory
+  codeStream.seekg(0, std::ios::beg);
+  std::copy(std::istreambuf_iterator<char>(codeStream), std::istreambuf_iterator<char>(), codeBuff);
+
+  // De-Serialize the code object that has been read into memory
+  status = hsa_code_object_deserialize(codeBuff, size, NULL, &code_object);
+  if (status != HSA_STATUS_SUCCESS) {
+    std::cout << "Failed to deserialize code object" << std::endl;
+    return false;
+  }
+
+  // Create executable.
+  hsa_executable_t hsaExecutable;
+  // status = hsa_executable_create(agent_info->profile,
+  status =
+      hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, "", &hsaExecutable);
+  check("Error in creating executable object", status);
+
+  // Load code object.
+  status = hsa_executable_load_code_object(hsaExecutable, agent_info->dev_id, code_object, "");
+  check("Error in loading executable object", status);
+
+  // Freeze executable.
+  status = hsa_executable_freeze(hsaExecutable, "");
+  check("Error in freezing executable object", status);
+
+  // Get symbol handle.
+  hsa_executable_symbol_t kernelSymbol;
+  status = hsa_executable_get_symbol(hsaExecutable, NULL, kernel_name, agent_info->dev_id, 0,
+                                     &kernelSymbol);
+  check("Error in looking up kernel symbol", status);
+
+  // Update output parameter
+  *code_desc = kernelSymbol;
+  return true;
+}
+
+// Add an instance of AgentInfo representing a Hsa Gpu agent
+void HsaRsrcFactory::AddAgentInfo(AgentInfo* agent_info, bool gpu) {
+  // Add input to Gpu list
+  if (gpu) {
+    gpu_list_.push_back(agent_info);
+    return;
+  }
+
+  // Add input to Cpu list
+  cpu_list_.push_back(agent_info);
+}
+
+// Print the various fields of Hsa Gpu Agents
+bool HsaRsrcFactory::PrintGpuAgents(const std::string& header) {
+  std::cout << header << " :" << std::endl;
+
+  AgentInfo* agent_info;
+  int size = uint32_t(gpu_list_.size());
+  for (int idx = 0; idx < size; idx++) {
+    agent_info = gpu_list_[idx];
+
+    std::cout << "> agent[" << idx << "] :" << std::endl;
+    std::cout << ">> Name : " << agent_info->name << std::endl;
+    std::cout << ">> Max Wave Size : " << agent_info->max_wave_size << std::endl;
+    std::cout << ">> Max Queue Size : " << agent_info->max_queue_size << std::endl;
+    std::cout << ">> Kernarg Region Id : " << agent_info->coarse_region.handle << std::endl;
+  }
+  return true;
+}
+
+// Returns the file path where brig files is located. Value is
+// available only after an instance has been built.
+char* HsaRsrcFactory::GetBrigPath() { return HsaRsrcFactory::brig_path_; }
+
+// Returns the number of compute units present on platform
+// Value is available only after an instance has been built.
+uint32_t HsaRsrcFactory::GetNumOfCUs() { return HsaRsrcFactory::num_cus_; }
+
+// Returns the maximum number of waves that can be launched
+// per compute unit. The actual number that can be launched
+// is affected by resource availability
+//
+// Value is available only after an instance has been built.
+uint32_t HsaRsrcFactory::GetNumOfWavesPerCU() { return HsaRsrcFactory::num_waves_; }
+
+// Returns the number of work-items that can execute per wave
+// Value is available only after an instance has been built.
+uint32_t HsaRsrcFactory::GetNumOfWorkItemsPerWave() { return HsaRsrcFactory::num_workitems_; }
+
+// Returns the number of times kernel loop body should execute.
+// Value is available only after an instance has been built.
+uint32_t HsaRsrcFactory::GetKernelLoopCount() { return HsaRsrcFactory::kernel_loop_count_; }
+
+// Returns boolean flag to indicate if debug info should be printed
+// Value is available only after an instance has been built.
+uint32_t HsaRsrcFactory::GetPrintDebugInfo() { return HsaRsrcFactory::print_debug_info_; }
+
+// Process command line arguments. The method will capture
+// various user command line parameters for tests to use
+void HsaRsrcFactory::ProcessCmdline() {
+  // Command line arguments are given
+  uint32_t idx;
+  uint32_t arg_idx;
+  for (idx = 1; idx < hsa_cmdline_arg_cnt; idx += 2) {
+    arg_idx = GetArgIndex((char*)hsa_cmdline_arg_list[idx]);
+    switch (arg_idx) {
+      case 0:
+        HsaRsrcFactory::brig_path_ = hsa_cmdline_arg_list[idx + 1];
+        break;
+      case 1:
+        HsaRsrcFactory::num_cus_ = atoi(hsa_cmdline_arg_list[idx + 1]);
+        break;
+      case 2:
+        HsaRsrcFactory::num_waves_ = atoi(hsa_cmdline_arg_list[idx + 1]);
+        break;
+      case 3:
+        HsaRsrcFactory::num_workitems_ = atoi(hsa_cmdline_arg_list[idx + 1]);
+        break;
+      case 4:
+        HsaRsrcFactory::kernel_loop_count_ = atoi(hsa_cmdline_arg_list[idx + 1]);
+        break;
+      case 5:
+        HsaRsrcFactory::print_debug_info_ = true;
+        break;
+    }
+  }
+}
+
+uint32_t HsaRsrcFactory::GetArgIndex(char* arg_value) {
+  // Map Brig file path to index zero
+  if (!strcmp(HsaRsrcFactory::brig_path_key_, arg_value)) {
+    return 0;
+  }
+
+  // Map Number of Compute Units to index one
+  if (!strcmp(HsaRsrcFactory::num_cus_key_, arg_value)) {
+    return 1;
+  }
+
+  // Map Number of Waves per CU to index two
+  if (!strcmp(HsaRsrcFactory::num_waves_key_, arg_value)) {
+    return 2;
+  }
+
+  // Map Number of Workitems per Wave to index three
+  if (!strcmp(HsaRsrcFactory::num_workitems_key_, arg_value)) {
+    return 3;
+  }
+
+  // Map Kernel Loop Count to index four
+  if (!strcmp(HsaRsrcFactory::kernel_loop_count_key_, arg_value)) {
+    return 4;
+  }
+
+  // Map print debug info parameter
+  if (!strcmp(HsaRsrcFactory::print_debug_key_, arg_value)) {
+    return 5;
+  }
+
+  return 108;
+}
+
+void HsaRsrcFactory::PrintHelpMsg() {
+  std::cout << "Key for passing Brig filepath: " << HsaRsrcFactory::brig_path_key_ << std::endl;
+  std::cout << "Key for passing Number of Compute Units: " << HsaRsrcFactory::num_cus_key_
+            << std::endl;
+  std::cout << "Key for passing Number of Waves per CU: " << HsaRsrcFactory::num_waves_key_
+            << std::endl;
+  std::cout << "Key for passing Number of Workitems per Wave: "
+            << HsaRsrcFactory::num_workitems_key_ << std::endl;
+  std::cout << "Key for passing Kernel Loop Count: " << HsaRsrcFactory::kernel_loop_count_key_
+            << std::endl;
+}

runtime/hsa-amd-aqlprofile/test/util/hsa_rsrc_factory.h

@@ -0,0 +1,262 @@
+#ifndef HSA_RSRC_FACTORY_H_
+#define HSA_RSRC_FACTORY_H_
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdint.h>
+#include <string.h>
+
+#include <iostream>
+#include <vector>
+#include <string>
+
+#include "perf_timer.h"
+#include "hsa.h"
+#include "hsa_ext_finalize.h"
+
+#define HSA_ARGUMENT_ALIGN_BYTES 16
+#define HSA_QUEUE_ALIGN_BYTES 64
+#define HSA_PACKET_ALIGN_BYTES 64
+
+#define check(msg, status)                                                                         \
+  if (status != HSA_STATUS_SUCCESS) {                                                              \
+    const char* emsg = 0;                                                                          \
+    hsa_status_string(status, &emsg);                                                              \
+    printf("%s: %s\n", msg, emsg ? emsg : "<unknown error>");                                      \
+    exit(1);                                                                                       \
+  }
+
+#define check_build(msg, status)                                                                   \
+  if (status != STATUS_SUCCESS) {                                                                  \
+    printf("%s\n", msg);                                                                           \
+    exit(1);                                                                                       \
+  }
+
+// Provide access to command line arguments passed in by user
+extern uint32_t hsa_cmdline_arg_cnt;
+extern char** hsa_cmdline_arg_list;
+
+// Encapsulates information about a Hsa Agent such as its
+// handle, name, max queue size, max wavefront size, etc.
+typedef struct {
+  // Handle of Agent
+  hsa_agent_t dev_id;
+
+  // Agent type - Cpu = 0, Gpu = 1 or Dsp = 2
+  uint32_t dev_type;
+
+  // Name of Agent whose length is less than 64
+  char name[64];
+
+  // Max size of Wavefront size
+  uint32_t max_wave_size;
+
+  // Max size of Queue buffer
+  uint32_t max_queue_size;
+
+  // Hsail profile supported by agent
+  hsa_profile_t profile;
+
+  // Memory region supporting kernel parameters
+  hsa_region_t coarse_region;
+
+  // Memory region supporting kernel arguments
+  hsa_region_t kernarg_region;
+
+} AgentInfo;
+
+class HsaRsrcFactory {
+ public:
+  // Constructor of the class. Will initialize the Hsa Runtime and
+  // query the system topology to get the list of Cpu and Gpu devices
+  HsaRsrcFactory();
+
+  // Destructor of the class
+  ~HsaRsrcFactory();
+
+  // Get the count of Hsa Gpu Agents available on the platform
+  //
+  // @return uint32_t Number of Gpu agents on platform
+  //
+  uint32_t GetCountOfGpuAgents();
+
+  // Get the count of Hsa Cpu Agents available on the platform
+  //
+  // @return uint32_t Number of Cpu agents on platform
+  //
+  uint32_t GetCountOfCpuAgents();
+
+  // Get the AgentInfo handle of a Gpu device
+  //
+  // @param idx Gpu Agent at specified index
+  //
+  // @param agent_info Output parameter updated with AgentInfo
+  //
+  // @return bool true if successful, false otherwise
+  //
+  bool GetGpuAgentInfo(uint32_t idx, AgentInfo** agent_info);
+
+  // Get the AgentInfo handle of a Cpu device
+  //
+  // @param idx Cpu Agent at specified index
+  //
+  // @param agent_info Output parameter updated with AgentInfo
+  //
+  // @return bool true if successful, false otherwise
+  //
+  bool GetCpuAgentInfo(uint32_t idx, AgentInfo** agent_info);
+
+  // Create a Queue object and return its handle. The queue object is expected
+  // to support user requested number of Aql dispatch packets.
+  //
+  // @param agent_info Gpu Agent on which to create a queue object
+  //
+  // @param num_Pkts Number of packets to be held by queue
+  //
+  // @param queue Output parameter updated with handle of queue object
+  //
+  // @return bool true if successful, false otherwise
+  //
+  bool CreateQueue(AgentInfo* agent_info, uint32_t num_pkts, hsa_queue_t** queue);
+
+  // Create a Signal object and return its handle.
+  //
+  // @param value Initial value of signal object
+  //
+  // @param signal Output parameter updated with handle of signal object
+  //
+  // @return bool true if successful, false otherwise
+  //
+  bool CreateSignal(uint32_t value, hsa_signal_t* signal);
+
+  // Allocate memory for use by a kernel of specified size in specified
+  // agent's memory region. Currently supports Global segment whose Kernarg
+  // flag set.
+  //
+  // @param agent_info Agent from whose memory region to allocate
+  //
+  // @param size Size of memory in terms of bytes
+  //
+  // @return uint8_t* Pointer to buffer, null if allocation fails.
+  //
+  uint8_t* AllocateLocalMemory(AgentInfo* agent_info, size_t size);
+  uint8_t* AllocateMemory(AgentInfo* agent_info, size_t size);
+
+  bool TransferData(uint8_t* dest_buff, uint8_t* src_buff, uint32_t length, bool host_to_dev);
+
+  // Allocate memory tp pass kernel parameters.
+  //
+  // @param agent_info Agent from whose memory region to allocate
+  //
+  // @param size Size of memory in terms of bytes
+  //
+  // @return uint8_t* Pointer to buffer, null if allocation fails.
+  //
+  uint8_t* AllocateSysMemory(AgentInfo* agent_info, size_t size);
+
+  // Loads an Assembled Brig file and Finalizes it into Device Isa
+  //
+  // @param agent_info Gpu device for which to finalize
+  //
+  // @param brig_path File path of the Assembled Brig file
+  //
+  // @param kernel_name Name of the kernel to finalize
+  //
+  // @param code_desc Handle of finalized Code Descriptor that could
+  // be used to submit for execution
+  //
+  // @return bool true if successful, false otherwise
+  //
+  bool LoadAndFinalize(AgentInfo* agent_info, const char* brig_path, char* kernel_name,
+                       hsa_executable_symbol_t* code_desc);
+
+  // Add an instance of AgentInfo representing a Hsa Gpu agent
+  void AddAgentInfo(AgentInfo* agent_info, bool gpu);
+
+  // Returns the file path where brig files is located
+  static char* GetBrigPath();
+
+  // Returns the number of compute units present on platform
+  static uint32_t GetNumOfCUs();
+
+  // Returns the maximum number of waves that can be launched
+  // per compute unit. The actual number that can be launched
+  // is affected by resource availability
+  static uint32_t GetNumOfWavesPerCU();
+
+  // Returns the number of work-items that can execute per wave
+  static uint32_t GetNumOfWorkItemsPerWave();
+
+  // Returns the number of times kernel loop body should execute.
+  static uint32_t GetKernelLoopCount();
+
+  // Returns boolean flag to indicate if debug info should be printed
+  static uint32_t GetPrintDebugInfo();
+
+  // Print the various fields of Hsa Gpu Agents
+  bool PrintGpuAgents(const std::string& header);
+
+ private:
+  // Number of queues to create
+  uint32_t num_queues_;
+
+  // Used to maintain a list of Hsa Queue handles
+  std::vector<hsa_queue_t*> queue_list_;
+
+  // Number of Signals to create
+  uint32_t num_signals_;
+
+  // Used to maintain a list of Hsa Signal handles
+  std::vector<hsa_signal_t*> signal_list_;
+
+  // Number of agents reported by platform
+  uint32_t num_agents_;
+
+  // Used to maintain a list of Hsa Gpu Agent Info
+  std::vector<AgentInfo*> gpu_list_;
+
+  // Used to maintain a list of Hsa Cpu Agent Info
+  std::vector<AgentInfo*> cpu_list_;
+
+  // Records the file path where Brig file is located.
+  // Value is available only after an instance has been built.
+  static char* brig_path_;
+  static char* brig_path_key_;
+
+  // Records the number of Compute units present on system.
+  // Value is available only after an instance has been built.
+  static uint32_t num_cus_;
+  static char* num_cus_key_;
+
+  // Records the number of waves that can be launched per Compute unit
+  // Value is available only after an instance has been built.
+  static uint32_t num_waves_;
+  static char* num_waves_key_;
+
+  // Records the number of work-items that can be packed into a wave
+  // Value is available only after an instance has been built.
+  static uint32_t num_workitems_;
+  static char* num_workitems_key_;
+
+  // Records the number of times kernel loop body should run. Value
+  // is available only after an instance has been built.
+  static uint32_t kernel_loop_count_;
+  static char* kernel_loop_count_key_;
+
+  // Records the number of times kernel loop body should run. Value
+  // is available only after an instance has been built.
+  static bool print_debug_info_;
+  static char* print_debug_key_;
+
+  // Process command line arguments. The method will capture
+  // various user command line parameters for tests to use
+  static void ProcessCmdline();
+
+  // Prints the help banner on user arg keys
+  static void PrintHelpMsg();
+
+  // Maps an index for the user argument
+  static uint32_t GetArgIndex(char* arg_value);
+};
+
+#endif  //  HSA_RSRC_FACTORY_H_

runtime/hsa-amd-aqlprofile/test/util/perf_timer.cpp

@@ -0,0 +1,157 @@
+#include "perf_timer.h"
+
+PerfTimer::PerfTimer() { freq_in_100mhz = MeasureTSCFreqHz(); }
+
+PerfTimer::~PerfTimer() {
+  while (!_timers.empty()) {
+    Timer* temp = _timers.back();
+    _timers.pop_back();
+    delete temp;
+  }
+}
+
+// a new cretaed timer instantance index will be returned
+int PerfTimer::CreateTimer() {
+  Timer* newTimer = new Timer;
+  newTimer->_start = 0;
+  newTimer->_clocks = 0;
+
+#ifdef _WIN32
+  QueryPerformanceFrequency((LARGE_INTEGER*)&newTimer->_freq);
+#else
+  newTimer->_freq = (long long)1.0E3;
+#endif
+
+  /* Push back the address of new Timer instance created */
+  _timers.push_back(newTimer);
+  return (int)(_timers.size() - 1);
+}
+
+int PerfTimer::StartTimer(int index) {
+  if (index >= (int)_timers.size()) {
+    Error("Cannot reset timer. Invalid handle.");
+    return FAILURE;
+  }
+
+#ifdef _WIN32
+// General Windows timing method
+#ifndef _AMD
+  long long tmpStart;
+  QueryPerformanceCounter((LARGE_INTEGER*)&(tmpStart));
+  _timers[index]->_start = (double)tmpStart;
+#else
+// AMD Windows timing method
+#endif
+#else
+// General Linux timing method
+#ifndef _AMD
+  struct timeval s;
+  gettimeofday(&s, 0);
+  _timers[index]->_start = s.tv_sec * 1.0E3 + ((double)(s.tv_usec / 1.0E3));
+#else
+  // AMD timing method
+  unsigned int unused;
+  _timers[index]->_start = __rdtscp(&unused);
+#endif
+#endif
+
+  return SUCCESS;
+}
+
+
+int PerfTimer::StopTimer(int index) {
+  double n = 0;
+  if (index >= (int)_timers.size()) {
+    Error("Cannot reset timer. Invalid handle.");
+    return FAILURE;
+  }
+#ifdef _WIN32
+#ifndef _AMD
+  long long n1;
+  QueryPerformanceCounter((LARGE_INTEGER*)&(n1));
+  n = (double)n1;
+#else
+// AMD Window Timing
+#endif
+
+#else
+// General Linux timing method
+#ifndef _AMD
+  struct timeval s;
+  gettimeofday(&s, 0);
+  n = s.tv_sec * 1.0E3 + (double)(s.tv_usec / 1.0E3);
+#else
+  // AMD Linux timing
+  unsigned int unused;
+  n = __rdtscp(&unused);
+#endif
+#endif
+
+  n -= _timers[index]->_start;
+  _timers[index]->_start = 0;
+
+#ifndef _AMD
+  _timers[index]->_clocks += n;
+#else
+  //_timers[index]->_clocks += 10 * n /freq_in_100mhz;      // unit is ns
+  _timers[index]->_clocks += 1.0E-6 * 10 * n / freq_in_100mhz;  // convert to ms
+#endif
+
+  return SUCCESS;
+}
+
+void PerfTimer::Error(string str) { cout << str << endl; }
+
+
+double PerfTimer::ReadTimer(int index) {
+  if (index >= (int)_timers.size()) {
+    Error("Cannot read timer. Invalid handle.");
+    return FAILURE;
+  }
+
+  double reading = double(_timers[index]->_clocks);
+
+  reading = double(reading / _timers[index]->_freq);
+
+  return reading;
+}
+
+
+uint64_t PerfTimer::CoarseTimestampUs() {
+#ifdef _WIN32
+  uint64_t freqHz, ticks;
+  QueryPerformanceFrequency((LARGE_INTEGER*)&freqHz);
+  QueryPerformanceCounter((LARGE_INTEGER*)&ticks);
+
+  // Scale numerator and divisor until (ticks * 1000000) fits in uint64_t.
+  while (ticks > (1ULL << 44)) {
+    ticks /= 16;
+    freqHz /= 16;
+  }
+
+  return (ticks * 1000000) / freqHz;
+#else
+  struct timespec ts;
+  clock_gettime(CLOCK_MONOTONIC_RAW, &ts);
+  return uint64_t(ts.tv_sec) * 1000000 + ts.tv_nsec / 1000;
+#endif
+}
+
+uint64_t PerfTimer::MeasureTSCFreqHz() {
+  // Make a coarse interval measurement of TSC ticks for 1 gigacycles.
+  unsigned int unused;
+  uint64_t tscTicksEnd;
+
+  uint64_t coarseBeginUs = CoarseTimestampUs();
+  uint64_t tscTicksBegin = __rdtscp(&unused);
+  do {
+    tscTicksEnd = __rdtscp(&unused);
+  } while (tscTicksEnd - tscTicksBegin < 1000000000);
+
+  uint64_t coarseEndUs = CoarseTimestampUs();
+
+  // Compute the TSC frequency and round to nearest 100MHz.
+  uint64_t coarseIntervalNs = (coarseEndUs - coarseBeginUs) * 1000;
+  uint64_t tscIntervalTicks = tscTicksEnd - tscTicksBegin;
+  return (tscIntervalTicks * 10 + (coarseIntervalNs / 2)) / coarseIntervalNs;
+}

runtime/hsa-amd-aqlprofile/test/util/perf_timer.h

@@ -0,0 +1,62 @@
+#ifndef _PERF_TIMER_H_
+#define _PERF_TIMER_H_
+
+// Will use AMD timer and general Linux timer based on users' need --> compilation flag
+// need to consider platform is Windows or Linux
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdint.h>
+#include <string.h>
+#include <iostream>
+#include <vector>
+#include <string>
+#if defined(_MSC_VER)
+#include <time.h>
+#include <windows.h>
+#include <intrin.h>
+#else
+#if defined(__GNUC__)
+#include <sys/time.h>
+#include <x86intrin.h>
+#endif  // __GNUC__
+#endif  //_MSC_VER
+
+using namespace std;
+
+class PerfTimer {
+ public:
+  enum { SUCCESS = 0, FAILURE = 1 };
+
+  PerfTimer();
+  ~PerfTimer();
+
+  // General Linux timing method
+  int CreateTimer();
+  int StartTimer(int index);
+  int StopTimer(int index);
+
+  // retrieve time
+  double ReadTimer(int index);
+  // write into a file
+  double WriteTimer(int index);
+
+ private:
+  struct Timer {
+    string name;     /* < name name of time object*/
+    long long _freq; /* < _freq frequency*/
+    double _clocks;  /* < _clocks number of ticks at end*/
+    double _start;   /* < _start start point ticks*/
+  };
+
+  std::vector<Timer*> _timers; /*< _timers vector to Timer objects */
+  double freq_in_100mhz;
+
+  // AMD timing method
+  uint64_t CoarseTimestampUs();
+  uint64_t MeasureTSCFreqHz();
+
+  void Error(string str);
+};
+
+#endif  // _PERF_TIMER_H_