Merge 'master' into 'amd-master'

Change-Id: I99fad287ff52fb4c34228ef5107118b6a038e3f9


[ROCm/hip commit: 7f9d78f7a1]

projects/hip/INSTALL.md

@@ -23,7 +23,11 @@ HIP code can be developed either on AMD ROCm platform using hcc or clang compile
 
 ## AMD-hcc
 
-* Install the [rocm](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/) packages.  ROCm will install all of the necessary components, including the kernel driver, runtime software, HCC compiler, and HIP.
+* Add the ROCm package server to your system as per the OS-specific guide available [here](https://rocm.github.io/ROCmInstall.html#installing-from-amd-rocm-repositories).
+* Install the "hip_hcc" package. This will install HCC and the HIP porting layer.
+```
+apt-get install hip_hcc
+```
 
 * Default paths and environment variables:
 
@@ -67,7 +71,7 @@ HIP code can be developed either on AMD ROCm platform using hcc or clang compile
    * Optionally, set HIPCC_VERBOSE=7 to output the command line for compilation to make sure clang is used instead of hcc.
 
 ## NVIDIA-nvcc
-* Configure the additional package server as described [here](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/).  
+* Add the ROCm package server to your system as per the OS-specific guide available [here](https://rocm.github.io/ROCmInstall.html#installing-from-amd-rocm-repositories).
 * Install the "hip_nvcc" package.  This will install CUDA SDK and the HIP porting layer.
 ```
 apt-get install hip_nvcc
@@ -122,20 +126,3 @@ make install
 
 * After installation, make sure HIP_PATH is pointed to `/where/to/install/hip`. 
 
-## HCC Options
-
-### Using HIP with the AMD Native-GCN compiler.
-AMD recently released a direct-to-GCN-ISA target.  This compiler generates GCN ISA directly from LLVM, without going through an intermediate compiler 
-IR such as HSAIL or PTX.
-The native GCN target is included with upstream LLVM, and has also been integrated with HCC compiler and can be used to compiler HIP programs for AMD.
-Binary packages for the direct-to-isa package are included with the [rocm](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/) package. 
-Alternatively, this sections describes how to build it from source: 
-
-1. Install the ROCm packages as described above.
-2. Follow the instructions [here](https://github.com/RadeonOpenCompute/HCC-Native-GCN-ISA/wiki)
-   * In the make step for HCC, we recommend setting -DCMAKE_INSTALL_PREFIX.  
-   * Set HCC_HOME environment variable before compiling HIP program to point to the native compiler:
-```shell
-export HCC_HOME=/path/to/native/hcc
-```
-

projects/hip/hipify-clang/CMakeLists.txt

@@ -1,4 +1,16 @@
-cmake_minimum_required(VERSION 3.12.3)
+if (CUDA_VERSION VERSION_GREATER "9.2")
+    cmake_minimum_required(VERSION 3.12.3)
+else()
+    if (CUDA_VERSION VERSION_LESS "9.0")
+        cmake_minimum_required(VERSION 3.5.0)
+    else()
+        cmake_minimum_required(VERSION 3.6.0)
+    endif()
+    if (MSVC AND MSVC_VERSION VERSION_GREATER "1900")
+        cmake_minimum_required(VERSION 3.7.2)
+    endif()
+endif()
+
 project(hipify-clang)
 
 if (MSVC AND MSVC_VERSION VERSION_LESS "1900")

projects/hip/hipify-clang/src/LLVMCompat.cpp

@@ -21,6 +21,7 @@ THE SOFTWARE.
 */
 
 #include "LLVMCompat.h"
+#include "llvm/Support/Path.h"
 
 namespace llcompat {
 
@@ -97,7 +98,13 @@ clang::SourceLocation getEndLoc(const clang::TypeLoc& typeLoc) {
 std::error_code real_path(const Twine &path, SmallVectorImpl<char> &output,
                           bool expand_tilde) {
 #if LLVM_VERSION_MAJOR < 5
-  return sys::fs::make_absolute(path, output);
+  output.clear();
+  std::string s = path.str();
+  output.append(s.begin(), s.end());
+  if (sys::path::is_relative(path)) {
+    return sys::fs::make_absolute(output);
+  }
+  return std::error_code();
 #else
   return sys::fs::real_path(path, output, expand_tilde);
 #endif

projects/hip/hipify-clang/src/main.cpp

@@ -138,20 +138,20 @@ int main(int argc, const char **argv) {
     // output (which may mean overwriting the input, if we're in-place).
     // Should we fail for some reason, we'll just leak this file and not corrupt the input.
     EC = llcompat::real_path(src, sourceAbsPath, true);
+    sourceFileName = sys::path::filename(sourceAbsPath);
     if (EC) {
       llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << src << "\n";
       Result = 1;
       continue;
     }
     if (TemporaryDir.empty()) {
-      EC = sys::fs::createTemporaryFile(src, ext, tmpFile);
+      EC = sys::fs::createTemporaryFile(sourceFileName, ext, tmpFile);
       if (EC) {
         llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << tmpFile << "\n";
         Result = 1;
         continue;
       }
     } else {
-      sourceFileName = sys::path::filename(sourceAbsPath);
       sTmpFileName = sTmpDirAbsParh + "/" + sourceFileName.str() + "." + ext.str();
       tmpFile = sTmpFileName;
     }

projects/hip/src/hip_hcc_internal.h

@@ -374,6 +374,7 @@ struct ihipModule_t {
     std::string fileName;
     hsa_executable_t executable = {};
     hsa_code_object_reader_t coReader = {};
+    std::string hash;
 
     ~ihipModule_t() {
         if (executable.handle) hsa_executable_destroy(executable);

projects/hip/src/hip_module.cpp

@@ -55,6 +55,21 @@ using namespace ELFIO;
 using namespace hip_impl;
 using namespace std;
 
+// calculate MD5 checksum
+inline std::string checksum(size_t size, const char *source) {
+    // FNV-1a hashing, 64-bit version
+    const uint64_t FNV_prime = 0x100000001b3;
+    const uint64_t FNV_basis = 0xcbf29ce484222325;
+    uint64_t hash = FNV_basis;
+
+    const char *str = static_cast<const char *>(source);
+    for (auto i = 0; i < size; ++i) {
+        hash ^= *str++;
+        hash *= FNV_prime;
+    }
+    return std::to_string(hash);
+}
+
 inline uint64_t alignTo(uint64_t Value, uint64_t Align, uint64_t Skew = 0) {
     assert(Align != 0u && "Align can't be 0.");
     Skew %= Align;
@@ -332,21 +347,21 @@ pair<hipDeviceptr_t, size_t> read_global_description(ForwardIterator f, ForwardI
 
 hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod,
                                          const char* name) {
-    static unordered_map<hipModule_t, vector<Agent_global>> agent_globals;
+    // the key of the map would the hash of code object associated with the
+    // hipModule_t instance
+    static unordered_map<std::string, vector<Agent_global>> agent_globals;
+    auto key = hmod->hash;
 
-    // TODO: this is not particularly robust.
-    if (agent_globals.count(hmod) == 0) {
+    if (agent_globals.count(key) == 0) {
         static mutex mtx;
         lock_guard<mutex> lck{mtx};
 
-        if (agent_globals.count(hmod) == 0) {
-            agent_globals.emplace(hmod, read_agent_globals(this_agent(), hmod->executable));
+        if (agent_globals.count(key) == 0) {
+            agent_globals.emplace(key, read_agent_globals(this_agent(), hmod->executable));
         }
     }
 
-    // TODO: This is unsafe iff some other emplacement triggers rehashing.
-    //       It will have to be properly fleshed out in the future.
-    const auto it0 = agent_globals.find(hmod);
+    const auto it0 = agent_globals.find(key);
     if (it0 == agent_globals.cend()) {
         throw runtime_error{"agent_globals data structure corrupted."};
     }
@@ -568,8 +583,14 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) {
 
     auto tmp = code_object_blob_for_agent(image, this_agent());
 
-    (*module)->executable = hip_impl::load_executable(
-        tmp.empty() ? read_elf_file_as_string(image) : tmp, (*module)->executable, this_agent());
+    auto content = tmp.empty() ? read_elf_file_as_string(image) : tmp;
+
+    (*module)->executable = hip_impl::load_executable(content,
+                                                      (*module)->executable,
+                                                      this_agent());
+
+    // compute the hash of the code object
+    (*module)->hash = checksum(content.length(), content.data());
 
     return (*module)->executable.handle ? hipSuccess : hipErrorUnknown;
 }

projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp

@@ -0,0 +1,713 @@
+/*
+Copyright (c) 2018 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.
+*/
+
+/*
+ * HIT_START
+ * BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM nvcc
+ * RUN: %t
+ * HIT_END
+ */
+
+// Includes HIP Runtime
+#include "hip/hip_runtime.h"
+#include <test_common.h>
+
+#include <set>
+#include <unordered_map>
+
+#define test_failed(test_name)                                                 \
+    printf("%s %s  FAILED!%s\n", KRED, test_name, KNRM);
+#define test_passed(test_name)                                                 \
+    printf("%s %s  PASSED!%s\n", KGRN, test_name, KNRM);
+
+/* BEGIN DUPLICATION */
+
+// TODO: The host code below is duplicated, and should ideally reside
+// in a common host-side library.
+
+typedef enum {
+    __OCKL_AS_PACKET_EMPTY = 0,
+    __OCKL_AS_PACKET_READY = 1
+} __ockl_as_packet_type_t;
+
+#define __OCKL_AS_PAYLOAD_ALIGNMENT 4
+#define __OCKL_AS_PAYLOAD_BYTES 48
+
+typedef enum {
+    __OCKL_AS_PACKET_HEADER_TYPE = 0, // corresponds to HSA_PACKET_HEADER_TYPE
+    __OCKL_AS_PACKET_HEADER_RESERVED0 = 8,
+    __OCKL_AS_PACKET_HEADER_FLAGS = 13,
+    __OCKL_AS_PACKET_HEADER_BYTES = 16,
+    __OCKL_AS_PACKET_HEADER_SERVICE = 24,
+} __ockl_as_packet_header_t;
+
+typedef enum {
+    __OCKL_AS_PACKET_HEADER_WIDTH_TYPE = 8,
+    __OCKL_AS_PACKET_HEADER_WIDTH_RESERVED0 = 5,
+    __OCKL_AS_PACKET_HEADER_WIDTH_FLAGS = 3,
+    __OCKL_AS_PACKET_HEADER_WIDTH_BYTES = 8,
+    __OCKL_AS_PACKET_HEADER_WIDTH_SERVICE = 8
+} __ockl_as_packet_header_width_t;
+
+// A packet is 64 bytes long, and the payload starts at index 16.
+struct __ockl_as_packet_t {
+    uint header;
+    uint reserved1;
+    ulong connection_id;
+
+    uchar payload[__OCKL_AS_PAYLOAD_BYTES];
+};
+
+typedef enum {
+    __OCKL_AS_STATUS_SUCCESS,
+    __OCKL_AS_STATUS_INVALID_REQUEST,
+    __OCKL_AS_STATUS_OUT_OF_RESOURCES,
+    __OCKL_AS_STATUS_BUSY,
+    __OCKL_AS_STATUS_UNKNOWN_ERROR
+} __ockl_as_status_t;
+
+typedef enum {
+    __OCKL_AS_CONNECTION_BEGIN = 1,
+    __OCKL_AS_CONNECTION_END = 2,
+} __ockl_as_flag_t;
+
+typedef enum { __OCKL_AS_FEATURE_ASYNCHRONOUS = 1 } __ockl_as_feature_t;
+
+typedef struct {
+    // Opaque handle. The value 0 is reserved.
+    ulong handle;
+} __ockl_as_signal_t;
+
+typedef struct __ockl_as_packet_t __ockl_as_packet_t;
+
+typedef struct {
+    ulong read_index;
+    ulong write_index;
+    __ockl_as_signal_t doorbell_signal;
+    __ockl_as_packet_t *base_address;
+    ulong size;
+} __ockl_as_stream_t;
+
+#define ATTR_GLOBAL __attribute__((address_space(1)))
+
+extern "C" __device__ __ockl_as_status_t
+__ockl_as_write_block(__ockl_as_stream_t ATTR_GLOBAL *stream, uchar service_id,
+                      ulong *connection_id, const uchar *str, uint32_t len,
+                      uchar flags);
+
+__device__ __ockl_as_status_t
+__hip_as_write_block(__ockl_as_stream_t *stream, uchar service_id,
+                     ulong *connection_id, const uchar *str, uint32_t len,
+                     uchar flags)
+{
+    __ockl_as_stream_t ATTR_GLOBAL *gstream =
+        reinterpret_cast<__ockl_as_stream_t ATTR_GLOBAL *>(stream);
+    return __ockl_as_write_block(gstream, service_id, connection_id, str, len,
+                                 flags);
+}
+
+/* END DUPLICATION */
+
+static __ockl_as_stream_t *
+createStream(void *ptr, uint buffer_size, uint num_packets)
+{
+    memset(ptr, 0, buffer_size);
+    __ockl_as_stream_t *r = (__ockl_as_stream_t *)ptr;
+    r->base_address = (__ockl_as_packet_t *)(&r[1]);
+    r->doorbell_signal = {0};
+    r->size = num_packets;
+
+    return r;
+}
+
+static uint8_t
+get_header_field(uint32_t header, uint8_t offset, uint8_t size)
+{
+    return (header >> offset) & ((1 << size) - 1);
+}
+
+static uint8_t
+get_packet_type(uint32_t header)
+{
+    return get_header_field(header, __OCKL_AS_PACKET_HEADER_TYPE,
+                            __OCKL_AS_PACKET_HEADER_WIDTH_TYPE);
+}
+
+static uint8_t
+get_packet_flags(uint32_t header)
+{
+    return get_header_field(header, __OCKL_AS_PACKET_HEADER_FLAGS,
+                            __OCKL_AS_PACKET_HEADER_WIDTH_FLAGS);
+}
+
+static uint8_t
+get_packet_bytes(uint32_t header)
+{
+    return get_header_field(header, __OCKL_AS_PACKET_HEADER_BYTES,
+                            __OCKL_AS_PACKET_HEADER_WIDTH_BYTES);
+}
+
+static uint8_t
+get_packet_service(uint32_t header)
+{
+    return get_header_field(header, __OCKL_AS_PACKET_HEADER_SERVICE,
+                            __OCKL_AS_PACKET_HEADER_WIDTH_SERVICE);
+}
+
+const unsigned int __OCKL_AS_PACKET_SIZE = sizeof(__ockl_as_packet_t);
+
+/* END DUPLICATION */
+
+using namespace std;
+
+#define STR_HELLO_WORLD "hello world"
+#define STRLEN_HELLO_WORLD 11
+
+const unsigned int THREADS_PER_BLOCK = 123; // include a partial warp
+const unsigned int NUM_BLOCKS = 3; // because powers of two are too convenient
+const unsigned int NUM_THREADS = NUM_BLOCKS * THREADS_PER_BLOCK;
+const unsigned int NUM_PACKETS_INSUFFICIENT = NUM_THREADS - 23;
+const unsigned int NUM_PACKETS_LARGE = NUM_THREADS * 4;
+const unsigned int NUM_SERVICES = 7;
+const unsigned int TEST_SERVICE = 42;
+
+unsigned int
+read_uint(const unsigned char *ptr)
+{
+    unsigned int value = 0;
+
+    for (int ii = sizeof(unsigned int) - 1; ii >= 0; --ii) {
+        value <<= 8;
+        value |= ptr[ii];
+    }
+
+    return value;
+}
+
+__global__ void
+singlePacketSingleProducer(__ockl_as_stream_t *stream)
+{
+    uint len = STRLEN_HELLO_WORLD;
+
+    uint64_t connection_id;
+
+    __hip_as_write_block(stream, TEST_SERVICE, &connection_id,
+                         (const uint8_t *)STR_HELLO_WORLD, STRLEN_HELLO_WORLD,
+                         __OCKL_AS_CONNECTION_BEGIN | __OCKL_AS_CONNECTION_END);
+}
+
+bool
+checkSinglePacketSingleProducer(__ockl_as_stream_t *stream)
+{
+    if (stream->write_index != 1)
+        return false;
+
+    __ockl_as_packet_t *packet = &stream->base_address[0];
+    uint header = packet->header;
+    if (get_packet_type(header) != __OCKL_AS_PACKET_READY)
+        return false;
+
+    if (get_packet_service(header) != TEST_SERVICE)
+        return false;
+
+    if (get_packet_bytes(header) != STRLEN_HELLO_WORLD)
+        return false;
+
+    if (get_packet_flags(header) !=
+        (__OCKL_AS_CONNECTION_BEGIN | __OCKL_AS_CONNECTION_END))
+        return false;
+
+    if (0 != strcmp(STR_HELLO_WORLD, (const char *)packet->payload))
+        return false;
+
+    return true;
+}
+
+bool
+ocklAsSinglePacketSingleProducer()
+{
+    bool success = true;
+    unsigned int numThreads = 1;
+    unsigned int numBlocks = 1;
+
+    unsigned int numPackets = 1;
+    unsigned int bufferSize =
+        sizeof(__ockl_as_stream_t) + numPackets * __OCKL_AS_PACKET_SIZE;
+
+    void *buffer;
+    HIPCHECK(hipHostMalloc(&buffer, bufferSize));
+
+    __ockl_as_stream_t *stream = createStream(buffer, bufferSize, numPackets);
+
+    hipLaunchKernelGGL(singlePacketSingleProducer, dim3(numBlocks),
+                       dim3(numThreads), 0, 0, stream);
+
+    HIPCHECK(hipDeviceSynchronize());
+
+    if (!checkSinglePacketSingleProducer(stream)) {
+        test_failed(__func__);
+        success = false;
+    }
+
+    HIPCHECK(hipHostFree(buffer));
+    return success;
+}
+
+__global__ void
+multipleProducers(__ockl_as_stream_t *stream)
+{
+    const unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
+    unsigned char data = (unsigned char)tid;
+
+    uint64_t connection_id;
+
+    __hip_as_write_block(stream, tid % NUM_SERVICES, &connection_id,
+                         (const unsigned char *)&tid, sizeof(unsigned int),
+                         __OCKL_AS_CONNECTION_BEGIN | __OCKL_AS_CONNECTION_END);
+}
+
+bool
+checkMultipleProducers(__ockl_as_stream_t *stream)
+{
+    int data[NUM_SERVICES] = {
+        0,
+    };
+
+    if (stream->write_index != NUM_THREADS)
+        return false;
+
+    for (int ii = 0; ii != NUM_THREADS; ++ii) {
+        __ockl_as_packet_t *packet = &stream->base_address[ii];
+        uint header = packet->header;
+        if (get_packet_type(header) != __OCKL_AS_PACKET_READY)
+            return false;
+
+        if (get_packet_bytes(header) != sizeof(unsigned int))
+            return false;
+
+        if (get_packet_flags(header) !=
+            (__OCKL_AS_CONNECTION_BEGIN | __OCKL_AS_CONNECTION_END))
+            return false;
+
+        unsigned char service = get_packet_service(header);
+        unsigned int payload = read_uint(packet->payload);
+
+        if (service != payload % NUM_SERVICES)
+            return false;
+        data[service]++;
+    }
+
+    int expected[NUM_SERVICES];
+    for (int ii = 0; ii != NUM_SERVICES; ++ii) {
+        expected[ii] = NUM_THREADS / NUM_SERVICES;
+        if (ii < NUM_THREADS % NUM_SERVICES) {
+            expected[ii]++;
+        }
+    }
+
+    for (int ii = 0; ii != NUM_SERVICES; ++ii) {
+        if (data[ii] != expected[ii])
+            return false;
+    }
+
+    return true;
+}
+
+bool
+ocklAsMultipleProducers()
+{
+    bool success = true;
+    unsigned int numPackets = NUM_THREADS;
+    unsigned int bufferSize =
+        sizeof(__ockl_as_stream_t) + numPackets * __OCKL_AS_PACKET_SIZE;
+
+    void *buffer;
+    HIPCHECK(hipHostMalloc(&buffer, bufferSize));
+
+    __ockl_as_stream_t *stream = createStream(buffer, bufferSize, numPackets);
+
+    hipLaunchKernelGGL(multipleProducers, dim3(NUM_BLOCKS),
+                       dim3(THREADS_PER_BLOCK), 0, 0, stream);
+
+    HIPCHECK(hipDeviceSynchronize());
+
+    if (!checkMultipleProducers(stream)) {
+        test_failed(__func__);
+        success = false;
+    }
+
+    HIPCHECK(hipHostFree(buffer));
+    return success;
+}
+
+__global__ void
+dropPackets(__ockl_as_stream_t *stream, unsigned int *status)
+{
+    const unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
+
+    uint64_t connection_id;
+
+    status[tid] =
+        __hip_as_write_block(stream, TEST_SERVICE, &connection_id,
+                             (const unsigned char *)&tid, sizeof(unsigned int),
+                             __OCKL_AS_CONNECTION_BEGIN |
+                                 __OCKL_AS_CONNECTION_END);
+}
+
+bool
+checkDropPackets(__ockl_as_stream_t *stream, unsigned int *status)
+{
+    unsigned int errorsExpected = NUM_THREADS - NUM_PACKETS_INSUFFICIENT;
+    for (int ii = 0; ii != NUM_THREADS; ++ii) {
+        switch (status[ii]) {
+        case __OCKL_AS_STATUS_OUT_OF_RESOURCES:
+            if (errorsExpected == 0)
+                return false;
+            --errorsExpected;
+            break;
+        case __OCKL_AS_STATUS_SUCCESS:
+            break;
+        default:
+            return false;
+        }
+    }
+    if (errorsExpected != 0)
+        return false;
+
+    if (stream->write_index != NUM_PACKETS_INSUFFICIENT)
+        return false;
+
+    for (int ii = 0; ii != NUM_PACKETS_INSUFFICIENT; ++ii) {
+        __ockl_as_packet_t *packet = &stream->base_address[ii];
+        uint header = packet->header;
+        if (get_packet_type(header) != __OCKL_AS_PACKET_READY)
+            return false;
+
+        if (get_packet_service(header) != TEST_SERVICE)
+            return false;
+
+        if (get_packet_bytes(header) != sizeof(unsigned int))
+            return false;
+
+        if (get_packet_flags(header) !=
+            (__OCKL_AS_CONNECTION_BEGIN | __OCKL_AS_CONNECTION_END))
+            return false;
+
+        unsigned int payload = read_uint(packet->payload);
+        if (payload >= NUM_THREADS)
+            return false;
+    }
+
+    return true;
+}
+
+bool
+ocklAsDropPackets()
+{
+    bool success = true;
+
+    unsigned int numPackets = NUM_PACKETS_INSUFFICIENT;
+    unsigned int bufferSize =
+        sizeof(__ockl_as_stream_t) + numPackets * __OCKL_AS_PACKET_SIZE;
+    void *buffer;
+    HIPCHECK(hipHostMalloc(&buffer, bufferSize));
+
+    void *status;
+    HIPCHECK(hipHostMalloc(&status, NUM_THREADS * sizeof(unsigned int)));
+
+    __ockl_as_stream_t *stream = createStream(buffer, bufferSize, numPackets);
+
+    hipEvent_t event;
+    HIPCHECK(hipEventCreate(&event));
+    hipLaunchKernelGGL(dropPackets, dim3(NUM_BLOCKS), dim3(THREADS_PER_BLOCK),
+                       0, 0, stream, (unsigned int *)status);
+
+    HIPCHECK(hipDeviceSynchronize());
+
+    if (!checkDropPackets(stream, (unsigned int *)status)) {
+        test_failed(__func__);
+        success = false;
+    }
+
+    HIPCHECK(hipHostFree(buffer));
+    return success;
+}
+
+#define STR30 "Cras nec volutpat mi, sed sed."
+#define STR47 "Lorem ipsum dolor sit amet, consectetur nullam."
+#define STR60 "Curabitur id maximus nibh. Donec quis porttitor nisl nullam."
+#define STR95                                                                  \
+    "In mollis imperdiet nibh nec ullamcorper."                                \
+    " Suspendisse placerat massa iaculis ipsum viverra sed."
+#define STR124                                                                 \
+    "Proin ut diam sit amet erat mollis gravida ac non sem."                   \
+    " Mauris viverra leo metus, id luctus metus feugiat sed. Morbi "           \
+    "posuere."
+
+#define DECLARE_TEST_DATA()                                                    \
+    const char *str30 = STR30;                                                 \
+    const char *str60 = STR60;                                                 \
+    const char *str47 = STR47;                                                 \
+    const char *str95 = STR95;                                                 \
+    const char *str124 = STR124;                                               \
+    const int numStr = 5;                                                      \
+    const char *strArray[5] = {str30, str60, str47, str95, str124};            \
+    unsigned char strLengths[5] = {30, 60, 47, 95, 124};
+
+__global__ void
+mixedProducers(__ockl_as_stream_t *stream)
+{
+    DECLARE_TEST_DATA();
+
+    const unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
+    const unsigned int idx = tid % 5;
+    uint64_t connection_id;
+
+    __hip_as_write_block(stream, idx, &connection_id,
+                         (unsigned const char *)strArray[idx], strLengths[idx],
+                         __OCKL_AS_CONNECTION_BEGIN | __OCKL_AS_CONNECTION_END);
+}
+
+bool
+checkMixedProducers(__ockl_as_stream_t *stream)
+{
+    typedef std::unordered_map<uint64_t, std::string> stream_buffer_map_t;
+    stream_buffer_map_t buffers;
+    std::unordered_map<std::string, int> strRecd;
+
+    DECLARE_TEST_DATA();
+
+    for (unsigned long read_index = 0; read_index != stream->write_index;
+         ++read_index) {
+        __ockl_as_packet_t *packet = stream->base_address + read_index;
+
+        uint header = packet->header;
+        if (get_packet_type(header) != __OCKL_AS_PACKET_READY)
+            return false;
+
+        uint bytes = get_packet_bytes(header);
+        unsigned char flags = get_packet_flags(header);
+        unsigned char service = get_packet_service(header);
+        unsigned long connection_id = packet->connection_id;
+        unsigned char *payload = packet->payload;
+
+        if ((flags & __OCKL_AS_CONNECTION_BEGIN) !=
+            (buffers.count(connection_id) == 0))
+            return false;
+
+        std::string &buf = buffers[connection_id];
+        buf.insert(buf.end(), payload, payload + bytes);
+
+        if (flags & __OCKL_AS_CONNECTION_END) {
+            if (buf != strArray[service])
+                return false;
+            strRecd[buf]++;
+            buffers.erase(connection_id);
+        }
+    }
+
+    int expected_counts[numStr];
+    for (int ii = 0; ii != numStr; ++ii) {
+        expected_counts[ii] = NUM_THREADS / numStr;
+        if (ii < (NUM_THREADS % numStr)) {
+            ++expected_counts[ii];
+        }
+    }
+
+    if (strRecd.size() != numStr)
+        return false;
+
+    for (int ii = 0; ii != numStr; ++ii) {
+        std::string mystr(strArray[ii]);
+        if (strRecd[mystr] != expected_counts[ii])
+            return false;
+    }
+    return true;
+}
+
+bool
+ocklAsMixedProducers()
+{
+    bool success = true;
+
+    unsigned int numPackets = NUM_PACKETS_LARGE;
+    unsigned int bufferSize =
+        sizeof(__ockl_as_stream_t) + numPackets * __OCKL_AS_PACKET_SIZE;
+
+    void *buffer;
+    HIPCHECK(hipHostMalloc(&buffer, bufferSize));
+
+    __ockl_as_stream_t *stream = createStream(buffer, bufferSize, numPackets);
+
+    hipLaunchKernelGGL(mixedProducers, dim3(NUM_BLOCKS),
+                       dim3(THREADS_PER_BLOCK), 0, 0, stream);
+
+    HIPCHECK(hipDeviceSynchronize());
+
+    if (!checkMixedProducers(stream)) {
+        test_failed(__func__);
+        success = false;
+    }
+
+    HIPCHECK(hipHostFree(buffer));
+    return success;
+}
+
+#define STR27 "In et consectetur mi metus."
+#define STR64 "Praesent tempus arcu id ligula blandit, eget congue justo metus."
+#define STR40 "Sed at dolor ipsum. Curabitur cras amet."
+
+__global__ void
+splitMessage(__ockl_as_stream_t *stream)
+{
+    const char *str27 = STR27;
+    const char *str64 = STR64;
+    const char *str40 = STR40;
+    const int numStr = 3;
+    const char *strArray[] = {str27, str64, str40};
+    int strLengths[] = {27, 64, 40};
+
+    const unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
+    int service = tid % 3;
+    int first = tid % 3;
+    int second = (tid + 1) % 3;
+    int third = (tid + 2) % 3;
+
+    uint64_t connection_id;
+    __hip_as_write_block(stream, service, &connection_id,
+                         (unsigned const char *)strArray[first],
+                         strLengths[first], __OCKL_AS_CONNECTION_BEGIN);
+    __hip_as_write_block(stream, service, &connection_id,
+                         (unsigned const char *)strArray[second],
+                         strLengths[second], 0);
+    __hip_as_write_block(stream, service, &connection_id,
+                         (unsigned const char *)strArray[third],
+                         strLengths[third], __OCKL_AS_CONNECTION_END);
+}
+
+bool
+checkSplitMessage(__ockl_as_stream_t *stream)
+{
+    typedef std::unordered_map<uint64_t, std::string> stream_buffer_map_t;
+    stream_buffer_map_t buffers;
+    std::unordered_map<std::string, int> strRecd;
+
+    static const int numExpected = 3;
+    const char *strExpected[numExpected] = {STR27 STR64 STR40,
+                                            STR64 STR40 STR27,
+                                            STR40 STR27 STR64};
+
+    for (unsigned long read_index = 0; read_index != stream->write_index;
+         ++read_index) {
+        __ockl_as_packet_t *packet = stream->base_address + read_index;
+
+        uint header = packet->header;
+        if (get_packet_type(header) != __OCKL_AS_PACKET_READY)
+            return false;
+
+        uint bytes = get_packet_bytes(header);
+        unsigned char flags = get_packet_flags(header);
+        unsigned long connection_id = packet->connection_id;
+        unsigned char *payload = packet->payload;
+        uint8_t service = get_packet_service(header);
+
+        if ((flags & __OCKL_AS_CONNECTION_BEGIN) !=
+            (buffers.count(connection_id) == 0))
+            return false;
+
+        std::string &buf = buffers[connection_id];
+        buf.insert(buf.end(), payload, payload + bytes);
+
+        if (flags & __OCKL_AS_CONNECTION_END) {
+            if (buf != strExpected[service])
+                return false;
+            strRecd[buf] += 1;
+            buffers.erase(connection_id);
+        }
+    }
+
+    int expected_counts[numExpected];
+    for (int ii = 0; ii != numExpected; ++ii) {
+        expected_counts[ii] = NUM_THREADS / numExpected;
+        if (ii < (NUM_THREADS % numExpected)) {
+            ++expected_counts[ii];
+        }
+    }
+
+    if (strRecd.size() != numExpected)
+        return false;
+
+    for (int ii = 0; ii != numExpected; ++ii) {
+        std::string mystr(strExpected[ii]);
+        if (strRecd[mystr] != expected_counts[ii])
+            return false;
+    }
+
+    return true;
+}
+
+bool
+ocklAsSplitMessage()
+{
+    bool success = true;
+
+    unsigned int numPackets = NUM_PACKETS_LARGE;
+    unsigned int bufferSize =
+        sizeof(__ockl_as_stream_t) + numPackets * __OCKL_AS_PACKET_SIZE;
+
+    void *buffer;
+    HIPCHECK(hipHostMalloc(&buffer, bufferSize));
+
+    __ockl_as_stream_t *stream = createStream(buffer, bufferSize, numPackets);
+
+    hipLaunchKernelGGL(splitMessage, dim3(NUM_BLOCKS), dim3(THREADS_PER_BLOCK),
+                       0, 0, stream);
+
+    HIPCHECK(hipDeviceSynchronize());
+
+    if (!checkSplitMessage(stream)) {
+        test_failed(__func__);
+        success = false;
+    }
+
+    HIPCHECK(hipHostFree(buffer));
+    return success;
+}
+
+#define TESTNAME "hipAsynchronousStreams"
+int
+main(int argc, char **argv)
+{
+    bool success = true;
+
+    success &= ocklAsSinglePacketSingleProducer();
+    success &= ocklAsMultipleProducers();
+    success &= ocklAsDropPackets();
+    success &= ocklAsMixedProducers();
+    success &= ocklAsSplitMessage();
+
+    hipDeviceReset();
+
+    if (success) {
+        test_passed(TESTNAME);
+        return 0;
+    }
+
+    failed(TESTNAME);
+}