From 98b7c724dc35f330f4d9f6698d0735067d7bcfdd Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 4 Jan 2019 11:04:47 +0530 Subject: [PATCH 1/8] Update INSTALL.md - Improved install instructions for AMD-hcc and NVIDIA-nvcc. - Removed redundant section related to HCC-Native-GCN-ISA. Change-Id: I33a9e71d95df51dbdcade4a48834840146bdbe6f [ROCm/hip commit: b8c78f90a681c04a99a9542d4e7045e78f27878c] --- projects/hip/INSTALL.md | 25 ++++++------------------- 1 file changed, 6 insertions(+), 19 deletions(-) diff --git a/projects/hip/INSTALL.md b/projects/hip/INSTALL.md index f773c8bde8..7b12bd42ef 100644 --- a/projects/hip/INSTALL.md +++ b/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 -``` - From d1a7eb0cb514717d2fe68c58b3ce4565e4fd0da5 Mon Sep 17 00:00:00 2001 From: Sameer Sahasrabuddhe Date: Tue, 23 Oct 2018 01:17:11 -0400 Subject: [PATCH 2/8] introduce test for asynchronous streams This test is identical to the HCC tests under tests/Unit/OCKL/AsyncStreams, but only for synchronous operation. [ROCm/hip commit: 8fe2a4b1671937cc8b7273418933ae93f839c867] --- .../src/deviceLib/hipAsynchronousStreams.cpp | 713 ++++++++++++++++++ 1 file changed, 713 insertions(+) create mode 100644 projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp diff --git a/projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp b/projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp new file mode 100644 index 0000000000..a047de322b --- /dev/null +++ b/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 + * RUN: %t + * HIT_END + */ + +// Includes HIP Runtime +#include "hip/hip_runtime.h" +#include + +#include +#include + +#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 stream_buffer_map_t; + stream_buffer_map_t buffers; + std::unordered_map 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 stream_buffer_map_t; + stream_buffer_map_t buffers; + std::unordered_map 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); +} From 60adc94cad230bb5ac756ccc3b7691a341984077 Mon Sep 17 00:00:00 2001 From: emankov Date: Sat, 5 Jan 2019 14:52:15 +0300 Subject: [PATCH 3/8] [HIPIFY][fix][#849] Create temporary file based on source file name only [ROCm/hip commit: 4a475a9db6edc988900b4a537bd72c09089d5723] --- projects/hip/hipify-clang/src/main.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp index 2a3b86f204..bd54975e59 100644 --- a/projects/hip/hipify-clang/src/main.cpp +++ b/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; } From 57b553324b39ddd1103bed4f9bd6be0eb3081b58 Mon Sep 17 00:00:00 2001 From: emankov Date: Sun, 6 Jan 2019 17:58:37 +0300 Subject: [PATCH 4/8] [HIPIFY][fix][#851] make_absolute only in case of relative path Otherwise return the input path. [ROCm/hip commit: bff9e1ebe24019fd1e78b7f9a2978a72890df8ec] --- projects/hip/hipify-clang/src/LLVMCompat.cpp | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/projects/hip/hipify-clang/src/LLVMCompat.cpp b/projects/hip/hipify-clang/src/LLVMCompat.cpp index 22a0b56d64..95c8ab7837 100644 --- a/projects/hip/hipify-clang/src/LLVMCompat.cpp +++ b/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 &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 From cf7c7f2db26948d3a94518304777ec9b89f85122 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 7 Jan 2019 09:26:21 +0530 Subject: [PATCH 5/8] Update hipAsynchronousStreams.cpp This test needs to be excluded from nvcc platform [ROCm/hip commit: 580c8084001136ae10d3966bf8a484211fcc53b1] --- projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp b/projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp index a047de322b..17e6254154 100644 --- a/projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp +++ b/projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp @@ -19,7 +19,7 @@ THE SOFTWARE. /* * HIT_START - * BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -std=c++11 + * BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM nvcc * RUN: %t * HIT_END */ From 2f9021afd3ca17bd9d5655425a82a9621d643356 Mon Sep 17 00:00:00 2001 From: "Wen-Heng (Jack) Chung" Date: Tue, 8 Jan 2019 17:18:06 +0000 Subject: [PATCH 6/8] Introduce hash key to HIP module implementation A hash calculated via FNV-1a algorithm is introduced in ihipModule_t, the internal of hipModule_t. The hash is used by HIP module APIs such as - read_agent_global_from_module to determine whether the agent-scope globals for a module have been iterated. This commit fixes one issue that applications which load / unload modules frequently would occasionally fail. After deep investigation of the issue it turns out the old implementation in read_agent_global_from_module uses hipModule_t as the key, which is not robust enough, as hipModule_t instances are allocated dynamically so there are cases that one memory address may be used by multiple hipModule_t instances. The real solution is to introduce a uniquely identifiable hash for the code object associated with the HIP module. And that's the rationale behind this commit. [ROCm/hip commit: b4d658a48f2c47568e9b333aa28b57f2478d7a76] --- projects/hip/src/hip_hcc_internal.h | 1 + projects/hip/src/hip_module.cpp | 41 ++++++++++++++++++++++------- 2 files changed, 32 insertions(+), 10 deletions(-) diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h index d71ad850e3..be794956f6 100644 --- a/projects/hip/src/hip_hcc_internal.h +++ b/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); diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index 657fb06b5e..780240c067 100644 --- a/projects/hip/src/hip_module.cpp +++ b/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(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 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> agent_globals; + // the key of the map would the hash of code object associated with the + // hipModule_t instance + static unordered_map> 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 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; } From e816f3e98a2eae44d6813d5fd04ef81be984c6fb Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 9 Jan 2019 13:59:48 +0300 Subject: [PATCH 7/8] [HIPIFY][cmake] Add cmake versioning depending on CUDA version [ROCm/hip commit: 8fb8ca02c85b12324c1cec4868b7480f9bb4eec6] --- projects/hip/hipify-clang/CMakeLists.txt | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/projects/hip/hipify-clang/CMakeLists.txt b/projects/hip/hipify-clang/CMakeLists.txt index 964e30d66a..bfd053984e 100644 --- a/projects/hip/hipify-clang/CMakeLists.txt +++ b/projects/hip/hipify-clang/CMakeLists.txt @@ -1,4 +1,13 @@ -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() +endif() + project(hipify-clang) if (MSVC AND MSVC_VERSION VERSION_LESS "1900") From 9314f882f3a095cbfce6e8f59f134466674b1652 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 9 Jan 2019 16:44:12 +0300 Subject: [PATCH 8/8] [HIPIFY][cmake][win] Add cmake versioning depending on Visual Studio version [ROCm/hip commit: 96f8bd7e2b287e4fe9faa11a22e7637246780eb9] --- projects/hip/hipify-clang/CMakeLists.txt | 3 +++ 1 file changed, 3 insertions(+) diff --git a/projects/hip/hipify-clang/CMakeLists.txt b/projects/hip/hipify-clang/CMakeLists.txt index bfd053984e..4b2d9cd1aa 100644 --- a/projects/hip/hipify-clang/CMakeLists.txt +++ b/projects/hip/hipify-clang/CMakeLists.txt @@ -6,6 +6,9 @@ else() 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)