diff --git a/README.md b/README.md
index 909dcda..71fa9ac 100755
--- a/README.md
+++ b/README.md
@@ -372,6 +372,31 @@
 
 <http://dist.gem5.org/dist/v20-1/test-progs/square/square.o>
 
+# Resource: HSA Agent Packet Example
+
+Based off of the Square resource in this repository, this resource serves as
+an example for using an HSA Agent Packet to send commands to the GPU command
+processor included in the GCN_X86 build of gem5.
+
+The example command extracts the kernel's completion signal from the domain
+of the command processor and the GPU's dispatcher. Initially this was a 
+workaround for the hipDeviceSynchronize bug, now fixed. The method of
+waiting on a signal can be applied to other agent packet commands though.
+
+Custom commands can be added to the command processor in gem5 to control
+the GPU in novel ways.
+
+## Compilation
+
+To compile:
+
+```
+cd src/hsa-agent-pkt
+docker run --rm -v ${PWD}:${PWD} -w ${PWD} -u $UID:$GID gcr.io/gem5-test/gcn-gpu make gfx8-apu
+```
+
+The compiled binary can be found in `src/hsa-agent-pkt/bin`
+
 # Resource: SPEC 2006
 
 The [Standard Performance Evaluation Corporation](
@@ -500,6 +525,10 @@
 `src/pthreads`.
 * **square**: Consult individual copyright notices of source files in
 `src/square`.
+* **hsa-agent-pkt**: `src/hsa-agent-pkt/square.cpp` is licensed under the
+same licence as 'src/square/square.cpp'.
+`src/hsa-agent-pkt/HSA_Interface.[h|.cpp]` are licensed under a BSD Lisense
+(A University of Maryland copyright).
 * **spec 2006**: SPEC CPU 2006 requires purchase of benchmark suite from
 [SPEC](https://www.spec.org/cpu2006/) thus, it cannot be freely distributed.
 Consult individual copyright notices of source files in `src/spec-2006`.
diff --git a/src/hsa-agent-pkt/HSA_Interface.cpp b/src/hsa-agent-pkt/HSA_Interface.cpp
new file mode 100644
index 0000000..e89cca4
--- /dev/null
+++ b/src/hsa-agent-pkt/HSA_Interface.cpp
@@ -0,0 +1,176 @@
+/*
+Copyright (c) 2020 University of Maryland
+All rights reserved.
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions are
+met: redistributions of source code must retain the above copyright
+notice, this list of conditions and the following disclaimer;
+redistributions in binary form must reproduce the above copyright
+notice, this list of conditions and the following disclaimer in the
+documentation and/or other materials provided with the distribution;
+neither the name of the copyright holders nor the names of its
+contributors may be used to endorse or promote products derived from
+this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+"AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
+LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
+A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
+OWNER 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 "HSA_Interface.h"
+
+void print_agent_dispatch_packet(hsa_agent_dispatch_packet_t* pkt)
+{
+
+    printf("Packet \t\t%p\n",
+        (void *)pkt);
+    printf("Packet 16t\t\t%p\n",
+        (uint16_t *)pkt);
+    printf("Packet 32t\t\t%p\n",
+        (uint32_t *)pkt);
+    printf("Packet void**\t\t%p\n",
+        (void **)pkt);
+    printf("%p header: \t\t%hu\n",
+        (void *)(&(pkt->header )),pkt->header );
+    printf("%p type: \t\t%hu\n",
+        (void *)(&(pkt->type )),pkt->type );
+    printf("%p reserved0: \t\t%u\n",
+        (void *)(&(pkt->reserved0 )),pkt->reserved0 );
+    printf("%p return_address: \t\t%p\n",
+        (void *)(&(pkt->return_address )),pkt->return_address );
+    printf("%p arg[0]: \t\t%lu\n",
+        (void *)(&(pkt->arg[0] )),pkt->arg[0] );
+    printf("%p arg[1]: \t\t%lu\n",
+        (void *)(&(pkt->arg[1] )),pkt->arg[1] );
+    printf("%p arg[2]: \t\t%lu\n",
+        (void *)(&(pkt->arg[2] )),pkt->arg[0] );
+    printf("%p arg[3]: \t\t%lu\n",
+        (void *)(&(pkt->arg[3] )),pkt->arg[1] );
+    printf("%p reserved2: \t\t%lu\n",
+        (void *)(&(pkt->reserved2 )),pkt->reserved2 );
+    printf("%p completion_signal: \t\t%lu\n",
+        (void *)(&(pkt->completion_signal )),pkt->completion_signal.handle );
+
+
+    fflush(stdout);
+}
+
+void agent_disp_packet_store_release(uint16_t* packet, uint16_t header) {
+    __atomic_store_n(packet, header, __ATOMIC_RELEASE);
+}
+
+uint16_t header(hsa_packet_type_t type) {
+    uint16_t header = type << HSA_PACKET_HEADER_TYPE;
+    header |=
+        HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE;
+    header |=
+        HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE;
+    return header;
+}
+
+hsa_status_t get_kernel_agent(hsa_agent_t agent, void* data) {
+    uint32_t features = 0;
+    hsa_agent_get_info(agent, HSA_AGENT_INFO_FEATURE, &features);
+    if (features & HSA_AGENT_FEATURE_KERNEL_DISPATCH) {
+        // Store kernel agent in the application-provided buffer and return
+        hsa_agent_t* ret = (hsa_agent_t*) data;
+        *ret = agent;
+        return HSA_STATUS_INFO_BREAK;
+    }
+    // Keep iterating
+    return HSA_STATUS_SUCCESS;
+}
+
+void signal_wait(hsa_signal_t signal)
+{
+    while (hsa_signal_wait_relaxed(signal, HSA_SIGNAL_CONDITION_EQ, 0,
+        UINT64_MAX, HSA_WAIT_STATE_ACTIVE) != 0);
+    // while (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_EQ, 0,
+    //  UINT64_MAX, HSA_WAIT_STATE_ACTIVE) != 0);
+}
+
+void initialize_agent_dispatch_packet(
+    hsa_agent_dispatch_packet_t* packet,
+    size_t header_size
+    )
+{
+    // Reserved fields, private and group memory,
+    // and completion signal are all set to 0.
+    memset(((uint8_t*) packet) + header_size, 0,
+        sizeof(hsa_agent_dispatch_packet_t) - header_size);
+}
+
+HSA_Interface::HSA_Interface(){
+
+    printf("INFO:: Setting up HSA Interface:\n");
+
+    CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/));
+    printf ("info: running on device %s\n", props.name); fflush(stdout);
+    #ifdef __HIP_PLATFORM_HCC__
+      printf ("info: architecture on AMD GPU device is: %d\n",
+        props.gcnArch); fflush(stdout);
+    #endif
+
+    printf ("INFO:: hsa_iterate_agents\n"); fflush(stdout);
+    hsa_agent_t kernel_agent;
+    hsa_iterate_agents(get_kernel_agent, &kernel_agent);
+    printf ("INFO:: hsa_queue_create\n"); fflush(stdout);
+    hsa_queue_create(kernel_agent, 4, HSA_QUEUE_TYPE_SINGLE,
+        NULL, NULL, 0, 0, &queue);
+    printf ("INFO:: hsa_queue_add_write_index_relaxed\n"); fflush(stdout);
+    hsa_queue_add_write_index_relaxed(queue, 1);
+
+    packet_id = 0;
+
+    printf("INFO:: Creating Stream\n");fflush(stdout);
+    stream = 0;
+    hipStreamCreate(&stream);
+}
+
+void HSA_Interface::steal_kernel_signal(uint32_t kid)
+{
+    hsa_agent_dispatch_packet_t * packet =
+        (hsa_agent_dispatch_packet_t*) queue->base_address + packet_id;
+    // Populate fields in kernel dispatch packet, except for the header,
+    // the setup, and the completion signal fields
+    initialize_agent_dispatch_packet(packet,sizeof(uint16_t));
+
+    uint64_t kernel_completion_signal_addr;
+    packet->type = AGENT_DISPATCH_PACKET_STEAL_KERNEL_SIGNAL;
+    packet->return_address = &kernel_completion_signal_addr;
+    packet->arg[0] = kid; //This field is for the kernel id.
+
+    //Create thief packet wait signal
+    hsa_signal_create(1, 0, NULL, &packet->completion_signal);
+
+    agent_disp_packet_store_release((uint16_t*) packet,
+        header(HSA_PACKET_TYPE_AGENT_DISPATCH));
+
+    print_agent_dispatch_packet(packet);
+
+    //Send thief packet
+    hsa_signal_store_screlease(queue->doorbell_signal, packet_id);
+
+    signal_wait(packet->completion_signal);
+    printf("INFO:: Done Waiting on Thief Signal\n");
+
+    hsa_signal_t * new_signal = new hsa_signal_t;
+    new_signal->handle = kernel_completion_signal_addr;
+    m_kernel_signals.push_back(new_signal);
+
+    packet_id++;
+}
+
+void HSA_Interface::wait_kernel(uint32_t kid)
+{
+    signal_wait(*(m_kernel_signals[kid]));
+}
\ No newline at end of file
diff --git a/src/hsa-agent-pkt/HSA_Interface.h b/src/hsa-agent-pkt/HSA_Interface.h
new file mode 100644
index 0000000..b0aab72
--- /dev/null
+++ b/src/hsa-agent-pkt/HSA_Interface.h
@@ -0,0 +1,83 @@
+/*
+Copyright (c) 2020 University of Maryland
+All rights reserved.
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions are
+met: redistributions of source code must retain the above copyright
+notice, this list of conditions and the following disclaimer;
+redistributions in binary form must reproduce the above copyright
+notice, this list of conditions and the following disclaimer in the
+documentation and/or other materials provided with the distribution;
+neither the name of the copyright holders nor the names of its
+contributors may be used to endorse or promote products derived from
+this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+"AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
+LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
+A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
+OWNER 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 "hip/hip_runtime.h"
+
+#include <include/hsa/hsa.h>
+
+#include <vector>
+
+#define AGENT_DISPATCH_PACKET_NOP 0
+#define AGENT_DISPATCH_PACKET_STEAL_KERNEL_SIGNAL 1
+
+#define CHECK(cmd) \
+{\
+    hipError_t error  = cmd;\
+    if (error != hipSuccess) {\
+      fprintf(stderr, "error: '%s'(%d) at %s:%d\n",\
+      hipGetErrorString(error), error,__FILE__, __LINE__);\
+    exit(EXIT_FAILURE);\
+    }\
+}
+
+void print_agent_dispatch_packet(hsa_agent_dispatch_packet_t* pkt);
+void agent_disp_packet_store_release(uint16_t* packet, uint16_t header);
+uint16_t header(hsa_packet_type_t type);
+hsa_status_t get_kernel_agent(hsa_agent_t agent, void* data);
+void signal_wait( hsa_signal_t signal);
+void initialize_agent_dispatch_packet(
+    hsa_agent_dispatch_packet_t* packet,
+    size_t header_size
+    );
+
+//Class for interacting with kernel agent and creating pipes
+class HSA_Interface {
+
+public:
+    HSA_Interface();
+    ~HSA_Interface(){};
+
+    void steal_kernel_signal(uint32_t kid);
+    void wait_kernel(uint32_t kid);
+
+    hipStream_t getStream() {return stream;}
+
+private:
+    hsa_queue_t * queue;
+    hipStream_t stream;
+    hipDeviceProp_t props;
+
+    //Store Kernel Signals for multuple launches
+    std::vector<hsa_signal_t *> m_kernel_signals;
+
+    //Each packet created will have an ID associated with it.
+    //It is used to index into the hsa queue.
+    uint64_t packet_id;
+
+};
+
diff --git a/src/hsa-agent-pkt/Makefile b/src/hsa-agent-pkt/Makefile
new file mode 100644
index 0000000..a5fec13
--- /dev/null
+++ b/src/hsa-agent-pkt/Makefile
@@ -0,0 +1,18 @@
+HIP_PATH?= /opt/rocm/hip
+HSA_PATH?= /opt/rocm/hsa
+HIPCC=$(HIP_PATH)/bin/hipcc
+
+BIN_DIR?= ./bin
+
+gfx8-apu: $(BIN_DIR)/square.o
+
+$(BIN_DIR)/square.o: square.cpp HSA_Interface.cpp $(BIN_DIR)
+	$(HIPCC) --amdgpu-target=gfx801 $(CXXFLAGS) -I$(HSA_PATH) square.cpp HSA_Interface.cpp -o $(BIN_DIR)/square.o
+
+$(BIN_DIR):
+	mkdir -p $(BIN_DIR)
+
+clean:
+	rm -rf $(BIN_DIR)
+
+.PHONY: gfx8-apu clean
diff --git a/src/hsa-agent-pkt/square.cpp b/src/hsa-agent-pkt/square.cpp
new file mode 100644
index 0000000..86653b1
--- /dev/null
+++ b/src/hsa-agent-pkt/square.cpp
@@ -0,0 +1,100 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in 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:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+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
+AUTHORS 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 IN
+THE SOFTWARE.
+*/
+
+#include <stdio.h>
+#include "hip/hip_runtime.h"
+#include "HSA_Interface.h"
+
+/*
+ * Square each element in the array A and write to array C.
+ */
+template <typename T>
+__global__ void
+vector_square(hipLaunchParm lp, T *C_d, const T *A_d, size_t N)
+{
+    size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
+    size_t stride = hipBlockDim_x * hipGridDim_x ;
+
+    for (size_t i=offset; i<N; i+=stride) {
+        C_d[i] = A_d[i] * A_d[i];
+    }
+}
+
+
+int main(int argc, char *argv[])
+{
+    float *A_h, *C_h;
+    size_t N = 1000000;
+
+    if (argc == 2)
+        N = atoi(argv[1]);
+
+    const unsigned threadsPerBlock = 256;
+    unsigned blocks = (N + threadsPerBlock - 1) / threadsPerBlock; 
+
+    size_t Nbytes = N * sizeof(float);
+
+    HSA_Interface * hsa= new HSA_Interface();
+
+    printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
+    A_h = (float*)malloc(Nbytes);
+    CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
+    C_h = (float*)malloc(Nbytes);
+    CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
+
+    // Fill with Phi + i
+    for (size_t i=0; i<N; i++)
+    {
+        A_h[i] = 1.618f + i;
+    }
+
+    int kernel_id = 0;
+    printf ("info: launch 'vector_square' kernel: "
+            "N = %lu | Blocks = %u | kernel_id %d\n",
+            N, blocks, kernel_id);
+    hipLaunchKernel(vector_square, dim3(blocks),
+                    dim3(threadsPerBlock), 0, hsa->getStream(),
+                    C_h, A_h, N);
+
+
+    //Kernel_id must match that of the launched kernel (ie launch order)
+    printf("info: Stealing kernel completion signal (kid: %d)\n",
+            kernel_id);
+    hsa->steal_kernel_signal(kernel_id);
+
+    //Theoretically Equivalent to hipDeviceSynchronize();
+    printf("info: Waiting on kernel completion signal (kid: %d)\n",
+            kernel_id);
+    hsa->wait_kernel(kernel_id);
+
+    //Increment the Kernel_id every time any kernel is launched.
+    kernel_id++;
+
+    printf ("info: check result\n");
+    for (size_t i=0; i<N; i++)  {
+        if (C_h[i] != A_h[i] * A_h[i]) {
+            CHECK(hipErrorUnknown);
+        }
+    }
+    printf ("PASSED!\n");
+	return 0;
+}