From patchwork Sat Jan 14 13:20:38 2017 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: =?utf-8?b?UGVra2EgSsOkw6Rza2Vsw6RpbmVu?= X-Patchwork-Id: 715355 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 3v10Wr6VThz9t17 for ; Sun, 15 Jan 2017 00:21:43 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="galZnx7E"; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :mime-version:in-reply-to:references:from:date:message-id :subject:to:cc:content-type; q=dns; s=default; b=TdN4jW6rchC66tx fhVkF0+udH9wqE54SWlgkaPE3KrQtvRNXKg1bmXERF6542lhGUpIt0i+hghsUJtz v7EaPyxnQm9Wa2R/uNOXblPQELMJEVGrLxMEQq2C0JMr7/MamrSuP+/Bv56uhKhi Gb/LWLfNNm277O8inpVMYycniKBk= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :mime-version:in-reply-to:references:from:date:message-id :subject:to:cc:content-type; s=default; bh=SshdgTnqDgW/HJ1NUe6Q9 JRnZY0=; b=galZnx7EBOERM8PyxXqq3UNXz5Y9elsfLSfw2EpnDxb0NNAN8wkEv CkO0NbC1336zsaezxdXPRIs13FXDHtQYfDCxE9i/JFQzJnWe4BefZA9UfIX/4LRd 7T7z6GcjPZI41s70+U4uKnRa8b/WIJlOgcIwnrRmlGW8F8wmSB9L20= Received: (qmail 57906 invoked by alias); 14 Jan 2017 13:21:13 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 57892 invoked by uid 89); 14 Jan 2017 13:21:11 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=0.4 required=5.0 tests=AWL, BAYES_50, KAM_STOCKGEN, RCVD_IN_DNSWL_LOW, RCVD_IN_SORBS_SPAM autolearn=ham version=3.3.2 spammy=tech, Tech, fiber, OTHER X-HELO: mail-qt0-f178.google.com Received: from mail-qt0-f178.google.com (HELO mail-qt0-f178.google.com) (209.85.216.178) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Sat, 14 Jan 2017 13:21:00 +0000 Received: by mail-qt0-f178.google.com with SMTP id k15so70319457qtg.3 for ; Sat, 14 Jan 2017 05:21:00 -0800 (PST) X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20161025; h=x-gm-message-state:mime-version:in-reply-to:references:from:date :message-id:subject:to:cc; bh=dWk5hUKXmj1oMvFVhR/hhXytS26ZHCAPDip3ZrbPtnA=; b=ZLu+nu95y/6Ix13gnKMrOE6Dnd8xGNKkkfbQscMKM2AEWA2rLRh5j2i2SghK6y5JQz TffoR23QzlFevcOMzxy2qjCr6YNJlTzQmHU6ce5fkKPoOTkmNMTP+ogPhTvmoYwQfZEF H6boOuNVv1JuKjral3+e38hCobWGbltOnAXEtnxaqnqZpx2PiWtyz8VTqSS8kIL971PT OmrMrP5rTcYxHF+K+TPzTSOruzX+hGeJTXXW2n7VAonuYobG/+EueIiGsSVPJsNJKKRl tVDyYYZvU7TazpCZopPC/MrYgWtYrxkj4l8qfn4nDO9OrXDpsxqoDF3F97y+RinSsRQn ASIA== X-Gm-Message-State: AIkVDXIhllpYPRHcxulr/ntjSr+H/FCzZi+YpwKdPLm92Q+5VsnyaKtnOvAm+eEUFn9PInFh9zRZ5Un6+gJ8Iw== X-Received: by 10.237.45.7 with SMTP id h7mr21312466qtd.280.1484400058605; Sat, 14 Jan 2017 05:20:58 -0800 (PST) MIME-Version: 1.0 Received: by 10.237.45.4 with HTTP; Sat, 14 Jan 2017 05:20:38 -0800 (PST) In-Reply-To: References: From: =?UTF-8?B?UGVra2EgSsOkw6Rza2Vsw6RpbmVu?= Date: Sat, 14 Jan 2017 15:20:38 +0200 Message-ID: Subject: Re: [PATCH] BRIG frontend: request for a global review To: Richard Biener Cc: GCC Patches , Martin Jambor Hi Richard, Attached is an updated BRIG FE patch set with everything else you found corrected except the builtin issue. It seems unclear if converting the HSAIL instruction-based builtins to direct runtime function declaration calls is an improvement to the current way. Changes to the previous in 005-diff-to-previous.patch Thanks again for the review, Pekka On Fri, Jan 13, 2017 at 5:43 PM, Pekka Jääskeläinen wrote: > Hi Richard, > > Thanks for the review! Replies/questions inline. > > On Fri, Jan 13, 2017 at 2:28 PM, Richard Biener > wrote: >> On Thu, Jan 12, 2017 at 3:55 PM, Pekka Jääskeläinen wrote: >>> Hi, >>> >>> A gentle ping... >> >> diff --git a/gcc/configure.ac b/gcc/configure.ac >> index 140b9f9..06941c5 100644 >> --- a/gcc/configure.ac >> +++ b/gcc/configure.ac >> @@ -996,6 +996,11 @@ if test x"$enable_hsa" = x1 ; then >> [Define this to enable support for generating HSAIL.]) >> fi >> >> +if echo "$enable_languages" | grep "brig" > /dev/null; then >> + AC_DEFINE(ENABLE_BRIG_FE, 1, >> + [Define this to enable the BRIG (HSAIL) frontend.]) >> +fi >> + >> >> this looks odd and I'd have expected this to be solely controlled via >> --enable-languages at configure time? > > Yes, it is. This AC_DEFINE just adds an autoconf variable ENABLE_BRIG_FE > using which I guard unnecessary inclusion of the brig-builtins.def in > builtins.def in > case BRIG FE is disabled. > >> @@ -1366,6 +1369,11 @@ assembler assembler-with-cpp >> ada >> f77 f77-cpp-input f95 f95-cpp-input >> go >> +<<<<<<< HEAD >> +java >> +brig >> +======= >> +>>>>>>> gcc-master >> @end smallexample >> >> @item -x none >> >> unmerged hunk (java is gone). > > Cleaned. > >> diff --git a/include/hsa-interface.h b/include/hsa-interface.h >> new file mode 100644 >> index 0000000..6765751 >> --- /dev/null >> +++ b/include/hsa-interface.h >> @@ -0,0 +1,630 @@ >> +/* HSA runtime API 1.0.1 representation description. >> + Copyright (C) 2016 Free Software Foundation, Inc. >> + >> ... >> >> this looks like a sub/superset of libgomp/plugin/hsa.h, please work on >> retaining only one version. > > Merged. > >> Did you check whether libhsail-rt builds for all GCC targets? If not >> you probably want to add >> a list where to disable the frontend and its runtime (see go / libgo >> for an example in the >> toplevel configure.ac). > > I have tested only x86 of the GCC upstream targets so far. > > Libgo seems to black list known broken ones with a separate --enable-libgo > switch to force enable. However, BRIG FE and libhsail-rt are disabled > by default. > Should I still add a separate switch to force enable BRIG and enable > it by default for x86? > > Or should I force disable it even when enabled with --enable-languages when > building for an untested target? > > > Thanks, > Pekka diff --git a/configure b/configure index 6ce390e..d66e219 100755 --- a/configure +++ b/configure @@ -3467,6 +3467,19 @@ if test x$enable_libgo = x; then esac fi +# Disable the BRIG frontend and libhsail-rt on untested or known +# broken systems. Currently it has been tested only on x86_64 Linux +# of the upstream gcc targets. More targets shall be added after testing. +case "${target}" in + x86_64-*-linux*) + ;; + *) + unsupported_languages="$unsupported_languages brig" + # This implicitly disables also target-libhsail-rt as it won't + # get added to the build without BRIG FE. + ;; +esac + # Default libgloss CPU subdirectory. libgloss_dir="$target_cpu" diff --git a/configure.ac b/configure.ac index 23e2277..aaa18f0 100644 --- a/configure.ac +++ b/configure.ac @@ -814,6 +814,19 @@ if test x$enable_libgo = x; then esac fi +# Disable the BRIG frontend and libhsail-rt on untested or known +# broken systems. Currently it has been tested only on x86_64 Linux +# of the upstream gcc targets. More targets shall be added after testing. +case "${target}" in + x86_64-*-linux*) + ;; + *) + unsupported_languages="$unsupported_languages brig" + # This implicitly disables also target-libhsail-rt as it won't + # get added to the build without BRIG FE. + ;; +esac + # Default libgloss CPU subdirectory. libgloss_dir="$target_cpu" diff --git a/gcc/doc/install.texi b/gcc/doc/install.texi index 0cf2f8a..1bec440 100644 --- a/gcc/doc/install.texi +++ b/gcc/doc/install.texi @@ -2594,7 +2594,10 @@ separately. Second, you must have the testing tools installed. This includes @uref{http://www.gnu.org/software/dejagnu/,,DejaGnu}, Tcl, and Expect; -the DejaGnu site has links to these. +the DejaGnu site has links to these. For running the BRIG frontend +tests, a tool to assemble the binary BRIGs from HSAIL text, +@uref{https://github.com/HSAFoundation/HSAIL-Tools/,,HSAILasm} must +be installed. If the directories where @command{runtest} and @command{expect} were installed are not in the @env{PATH}, you may need to set the following diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi index 4cbca4f..4c18397 100644 --- a/gcc/doc/invoke.texi +++ b/gcc/doc/invoke.texi @@ -1369,11 +1369,7 @@ assembler assembler-with-cpp ada f77 f77-cpp-input f95 f95-cpp-input go -<<<<<<< HEAD -java brig -======= ->>>>>>> gcc-master @end smallexample @item -x none diff --git a/gcc/testsuite/brig.dg/README b/gcc/testsuite/brig.dg/README index cc313c4..2ad5b24 100644 --- a/gcc/testsuite/brig.dg/README +++ b/gcc/testsuite/brig.dg/README @@ -7,4 +7,6 @@ suite for all HSAIL instructions. The HSA PRM conformance suite is supposed to be used for that. HSAILasm is required for converting the text HSAIL files to BRIGs -which the compiler consumes. +which the compiler consumes. It can be built from +https://github.com/HSAFoundation/HSAIL-Tools + diff --git a/gcc/testsuite/brig.dg/dg.exp b/gcc/testsuite/brig.dg/dg.exp index fd75cae..d4f37e6 100644 --- a/gcc/testsuite/brig.dg/dg.exp +++ b/gcc/testsuite/brig.dg/dg.exp @@ -4,12 +4,12 @@ # it under the terms of the GNU General Public License as published by # the Free Software Foundation; either version 3 of the License, or # (at your option) any later version. -# +# # This program is distributed in the hope that it will be useful, # but WITHOUT ANY WARRANTY; without even the implied warranty of # MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the # GNU General Public License for more details. -# +# # You should have received a copy of the GNU General Public License # along with GCC; see the file COPYING3. If not see # . @@ -21,7 +21,11 @@ load_lib brig-dg.exp # Initialize `dg'. dg-init -dg-runtest [find $srcdir/$subdir *.hsail] "" "" +if [expr [llength [auto_execok HSAILasm]] > 0] { + dg-runtest [find $srcdir/$subdir *.hsail] "" "" +} else { + unsupported "All BRIG FE tests require HSAILasm in PATH." +} # All done. dg-finish diff --git a/libgomp/plugin/hsa.h b/libgomp/plugin/hsa.h deleted file mode 100644 index 6765751..0000000 --- a/libgomp/plugin/hsa.h +++ /dev/null @@ -1,630 +0,0 @@ -/* HSA runtime API 1.0.1 representation description. - Copyright (C) 2016 Free Software Foundation, Inc. - -This file is part of GCC. - -GCC is free software; you can redistribute it and/or modify -it under the terms of the GNU General Public License as published by -the Free Software Foundation; either version 3, or (at your option) -any later version. - -GCC is distributed in the hope that it will be useful, -but WITHOUT ANY WARRANTY; without even the implied warranty of -MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the -GNU General Public License for more details. - -You should have received a copy of the GNU General Public License -along with GCC; see the file COPYING3. If not see -. - -The contents of the file was created by extracting data structures, enum, -typedef and other definitions from HSA Runtime Programmer’s Reference Manual -Version 1.0 (http://www.hsafoundation.com/standards/). - -HTML version is provided on the following link: -http://www.hsafoundation.com/html/Content/Runtime/Topics/Runtime_title_page.htm -*/ - -#ifndef _HSA_H -#define _HSA_H 1 - -#define HSA_LARGE_MODEL 1 - -typedef struct hsa_signal_s { uint64_t handle; } hsa_signal_t; -typedef enum { - HSA_QUEUE_TYPE_MULTI = 0, - HSA_QUEUE_TYPE_SINGLE = 1 -} hsa_queue_type_t; - -typedef enum { HSA_PROFILE_BASE = 0, HSA_PROFILE_FULL = 1 } hsa_profile_t; -typedef struct hsa_region_s { uint64_t handle; } hsa_region_t; -typedef enum { - HSA_EXECUTABLE_SYMBOL_INFO_TYPE = 0, - HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH = 1, - HSA_EXECUTABLE_SYMBOL_INFO_NAME = 2, - HSA_EXECUTABLE_SYMBOL_INFO_MODULE_NAME_LENGTH = 3, - HSA_EXECUTABLE_SYMBOL_INFO_MODULE_NAME = 4, - HSA_EXECUTABLE_SYMBOL_INFO_AGENT = 20, - HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS = 21, - HSA_EXECUTABLE_SYMBOL_INFO_LINKAGE = 5, - HSA_EXECUTABLE_SYMBOL_INFO_IS_DEFINITION = 17, - HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ALLOCATION = 6, - HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SEGMENT = 7, - HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ALIGNMENT = 8, - HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE = 9, - HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_IS_CONST = 10, - HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT = 22, - HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE = 11, - HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT = 12, - HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE = 13, - HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE = 14, - HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK = 15, - HSA_EXECUTABLE_SYMBOL_INFO_INDIRECT_FUNCTION_OBJECT = 23, - HSA_EXECUTABLE_SYMBOL_INFO_INDIRECT_FUNCTION_CALL_CONVENTION = 16 -} hsa_executable_symbol_info_t; -typedef enum { - HSA_REGION_GLOBAL_FLAG_KERNARG = 1, - HSA_REGION_GLOBAL_FLAG_FINE_GRAINED = 2, - HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED = 4 -} hsa_region_global_flag_t; -typedef struct hsa_code_object_s { uint64_t handle; } hsa_code_object_t; -typedef enum { - HSA_KERNEL_DISPATCH_PACKET_SETUP_WIDTH_DIMENSIONS = 2 -} hsa_kernel_dispatch_packet_setup_width_t; -typedef enum { - HSA_DEVICE_TYPE_CPU = 0, - HSA_DEVICE_TYPE_GPU = 1, - HSA_DEVICE_TYPE_DSP = 2 -} hsa_device_type_t; -typedef enum { - HSA_STATUS_SUCCESS = 0x0, - HSA_STATUS_INFO_BREAK = 0x1, - HSA_STATUS_ERROR = 0x1000, - HSA_STATUS_ERROR_INVALID_ARGUMENT = 0x1001, - HSA_STATUS_ERROR_INVALID_QUEUE_CREATION = 0x1002, - HSA_STATUS_ERROR_INVALID_ALLOCATION = 0x1003, - HSA_STATUS_ERROR_INVALID_AGENT = 0x1004, - HSA_STATUS_ERROR_INVALID_REGION = 0x1005, - HSA_STATUS_ERROR_INVALID_SIGNAL = 0x1006, - HSA_STATUS_ERROR_INVALID_QUEUE = 0x1007, - HSA_STATUS_ERROR_OUT_OF_RESOURCES = 0x1008, - HSA_STATUS_ERROR_INVALID_PACKET_FORMAT = 0x1009, - HSA_STATUS_ERROR_RESOURCE_FREE = 0x100A, - HSA_STATUS_ERROR_NOT_INITIALIZED = 0x100B, - HSA_STATUS_ERROR_REFCOUNT_OVERFLOW = 0x100C, - HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS = 0x100D, - HSA_STATUS_ERROR_INVALID_INDEX = 0x100E, - HSA_STATUS_ERROR_INVALID_ISA = 0x100F, - HSA_STATUS_ERROR_INVALID_ISA_NAME = 0x1017, - HSA_STATUS_ERROR_INVALID_CODE_OBJECT = 0x1010, - HSA_STATUS_ERROR_INVALID_EXECUTABLE = 0x1011, - HSA_STATUS_ERROR_FROZEN_EXECUTABLE = 0x1012, - HSA_STATUS_ERROR_INVALID_SYMBOL_NAME = 0x1013, - HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED = 0x1014, - HSA_STATUS_ERROR_VARIABLE_UNDEFINED = 0x1015, - HSA_STATUS_ERROR_EXCEPTION = 0x1016 -} hsa_status_t; -typedef enum { - HSA_EXTENSION_FINALIZER = 0, - HSA_EXTENSION_IMAGES = 1 -} hsa_extension_t; -typedef struct hsa_queue_s { - hsa_queue_type_t type; - uint32_t features; - -#ifdef HSA_LARGE_MODEL - void *base_address; -#elif defined HSA_LITTLE_ENDIAN - void *base_address; - uint32_t reserved0; -#else - uint32_t reserved0; - void *base_address; -#endif - - hsa_signal_t doorbell_signal; - uint32_t size; - uint32_t reserved1; - uint64_t id; -} hsa_queue_t; -typedef struct hsa_agent_dispatch_packet_s { - uint16_t header; - uint16_t type; - uint32_t reserved0; - -#ifdef HSA_LARGE_MODEL - void *return_address; -#elif defined HSA_LITTLE_ENDIAN - void *return_address; - uint32_t reserved1; -#else - uint32_t reserved1; - void *return_address; -#endif - uint64_t arg[4]; - uint64_t reserved2; - hsa_signal_t completion_signal; -} hsa_agent_dispatch_packet_t; -typedef enum { - HSA_CODE_SYMBOL_INFO_TYPE = 0, - HSA_CODE_SYMBOL_INFO_NAME_LENGTH = 1, - HSA_CODE_SYMBOL_INFO_NAME = 2, - HSA_CODE_SYMBOL_INFO_MODULE_NAME_LENGTH = 3, - HSA_CODE_SYMBOL_INFO_MODULE_NAME = 4, - HSA_CODE_SYMBOL_INFO_LINKAGE = 5, - HSA_CODE_SYMBOL_INFO_IS_DEFINITION = 17, - HSA_CODE_SYMBOL_INFO_VARIABLE_ALLOCATION = 6, - HSA_CODE_SYMBOL_INFO_VARIABLE_SEGMENT = 7, - HSA_CODE_SYMBOL_INFO_VARIABLE_ALIGNMENT = 8, - HSA_CODE_SYMBOL_INFO_VARIABLE_SIZE = 9, - HSA_CODE_SYMBOL_INFO_VARIABLE_IS_CONST = 10, - HSA_CODE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE = 11, - HSA_CODE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT = 12, - HSA_CODE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE = 13, - HSA_CODE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE = 14, - HSA_CODE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK = 15, - HSA_CODE_SYMBOL_INFO_INDIRECT_FUNCTION_CALL_CONVENTION = 16 -} hsa_code_symbol_info_t; -typedef enum { - HSA_QUEUE_FEATURE_KERNEL_DISPATCH = 1, - HSA_QUEUE_FEATURE_AGENT_DISPATCH = 2 -} hsa_queue_feature_t; -typedef enum { - HSA_VARIABLE_ALLOCATION_AGENT = 0, - HSA_VARIABLE_ALLOCATION_PROGRAM = 1 -} hsa_variable_allocation_t; -typedef enum { - HSA_FENCE_SCOPE_NONE = 0, - HSA_FENCE_SCOPE_AGENT = 1, - HSA_FENCE_SCOPE_SYSTEM = 2 -} hsa_fence_scope_t; -typedef struct hsa_agent_s { uint64_t handle; } hsa_agent_t; -typedef enum { HSA_CODE_OBJECT_TYPE_PROGRAM = 0 } hsa_code_object_type_t; -typedef enum { - HSA_SIGNAL_CONDITION_EQ = 0, - HSA_SIGNAL_CONDITION_NE = 1, - HSA_SIGNAL_CONDITION_LT = 2, - HSA_SIGNAL_CONDITION_GTE = 3 -} hsa_signal_condition_t; -typedef enum { - HSA_EXECUTABLE_STATE_UNFROZEN = 0, - HSA_EXECUTABLE_STATE_FROZEN = 1 -} hsa_executable_state_t; -typedef enum { - HSA_ENDIANNESS_LITTLE = 0, - HSA_ENDIANNESS_BIG = 1 -} hsa_endianness_t; -typedef enum { - HSA_MACHINE_MODEL_SMALL = 0, - HSA_MACHINE_MODEL_LARGE = 1 -} hsa_machine_model_t; -typedef enum { - HSA_AGENT_INFO_NAME = 0, - HSA_AGENT_INFO_VENDOR_NAME = 1, - HSA_AGENT_INFO_FEATURE = 2, - HSA_AGENT_INFO_MACHINE_MODEL = 3, - HSA_AGENT_INFO_PROFILE = 4, - HSA_AGENT_INFO_DEFAULT_FLOAT_ROUNDING_MODE = 5, - HSA_AGENT_INFO_BASE_PROFILE_DEFAULT_FLOAT_ROUNDING_MODES = 23, - HSA_AGENT_INFO_FAST_F16_OPERATION = 24, - HSA_AGENT_INFO_WAVEFRONT_SIZE = 6, - HSA_AGENT_INFO_WORKGROUP_MAX_DIM = 7, - HSA_AGENT_INFO_WORKGROUP_MAX_SIZE = 8, - HSA_AGENT_INFO_GRID_MAX_DIM = 9, - HSA_AGENT_INFO_GRID_MAX_SIZE = 10, - HSA_AGENT_INFO_FBARRIER_MAX_SIZE = 11, - HSA_AGENT_INFO_QUEUES_MAX = 12, - HSA_AGENT_INFO_QUEUE_MIN_SIZE = 13, - HSA_AGENT_INFO_QUEUE_MAX_SIZE = 14, - HSA_AGENT_INFO_QUEUE_TYPE = 15, - HSA_AGENT_INFO_NODE = 16, - HSA_AGENT_INFO_DEVICE = 17, - HSA_AGENT_INFO_CACHE_SIZE = 18, - HSA_AGENT_INFO_ISA = 19, - HSA_AGENT_INFO_EXTENSIONS = 20, - HSA_AGENT_INFO_VERSION_MAJOR = 21, - HSA_AGENT_INFO_VERSION_MINOR = 22 -} hsa_agent_info_t; -typedef struct hsa_barrier_and_packet_s { - uint16_t header; - uint16_t reserved0; - uint32_t reserved1; - hsa_signal_t dep_signal[5]; - uint64_t reserved2; - hsa_signal_t completion_signal; -} hsa_barrier_and_packet_t; -typedef struct hsa_dim3_s { - uint32_t x; - uint32_t y; - uint32_t z; -} hsa_dim3_t; -typedef enum { - HSA_ACCESS_PERMISSION_RO = 1, - HSA_ACCESS_PERMISSION_WO = 2, - HSA_ACCESS_PERMISSION_RW = 3 -} hsa_access_permission_t; -typedef enum { - HSA_AGENT_FEATURE_KERNEL_DISPATCH = 1, - HSA_AGENT_FEATURE_AGENT_DISPATCH = 2 -} hsa_agent_feature_t; -typedef enum { - HSA_WAIT_STATE_BLOCKED = 0, - HSA_WAIT_STATE_ACTIVE = 1 -} hsa_wait_state_t; -typedef struct hsa_executable_s { uint64_t handle; } hsa_executable_t; -typedef enum { - HSA_REGION_SEGMENT_GLOBAL = 0, - HSA_REGION_SEGMENT_READONLY = 1, - HSA_REGION_SEGMENT_PRIVATE = 2, - HSA_REGION_SEGMENT_GROUP = 3 -} hsa_region_segment_t; -typedef enum { - HSA_REGION_INFO_SEGMENT = 0, - HSA_REGION_INFO_GLOBAL_FLAGS = 1, - HSA_REGION_INFO_SIZE = 2, - HSA_REGION_INFO_ALLOC_MAX_SIZE = 4, - HSA_REGION_INFO_RUNTIME_ALLOC_ALLOWED = 5, - HSA_REGION_INFO_RUNTIME_ALLOC_GRANULE = 6, - HSA_REGION_INFO_RUNTIME_ALLOC_ALIGNMENT = 7 -} hsa_region_info_t; -typedef enum { - HSA_ISA_INFO_NAME_LENGTH = 0, - HSA_ISA_INFO_NAME = 1, - HSA_ISA_INFO_CALL_CONVENTION_COUNT = 2, - HSA_ISA_INFO_CALL_CONVENTION_INFO_WAVEFRONT_SIZE = 3, - HSA_ISA_INFO_CALL_CONVENTION_INFO_WAVEFRONTS_PER_COMPUTE_UNIT = 4 -} hsa_isa_info_t; -typedef enum { - HSA_VARIABLE_SEGMENT_GLOBAL = 0, - HSA_VARIABLE_SEGMENT_READONLY = 1 -} hsa_variable_segment_t; -typedef struct hsa_callback_data_s { uint64_t handle; } hsa_callback_data_t; -typedef enum { - HSA_SYMBOL_KIND_VARIABLE = 0, - HSA_SYMBOL_KIND_KERNEL = 1, - HSA_SYMBOL_KIND_INDIRECT_FUNCTION = 2 -} hsa_symbol_kind_t; -typedef struct hsa_kernel_dispatch_packet_s { - uint16_t header; - uint16_t setup; - uint16_t workgroup_size_x; - uint16_t workgroup_size_y; - uint16_t workgroup_size_z; - uint16_t reserved0; - uint32_t grid_size_x; - uint32_t grid_size_y; - uint32_t grid_size_z; - uint32_t private_segment_size; - uint32_t group_segment_size; - uint64_t kernel_object; - -#ifdef HSA_LARGE_MODEL - void *kernarg_address; -#elif defined HSA_LITTLE_ENDIAN - void *kernarg_address; - uint32_t reserved1; -#else - uint32_t reserved1; - void *kernarg_address; -#endif - uint64_t reserved2; - hsa_signal_t completion_signal; -} hsa_kernel_dispatch_packet_t; -typedef enum { - HSA_PACKET_TYPE_VENDOR_SPECIFIC = 0, - HSA_PACKET_TYPE_INVALID = 1, - HSA_PACKET_TYPE_KERNEL_DISPATCH = 2, - HSA_PACKET_TYPE_BARRIER_AND = 3, - HSA_PACKET_TYPE_AGENT_DISPATCH = 4, - HSA_PACKET_TYPE_BARRIER_OR = 5 -} hsa_packet_type_t; -typedef enum { - HSA_PACKET_HEADER_TYPE = 0, - HSA_PACKET_HEADER_BARRIER = 8, - HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE = 9, - HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE = 11 -} hsa_packet_header_t; -typedef struct hsa_isa_s { uint64_t handle; } hsa_isa_t; -typedef enum { - HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT = 0, - HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO = 1, - HSA_DEFAULT_FLOAT_ROUNDING_MODE_NEAR = 2 -} hsa_default_float_rounding_mode_t; -typedef struct hsa_code_symbol_s { uint64_t handle; } hsa_code_symbol_t; -typedef struct hsa_executable_symbol_s { - uint64_t handle; -} hsa_executable_symbol_t; -#ifdef HSA_LARGE_MODEL -typedef int64_t hsa_signal_value_t; -#else -typedef int32_t hsa_signal_value_t; -#endif -typedef enum { - HSA_EXCEPTION_POLICY_BREAK = 1, - HSA_EXCEPTION_POLICY_DETECT = 2 -} hsa_exception_policy_t; -typedef enum { - HSA_SYSTEM_INFO_VERSION_MAJOR = 0, - HSA_SYSTEM_INFO_VERSION_MINOR = 1, - HSA_SYSTEM_INFO_TIMESTAMP = 2, - HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY = 3, - HSA_SYSTEM_INFO_SIGNAL_MAX_WAIT = 4, - HSA_SYSTEM_INFO_ENDIANNESS = 5, - HSA_SYSTEM_INFO_MACHINE_MODEL = 6, - HSA_SYSTEM_INFO_EXTENSIONS = 7 -} hsa_system_info_t; -typedef enum { - HSA_EXECUTABLE_INFO_PROFILE = 1, - HSA_EXECUTABLE_INFO_STATE = 2 -} hsa_executable_info_t; -typedef enum { - HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS = 0 -} hsa_kernel_dispatch_packet_setup_t; -typedef enum { - HSA_PACKET_HEADER_WIDTH_TYPE = 8, - HSA_PACKET_HEADER_WIDTH_BARRIER = 1, - HSA_PACKET_HEADER_WIDTH_ACQUIRE_FENCE_SCOPE = 2, - HSA_PACKET_HEADER_WIDTH_RELEASE_FENCE_SCOPE = 2 -} hsa_packet_header_width_t; -typedef enum { - HSA_CODE_OBJECT_INFO_VERSION = 0, - HSA_CODE_OBJECT_INFO_TYPE = 1, - HSA_CODE_OBJECT_INFO_ISA = 2, - HSA_CODE_OBJECT_INFO_MACHINE_MODEL = 3, - HSA_CODE_OBJECT_INFO_PROFILE = 4, - HSA_CODE_OBJECT_INFO_DEFAULT_FLOAT_ROUNDING_MODE = 5 -} hsa_code_object_info_t; -typedef struct hsa_barrier_or_packet_s { - uint16_t header; - uint16_t reserved0; - uint32_t reserved1; - hsa_signal_t dep_signal[5]; - uint64_t reserved2; - hsa_signal_t completion_signal; -} hsa_barrier_or_packet_t; -typedef enum { - HSA_SYMBOL_KIND_LINKAGE_MODULE = 0, - HSA_SYMBOL_KIND_LINKAGE_PROGRAM = 1, -} hsa_symbol_kind_linkage_t; -hsa_status_t hsa_executable_validate(hsa_executable_t executable, - uint32_t *result); -uint64_t hsa_queue_add_write_index_acq_rel(const hsa_queue_t *queue, - uint64_t value); - -uint64_t hsa_queue_add_write_index_acquire(const hsa_queue_t *queue, - uint64_t value); - -uint64_t hsa_queue_add_write_index_relaxed(const hsa_queue_t *queue, - uint64_t value); - -uint64_t hsa_queue_add_write_index_release(const hsa_queue_t *queue, - uint64_t value); -hsa_status_t hsa_shut_down(); -void hsa_signal_add_acq_rel(hsa_signal_t signal, hsa_signal_value_t value); - -void hsa_signal_add_acquire(hsa_signal_t signal, hsa_signal_value_t value); - -void hsa_signal_add_relaxed(hsa_signal_t signal, hsa_signal_value_t value); - -void hsa_signal_add_release(hsa_signal_t signal, hsa_signal_value_t value); -hsa_status_t hsa_executable_readonly_variable_define( - hsa_executable_t executable, hsa_agent_t agent, const char *variable_name, - void *address); -hsa_status_t hsa_agent_extension_supported(uint16_t extension, - hsa_agent_t agent, - uint16_t version_major, - uint16_t version_minor, - bool *result); -hsa_signal_value_t hsa_signal_load_acquire(hsa_signal_t signal); - -hsa_signal_value_t hsa_signal_load_relaxed(hsa_signal_t signal); -hsa_status_t hsa_executable_get_info(hsa_executable_t executable, - hsa_executable_info_t attribute, - void *value); -hsa_status_t hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent, - void *data), - void *data); -void hsa_signal_subtract_acq_rel(hsa_signal_t signal, hsa_signal_value_t value); - -void hsa_signal_subtract_acquire(hsa_signal_t signal, hsa_signal_value_t value); - -void hsa_signal_subtract_relaxed(hsa_signal_t signal, hsa_signal_value_t value); - -void hsa_signal_subtract_release(hsa_signal_t signal, hsa_signal_value_t value); -hsa_status_t -hsa_executable_symbol_get_info(hsa_executable_symbol_t executable_symbol, - hsa_executable_symbol_info_t attribute, - void *value); -void hsa_signal_xor_acq_rel(hsa_signal_t signal, hsa_signal_value_t value); - -void hsa_signal_xor_acquire(hsa_signal_t signal, hsa_signal_value_t value); - -void hsa_signal_xor_relaxed(hsa_signal_t signal, hsa_signal_value_t value); - -void hsa_signal_xor_release(hsa_signal_t signal, hsa_signal_value_t value); -hsa_status_t hsa_code_object_get_info(hsa_code_object_t code_object, - hsa_code_object_info_t attribute, - void *value); -hsa_status_t hsa_code_object_deserialize(void *serialized_code_object, - size_t serialized_code_object_size, - const char *options, - hsa_code_object_t *code_object); -hsa_status_t hsa_status_string(hsa_status_t status, const char **status_string); -hsa_status_t hsa_code_object_get_symbol(hsa_code_object_t code_object, - const char *symbol_name, - hsa_code_symbol_t *symbol); -void hsa_signal_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value); - -void hsa_signal_store_release(hsa_signal_t signal, hsa_signal_value_t value); -hsa_status_t hsa_signal_destroy(hsa_signal_t signal); -hsa_status_t hsa_system_get_extension_table(uint16_t extension, - uint16_t version_major, - uint16_t version_minor, - void *table); -hsa_status_t hsa_agent_iterate_regions( - hsa_agent_t agent, - hsa_status_t (*callback)(hsa_region_t region, void *data), void *data); -hsa_status_t hsa_executable_agent_global_variable_define( - hsa_executable_t executable, hsa_agent_t agent, const char *variable_name, - void *address); -hsa_status_t hsa_queue_create(hsa_agent_t agent, uint32_t size, - hsa_queue_type_t type, - void (*callback)(hsa_status_t status, - hsa_queue_t *source, void *data), - void *data, uint32_t private_segment_size, - uint32_t group_segment_size, hsa_queue_t **queue); -hsa_status_t hsa_isa_compatible(hsa_isa_t code_object_isa, hsa_isa_t agent_isa, - bool *result); -hsa_status_t hsa_code_object_serialize( - hsa_code_object_t code_object, - hsa_status_t (*alloc_callback)(size_t size, hsa_callback_data_t data, - void **address), - hsa_callback_data_t callback_data, const char *options, - void **serialized_code_object, size_t *serialized_code_object_size); -hsa_status_t hsa_region_get_info(hsa_region_t region, - hsa_region_info_t attribute, void *value); -hsa_status_t hsa_executable_freeze(hsa_extension_t executable, - const char *options); -hsa_status_t hsa_system_extension_supported(uint16_t extension, - uint16_t version_major, - uint16_t version_minor, - bool *result); -hsa_signal_value_t hsa_signal_wait_acquire(hsa_signal_t signal, - hsa_signal_condition_t condition, - hsa_signal_value_t compare_value, - uint64_t timeout_hint, - hsa_wait_state_t wait_state_hint); - -hsa_signal_value_t hsa_signal_wait_relaxed(hsa_signal_t signal, - hsa_signal_condition_t condition, - hsa_signal_value_t compare_value, - uint64_t timeout_hint, - hsa_wait_state_t wait_state_hint); -hsa_status_t hsa_memory_copy(void *dst, const void *src, size_t size); -hsa_status_t hsa_memory_free(void *ptr); -hsa_status_t hsa_queue_destroy(hsa_queue_t *queue); -hsa_status_t hsa_isa_from_name(const char *name, hsa_isa_t *isa); -hsa_status_t hsa_isa_get_info(hsa_isa_t isa, hsa_isa_info_t attribute, - uint32_t index, void *value); -hsa_status_t hsa_signal_create(hsa_signal_value_t initial_value, - uint32_t num_consumers, - const hsa_agent_t *consumers, - hsa_signal_t *signal); -hsa_status_t hsa_code_symbol_get_info(hsa_code_symbol_t code_symbol, - hsa_code_symbol_info_t attribute, - void *value); -hsa_signal_value_t hsa_signal_cas_acq_rel(hsa_signal_t signal, - hsa_signal_value_t expected, - hsa_signal_value_t value); - -hsa_signal_value_t hsa_signal_cas_acquire(hsa_signal_t signal, - hsa_signal_value_t expected, - hsa_signal_value_t value); - -hsa_signal_value_t hsa_signal_cas_relaxed(hsa_signal_t signal, - hsa_signal_value_t expected, - hsa_signal_value_t value); - -hsa_signal_value_t hsa_signal_cas_release(hsa_signal_t signal, - hsa_signal_value_t expected, - hsa_signal_value_t value); -hsa_status_t hsa_code_object_iterate_symbols( - hsa_code_object_t code_object, - hsa_status_t (*callback)(hsa_code_object_t code_object, - hsa_code_symbol_t symbol, void *data), - void *data); -void hsa_queue_store_read_index_relaxed(const hsa_queue_t *queue, - uint64_t value); - -void hsa_queue_store_read_index_release(const hsa_queue_t *queue, - uint64_t value); -hsa_status_t hsa_memory_assign_agent(void *ptr, hsa_agent_t agent, - hsa_access_permission_t access); -hsa_status_t hsa_queue_inactivate(hsa_queue_t *queue); -hsa_status_t hsa_executable_get_symbol(hsa_executable_t executable, - const char *module_name, - const char *symbol_name, - hsa_agent_t agent, - int32_t call_convention, - hsa_executable_symbol_t *symbol); -uint64_t hsa_queue_cas_write_index_acq_rel(const hsa_queue_t *queue, - uint64_t expected, uint64_t value); - -uint64_t hsa_queue_cas_write_index_acquire(const hsa_queue_t *queue, - uint64_t expected, uint64_t value); - -uint64_t hsa_queue_cas_write_index_relaxed(const hsa_queue_t *queue, - uint64_t expected, uint64_t value); - -uint64_t hsa_queue_cas_write_index_release(const hsa_queue_t *queue, - uint64_t expected, uint64_t value); -void hsa_signal_and_acq_rel(hsa_signal_t signal, hsa_signal_value_t value); - -void hsa_signal_and_acquire(hsa_signal_t signal, hsa_signal_value_t value); - -void hsa_signal_and_relaxed(hsa_signal_t signal, hsa_signal_value_t value); - -void hsa_signal_and_release(hsa_signal_t signal, hsa_signal_value_t value); -uint64_t hsa_queue_load_read_index_acquire(const hsa_queue_t *queue); - -uint64_t hsa_queue_load_read_index_relaxed(const hsa_queue_t *queue); -hsa_status_t hsa_executable_load_code_object(hsa_executable_t executable, - hsa_agent_t agent, - hsa_code_object_t code_object, - const char *options); -uint64_t hsa_queue_load_write_index_acquire(const hsa_queue_t *queue); - -uint64_t hsa_queue_load_write_index_relaxed(const hsa_queue_t *queue); -hsa_status_t hsa_agent_get_exception_policies(hsa_agent_t agent, - hsa_profile_t profile, - uint16_t *mask); -hsa_status_t hsa_memory_deregister(void *ptr, size_t size); -void hsa_signal_or_acq_rel(hsa_signal_t signal, hsa_signal_value_t value); - -void hsa_signal_or_acquire(hsa_signal_t signal, hsa_signal_value_t value); - -void hsa_signal_or_relaxed(hsa_signal_t signal, hsa_signal_value_t value); - -void hsa_signal_or_release(hsa_signal_t signal, hsa_signal_value_t value); -hsa_status_t hsa_soft_queue_create(hsa_region_t region, uint32_t size, - hsa_queue_type_t type, uint32_t features, - hsa_signal_t doorbell_signal, - hsa_queue_t **queue); -hsa_status_t hsa_executable_iterate_symbols( - hsa_executable_t executable, - hsa_status_t (*callback)(hsa_executable_t executable, - hsa_executable_symbol_t symbol, void *data), - void *data); -hsa_status_t hsa_memory_register(void *ptr, size_t size); -void hsa_queue_store_write_index_relaxed(const hsa_queue_t *queue, - uint64_t value); - -void hsa_queue_store_write_index_release(const hsa_queue_t *queue, - uint64_t value); -hsa_status_t hsa_executable_global_variable_define(hsa_executable_t executable, - const char *variable_name, - void *address); -hsa_status_t hsa_executable_destroy(hsa_executable_t executable); -hsa_status_t hsa_code_object_destroy(hsa_code_object_t code_object); -hsa_status_t hsa_memory_allocate(hsa_region_t region, size_t size, void **ptr); -hsa_signal_value_t hsa_signal_exchange_acq_rel(hsa_signal_t signal, - hsa_signal_value_t value); - -hsa_signal_value_t hsa_signal_exchange_acquire(hsa_signal_t signal, - hsa_signal_value_t value); - -hsa_signal_value_t hsa_signal_exchange_relaxed(hsa_signal_t signal, - hsa_signal_value_t value); - -hsa_signal_value_t hsa_signal_exchange_release(hsa_signal_t signal, - hsa_signal_value_t value); -hsa_status_t hsa_agent_get_info(hsa_agent_t agent, hsa_agent_info_t attribute, - void *value); -hsa_status_t hsa_init(); -hsa_status_t hsa_system_get_info(hsa_system_info_t attribute, void *value); -hsa_status_t hsa_executable_create(hsa_profile_t profile, - hsa_executable_state_t executable_state, - const char *options, - hsa_executable_t *executable); - -#endif /* _HSA_H */ diff --git a/libgomp/plugin/plugin-hsa.c b/libgomp/plugin/plugin-hsa.c index b829c8c..469a90d 100644 --- a/libgomp/plugin/plugin-hsa.c +++ b/libgomp/plugin/plugin-hsa.c @@ -34,7 +34,7 @@ #include #include #include -#include +#include "hsa-interface.h" #include #include #include "libgomp-plugin.h" diff --git a/libhsail-rt/include/internal/fibers.h b/libhsail-rt/include/internal/fibers.h index c0682e1..033146c 100644 --- a/libhsail-rt/include/internal/fibers.h +++ b/libhsail-rt/include/internal/fibers.h @@ -1,23 +1,27 @@ /* fibers.h -- an extremely simple lightweight thread (fiber) implementation - Copyright (C) 2016 Free Software Foundation, Inc. + Copyright (C) 2015-2016 Free Software Foundation, Inc. Contributed by Pekka Jaaskelainen for General Processor Tech. - This file is part of GCC. - - GCC is free software; you can redistribute it and/or modify it under - the terms of the GNU General Public License as published by the Free - Software Foundation; either version 3, or (at your option) any later - version. - - GCC is distributed in the hope that it will be useful, but WITHOUT ANY - WARRANTY; without even the implied warranty of MERCHANTABILITY or - FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License - for more details. - - You should have received a copy of the GNU General Public License - along with GCC; see the file COPYING3. If not see - . */ + 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. +*/ #ifndef PHSA_RT_FIBERS_H #define PHSA_RT_FIBERS_H diff --git a/libhsail-rt/rt/fibers.c b/libhsail-rt/rt/fibers.c index 6de5dd2..a3056a9 100644 --- a/libhsail-rt/rt/fibers.c +++ b/libhsail-rt/rt/fibers.c @@ -3,21 +3,29 @@ Contributed by Pekka Jaaskelainen for General Processor Tech. - This file is part of GCC. - - GCC is free software; you can redistribute it and/or modify it under - the terms of the GNU General Public License as published by the Free - Software Foundation; either version 3, or (at your option) any later - version. - - GCC is distributed in the hope that it will be useful, but WITHOUT ANY - WARRANTY; without even the implied warranty of MERCHANTABILITY or - FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License - for more details. + Copyright (C) 2015-2016 Free Software Foundation, Inc. + Contributed by Pekka Jaaskelainen + for General Processor Tech. - You should have received a copy of the GNU General Public License - along with GCC; see the file COPYING3. If not see - . */ + 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 #include