From d0951a9c205cfdb7fdcd9746bd7267fac1009f75 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 18 Apr 2019 07:35:06 +0530 Subject: [PATCH 01/17] Fix D2DAsync test [ROCm/hip commit: e5e7651a4aeee3b9a4d6a9010105df82d735ab19] --- projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp index 2d1a591e73..13eca9c24d 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp @@ -58,8 +58,8 @@ int main() { HIPCHECK(hipDeviceSynchronize()); HipTest::checkVectorADD(A_h, B_h, C_h, N); - HIPCHECK(hipStreamCreate(&s)); HIPCHECK(hipSetDevice(1)); + HIPCHECK(hipStreamCreate(&s)); HIPCHECK(hipMemcpyDtoDAsync((hipDeviceptr_t)X_d, (hipDeviceptr_t)A_d, Nbytes, s)); HIPCHECK(hipMemcpyDtoDAsync((hipDeviceptr_t)Y_d, (hipDeviceptr_t)B_d, Nbytes, s)); From 53b232fb2ab2ea1e0cf615f0075d174790806e58 Mon Sep 17 00:00:00 2001 From: Konstantin Pyzhov Date: Thu, 18 Apr 2019 13:07:24 -0400 Subject: [PATCH 02/17] Fix for __ffsll() device functions. [ROCm/hip commit: b7bd29924a6bdecb873386cc95ae77f2a15c15c7] --- projects/hip/include/hip/hcc_detail/device_functions.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/device_functions.h b/projects/hip/include/hip/hcc_detail/device_functions.h index b88df2e168..7cb9406856 100644 --- a/projects/hip/include/hip/hcc_detail/device_functions.h +++ b/projects/hip/include/hip/hcc_detail/device_functions.h @@ -58,7 +58,7 @@ __device__ static inline unsigned int __ffs(unsigned int input) { } __device__ static inline unsigned int __ffsll(unsigned long long int input) { - return ( input == 0 ? -1 : __builtin_ctzl(input) ) + 1; + return ( input == 0 ? -1 : __builtin_ctzll(input) ) + 1; } __device__ static inline unsigned int __ffs(int input) { @@ -66,7 +66,7 @@ __device__ static inline unsigned int __ffs(int input) { } __device__ static inline unsigned int __ffsll(long long int input) { - return ( input == 0 ? -1 : __builtin_ctzl(input) ) + 1; + return ( input == 0 ? -1 : __builtin_ctzll(input) ) + 1; } __device__ static inline unsigned int __brev(unsigned int input) { From d4bce6c019fd88a022d78c438629e52bd0d0786e Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Thu, 18 Apr 2019 16:18:31 -0400 Subject: [PATCH 03/17] Fix missing arg in HIP_INIT_API [ROCm/hip commit: bb5c620b139b3f740d2e33e95759c03853339d4f] --- projects/hip/src/hip_module.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index 05daca6055..e1a3b1ff65 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -474,7 +474,7 @@ hipError_t hipModuleGetFunction(hipFunction_t* hfunc, hipModule_t hmod, const ch // Get kernel for the given hsa agent. Internal use only. hipError_t hipModuleGetFunctionEx(hipFunction_t* hfunc, hipModule_t hmod, const char* name, hsa_agent_t *agent) { - HIP_INIT_API(hipModuleGetFunctionEx, hfunc, hmod, name); + HIP_INIT_API(hipModuleGetFunctionEx, hfunc, hmod, name, agent); return ihipLogStatus(ihipModuleGetFunction(hfunc, hmod, name, agent)); } From 2e0c67bbd3b987e62aee7334824cc6bffa3ba9dd Mon Sep 17 00:00:00 2001 From: Konstantin Pyzhov Date: Fri, 19 Apr 2019 08:53:22 -0400 Subject: [PATCH 04/17] Fix for __popcll() device function implementation. [ROCm/hip commit: beadaab661798ae0b44fb8cc7810725402474c91] --- projects/hip/include/hip/hcc_detail/device_functions.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hcc_detail/device_functions.h b/projects/hip/include/hip/hcc_detail/device_functions.h index 7cb9406856..044ca1e634 100644 --- a/projects/hip/include/hip/hcc_detail/device_functions.h +++ b/projects/hip/include/hip/hcc_detail/device_functions.h @@ -42,7 +42,7 @@ __device__ static inline unsigned int __popc(unsigned int input) { return __builtin_popcount(input); } __device__ static inline unsigned int __popcll(unsigned long long int input) { - return __builtin_popcountl(input); + return __builtin_popcountll(input); } __device__ static inline int __clz(int input) { From be49602b292e039db21b56750e4d9752f71f5109 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 23 Apr 2019 03:22:25 +0530 Subject: [PATCH 05/17] Fix hipPointerGetAttributes for NVCC [ROCm/hip commit: 69a3d6b72ad240561e76103b1de396ec14c154e5] --- projects/hip/include/hip/nvcc_detail/hip_runtime_api.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index 4d84cfbc51..e5c8488079 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -866,7 +866,7 @@ inline static hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(int* numBl return hipCUDAErrorTohipError(cerror); } -inline static hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, void* ptr) { +inline static hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void* ptr) { struct cudaPointerAttributes cPA; hipError_t err = hipCUDAErrorTohipError(cudaPointerGetAttributes(&cPA, ptr)); if (err == hipSuccess) { From 17474d2b5868173acfa2144a3220ded0769d459f Mon Sep 17 00:00:00 2001 From: Evgeny Date: Mon, 22 Apr 2019 21:14:18 -0500 Subject: [PATCH 06/17] hip_prof_api.h include under __cplusplus [ROCm/hip commit: af3f3ccb2b4fc4b519dcdd363f295a26962a5ac3] --- projects/hip/include/hip/hcc_detail/hip_runtime_api.h | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index ccfe7aefc3..bfc9938171 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -2966,7 +2966,9 @@ hipError_t hipLaunchByPtr(const void* func); } /* extern "c" */ #endif +#ifdef __cplusplus #include +#endif #ifdef __cplusplus extern "C" { From 27aed8c1728cf461c4900d12960b9181c6dbf6e9 Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Mon, 22 Apr 2019 23:13:02 -0400 Subject: [PATCH 07/17] [hip] Fix use of `__HIP_CLANG_ONLY__` in `hip_ldg.h`. - Check its value instead of whether it's defined or not. [ROCm/hip commit: 6fb07acc8ce392f03e87abf204d2dcb0b085e5ff] --- projects/hip/include/hip/hcc_detail/hip_ldg.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_ldg.h b/projects/hip/include/hip/hcc_detail/hip_ldg.h index a5b80b0a1b..ab86955cd6 100644 --- a/projects/hip/include/hip/hcc_detail/hip_ldg.h +++ b/projects/hip/include/hip/hcc_detail/hip_ldg.h @@ -24,7 +24,7 @@ THE SOFTWARE. #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_LDG_H #if defined(__HCC_OR_HIP_CLANG__) -#if __hcc_workweek__ >= 16164 || defined(__HIP_CLANG_ONLY__) +#if __hcc_workweek__ >= 16164 || __HIP_CLANG_ONLY__ #include "hip_vector_types.h" #include "host_defines.h" @@ -96,7 +96,7 @@ __device__ inline static double __ldg(const double* ptr) { return ptr[0]; } __device__ inline static double2 __ldg(const double2* ptr) { return ptr[0]; } -#endif // __hcc_workweek__ || defined(__HIP_CLANG_ONLY__) +#endif // __hcc_workweek__ || __HIP_CLANG_ONLY__ #endif // defined(__HCC_OR_HIP_CLANG__) From 6eec9d4ac82ca0936577c00c446090c450b24f70 Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Mon, 22 Apr 2019 23:21:00 -0400 Subject: [PATCH 08/17] [hip] Fix including of hip_fp16.h - Separate the definition of `__HCC_OR_HIP_CLANG__`, `__HCC_ONLY__`, and `__HIP_CLANG_ONLY__` into hip_common.h so that it could be included in hip_fp16.h, which may be included separately in app. [ROCm/hip commit: dc0d7bd5ce586c9dc597b9c4c9ed99b8d2a863f1] --- .../hip/include/hip/hcc_detail/hip_common.h | 40 +++++++++++++++++++ .../hip/include/hip/hcc_detail/hip_fp16.h | 7 ++++ .../hip/include/hip/hcc_detail/hip_runtime.h | 14 +------ 3 files changed, 48 insertions(+), 13 deletions(-) create mode 100644 projects/hip/include/hip/hcc_detail/hip_common.h diff --git a/projects/hip/include/hip/hcc_detail/hip_common.h b/projects/hip/include/hip/hcc_detail/hip_common.h new file mode 100644 index 0000000000..2e2abace0c --- /dev/null +++ b/projects/hip/include/hip/hcc_detail/hip_common.h @@ -0,0 +1,40 @@ +/* +Copyright (c) 2019 - present 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. +*/ + +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMMON_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMMON_H + +#if defined(__HCC__) +#define __HCC_OR_HIP_CLANG__ 1 +#define __HCC_ONLY__ 1 +#define __HIP_CLANG_ONLY__ 0 +#elif defined(__clang__) && defined(__HIP__) +#define __HCC_OR_HIP_CLANG__ 1 +#define __HCC_ONLY__ 0 +#define __HIP_CLANG_ONLY__ 1 +#else +#define __HCC_OR_HIP_CLANG__ 0 +#define __HCC_ONLY__ 0 +#define __HIP_CLANG_ONLY__ 0 +#endif + +#endif // HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMMON_H diff --git a/projects/hip/include/hip/hcc_detail/hip_fp16.h b/projects/hip/include/hip/hcc_detail/hip_fp16.h index 849b7278d8..93ede207c2 100644 --- a/projects/hip/include/hip/hcc_detail/hip_fp16.h +++ b/projects/hip/include/hip/hcc_detail/hip_fp16.h @@ -21,6 +21,11 @@ THE SOFTWARE. */ #pragma once +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H + +#include + #include "hip/hcc_detail/host_defines.h" #include #if defined(__cplusplus) @@ -1643,3 +1648,5 @@ THE SOFTWARE. #elif defined(__GNUC__) #include "hip_fp16_gcc.h" #endif // !defined(__clang__) && defined(__GNUC__) + +#endif // HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h index 212502ea5e..e6f03ab797 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h @@ -29,19 +29,7 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_H #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_H -#if defined(__HCC__) -#define __HCC_OR_HIP_CLANG__ 1 -#define __HCC_ONLY__ 1 -#define __HIP_CLANG_ONLY__ 0 -#elif defined(__clang__) && defined(__HIP__) -#define __HCC_OR_HIP_CLANG__ 1 -#define __HCC_ONLY__ 0 -#define __HIP_CLANG_ONLY__ 1 -#else -#define __HCC_OR_HIP_CLANG__ 0 -#define __HCC_ONLY__ 0 -#define __HIP_CLANG_ONLY__ 0 -#endif +#include //--- // Top part of file can be compiled with any compiler From a2b6b1e7e95b959efa69d9b996f800342c297abb Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 23 Apr 2019 17:55:47 +0300 Subject: [PATCH 09/17] [HIPIFY][hipify-perl] Formatting [ROCm/hip commit: defc6f81552200b8acae0a33048576bfe2a7d443] --- projects/hip/bin/hipify-perl | 265 ++++++++--------------------------- 1 file changed, 62 insertions(+), 203 deletions(-) diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index 8bb7bdf31f..3456f9fa41 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -1,6 +1,6 @@ #!/usr/bin/perl -w ## -# Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +# Copyright (c) 2015-present 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 @@ -25,16 +25,16 @@ use Getopt::Long; my $warn_whitelist =""; GetOptions( - "print-stats" => \$print_stats # print the command-line, like a header. - , "count-conversions" => \$count_conversions # count conversions. - , "quiet-warnings" => \$quiet_warnings # don't print warnings on unknown CUDA functions. + "print-stats" => \$print_stats # print the command-line, like a header. + , "count-conversions" => \$count_conversions # count conversions. + , "quiet-warnings" => \$quiet_warnings # don't print warnings on unknown CUDA functions. , "warn-whitelist=s"=> \$warn_whitelist - , "no-translate-builtins" => \$no_translate_builtins # don't translate math functions. - , "no-translate-textures" => \$no_translate_textures - , "no-output" => \$no_output # don't write any translated output to stdout. - , "inplace" => \$inplace # modify input file inplace, replacing input with hipified output, save backup in ".prehip" file. - # If .prehip file exists, use that as input to hip. - , "n" => \$n # combination of print_stats + no-output. + , "no-translate-builtins" => \$no_translate_builtins # don't translate math functions. + , "no-translate-textures" => \$no_translate_textures # don't translate texture functions. + , "no-output" => \$no_output # don't write any translated output to stdout. + , "inplace" => \$inplace # modify input file inplace, replacing input with hipified output, save backup in ".prehip" file. + # If .prehip file exists, use that as input to hip. + , "n" => \$n # combination of print_stats + no-output. ); $print_stats = 1 if $n; @@ -47,20 +47,17 @@ $no_output = 1 if $n; ,"cudaStatus" ,"cudaDevice" ,"cudaDevice_t" - ,"cudaIDs" ,"cudaGridDim" ,"cudaDimGrid" ,"cudaDimBlock" ,"cudaDeviceId" ,"cudaDevices", - ,"cudaGradOutput", ,"cudaInput", ,"cudaOutput", ,"cudaGradInput", ,"cudaIndices", - ,"cudaColorSpinorField" ,"cudaGaugeField" ,"cudaMom" @@ -77,126 +74,96 @@ $no_output = 1 if $n; ); #print "WW=@warn_whitelist\n"; - # Allow users to add their own functions. push (@warn_whitelist, split(',',$warn_whitelist)); - #--- #Stats tracking code: @statNames = ("dev", "mem", "kern", 'coord_func', "math_func", "special_func", "stream", "event", "err", "def", "tex", "extern_shared", "other"); - #--- #Compute total of all individual counts: sub totalStats { my %count = %{ shift() }; - my $total = 0; foreach $key (keys %count) { $total += $count{$key}; } - return $total; }; #--- sub printStats { - my $label = shift(); + my $label = shift(); my @statNames = @{ shift() }; - my %counts = %{ shift() }; - my $warnings = shift(); + my %counts = %{ shift() }; + my $warnings = shift(); my $loc = shift(); - - my $total = totalStats(\%counts); - + my $total = totalStats(\%counts); printf STDERR "%s %d CUDA->HIP refs( ", $label, $total; - foreach $stat (@statNames) { printf STDERR "%s:%d ", $stat, $counts{$stat}; } - printf STDERR ") warn:%d LOC:%d", $warnings, $loc; } - #--- -# Add adder stats to dest. Used to add stats for current file to a running total for all files: +# Add adder stats to dest. Used to add stats for current file to a running total for all files: sub addStats { - my $dest_ref = shift() ; + my $dest_ref = shift(); my %adder = %{ shift() }; - foreach $key (keys %adder) { $dest_ref->{$key} += $adder{$key}; #printf ("D{$key} += %d => %d\n", $adder{$key}, $dest{$key}); } } - #--- sub clearStats { my $dest_ref = shift() ; my @statNames = @{ shift() }; - foreach $stat (@statNames) { $dest_ref->{$stat} = 0; } } - - - +#--- # count of transforms in all files: my %tt; clearStats(\%tt, \@statNames); - - - $Twarnings = 0; $TlineCount = 0; -my %TwarningTags ; -my %Tkernels ; - +my %TwarningTags; +my %Tkernels; my $fileCount = @ARGV; my $fileName = ""; while (@ARGV) { $fileName=shift (@ARGV); - - if ($inplace) { my $file_prehip = "$fileName" . ".prehip"; - my $infile; - my $outfile; - if (-e $file_prehip) { - $infile = $file_prehip; - $outfile = $fileName; - } else { - system ("cp $fileName $file_prehip"); - $infile = $file_prehip; - $outfile = $fileName; - } - + my $infile; + my $outfile; + if (-e $file_prehip) { + $infile = $file_prehip; + $outfile = $fileName; + } else { + system ("cp $fileName $file_prehip"); + $infile = $file_prehip; + $outfile = $fileName; + } open(INFILE,"<", $infile) or die "error: could not open $infile"; open(OUTFILE,">", $outfile) or die "error: could not open $outfile"; $OUTFILE = OUTFILE; - } else { open(INFILE,"<", $fileName) or die "error: could not open $fileName"; $OUTFILE = STDOUT; } - - - - - # Note : \b is used in perl to indicate the start of a word - typically that is what we want in this case: - # - # count of transforms in this file, init to 0 here: my %ft; clearStats(\%ft, \@statNames); @@ -206,44 +173,34 @@ while (@ARGV) { my $warningsCublas = 0; my $warningsCurand = 0; my %warningTags; # hash with counts of particular unknown keywords. - my $lineCount = 0; - undef $/; # Read whole file at once, so we can match newlines. while () { - #-------- # Compiler Defines # __CUDACC__ is set by NVCC to indicate it is treating the input file as CUDA code (as opposed to host) # Typically we want any code treated as CUDA code to be treated as accelerator code by Kalmar too # __HIPCC__ will set KALMARCC $ft{'def'} += s/\b__CUDACC__\b/__HIPCC__/g; - # __CUDA_ARCH is often used to detect when a function or kernel is being compiled for the device. # Don't automaticall convert this - likely these will need special attention with HIP_ARCH_HAS_* macros #$ft{'def'} += s/\b__CUDA_ARCH__\b/__HIP_ARCH__/g; - - - #-------- #Includes: $countIncludes += s/(\s*#\s*include\s+)[<"]cuda_runtime\.h[>"]/$1/; $countIncludes += s/(\s*#\s*include\s+)[<"]cuda_runtime_api\.h[>"]/$1/; $countIncludes += s/(\s*#\s*include\s+)[<"]cuda_fp16\.h[>"]/$1/; - - #-------- # Error codes and return types: $ft{'err'} += s/\bcudaError_t\b/hipError_t/g; $ft{'err'} += s/\bcudaError\b/hipError_t/g; $ft{'err'} += s/\bcudaSuccess\b/hipSuccess/g; - $ft{'err'} += s/\bcudaErrorUnknown\b/hipErrorUnknown/g; $ft{'err'} += s/\bcudaErrorMemoryAllocation\b/hipErrorMemoryAllocation/g; $ft{'err'} += s/\bcudaErrorMemoryFree\b/hipErrorMemoryFree/g; $ft{'err'} += s/\bcudaErrorUnknownSymbol\b/hipErrorUnknownSymbol/g; - $ft{'err'} += s/\bcudaErrorInvalidSymbol\b/hipErrorInvalidSymbol/g; + $ft{'err'} += s/\bcudaErrorInvalidSymbol\b/hipErrorInvalidSymbol/g; $ft{'err'} += s/\bcudaErrorOutOfResources\b/hipErrorOutOfResources/g; $ft{'err'} += s/\bcudaErrorInvalidValue\b/hipErrorInvalidValue/g; $ft{'err'} += s/\bcudaErrorInvalidResourceHandle\b/hipErrorInvalidResourceHandle/g; @@ -253,15 +210,12 @@ while (@ARGV) { $ft{'err'} += s/\bcudaErrorNotReady\b/hipErrorNotReady/g; $ft{'err'} += s/\bcudaErrorUnknown\b/hipErrorUnknown/g; $ft{'err'} += s/\bcudaErrorPeerAccessAlreadyEnabled\b/hipErrorPeerAccessAlreadyEnabled/g; - - # error APIs: + #-------- + # error APIs: $ft{'err'} += s/\bcudaGetLastError\b/hipGetLastError/g; $ft{'err'} += s/\bcudaPeekAtLastError\b/hipPeekAtLastError/g; $ft{'err'} += s/\bcudaGetErrorName\b/hipGetErrorName/g; $ft{'err'} += s/\bcudaGetErrorString\b/hipGetErrorString/g; - - - #-------- # Memcpy $ft{'mem'} += s/\bcudaMemcpy\b/hipMemcpy/g; @@ -272,48 +226,39 @@ while (@ARGV) { $ft{'mem'} += s/\bcudaMemcpyDefault\b/hipMemcpyDefault/g; $ft{'mem'} += s/\bcudaMemcpyToSymbol\s*\(\s*(\w+)\b/hipMemcpyToSymbol\(HIP_SYMBOL\($1\)/g; $ft{'mem'} += s/\bcudaMemcpyFromSymbol\s*\(\s*(.+?)\s*,\s*(.+?)\b/hipMemcpyFromSymbol\($1, HIP_SYMBOL\($2\)/g; - $ft{'mem'} += s/\bcudaMemset\b/hipMemset/g; $ft{'mem'} += s/\bcudaMemsetAsync\b/hipMemsetAsync/g; - $ft{'mem'} += s/\bcudaMemcpyAsync\b/hipMemcpyAsync/g; - $ft{'mem'} += s/\bcudaMemGetInfo\b/hipMemGetInfo/g; - $ft{'mem'} += s/\bcudaMemcpyKind\b/hipMemcpyKind/g; - $ft{'mem'} += s/\bcudaPointerAttributes\b/hipPointerAttribute_t/g; $ft{'mem'} += s/\bcudaPointerGetAttributes\b/hipPointerGetAttributes/g; - $ft{'mem'} += s/\bcudaMemcpy2D\b/hipMemcpy2D/g; $ft{'mem'} += s/\bcudaMemcpy2DToArray\b/hipMemcpy2DToArray/g; $ft{'mem'} += s/\bcudaMemcpyToArray\b/hipMemcpyToArray/g; - $ft{'mem'} += s/\bcudaGetSymbolAddress\s*\(\s*(.+?)\s*,\s*(.+?)\b/hipGetSymbolAddress\($1, HIP_SYMBOL\($2\)/g; $ft{'mem'} += s/\bcudaGetSymbolSize\s*\(\s*&(\w+)\s*,\s*(.+?)\b/hipGetSymbolSize(&$1, HIP_SYMBOL\($2\)/g; - #-------- # Memory management: $ft{'mem'} += s/\bcudaMalloc\b/hipMalloc/g; - $ft{'mem'} += s/\bcudaMallocHost\b/hipHostMalloc/g; # note conversion to standard hipHost* naming convention + # note conversion to standard hipHost* naming convention + $ft{'mem'} += s/\bcudaMallocHost\b/hipHostMalloc/g; $ft{'mem'} += s/\bcudaFree\b/hipFree/g; - $ft{'mem'} += s/\bcudaFreeHost\b/hipHostFree/g; # note conversion to standard hipHost* naming convention - $ft{'mem'} += s/\bcudaHostAlloc\b/hipHostMalloc/g; - $ft{'mem'} += s/\bcudaHostGetDevicePointer\b/hipHostGetDevicePointer/g; - $ft{'mem'} += s/\bcudaHostAllocDefault\b/hipHostMallocDefault/g; - $ft{'mem'} += s/\bcudaHostAllocPortable\b/hipHostMallocPortable/g; - $ft{'mem'} += s/\bcudaHostAllocMapped\b/hipHostMallocMapped/g; - $ft{'mem'} += s/\bcudaHostAllocWriteCombined\b/hipHostMallocWriteCombined/g; - $ft{'mem'} += s/\bcudaHostRegisterMapped\b/hipHostRegisterMapped/g; + # note conversion to standard hipHost* naming convention + $ft{'mem'} += s/\bcudaFreeHost\b/hipHostFree/g; + $ft{'mem'} += s/\bcudaHostAlloc\b/hipHostMalloc/g; + $ft{'mem'} += s/\bcudaHostGetDevicePointer\b/hipHostGetDevicePointer/g; + $ft{'mem'} += s/\bcudaHostAllocDefault\b/hipHostMallocDefault/g; + $ft{'mem'} += s/\bcudaHostAllocPortable\b/hipHostMallocPortable/g; + $ft{'mem'} += s/\bcudaHostAllocMapped\b/hipHostMallocMapped/g; + $ft{'mem'} += s/\bcudaHostAllocWriteCombined\b/hipHostMallocWriteCombined/g; + $ft{'mem'} += s/\bcudaHostRegisterMapped\b/hipHostRegisterMapped/g; $ft{'mem'} += s/\bcudaHostRegister\b/hipHostRegister/g; $ft{'mem'} += s/\bcudaHostUnregister\b/hipHostUnregister/g; $ft{'mem'} += s/\bcudaHostGetDevicePointer\b/hipHostGetDevicePointer/g; - $ft{'mem'} += s/\bcudaMallocArray\b/hipMallocArray/g; $ft{'mem'} += s/\bcudaFreeArray\b/hipFreeArray/g; $ft{'mem'} += s/\bcudaMallocPitch\b/hipMallocPitch/g; - - #-------- # Events $ft{'event'} += s/\bcudaEvent_t\b/hipEvent_t/g; @@ -325,7 +270,6 @@ while (@ARGV) { $ft{'event'} += s/\bcudaEventSynchronize\b/hipEventSynchronize/g; $ft{'event'} += s/\bcudaEventDisableTiming\b/hipEventDisableTiming/g; $ft{'event'} += s/\bcudaEventQuery\b/hipEventQuery/g; - #-------- # Streams $ft{'stream'} += s/\bcudaStream_t\b/hipStream_t/g; @@ -336,23 +280,22 @@ while (@ARGV) { $ft{'stream'} += s/\bcudaStreamSynchronize\b/hipStreamSynchronize/g; $ft{'stream'} += s/\bcudaStreamDefault\b/hipStreamDefault/g; $ft{'stream'} += s/\bcudaStreamNonBlocking\b/hipStreamNonBlocking/g; - - #-------- # Other synchronization $ft{'dev'} += s/\bcudaDeviceSynchronize\b/hipDeviceSynchronize/g; - $ft{'dev'} += s/\bcudaThreadSynchronize\b/hipDeviceSynchronize/g; # translate deprecated cudaThreadSynchronize + # translate deprecated cudaThreadSynchronize + $ft{'dev'} += s/\bcudaThreadSynchronize\b/hipDeviceSynchronize/g; $ft{'dev'} += s/\bcudaDeviceReset\b/hipDeviceReset/g; - $ft{'dev'} += s/\bcudaThreadExit\b/hipDeviceReset/g; # translate deprecated cudaThreadExit + # translate deprecated cudaThreadExit + $ft{'dev'} += s/\bcudaThreadExit\b/hipDeviceReset/g; $ft{'dev'} += s/\bcudaSetDevice\b/hipSetDevice/g; $ft{'dev'} += s/\bcudaGetDevice\b/hipGetDevice/g; - #-------- # Device $ft{'dev'} += s/\bcudaDeviceProp\b/hipDeviceProp_t/g; $ft{'dev'} += s/\bcudaGetDeviceProperties\b/hipGetDeviceProperties/g; $ft{'dev'} += s/\bcudaDeviceGetPCIBusId\b/hipDeviceGetPCIBusId/g; - + #-------- # Attribute $ft{'err'} += s/\bcudaDevAttrMaxThreadsPerBlock\b/hipDeviceAttributeMaxThreadsPerBlock/g; $ft{'err'} += s/\bcudaDevAttrMaxBlockDimX\b/hipDeviceAttributeMaxBlockDimX/g; @@ -378,9 +321,10 @@ while (@ARGV) { $ft{'err'} += s/\bcudaDevAttrMaxSharedMemoryPerMultiprocessor\b/hipDeviceAttributeMaxSharedMemoryPerMultiprocessor/g; $ft{'err'} += s/\bcudaDevAttrMemoryClockRate\b/hipDeviceAttributeMemoryClockRate/g; $ft{'err'} += s/\bcudaDevAttrGlobalMemoryBusWidth\b/hipDeviceAttributeMemoryBusWidth/g; + #-------- $ft{'dev'} += s/\bcudaDeviceAttr\b/hipDeviceAttribute_t/g; $ft{'dev'} += s/\bcudaDeviceGetAttribute\b/hipDeviceGetAttribute/g; - + #-------- # Cache config $ft{'dev'} += s/\bcudaDeviceSetCacheConfig\b/hipDeviceSetCacheConfig/g; $ft{'dev'} += s/\bcudaThreadSetCacheConfig\b/hipDeviceSetCacheConfig/g; # translate deprecated @@ -393,10 +337,7 @@ while (@ARGV) { $ft{'dev'} += s/\bcudaFuncCachePreferEqual\b/hipFuncCachePreferEqual/g; # function $ft{'dev'} += s/\bcudaFuncSetCacheConfig\b/hipFuncSetCacheConfig/g; - - $ft{'dev'} += s/\bcudaDriverGetVersion\b/hipDriverGetVersion/g; - #-------- # Peer2Peer $ft{'dev'} += s/\bcudaDeviceCanAccessPeer\b/hipDeviceCanAccessPeer/g; @@ -409,8 +350,7 @@ while (@ARGV) { $ft{'mem'} += s/\bcudaIpcGetMemHandle\b/hipIpcGetMemHandle/g; $ft{'mem'} += s/\bcudaIpcMemHandle_t\b/hipIpcMemHandle_t/g; $ft{'mem'} += s/\bcudaIpcMemLazyEnablePeerAccess\b/hipIpcMemLazyEnablePeerAccess/g; - - + #-------- # Shared mem: $ft{'dev'} += s/\bcudaDeviceSetSharedMemConfig\b/hipDeviceSetSharedMemConfig/g; $ft{'dev'} += s/\bcudaThreadSetSharedMemConfig\b/hipDeviceSetSharedMemConfig/g; # translate deprecated @@ -420,33 +360,25 @@ while (@ARGV) { $ft{'dev'} += s/\bcudaSharedMemBankSizeDefault\b/hipSharedMemBankSizeDefault/g; $ft{'dev'} += s/\bcudaSharedMemBankSizeFourByte\b/hipSharedMemBankSizeFourByte/g; $ft{'dev'} += s/\bcudaSharedMemBankSizeEightByte\b/hipSharedMemBankSizeEightByte/g; - $ft{'dev'} += s/\bcudaGetDeviceCount\b/hipGetDeviceCount/g; - #-------- # Profiler - #$aOt += s/\bcudaProfilerInitialize\b/hipProfilerInitialize/g; // see if these are called anywhere. + #$aOt += s/\bcudaProfilerInitialize\b/hipProfilerInitialize/g; $ft{'other'} += s/\bcudaProfilerStart\b/hipProfilerStart/g; $ft{'other'} += s/\bcudaProfilerStop\b/hipProfilerStop/g; - - - + #-------- $countKeywords += m/__global__/; $countKeywords += m/__shared__/; - #-------- # CUDA extern __shared__ syntax # Note these only work if declaration is on a single line. { # match uses ? for <.*> which will be unitialized if this is not present in launch syntax. no warnings qw/uninitialized/; - my $k = 0; - # Match extern __shared__ type foo[]; syntax # Replace as HIP_DYNAMIC_SHARED() macro $k += s/extern\s+([\w\(\)]+)?\s*__shared__\s+([\w:<>\s]+)\s+(\w+)\s*\[\s*\]\s*;/HIP_DYNAMIC_SHARED($1 $2, $3)/g; - # test patterns for the regular expression above: #'extern __shared__ double foo[];' #'extern __shared__ unsigned int foo[];' @@ -458,47 +390,35 @@ while (@ARGV) { #'extern __shared__ blah::type s[];' #'extern __shared__ typename mapper::type s_data[];' #'extern __attribute__((used)) __shared__ typename mapper::type s_data[];' - $ft{'extern_shared'} += $k; } - #-------- # CUDA Launch Syntax # Note these only work if launch is on a single line. - # Handle the <>> syntax: { # match uses ? for <.*> which will be unitialized if this is not present in launch syntax. no warnings qw/uninitialized/; - my $k = 0; my $kernelName; - # Handle the <>> syntax: $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), $5, $6, /g; $kernelName = $1 if $k; - # Handle the <>> syntax: $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), $5, 0, /g; $kernelName = $1 if $k; - # Handle the <>> syntax: $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>([\s\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), 0, 0, /g; $kernelName = $1 if $k; - $ft{'kern'} += $k; if ($k) { $Tkernels{$kernelName} ++; } - } - - - + } unless ($no_translate_textures) { $ft{'tex'} += s/\bcudaChannelFormatDesc\b/hipChannelFormatDesc/g; $ft{'tex'} += s/\bcudaFilterModePoint\b/hipFilterModePoint/g; $ft{'tex'} += s/\bcudaReadModeElementType\b/hipReadModeElementType/g; - $ft{'tex'} += s/\bcudaArray\b/hipArray/g; $ft{'tex'} += s/\bcudaCreateChannelDesc\b/hipCreateChannelDesc/g; $ft{'tex'} += s/\bcudaBindTexture\b/hipBindTexture/g; @@ -508,23 +428,15 @@ while (@ARGV) { $ft{'tex'} += s/\bcudaAddressMode/hipAddressMode/g; $ft{'tex'} += s/\bcudaFilterMode/hipFilterMode/g; } - - if ($count_conversions) { while (/(\bhip[A-Z]\w+\b)/g) { $convertedTags{$1}++; - #print STDERR "HIP: $1 : ", $translateTags{$1}, "\n"; } } - - # guess that we are in device code , or at least in a file that calls device code. # will almost certainly call one of the coordiante functions - could be fooled by clever macros but usually works: my $hasDeviceCode = $countKeywords + $ft{'coord_func'} + $ft{'math_func'} + $ft{'special_func'}; - - - unless ($quiet_warnings) { #print STDERR "Check WARNINGs\n"; # copy into array of lines, process line-by-line to show warnings: @@ -532,58 +444,41 @@ while (@ARGV) { my @lines = split /\n/, $_; my $tmp = $_; # copies the whole file, could be a little smarter here... my $line_num = 0; - foreach (@lines) { - $line_num ++; - # remove any whitelisted words: foreach $w (@warn_whitelist) { s/\b$w\b/ZAP/ } - - my $tag ; + my $tag; if ((/(\bcuda[A-Z]\w+)/) or (/<<<.*>>>/)) { # flag any remaining code that look like cuda API calls, may want to add these to hipify $tag = (defined $1) ? $1 : "Launch"; - } elsif (/(\bcublas[A-Z]\w+)/) { + } elsif (/(\bcublas[A-Z]\w+)/) { $warningsCublas++; - $tag = $1; - } elsif (/(\bcurand[A-Z]\w+)/) { + $tag = $1; + } elsif (/(\bcurand[A-Z]\w+)/) { $warningsCurand++; - $tag = $1; - } - - if (defined $tag) { + $tag = $1; + } + if (defined $tag) { $warnings++; $warningTags{$tag}++; print STDERR " warning: $fileName:#$line_num : $_"; print STDERR "\n"; } - $s = warnUnsupportedSpecialFunctions($line_num); $warnings += $s; - - } - $_ = $tmp; } } - - - - #-------- # Math libraries # To limit bogus translations, try to make sure we are in a kernel (ft{'builtin'} != 0): if (not $no_translate_builtins and ($hasDeviceCode > 0)) { $ft{'special_func'} += countSupportedSpecialFunctions(); } - - - - #-------- # Print it! # TODO - would like to move this code outside loop but it uses $_ which contains the whole file. @@ -591,33 +486,23 @@ while (@ARGV) { my $apiCalls = $ft{'err'} + $ft{'event'} + $ft{'mem'} + $ft{'stream'} + $ft{'dev'} + $ft{'def'} + $ft{'tex'} + $ft{'other'} + $ft{'math_func'}; my $kernStuff = $hasDeviceCode + $ft{'kern'}; my $totalCalls = $apiCalls + $kernStuff; - $is_dos = m/\r\n$/; - if ($totalCalls and ($countIncludes == 0) and ($kernStuff != 0)) { # If this file makes kernel builtin calls, and does not include the cuda_runtime.h, # then add an #include to match "magic" includes provided by NVCC. - # This logic can miss cases where cuda_runtime.h is included by another include file. + # This logic can miss cases where cuda_runtime.h is included by another include file. print $OUTFILE '#include "hip/hip_runtime.h"' . ($is_dos ? "\r\n" : "\n"); } print $OUTFILE "$_"; } - $lineCount = $_ =~ tr/\n//; - } - my $totalConverted = totalStats(\%ft); - #printf "TOTAL-CONV=%d\n", $totalConverted; - - if (($totalConverted+$warnings) and $print_stats) { printStats("info: converted", \@statNames, \%ft, $warnings, $lineCount); print STDERR " in '$fileName'\n"; } - - # Update totals for all files: addStats(\%tt, \%ft); $Twarnings += $warnings; @@ -626,41 +511,31 @@ while (@ARGV) { $TwarningTags{$key} += $warningTags{$key}; } } - #-- Print total stats for all files processed: if ($print_stats and ($fileCount > 1)) { print STDERR "\n"; printStats("info: TOTAL-converted", \@statNames, \%tt, $Twarnings, $TlineCount); print STDERR "\n"; - foreach my $key (sort { $TwarningTags{$b} <=> $TwarningTags{$a} } keys %TwarningTags) { printf STDERR " warning: unconverted %s : %d\n", $key, $TwarningTags{$key}; } - my $kernelCnt = keys %Tkernels; printf STDERR " kernels (%d total) : ", $kernelCnt; foreach my $key (sort { $Tkernels{$b} <=> $Tkernels{$a} } keys %Tkernels) { printf STDERR " %s(%d)", $key, $Tkernels{$key}; } print STDERR "\n"; - - print STDERR "\n"; } - if ($count_conversions) { foreach my $key (sort { $convertedTags{$b} <=> $convertedTags{$a} } keys %convertedTags) { printf STDERR " %s %d\n", $key, $convertedTags{$key}; } } - - - sub countSupportedSpecialFunctions { my $m = 0; - #supported special functions: foreach $func ( # Synchronization: @@ -669,53 +544,39 @@ sub countSupportedSpecialFunctions { # match math at the beginning of a word, but not if it already has a namespace qualifier ('::') : $m += m/[:]?[:]?\b($func)\b(\w*\()/g; - } - return $m; } sub warnUnsupportedSpecialFunctions { my $line_num = shift; - my $m = 0; - - - foreach $func ( # memory fence: "__threadfence_block", "__threadfence", "__threadfence_system", - # Synchronization: "__syncthreads_count", "__syncthreads_and", "__syncthreads_or", - # Read-only cache function: "__ldg", - # Cross-lane and warp-vote instructions: #"__all", #"__any", #"__ballot", - #"__popc", #"__clz", - #"__shfl", #"__shfl_up", #"__shfl_down", #"__shfl_xor", - "__prof_trigger", - # too popular, and we can't tell if we are in device or host code. #"assert", #"printf", - #"malloc", #"free", #"memset", @@ -728,8 +589,6 @@ sub warnUnsupportedSpecialFunctions $m += $mt; print STDERR " warning: $fileName:#$line_num : unsupported device function : $_\n"; } - } - return $m; } From 3fb77dbef535db108b48d5dc523aac314bbceab1 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 23 Apr 2019 21:13:00 +0300 Subject: [PATCH 10/17] [HIPIFY][doc] Provide patches for clang's bug 38811 + Update Readme.md accordingly [ROCm/hip commit: 57931b30560d6791178172564362aec63ac14afc] --- projects/hip/hipify-clang/README.md | 8 +++++--- .../patches/patch_for_clang_7.0.0_bug_38811.zip | Bin 0 -> 14641 bytes .../patches/patch_for_clang_7.0.1_bug_38811.zip | Bin 0 -> 14641 bytes .../patches/patch_for_clang_8.0.0_bug_38811.zip | Bin 0 -> 14641 bytes 4 files changed, 5 insertions(+), 3 deletions(-) create mode 100644 projects/hip/hipify-clang/patches/patch_for_clang_7.0.0_bug_38811.zip create mode 100644 projects/hip/hipify-clang/patches/patch_for_clang_7.0.1_bug_38811.zip create mode 100644 projects/hip/hipify-clang/patches/patch_for_clang_8.0.0_bug_38811.zip diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md index e3e26af390..37a1645c4b 100644 --- a/projects/hip/hipify-clang/README.md +++ b/projects/hip/hipify-clang/README.md @@ -49,12 +49,14 @@ | 5.0.2 | 8.0 | + | + | | 6.0.0 | 9.0 | + | + | | **6.0.1** | **9.0** | +
**LATEST STABLE RELEASE** | + | -| 7.0.0 | 9.2 | -
not working due to
the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | -
not working due to
the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) | -| 7.0.1 | 9.2 | -
not working due to
the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | -
not working due to
the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) | +| 7.0.0 | 9.2 | -
not working due to
the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811)
+
[patch](patches/patch_for_clang_7.0.0_bug_38811.zip)*
| -
not working due to
the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) | +| 7.0.1 | 9.2 | -
not working due to
the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811)
+
[patch](patches/patch_for_clang_7.0.1_bug_38811.zip)*
| -
not working due to
the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) | | 7.1.0 | 9.2 (?) | -
LLVM 7.1.0
is not yet released | -
LLVM 7.1.0
is not yet released | -| **8.0.0** | **10.0** | -
not working due to
the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | +
**LATEST STABLE RELEASE** | +| **8.0.0** | **10.0** | -
not working due to
the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811)
+
[patch](patches/patch_for_clang_8.0.0_bug_38811.zip)*
| +
**LATEST STABLE RELEASE** | | 8.0.1 | 10.1 (?) | -
LLVM 8.0.1
is not yet released | -
LLVM 8.0.1
is not yet released | +`*` Download the patch and unpack it into your LLVM distributive directory; a few header files will be overwritten; rebuilding of LLVM is not needed. + In most cases, you can get a suitable version of LLVM+CLANG with your package manager. Failing that or having multiple versions of LLVM, you can [download a release archive](http://releases.llvm.org/), build or install it, and set diff --git a/projects/hip/hipify-clang/patches/patch_for_clang_7.0.0_bug_38811.zip b/projects/hip/hipify-clang/patches/patch_for_clang_7.0.0_bug_38811.zip new file mode 100644 index 0000000000000000000000000000000000000000..28205a2645b83be0a278f2132b603ceb697b5d53 GIT binary patch literal 14641 zcma)j19ariwrwK&U?4IcUA9S?{Dww zTD@u&_9rU=0*V3v1^@saIVsX#&3_fpuTPd{`gF4Y(E|D(76z7j)+YZCJB0sX$3{y} zOaI@jaq|9WT#SEXZDws?>11g1UnodGX>N-2TGCYpGXVgAg8>9!`+uNxbpF;u9Rnvr zJsm?MXEOsM9b+eJ14lC(YX@4>aV-yr*`c~mJ|)i(j=W@hcIh3^@HQBjruIM7RoF;W zP&FD(r==(obzpUS^TZ@;^@sJF$J0kf7VdaRaGAYG9<4s_soB{o5*(gyH~XVrFCnT6 z)HS^=ETcH9_G{Qa2kP2r6Y_k}2?1cssY8V0A}UIT{)g zn1W46wX`8U9h?%Wvz}k4DQlyJDAtTFWrnuK48NaE$UlLzw^KmoSL^OlL*wQ+JveMc zsQg~Cm3W=Dr(Bk77$w$j+p4=+=TwzNhZCu}KBmq{Ri zC)v0{+j+Av&QOW;eBqX9h>WKryhR}-ka(dM038b``;OL zWP4_2OIEPqPo#dNENG0bxQPzGO)2Gq$|gfp#{8D)2I`8}5iXdX5MYZoQ`PyC>iY)b zHd%5^>v%IQG6*7t05?kVP4_d!C!XTl1suh9T-S|>iETn9Ux^0a&EhnsT#!iWb2Xg( zih9$QLIWB*1q>SJky|KK@uJ?XXROulqOgF+q_SE~FxtWa?tQy+xFPbCK-zh7KS$Pp z0Sf0*ilKhCxJ!hTgH)88pJ@I&0F?=F%O?#qR1SxC`rN7l=zuC+hgfo(G6{Jz0T$BN zLy@^9CCkYa*73069(`w6=kMh%kLF%-t^TK`sP*v;v>DH|SJk)9b>Rirn?N~BP~*$f zDlbu7w`%dg^5V5*K4`TT>+YOj;~Z@iH(zjk)Uw@hur-u{CB?LC=F) z)6#`d)MFmCXu`XX%v$ey(|OL$Zj2lrQV{P?ckwG%NqP*KYlb0HV)5gD0)s*Nmmj-* z-;1PPv#9DX|ID?^;9`k1y;5dLC1CcUD{S(JSqA%LYsQ=e^~3mUmBhz<_Gb>Ap79ewZJr2|4O7t`e8=|_Is%Vv%;$hE*)7< zm0u7f=$%}o(A(4sWDmE_PpjFg#BU;)cgsty2t;;RI{9wY7cQb3_pI(x=d9{-*U^J4@`PjIsmQtZ`CtSgLx|KTPUse=!EfVrBzi49X%%DJ5n(qiS&N+DJ`*(vBu+3pa?M&1x01U3T{DF*_6AZ!847D2)T z^Si?QNZR3oRP^TxP4wp?Pq<`bcUwR^I5#)G7u}Q_P(G=sztwaTMLDk;DT}byYC55JCeRY0X)LF47EcO;nJP+Xi({_5 z%h(4m2*NhV_N&!|rJLo2E9fHMXP}jFf4aSqNFplPTQ3MIJ(TZFRfLtDsSJmyAuBCc z$M4p0mR&Yy+q+C-=I#o@GoGUFrs=7~KB>gFcs`^e6&$>Pjj0%02ym?T583a9JmA(+4Q>~`*>xN#^ockAvHmR?k&S!@YNiQBQF|f znpJ~ex|mqx_o;I1Eo_QHRdRBp$#d0dcyj)H%y~SOwU*m?grbsL{D*nmVyCB468Uix zn%g;yhLRiSyUCi9dWZ3vp1POOnv^<^;Tp%udpsCc8`p4@7qXI@3wcYGO+MA-1&PbO zl;Zw+){aoeV9aShQ}z>gSC}7-XMf4F)Hlo8@4TZn`noWh*J9p5ayZq*6hly+VG5ZK z5qhz&zAPw`m`+;sheR$w78I!-aT&hQqit-6nNnFzFl*bHuP*6{y0m49Tthj9e=u-;WU; zlxkYpG9DJ!iBv0#ANiG&`@%14R8dBVu0oz@Lg&0(OdiKIBWcZ3ls0s*!Gf$k>9Epc z*2Cn;GKmI`)YjG-aCT9Z3`5sMH!lt9?CL+dPli~h0oL7C3#-d zd_bM^9Oz1PdiRQxWbU_6D0p=UV}wM3iDim)RW5ncDg6{ZmL6G<-}k`vu*|antBz(% zkoRJs6?Z)~N7AD2$nfgl6vzdIZG;++C4unLi>T_P_@J+*B$eWVTt5=T(zcP+xLBMjawcuKyDr%cr z%3|WAoQorxi^DllMLbc3yB9<}83soY2~}pA=JjiUQfZi?DhE1X9>P!;(rGwJ$MhnO z-Qp3tJYUjorED*yl8~>fvC-5|pwqi!F}&we+fy?;n>72b5Z7TI ze1P{+hdGLMR+W4K+unNOGn*q!=RACS*!DqNu)gXjTHy$RQIO(m@t(1i;frlkXhwMk z;4R9J?kkpLRTt2TgA@RoYk5$xEK!Xg`av%R9gg_n6b=h6|&AY^c3s3EcwjPcpq7K1L4F#lE#057B z|1dl+9|L=bvh~?WTKwFZ{91b)^?J?rAo)0YRHs#<%%26G^=6ffPKUUGc!~ufz?JaN zRz#LF!U_d=*tyd1g6A((W~xs&Xt!fUs#ZF7L1-(@}mw5$kr=Iw@>rxgJxv?(HiqNHLxfj=U zSb4viN-6T-+7(oEFet2qkAwt$IH5rbGUCVU;Jir7*aj`hHWh>Q zU=dKuDh6*?9&fy5vR^F+d5!Ti_?7rzmeSbh(^$2B}CUh$H;06z%TizSCR(q z*J;Ty0xW;+bN|W<6J;}jYqKD&@X0GnaB_)i-NT^j5*@5b7s=c1-j2Pyx3G2xNmryW z@y(Ik?W@l5;vjKfGWIJcNAd^JS!UV(bs*aB@g*ApkFwB%r|iZ2T=9~HB~-4)J$mde z7H26xH8vO~E0}s~G-Gu%vU9`*|11=_V&7#flgd)JKdHz~gpkj+iEX2kbDdk@Y_uGQqS4)N(l{{AqWor*KU8 zQ|!=8qw8}=*b2u zsV5A(NP)e}K+)VZ+=MSvemgT*jwDQai5dKul!ASu zj1$?S97?mZXlVfVCM-@(G}1DE4IdE-qBlDypmg}p8^BWpGdjij#)!TZ7dJ}y%H`=6 zZ`nskxG9?r>yJ*ZiO<3jCQA{YyD}r-8S;D5F953tp0_AkX9{a8PHWoRi^X&u{O_9? z@af=|YpXYZZv*ntKh|;ycAl}^l%MfA!o4}s!F;&zKJGXADd-_C1UeUKM z&v45!1vi@==t+pRng-L(n)%nWud}N46-RAtcqC9#$4&ajZ7$H}64ec|GuFe|i9VP8 ziFh9Scv6Tt+sBa+0e#Xm@nzX9Y8z}CB@R$a8ZM2SZFs;f~~_d zjuJ*%O&*lI_`%APoCc}}kE~1_V>NEn0jgGvZ^)@)XP!pA9zGaZg}}z0oNmaN1tZiT z*+`Rek7)RdpKP`TRcX{v<3^P(jhbcHkkh5jJk3&r_LOw!jwwx)nxMAP~mNaFfIC-+w^ig*x#sJBM*L3P7 z@e8)^rz=;f!|9UcLlQIzR-;B-o|(>K+*A4)1zoXrv%RcNRC1=p0#Y#_Y~j zKJS}RhdR;O`P*Iraqd!OO$E-ay?^Rb_TYF zf|iF&EJ>}u1s&^_;16Y2Qc??7~b-TkiK9bn>lKF@}G1ou1v7)wY2DpKkV=%^D z+=x)8i9U5Q(B63}TRC<-W4<*vuc3NemC<5|&$DsR`Td-dk$VeTvP%qtgKa;xgv)Ce zx3R)P{W|gTMUo}Ys70E>=DNZ)cl{Xvxd%GEm_G=^PO_a&A*T8+4r*>$2Tbjn2CJ~9+Fg|r8HP(jzt6Vm#; z3s^6i4q_(NTZM`mCPr)9q5qi4T8IzTYn56S-xI@}do0^Nky`cxJ z)Kg$4c9A$HV5Vy{-)`DKRc_=yX{A(-WeYz1B=Yo=&=KvUlYI`u*Wbxgqt0s7?@?o- z#cs%ID43C3y>%A!HZN!tV4CpKG{9HxGI?l)E%SbvF^8L)8^cWQ`q6+pq0rsi-kWZB zJ2jzqzBA!x9OnAl!FbUwJk(_T=kZiHVIUpj4d4y$dLzH z0Sm}G$Q#E15bUg#8$?JF& zy3woOUi-T}D!f!*%}yR}oDaI{qIG<;-+l}|50<08*NI`OwiGAJdo;>Dm;CIds9YDZ zBp=%YLS-)Me%AIFQ!VLKI7TenZr%HpFcBVwEG&heAe6A*sonSK9kj)5<@`TZ9{wdsyE~ZL}G^cs3Xn;HgzPWD;8&CEKW%pDJbJC zEJz=q^EvJEur_s@Xw^zN-cJUdZQ(_#e|IUt>Jub^GLXuB)Bl=7c8X%CPZqD>3!Jdc3GQKCnHns?HuA`bHy4o)++NAr(HN`s_+ZsZt=)b zxBaxd5yq?I#7J$>EhE*48Gl_lUsv2kDS$LBxwC$n(y%W7DlD5O!{zAW2tY+Jq1*scgWOB5akK6UFcV}odblHo4? zS<_}U?uAXw{IQhb#)YH+#t-w;sRBjO8tu7tT7w8B^)kI{Q1L-S%mHuuSOz3RT4iOa z^Qo^a8h2K@kCTu*76YDHRm2t;6B-46r?csMyeW%*J04nKnI4o9+lDK0R>Pci=@=FSn;jA^m^8zt6djx51yRe2kv zT`=<6!X1+2q|CmHnFj$8;%OBe*L=ETMP#oi{`e!xeH}HDG5t5@dQv5O;ULRQg>d;~ zNq=?5Xi|uh)2UOFl1lg8kdn$+d!JnxKn~7s9;cGEmwqL!O zc;zxDwbvDRr;?|3GPmv(9JPa9V`sH5bJ_`I96tTw$f;+663sAsWn$8eg@@{PH|V_h zHDA+)>A>FMx&HI;7+jS=Qd+#ij)ReOA)^LbV*GP_9GDbXMo7h#BJ=wX?YLviHT}=H z6yHVKJe@eiXPcXc-b>`u-KRuu^4I=JI`g-#gSS~GpwA;+(2Bbp^=H?0N`zdG9>H-~ zv*zLqLYtH1_Gza6QK@TI$!lAUaBcpG7O;l3vnAg~mQnIkuG_*7Ap5<52D$MTw=oiK zA>O6uSHB;!uh~4m-rEUL3MJXH{+p%^*D4VnFVQLeNg-E z$pMBF08asc_(OnU3{YWR+oYoXW7^0Cufp}ZNuEGvkSeCt2JHeModN9T0qx`Fxcp>> zcJ%V6hB^U5qRq;W&)yjU+qruxyA1<^#JiK_p(r?m)MmaffbKbKTE$tJ(t~{lyiVscx zcI)}<^UtIRNWjWhB7)_Pbwf1_08pR?0C4d*h zTqBy-#P(S1j)CBQ@9Ja@Uzf*VvIL!D{`}0SmkB2H;$ckBu0)Q2UC->J#Y6OPL?W17 z4;%n(FD@mFU3|ICCD(CaqX)iicUZWPGLb2+jzy-q&7iz5jPU|qHP&`w$Jfn*TVTSo35F~yqro*$hU@We< z=0_m65&Cj^{jJ+D`={3Yt^;|ZKau{;>3h?enBgh=Bt~;O=}$LiKrH$}=WoMLCkod5 z3c(rVT6L3Ob0EjuU@C=wKpMlm5QjcnZfr8z@WH zC)p7Ybp!3+O$&o1pjsK()-!7UqIOds%xLBNsfM~I28!w+bZY>aQ}*cRkwzUCNdVXT zB=&ODA4jLC@KI^m59WEj%Z8RAQ;ohE!W6OW;LMicT`#f_ngxGo>&2_PdEP&M z=luu3fGv&2w{({dOU}2{bvA6;UE#AiUh$kMWg6A)gZVNvG_k1Bv!-}whDWAodffm%i)_C^3UlNm)@`o??evW^Yf znK37xn(R|Q@N2NnfU+F6s|yR>N)rex)&^XcdN>Oh;eMeYnsMDFTQYT01KjqZxtBT);-|VN<6O03rqbh{+bZ$y8f_O8tIuQByZUMopkaK&dHz(fAn$hQU#HSfvMl>R5 zC<0xo8`LQ#HUJ8u=$`EB!$@*Ot}afKPYx~(+R<8jC|n+ppd-8E=3^nd?YECds-Bqb z#zWGfh}j4Meh`en+5}v?KmDD#0=7{~n!6+lht3X63=}FGNL1bLCs0w-zTJ=3_J#53@~=)`5<&Lv=t>9(%T~G_UXIO@e(&% zUg5Jh_$Q?pBYluU!czx@LXkiFDs2Y6yu@Z&%ro!`3Q z@Q_e|~iW*0pEoSbgUKitxwjg+EE5wTxU z)EMKWxisyC%jbo|74)L#41Ed32|eE)8@GzmBW`D@L8QfA@_G+`|0AC~+@&72BAe=} zp^z(G`%;ssf6vd}f1*{TH=7ey3)AEqhlZ|OI1Gki;c)Sgj)q!2_`B4Vr&y@JK<1Sj zihGQi)UMPi7c~3`u^dYF_ZgG&9p?8il)jPecj318_IHp5a@ijhd!+K$NZ&)}dPl5e zJy5^|#lRWef@BlfD|KO2&YW7C8p#Vw;VCZ#-1FjR#+oaw^~ZUtz@i4gYl0AY2MgiS z=jgQ&BY}T+Ml|_OFIl|>)_yRcc?Tthmu>(&c2@TpD?J~^k}__RrKnlUA*AZ^tXoJz zR)Y7*+91N*4U$$*@kqv^oKz7Z{K%}dg$QPo=`BJdEmlu!S|u0l>Pke9_E(?+)B_z+ zDI8lS34FyyIKbZWLD-b1z>1dGWHHrG8@^S-s0S~^%{g)rg??#tRb`jHD-~}ADUjh3w z%||c?unnRSQH(owC9y3?)K#x%QKH1;KZk{e+DAdCo?m@O-*|Ip%HUrz%~8P~@Grji zX?`N$Zl@k1eL%9-+Y&y;Ad# z-Ip%OKLr<*234JuVeSKkzN%t#3iP|g{=vE4;?*!a2aW)u`&j$_VPKAk z%)r2{Y0g(1?}pd6S9%Y zlq*p}(?*J+a_Gqvd0&Mcc%3pF-i|V8ZawGz1S5r8>;p$|Z`leOB*!CZsY#-FX}B1$ zMVa6HE*RRzTOZJ@`d?LPCeF|H7jGGL>Zh)L|S4A=%-3EO%7x4FPBD`*-PtKQC zC-vM79V&q*fgoZj{PC7HxYE`<#(5kP=U?4svw9rDn@qq+$YS8JPln`ZpRbH=Urb-c zZsA{zU!R}?zcO5n+mNsbw<=YsZJ_%T3)Dh#7d+P-$Njjvu7*Mk>TIH{!HNhv{L5^6n^xdtuqBJchPykb@}ju1(JRNmtGx?9#xS+vQwO54>KYTRDk zXRVa72DPg@ypmv5&O&JId)5*gXYJP-Y4Z@2;G&WNtWS){hGjQP8C9?oGMtAY*b8>< z3C&arA**8aYPgkpn0W5l1ndm2)}E!DZNm9Sh;<@H@P2#ts0~8ULU()eyp6hVe%4%Y z!T!4`9rFsqt4;B|zS?MiBi=cGU4csNtxBzq?3sCU8OQ2}X6EF%X6E3RbhQSkyT$#s zO2NovxXV!nL$?W>>c9Ln?z2NF69RBRHVlW(66VE3j{LbtGaqoy$?@=Vd>vHTfA;1# zjn}>viip*uav9b2pWPZ#TEs@Azdse|RI6`!L#6zD8~b><-W7RRwYk6X`Ek8JxR7=| zZ8tj_-5$~B`BT>u5mcnMYX8x1IthYTZ*Py&ULKc{iinbktW$=5{t-&~#akSzAuwUi zA>CbC&PP+O4@hdJ_zv#JZbQd-@@HSfcXlzaekeQ2D?zWHmeAZ%e8ej`i06}`4<^j7 z0{UzYkaz=;JuJjHYqw6ua z6awD>TTScLB!KI?x-H#xE<5P#Qfkn~wNtyEb+L5eV?+_%#o~xhK+*TsIsFOruU)`j zNzDH95>2{L001Ki0Komf-33_bIhyJi+t|D4*&BZ4CjWCIP^o5Pv)Yd6W3%-PC-|Z- zcE~PsTQ|x1BP>bK6=$_3s~8L%GEq=%%}|bXpL3w*vXe`^jzT0~%Tj1+&()>tRp8o{ zbNz4W$Sl|I^hL`nn^Q#B281q(!rpzla*QnELo8_$a-sS;OhSU+ZV+FG1rA<8l|(Zs z?Y&@SL2eMehcbs@Rl*}8ZH17?)N05r#OnQp1dt>K5J)9UHR}l}LGzZg_av!g5%~Gt z2I^VzNKl^fu75O(0>r|a0h<^@B7sz-Kc#`<1&P+IRh6fJr3Sdc*ythx<06cT?Cgrp z1tW_ghnW1&(j}bRMf(~F5UdH^Lg@t%pl7WDv**qdZCs5&=Sg+QgMM=RZcssj;o`+d zg}O#|rA5Wvt#|q+12Ur0`K%)yJ=b;HY5R=GJ&rbkB(%o9Vl=e@_Mi0}JIZ|pa zlj_e#hgekaRE~Ai&zi^i-kG(18tVCdFGY-CM&*iRKvC;sV7_lz=|{aXBk*R#O@>Q0 z%m-9UsMw0!7?*rc7B@|g_fI0~lPa@%4iZHZmEu<@p-%DzP(q}Z39xOA#_TR^B`fCW)aRj&T97q3zH6q?Hw)A~vlqxN z?Sbgevx;m_QG!&wR|9m_E7K>LPQE8V54XYLDBBYSB#flEYaN}kD{T$9-g$t&cr1L( z2ic*S8kb=5g$Q~sfSeZZo3*uQfr8TJY3s{Qo~(Xv+0VA|o|}p@j22K!ZaOSJJ!b!K z+aOaeIFm|f2uVf`qP?WJNL(tQ?De`_08l@Pt}gmS3oCHdoQX@QEf5T)BAXQqycOsN zKjk-)BoA8)8Kg1`!D8r#K_ChRKp?A~UT3 zn-E}!_Oo#)%4!#G`O5a|NMO*L$OhYs)_3XVf(lckq@Eju5vLWcJ6yH*YEoyGR+J1gMS#Rs~u`}Yfx+i+p z`FJ<@VQY~K03n)`_XQW+ED>_re)`xam2z>crh1s4XK4F&abiwSpBWCirr-)&I(jo% z&L`>oO^kwB<#XQnro!;gB`YziYa0l}*LhOI(_V+{xgte{Hep*YjQ|%%UiQ)*>ZIk# zP@{!LrA4O@{=~Lb5g!Q}4o%4?9uimPIJV(xL**Y+?=EkQp`QGzWi+b+g$&j@nw6;) ziq46Np?&^4L(X&f;`bKrYt?L5sSyO5bH?HteUj(s zEXyfBPKPIR6F9h@)nrr*YhHUt0@tQ4CJJe$;cLO5T+0m}ULK^I`F4#2*4N!t_Q|56 z!&Ou&DTC5y=A(WWXY_i_g~>zTz5M!=prHdUP6cJ-wCX@Up<9#v5 z)3~?dN&FhuHj7t&t{RKgq}Rt|!IRf`ZM%Nf2+NgppB+pyC`8wcobXew)vpDaiPj7yKCjXybu?rbauHlVR`s47{sX*}wV zNBSqAmk1XW#AORN*SOL;^@*2*Z>TxHPY6F4W)i#@eBZdA>)_EM-MlD!^=R8jP{><; zZ2UY@Jn<4KnNcB>3iy58w~4-lRFYiYbKJrCThrOR_xrET_qis_PGg96E*2Uf3EPG1 zNH0zYeRcSB?t!c1A6iS`u<<9tw$Zz7N9NTmb)SGI;ZfCS+u6H z7hxZvg(cok9bwWt_YV&xg~9PaG^4{l?xQ&2mupjW{5_^zceq{2-TXSC9reGXO87A1 zy!f|r7}0f5oAi#iV_Yp^!`higHDj*RC(X(d;})MnD#0%iVhGV@*u;gv#}wl1u_CAl z5cz(0X~$W(9F|UGvgsRjd!oCt%t+AH8NuwmrW$dXeV1TWk5RXWyk~n$jjLHWg%6Ks4N6Aampn1n&{X01SmjI>;Ay;?{Z zu`3mz&&^V5zGqrhilzr(z-$Vg%H6m@vs=w)&B0={HUg9=g9pu$2dqt@T^q|A+uPph z%Z4C^*OF=AoyZ9boU{b;D=Ppy6uA?gsm@fF9dNcYW{`AB#vB{c*ik{YU#ngh^jx%fYU6{?95(TG1|m6(Ts>5F_c|kt1GXX3s=UwQ}~9^B@8inepvy`C0iyu^ zGx6a{!WAY&xpT>aDOj%P{|84K6DI2L6{fh_4sP7en%q^)K!J0K7J^f&c&j literal 0 HcmV?d00001 diff --git a/projects/hip/hipify-clang/patches/patch_for_clang_7.0.1_bug_38811.zip b/projects/hip/hipify-clang/patches/patch_for_clang_7.0.1_bug_38811.zip new file mode 100644 index 0000000000000000000000000000000000000000..fca5bf52e36ea96d4e9d5a32ced95e8a2d4e3588 GIT binary patch literal 14641 zcma)j1C(XWlJzaC%eL8NySln;+qP}nw%KKLnO(MR+jjl+J98YW|~uetokv)2EaD4~v7+|7rpK4+{fJJ!_NyWcS}a zA^aCRHd=aGhX2JHYa!D9)<=8n(}Yx;EN`JRfvI0N64)Ilgv`SZ2_#E=RHWN@cms`1ccdp%ucI z~+Zalfv+76jdx_h%w0K+I1Q zC>Agjym+I%C%m^)##iNKr!DdbLOahNT@&&zM!-4~YY5GyEFp4|NNl9&n34nK#2Y4*{P5J2c*Zm0gK$CjR`xFndfpMHwpM%UFy=(%vEVy# z5kiS$+pMqo^9EO;TDW~2?#(zRM+W5IZ~4wnzH-}E(!JJAXKT|B4qf-5cbc`p6<{y0P8oubUUE%_=MX> zA8kn=9}l68?3~W*DLS7F+4q62ENOeQxz>c)>$~;XcDyd^PhLIELzGZVsKl1^z1dYL zPZ;^ZyH9eJ*;}TX!@ak&6Mt_Uc?Y^G@e7&R(c#!gy5f|rG&c0Z%haMRtr5owTZ+?V z5(wZ)Hm=b4?nC&L1f+!a>_{0(ob$^1d$|3qv3gl}pBrk&=K9;I$Df1uruqfZoUiNt zcSarAo|)N_6>RtusUIl|8lx+2qQh@fO1Yr2$qzh%0Cy5e<&3#KOo*rLr;b^fIK zzJa(+mK@VM-b{-Of=D62jgoxR{Y>$Rr}%aONAVrkbz@>;n^4JDqQQ5wIE^V6B$E1E z4QIci-n6CAfW}S%gT{H}77A6osCVlbYxTP*EZ{MztX30@ws3%Z-|if4h&&~bcAnhN zku_j|!nu@UsGlwF5+UUv73Jn9n*R?)pJ|O9okNjTu$Y z^PtwWbRiV=m`5#|@a`kC*1O(xp0l$XBZr3+#QW1-{K{359z*7uVaSwN{P>^1V37Xh z$8O*EBB|Fbs`|@6bL}#?SRzfYlvz>aB3;V~< zppeNTd0t#kvYPKrzn+WFJ@=_}TksX#!z8m$0i|EEZtGm=Wt2&+A zDqB_n7JIW3>41OZxMe9^hFQ1qNzEp+Zdb#f0|zQ{&8Is9>P#40B{1)V9|9Rk_N8v=yDW`H8)K)?@#EnwLq zNO)j=SC}73J6w>8{#>Dn{#@hoE*ij=@_f_9%|}f2(=7P$)l>Hejp(6p?t&QFUvEB1uN-lB=;0sMcu&f7_R zD0vgeF^|G)aBr*i6+?P2AMG*@IfI>7$W;#Woo4_N>2Oj;#)_W~c(hF0-1;yXPV&KreV`$#4C zhDdK2```sZ*aq2twVJSWv%GKxUF7==v@-5bw>J_=L?wIc1wo~U^1Z2wu(C6i;V?C1 zrRD1Q-8#;)%jRr*mubxWJw(6541`Gu!AIyRlC|Mk;pL8>(nP{n0(?$!N>Mp`BtvE* z(!Ns3nYlk)k37zJKi2o5G;{b6`0Z!SM>+662nzl4f<-bR92n-(o|QC@=K^VoEDJWPBHWN zDfrS$XG6-msSYWfUL;Xf_vwqIKc!Gr-${ZTtb{GNS~nQPqs01`-AyQ1$WTL0adYQV z6*Fl$+%2iMt_bC6ZW9L}V@*y=tvHhdm!Kk>e)nS^uT0Xod}KDHCTP&TW!MY8nuBrV zMdM7fYVb=J6N~&lRgS%dO;M;yPHr@Lt~w1*&YzDtkEgQMayySuRC0^|Fppd8^i)bB zKW;*EJBQIwa^rkAS#whFFkaJB_cB_OQs*&T<2ZSb2g7RP8jkWpR&sM8Z>h4$r@Fi# zak-aL++WYy5$YI>Iqhf4e&X&5^P}wW?B24chp8-7e@11%sWU9r<#~z2+A`| zA@dg$7$#a#{^@#eC>72hsqr3@#goBjbgNtotSW}B@R#{{02jXzM5$ZC zn|e2Tc2;%3?X8&~Y@Qi4zoal+T zyUC`xsOm~n`j+qoBZ9#mwHAkW>F#&FIZ(FSq)rSbo#SDS7?!I7Rrr-5xs{)hYlYgLq(ZG?~+FAq7E~=7Y=$h!}r9qut{YUr75bHF+dfrTVo#|`hg<)ZoDhlO# zM^I{}yu6~MpAvNfzz569O#4((qpn7pT751ZMJ~y`kManYC?y#c zy8}Z;UN{I$|F2rIn%jyse#b+xWm(qs-qhhe7!51CtSn3F=K6?nlHQL&$%D$;?g{A& zg4X&BRKpZZiu6&7GY-V$q5S7s;E~#Ife1j=#<462SnjnJTnki1 zZBt8GOq`T+aYS=*I47!zC#rDwf`})>;3y)Y%1qO|ehpA64O3L*KnKi280tbg4JYZC zUc|9mJVKY}OWLiJ?WI%_@^v*fn)(TJdUq^_m)+^J$=F;&YXK`Te#&-xYG!AXX5SU! zI?RI)@ILA=N3qVTk}qJ}TTgsubA;)fhi?ztK4=TpR~ale?m72i5PuW=&OF zaHC)G@M!RhhhmA}WW#DbQF*ad%|W-=X{olLw`%7yub}PJ6CZb7%7P*{)`eLSx->2K z;<^qi?^jbPMIKzcf{G3Xg_ZD;kf0AIG)O^4{CFLl7il@W>-hvHtm|MF0}bQ#e1YuM zLx?-M)umo6_UUMuOY2?vRh|hfnESC#2i%3wxi_bJbx9i)FKkuz&G z<0E}^)O`?i6z#ft;`da@4euuBnFqYqc}*R@vxUlb<{rT-9!CpwtIGlnse7Inu6uK} zJ|kw3L>Fx>MN`Gcw6bZ56tjC(wxsQoidpqv^-t$9!+4IJA#r9EEL@+-UB+IwvRoQm z9O!dy-M6a^oC(QIYt_3i8U3c`gB_7|6}z&6)(c)keo0E!Xud;!Oj5dc6THtF4_ zVz3@80%}>s;O)xejkjzZ+?Fv>!deQaT(B#G6eNcbz5A6wz#H=+{D@1gg&h6lwfUUHXJD=XdfPRYacue2NH+OCeb3B^+T z(_VaT1?LBQcu4xYtVJTDog$(cf_jsUEN~}HEkhKOxpADvzfSlDdh+oVykq+X;Rl{7 zHAZLx741IOh3@f53f&p@MWB9749w9{v|{A_^1Lq;f=_e4%`Wh8!4Op5Wc3QcV1lgG zt~6gOw`B5Lf3O;ei5q)iJB1Ll90W#}7*t)NgEi?QdE4FFv3K_t*6twb ziWDZkIg-15)j3`qB<@Sbe&ysy{vbNbEZe^hMEgCyWFz2F7JBfMy_laXUb3)+%GJ0> zkKM)MEaj)h2E$|pQ*VuCtd2%@j=12Tg(6q%yNqR0nJSY+k30Aq*;ip1H`?IDeQOS3 z_P|MrKWriu$lf6XDQkL%T@9j~k*B;4PT7c4d3`J@O22~BUs6iH0?R*=3Cmy}(a<_h z7%hWFM{!Uboz7JiE`4BDXNKbDs!{V z^ysQgC6U-XF^=+Jf?4YkH7C$Cjlx+GZ__oA$gc;#gd&v$V(su5lyUhF(9G`SjwVni zE-{jN!mx`J*t-lA%}v8i_%h|UGlS)*vfgdQ!N&TWC3m?^l+LMhStG+UztCoxhhM!i zMl%z7+b0#}vX!YKk@!oFqdJ)47T5Jp1^L(-eI+h?0b4ZxL>V2S7 zjiiZ5_fgwj+QbAw$>moE<)mG9!~E58$!6!>m|h{gvnzgwcUsi5XZ5tZk`@rzvo1A~ z4w?V3`nrG0w2RArIFDpxojd`@JzOPa1uPBAQ2YQW3C7rMukYw$OPgA_sTQ6Xj!95%euiNf3$NGsKB`3 z4#|j`zFxU6$6}=ev@wcE)Qdi5K82$mOVPfw6@OBn+aBQ6JkEAHfxKu@Vg`=qUkyYl z*f+{JkuA!hG&_ry25@h};?zVVE%Vp#5uqS@vvUGUhyT0*JVh|0Q=D&%=v#4dqlB+q zo^J7$eT0OYve~fy=;WICEF5976!Ez$GXkC=zbE|yuzKKmi=uU=u(sl~roFvbOxMBx zzL^1^4sN-&dh_=-ARql>Etg>D8Ou%i8J{EEn-d+(hYRoHexsj)PQ*jrMz_1p->%{n zee3cJw=7d|v)O^3gjlO-Fzu|Fe?9v;t6E=i)YgVa0ws0aq<`Gz0&Ol)-5@(-J)E8B zbJ?GW=dq6`g_yH_92pVNC*8y~)PESrX`s96ema*2b)HMG6Z(t!l1|)lWCfuzj{TEY z{j?%b&gK!3A?xgD`18>+WCBZ@<>hiP`eU&7;e4u05+&UZc`^=l<`H92Jt#VLW61pI z^=N(YB;?tn+)c6k8}j2L>-|Ez0urUP9+gjRD@hPtfi5M8v+A~Bi-6%KgLdV ziN<5fAGh)1uIz>Pdxy;w4^vN6H`5F2ULE6<_)|EG4QVRj)CE*hEPV|ZHF$=Fb^|5Y zIy~bjVWidMLCK3BtUSqSpla~Q%EU2N<3=5zYQ^}5oGNzaY1HfCgOODTY~0D|hKyM- zLJg9QG%5FphQIj9W?N8|Mh!J?RO!;FS%wWcUE0jkEH!9PNr(v%59<;h?g<6p<2 zRTp^BQY(gkkEkUwE-z><2*iScTQF9p;}m`Qin)&)22Nl}Q#Oi|CtFP)b%$aMkX(38 zr(P1jVEcZ$a+NxqE?GV#L6cxLYUBl)_UB8Jm-^C%4LS@N{{nVosgtcH{{?)PELLDH zMKRb`mhY0Q+->c`5{d~&ELM0fMbTZoQGvgnFI}VVfr0Muu1AJO3NmzOaTAHok>zI0 z?p)>bz8Q6>6P=yE?G+$bTV{IHJuZQ$O2DN(n}S;FE%@Z*f9sEr$3!8wZ`=&nX$Xx1c4v#2`4>_ESr^ zymoOLD=gHn9WP%bS@Mipq$zB!D_nEep8=42pwo-_gD~tA%Z`W6C!5U>Tv2Q>wW{uF zd{pI1#lS}2{4&48{q}HVx2^o_zlxv~_xKa_^b;xKT+Y zqa1Oz@w?9^7M|OaT~2akU+yIz}{^v6Q@T1`6giw7?pY|v{(v69EDu%7jHhBp3@0P)_2 z7qM4_s9sf_%}c4awhW=hW7NTJRGHe=>(ccU?C~0*r)!-tS*eIgI=_ZHqz$#Hxft9V zdcaCO1!iIwiDLq0x<>QurVUi(M(&eVO66F#;KNTMPd^DA(LOra=P-Qzojf(_tVaDF zH6~i@hOCBy8M)P4XF+fCf<^(R2`^0peB~~ahgR4!@0S^KxT(1@%;c^g4Y(5u-M#I- z>2|kM6ME-66Mn{FuD>0O7wy7BO~!v7PlZEHW*{zC;p^6YdCzFnuTyh?VJob{eBH;l zpk|Z!?rb%vUeu(F)aDSOPN0!Uj(gA@S zd7u@rfV`6oZD58ozJo2+Ue%j%y`SXdt*4cp?Sn;6Jt-{@xCNjE`VgumMpQ|sNO zwLON^k)ldFvdXY5LH&5;T`{syBzHb zwNa%@(L$lCX#L{-WYF0bUZnbWmlCW#K@unfsoXdHuQ6n&D2DoE@#;;sER#-q#KxDB z0-a-IL?~LZWk^m5j>2f9e|(wVhiz_`wK;V%GS%MBAwD)&tU+U~GGB4pg@dLFzd-I5 zj~sQ|PsK~62$kcs2{HD*Q+AMx#@iRH~l z7h|EG&3@PX4;r;^xu8`J?L7w##Z~~ zV8>G*+_-4*S62_tOFb=P23O^tmo5@n3}zOcOQ1!PBh6I*6oe$lKdUoqrJU&9HrTe( zwn$l)rzzJssna7y-i<)&;2AnOw4l`m0C0sZ()?shpBOTF##Kgntjt&^EyIPc2 zd2cyy1LojqqW!GdB+SAx7K-fp=L-RJUndUz@cZ8UtP$;W%e3W25i^aa{4R!W5Flcs za6*A(TINAjxH_BL$X-IkV>+Zz@xT_$JCsC@HVUf}>#FUQ8XA*|IQcA6c|;`&9l_-c zp|LPU1+K#e{nsf%M=H2kjlLi-R$BgpP>B+HlL0yF#+^DkoV9AV9v7Q4SuVqNCQ60BC0 zw?WzkBd;yoAxTck?7NtG5D+1rR>5)2r#n_e_KM<t4Z8JLolbR{JuiolwT%(;tqUdKM_r46|1zCf!(gsBU+I z&Wm63HEfs;>>Zx#KM#+=RS6`e#VhPM7)cj0YM>>?KgY*`Nr7dAR9q=CzyHvVJH}ko z|BOrVU8K#^i9>w0xq0ZlL_XbpO5`Sg?VqGGf9pDUn`Hv}JkkZNxXV$0c3r1L$o1$E z9EUY)F3updIZ1AxX6hf6x@MKUw$%vN=8tFrYgjv5@@-@pB|qi5E&Kqo-wSAv8*gzN zBjFa}U3z}?`yu<9&GYNMolw~JZ16hRN$94RPe@7Go2^`>GnpP_)FQtfO^t@3Lp!

j10_E1wk5$p#$(ov;US^kS5j# zweOxBU^oHr6aa`n1Q^Bu71p&)D%wA$ja=|5T(6tt31kMTVp?s`E&$RQz-}JUK5mZ7 zPiAOGFMn#N6EGy&to-==&Dcb}#q_qJl#82X*jS8f@387>Ph;^Uc{8RV5E>6+kLc9f zAV{EmmD|=^C3YFhHN$o0e+Sor+tFL5auIuy#8u(G4_B$(Qoha*8dGs$Sias?=MKhI z)vBShf5ENU|KMqsp!B;g`GGnbybbl?LfgZ`ix({YCQWFRlLI&BJ$up8KJn|wQ@X17 z(ByBYp3gr2Op1U6tb8RRSngOiRKoxO1!@2Q$Nx1E;VZY|AGrvIMql{}_C{Z+2o0JR zHmjUS@7X&1h228d+q~A+wcWg}zS!TRYNBzMq}$#NkNB6r84k}Aht)5Ab_$3$lByLN zCtAlfqIpehkHzj72=4c;PUi4+c?>2?&^hMM&y0GRU_vh*#`Nq;d0##u6P}0m>&W`;>Th-JbMeq z;)-j21acdpFQ?bvx(%~`YR&IDkSF>R>EE2bH=T(Yp0ZD3G^dmPbYlj@q91hrHvDv= zV9l=(oI$QtH~BROa?A~;Quqg?G0Y2b=(FXU?d`n$t!=~L8KAYne&zVvt^5dNy*sEwLs|n@T??c(1MmoNn zX|zjAoek)@ZK-9P5-h9f5*->A&Q2gHy>t&t+HV-5l@=ZhaU+YfnAmk~+m?Dd>oFlr z&Mtgy851=0`k*7o9`-x|;T0Mdb*}KJ0+{ukf50*}zYQ9wWh7#61Yk3nQM9FRyr(Yf z*r1&mbKo*IlwDQztdRZ6BKJ z$&k+N5X=fWw|9DT;?1lXeNIPwn&D+c zBZ7t^(51RTonm4GpdgCw$-X{}BuC`x;xzf>;L@NSt+j{3&EBpr&FjS%1m!3eBPz_t6+-OVVOQL|7hXNFW*0)UI$1a8;Li;^FV7IAC zTRWcWCZIQ^s)a1^+*cqYIwBaqgHjv*6Qw%DCw#HDOA7DdM&{c6sl=7%O@TYpCkamL2d$FWx)cc*i^KvY zN7hY#5l--Bd<~kczi@BE@Tqe1#_pG_e#z`WyLNof)Shp4!Q;=#>9+jCEe+a8DVh`! z`xQlvF;1FG(_XlIUN~GqFM7_R~Ig zsjeCdxze>SHJSSN{OtWFT2*?pIbpRhO}=qx=(>f&U>Ft-7Z2%ZsMUkNOI>-2h58F* zUb&&T$Cyd&N}X~+!;cWlp=5ubF)80+eh)+G8`*vrZhLQk2WcRe{ZX+;Du0diJ#?;j z#7fo!1x!#3oY5^vHj%wj7gpuWskNz*ys#9W@>0M(FMejMxzbvHoTmybY5=?@2$6TN z5H5X=UK=qI_;+VSlkfDB)mvcg2LqaSP*Qm52Eb!yb)T`)^I>@6RJO?e8eXo*c0Q~k8zTP2Ko@IuTik0Srbc%^Q!84u&^Z|gL( z_0EQ1Qnp4k|JC!N?A7zK?DG32JlvM71Z*#LNz(6ZQ|1Bo7=uveTEu?McHXIn)DAYAwuRR1IUHWcohwH@-px!@JpNNUgb7ua*1Az2zqb4Fk!(+igMGd zhxiOrM6aT!*hVL{5Me7$qlE@+$2CK0T%A7lCqFg()|ydb256<7*(|(DG+^!{_viCnHe4HCS?Se#h=~>QZA+&5!6e-cjf#+0O)Cz)|}Z zus_p$1akn}AQ}LF6zCaX=VwLm|68tYh6MW&QhZvncBsA6L8mF@SJB?VA8S}$F25M zcBRSQT~&S~3@5}|Ya6xHM)kzuVK2YeZDmJG;zf0^Noa-%K4YPCL2q3zTweKGW=7>l>C6b&_ z+W;Kg&mvM0^~g({@2V@S4>)y_H1Oe1V+vjW5+yWkq!=oPo=lPVRoH>oDZ}CID1+wKbM8+tQn4eN0gY+{2tvXq;+xf{u$9NS$q?)Y$QgH4B-7Dt&{uN-fA1#3>qh$I zd}(!3&)v|W5_l2_B9_7*Z)t-oZOvnx$02e4)onJb$05AQ1dN0%1|IulNRIaT%INmR z^i}K@{?+*P2`ca_!_~MA35#&6QkB{Ux<9c%EhKlrbIozwkE`oyD8!)7CQ1q?j&bwd z1j>oBIn}b+g9DJ15)>w()@_iznjuAuQ0sY6wm9cjrKRYD2d2AW$5Q0p_E^|#i1Gk z6XqP!-KFJxH0An$q-Kim;C}2jbc`o|_CJ{kI8 z!u%?r|MqV`#M1FG?AlrrhoHNtIMr~coIm$S(+j2OHS?iXi1L?UP2*Ynt-m4=cb;VC z0nKvofq*kCepf6(tGI)#O*C!q@+Iq&|FY#S7&C{KmM0%K2g@#%35)8Py`wBJ_ai?+ zJ}-$Z^Zm#D`W)tRF>Z#BLN%%Fum|K_nY4M7%-rlY(DOA?N`(jXFApVWbSUjg@31kt z9)n9E@C~rlv|ddDxW22~(p~4WgU&9c25nqBwd+|IOBX&y6wzHQj`#!=eQ%x9pFsaw z1>oeRNc+!AH0eSC0E{F60QdiO6=0?3XsTmuWACD8Z}^p){Lh6zrJ9Y+YCEEj&DJxV z;ETT4A-l|N-6ZFaup~iOoYk7FVlZ&XL_xJRLpjoY&VicCPA>5}3Xyy*OQESfSC_6= zfooUJ^}nSfvs}N^7cH-BP7z%j5V|M|d-v(eF|vpcv7|}Jh3e-p2?>6?L3|w+ICuqB z63wKv_kxuLxk2KfUV7EQMaPN^~+D~K~4U};KLG?ucWe1l+CL>5Jxo*P zNU6C@sy`bYVo|+QIo3@-YaZu&XV&&fET=G3x+%!GjSLqmJms@~YvywRBD41%?ZHt+-C3fY)uCHJvhCGGZ z1ioCwWE;sw>>y%U#pWz}IH(@=gtjeK2j|XMnwsf-$)bvAI#LH!^H9nx@H8*6N zd#_Mc_zm;$A^3VQyP-9;iuTu(uzG^qlzz`f$3I~V1Tr1Ry};jECNCUWfO~6%*Q$TS z?k6y%<+21iXcI(&af)A<++M<5UYiMJ8zqP85_k z7y}9e1czs&ZeI3wsMMXd4!!rsi!mnOEnC$A9D%mz{G636WYRz_v9Sv%9dBteB%ypNBeXLDt~-u9-sLEKv8% zULd=)2ckdEDzZIA2~zQ14bWAuOrK;r`JMnh+y;lEY)=%BFp}b~b#%(Ev^C&*=K=cS zvG6e;WQS&IT!P6LBIvmQa$3A^*4Cm03QC)&tuH%yviiMcKikH8ZYs_&T0kwi>9F|p znEk_TgG{;LOe&!vBpEq~_LAZvajAf^*XwctK>Z}Ty66)vtiV-sCN812KrobwY*sMv zR-hmJl;22_JZvpwkjg9si=iKqM~P+rG?t8#{+#|%OFcKA6R&9>+TAQcZ%i>8CQGJ? z%(Mb*LVz9G&&HuBt6jL|E8DLlfkAH~8*DFH-=&)i3Y{(q)_+7$uN#fVX{MhcyoEd) zcdE-^dZlPdir6-G_j`bX4Z|==&zu>0_T%%Ehso>S2DKq3zqni8(!eW;p1Yf-7w4 z=*?s~pQQ6QF$!jt&w1mU3d28_ti-6UZ6FX|`$-K?dmXapiWC*vgl)Yv0$dz<*-Lk* zla?n#jTRb}7M())6Wdlrd?aW%G$o&SNL-oY*oLbOm48gVySy=mdh)B5(X0j(GFaeo(&hl zznoZBTieOFm4A);9J|}o`sDaHTaaeALbo(8*zS7~^W?FtT6OM?c{YcEVfrS%*A8%* zXj~wYf+l``s(cvQ%64!(=DuM;%vNVTpc8G~-6^SsXPB=D0?C5l1GXAW|ih`N*} zF|j53xOdp$%w1U4^o>quT${R>D5RN&uLXm0EjM_0d5~`A+cgqcUw2p8 zCyR;>S5c{?3`(DwkNRDl(d#uACJ%l0^6OKAh7Pzm6_kzBsss6iZcX-+_ubk-_4fRw zk+*nHOE7RgJB3w`qmo3~}<4Wt)CteP|q2~NPA^c#NN$_IuedB(vgGY;W^P=q4qirKW zA#eGy@$*RW#7m@PMuku+;P-LgCi)UmNpgA5aR=vbO=t7o@4q_V=bA7(jUn2(SZIJG zY!|L0y*M58*&Xy*3%peaG1J$pt#bO7AQ;FG=*b7@`+vQ*9(7b(fVPZV?uNwcoWeV1 z(VE6ygnfh-mUusPgnieFfc7i6B3vfrmbBf$h05GkE>MUg1WcHft5xqHa(ydTz*8KJ z(}kX2cq5N?gc`p)O2y^HNDjEl!lOz-b7r2^8M;%Uwjx^c_n30s;dUi=^Xr6m)c=ku z;lqgY;@`?)MAt!W(mURcakYdEYiAzSjJZyqG%HJtTYL(s1iwUxAw-*D6Bhy>Q;4(2 zil8DubR6mzPHgl-B1=h7<%jSXm6+JzHv}_Y1F2u<>n7Fl0Ys5(deewE`J3(sHr& zY9V37u2g_NH%qDco@rGnnjU}wvng~ccjE@lZZ)4Z2aD0#2vDL79yCiHur`HuZ7gqW zZ+oXN8-f^KOQwN$A}1_x(h|t8tN`p#Ws{S|wPP0I>Gka;pXcN~&Ga`D#PuDlAf1k0xQ z9cpO}OL09nT|y~lnVR2a5g}xwq5+gAiN6UD<5}*WQ~;Jh%LU<(*Uu6|racck&pCU} zq=|E$(Rs8L%vI)h;kc#%wCOPN@;G{51?#AI{t-|uTp909;TuAiOmz9Cs|NFdE-L{9 zi~{(diEm%ix3BUSaqX}2cN*Nk5dV${_jk*`5Jdk#{Ec+>MOph7?%xsP{)H3&FSviw zT3(BE6P$mBCB^0N`K%0oeW*l#b5d-cZND z$xu(n(8$@$z(~i~$=blt%*NV*)^uFU!(n!)?vqc+GlU~A*`8f`2Q<752BxX~4|Nqb z5*1X9hSO;&ibNe)-QGMg$y)tk{pRuXk&%Tv9uiz;FOo;A&wFZiwu%IY=iANxsMkw~ z>H>95Zwt#P&Z_+yw$FjOHrj+dA9O+h*fKdezIKdQX3(!LN3r-yWx34w_Y-)b6~dU~ z&tVTX?H9419AP=JLD+4JE{1rkdMBI$JN;D>L@{G%&fG2`BvGNb8*xe4(pG$T zy&Sc})avU%c%;(kD3oFCcjdTki5rB@@q$d7#j$;1#>+FX3fPIH+4#*+n1wJYXWI7B z>pWzGcIN=Vd+ZN~Y4OgNW1P*J;YyXd#Lo>o2JD~9Rgxg0S zZAl*=5220hoX+eiI-d;L_kpe~X?wG|)`Z#XyY<+1ye{ldUOml2lu%5l#Fq5E*;Ob{ z82Q1wPjZ#nTc(=By|=Rye{URl2f8Zp3z^x`;n+yJ;*_m4HuS^G)S@k|5yuHziqmBh z2;fOJuF&}IL->>gq=feDNEu3;^UC^rxc#iLdRcg%8*0bq`rE0;pM&?N`UTOPuj~GI zMjhFnnc0#RZ1@wYA1MnOqbqKr!*5ecxuCMi5S1~%Wx9d7;&p@zrY8j0qRmuw{-pZ8 zfw)bU9Md}9Op6SHNFl(Dl6=$sO!0}Q_;vwD@g3K7V`5^PP{~)K!FRJbjVTu-lKNZ? zXTPG}w58C1#!dl)#(Crx3RS$Qck3Bz^}8r6;4!JJRuhc2aDaQ??i_B2JSC8Jp4`uo zHDG|kxs+n4pDpeZA>|+y<>n`v{|-Q90^IUR0}YkK;hjFWssK8mO4lKl+@?%I-b{dn z^z~3=Zb`{Ar$qPM=hH0?jy6-yWVu3v$Go`hldoz`_ov&IjvXE+svxqCLpW`N?a~39qME4YPc|H&C)Y9p3w1|@3|oBXai0aV4&o)u zFO=;qYU*rqCydahUi=@{dhnNqXNmbiJTS;Ucy%o>&DFmWDUyEJ5rO?)>cFhGvd1POX47b)~MwF240t@G1rwkq+P2O20>bIF_r!)(+-Q~eR;_a{H<*;xdBTVlnom$pWjM> zY}6UY;vyK0By>JDuBxR}DBKBGu%2$E&iI!V1zn4Tyvr||*g7*P(3R#pLYT~(!_#q$ z??(}k34oCF?mQ~nVyMhtz&^8}AH{&PCpLB7gS+rSS-0^?%_g&MSHqtJ2P$&Sr#l1WT39-AWa+8Q zaPOs~LRMz!5>Rxio*;AEG$PGn@JA`cDknPyohjQL0@=tL0))V3fFk8Uzz>8iVA&!_ zcwl~4m>)?yT#$Prx5kx@#kcLwYadlb3qWGK>_jpp{nBh zb3Dz_r+eZW@6wJmS7bRl!2ZzUS4Ia9SpV-#S|xb*0w2_dR`D|8J3SfB8--f?NG11% zNN*YY-~~b02HAeKny_@Uyl@3wg(N9ZY%wc%Oe<&K}yM8a1Bd`@vnQ8{}gLuMk< zzEa7Vxj$TwJkEGO*7u<_bNCSW?PtwLI$UKn5v}Wd?{*A@E5x(JE#-tzsq_d#*jUOW zp?)`oaRsdFAO)A?G!&O-g`f9RNR-OLRR9B7iOQ<-U7LeRdG4uE- z_|i*fL(00T4k?{pBvDoO>5HU4rBGGhNrD`#ge|yQHyFgD#QK-rO(8EvdJz2<2&R69*t;O-@U#IFkdHpdy=o_hTQgOwzb~WHzKGXwbc7*bBazgK^|V z<4m(^@JkmHi~K%Sj=hCVQK(8zZZvtWIt@?GpN~0@r?S>^JC9IQa*O{kk6Y~YR7xU0 zZbEZAhtW`S<9s(+b5idxUeiT=P_L4IC+l;!)oIij`BiQa&sYXsj|tZy1XE9 zxtCJhU(ebR>KKeU?Ptn<;_eFbqw(x7d6xQSS^J%L)J9(yM)O+CJ4gzM z^C3bn_SKgKMH16Vi~f+v1;~OT)gvy$7n+oBNhMN$nB~V;3>=9IKzixOuBtuIl%Nvc zVEpW-XI4uMZV9^>CR$Pc>3VM{70w^2@g9@Klfh?nt6U4LDu%7_m-%@B7rzlisawLE zdN+D@R&~Jbt(hO>w3IJ6FQ29|f-K*g3MzuR8~GgAU+LMA9WAlO}V_T^FW$y}Um^{QnRl*(cjA*Q!=!v+y z$)>re>Pl1kmhc55g25iO7KeA~?svX9P`2BoP7EfU<6({%ma76)_?02Km7kGoh2;A& zqJvURD_h3H;yRIPW$`1wa&lkzWsNGz2+>u@6HVxxmy5~cxMn1+d5Y484mMbjwI>}` zddzy599bsOz>(V8S_955s*+*on&{@GL7iRwNB7AP>omZ6-b{I&>1*PJVPTXi3gvo7 zP->>UyrQI^5_JN=2g}M#`&3b*u11<#eJ&kEF3G)*@(_M&4a8_(-#Xu8dk^F~{kbI1 zi<%FpbDjfTiB9icagxma777Kg4q=RtC@`^1v98J`Z#t!)qQ}xB3-bFOxE_{y7GTxU zYzgvS47B2|r{+jn^c@*q{hI>0psr=FF6e;B^edF z14Bk$I0#JtuUfL2+ln=Q$3wAYS=RO5)ZslC4J*5>EKBO<`iOCo-j6}agUZ_O3F!-h z*7^)o!xT)4^ihj54#ec4{O4NWk=k#9Kp;L0z)uh#J^qn7o1+KCQgC`=TmVx(co9cu`CE!?zI+N3sgmI zQ%hM)oRo8MM00UCC#r}is&My$h$q9~C?cWCOw+u64Nxi#Q&iOwjVC+V18 z#IajELYL=D+O3rBrBo8~bu~7c`U!M;cPxgN-RZN**jz(v0V^vCgWJFJRkSPkd%`gz21zZx7o(XbaX?9Yre~AutM3d@bHHmNI;?Z3@jO z&j7qd`O$sFlC0_iT5*s9Kyxh*3YI0R@k2l8rJ%#nzI^}_v^H6ERuzX*aFUWXtGfG4deBEf$Y^o zh&#E}rCu!d>1dft>s|R(o(U|N`>{?3+=bA&H>Y}aNgEX}Y+IB;D+ogbGnMP>ZgWE8 zBYkw#eGqgM?Yes6_f*IY?FUfM+D{Jc zupTS|YFWkL?aJegw`?5TmN8MnS_-FJuq%QTDgs7PNaV<#?etU@id<=1yDai*1y{qM z_?N5hn2B($GZ_bq$OM>n;RT+s62dTu+rsFBTj1 zdq~#Hd}N+T_$!znTjAn2q7uvRq4&6k2gNI1a+g&rE8X-?$;0%ov>WEyu8y|}#Zvpz zUVLr^=LdUuNcy|1MIxh}BBB|BdXtSTa3@VILll#_ah%4#PWT3T^6?eCWBUc+2c9Z5 zMrZ;R?LO9p?(sW~IAUew|+rJJ(`#rv7Bj8aMdhnFJn4c?Nvap28)woBG z-NoW8<)_95!(;_hZ;fWGjz)HlxZt0KB3JCYjAc@pDw9NyJNO&fS78}9+Tg@}YYt)d zz)6WeY$6rN-XQ}iYkG%W4WgWpr@RhM*@#nleJm?kGD9^Pt+4Q6^aWF16!0EXY^j17o1-G79!|`g7Wea?+70s)=$cbF0OEuLr+`B9#PU?eH0tarqC>%HkcezcJ&Z%=*Bf~Sl&}Ny3U%fI$ zGZT8-Cl%$gm8l|;_)CtXI+)@X*Y!^Y`PdqLB`$jbTQvVf1$jk;k$aqJCkS2YeV|i~ zq=`xQQQKYG!~{Xf=^|ZT^77*F9E;W)4 zng6i*x_`>Fi_3mEk7Q(>JORf&TqR`%EDg(0`~WBk#@KD+|Enw3$gv@uG`nx>Y{7NN1mXqv$}_rXS))qJy1^uWv~v@vz_{QJ z$%vZ1Ub!#FVx+%e@EK_i^*@2#fSgUC;?W~!9J^MPVT3>P0)`mv{C3W1Sf86E*Z7xyWAUk6{oSo=% z*`J8#v5zN(n6rHx84=JY-NZH2e;CMVpu6gRI+q7^o=dP3`iuFJPTX;11)(yI{gYSy zv?5T><`I!0>+ERw^U*S70!y3a<#I6kW3cz(e5y$%`MXJjrRGYVgR)#4%RmMjfDP#rTGtDt6{+)a&7ckyQw6+{x*Nj9D;3 z4U&yCDfftmzxc^!TTqoo4K;35>C&iKh7CDg+RW1|HE2&shwhltlnF!S$zU4eU&oHj0xcTTLHzhhhwnTzE~V zUJ}1x`+mA|l{%a*Sw18|lVCM!t*R$wkg zG1yj??~<$BZSBGmiU~(7R(LK&(Otb!fxn(FU8C-Sf$s3GM}|fUGIVEg6N%1|_~hi}Waso`eR;+Nz*^pf%CdYw z@x5bvyNg-A9H)OsbUmwVI7{NCB+}{Fb^7#_2&@QG!px{t4okliKR(Xmf(Ox*Zq zRin~o%Dqah(Ak24JS(|LJ-qO7?{{Bf$)nYXOPO6Kxa9M*SW& zCR*%(+gF&uG-IQ*(e}E3CqN-N(0} zW|R2tY&ED})TE5m<`AJ!d9Bj5?AO5OQ5Fp12uvqf&NcM0gHbJ0qO1N#mdn`E0f8KO zpcSxyyps%VV1_fkgDuux)thm>pXB7Nr7U3np!BVl!0}R!rA3fSSCH zN1+?N`t7y9+oQru_0{a;;l}x(t1eo{H~a0!(DPt9>U*6SrfN%Zvb;y5+;hp#UW&?f z5liy1Js?!(qV8vHk1^GfPK9H{vhCKrUkMZ8QOLql_z6M@`<>c-pWZ=R+*Z#2V^#hZ zkAIkaMC+6{79bBX^5)XkgOXcGaf!%!aO+(AFcnFe(^PO1V|QIiH$h#5zuOa2>)oZb zJ%-egqDnim%CIa!{dnbE>j}>)sehfV%a;Tbs4VEZ5jo`iR9=B7jOCYDAc$i4^ zTb#p{D<<5Vl_T&{P#q}J6%4csB%F?WBdU5sen2EtqrJYZ8tQo3SsHpb$Vq>+L$ z&ccH90Xm=4E)Q!{w~1D*lv7S~`PiyCqCr5QEoe$}d2!;R;TruWrLlXoAxFcz9PJ9V zQKd`KLZPc@{o?&((AgGVr22Q460AN!5-0P@yRlTLfY#+Q-; zonvG~C|a>)NKOfk!f2#_e3{;dZElyfIdw8J)!xn_J~mgZL1V2lUvb)ngQg0r0r|<$nPA%AwiR1k>W=0Pm@$I9D<;_Rr z#9EGkH#pXUyALu5xYLrmUJpwyNJAfDc~7tdEJQI4=krq@49M->dwofAgUHAG<;keb zY`uxhBc3ImW_Nu4Q-3lG2fwU_1&Tr%)$Pl&4ac_C3ytlHu(L$rVc=7@9y2zG1}GWs z;-58bR^wjS)XX1C8E#xi3Sj&&KbG{hY6rjKPnGNe^j zmO7vM%A#>+rTaJu$zw6#nN>w>fia;`;CEW-p;0k3Gctc>+L&(i-+lc(=vukPR{P~( z$5S8NxM=ZLR}aohJuPDfSLL3UE)rP`W)_`Gphc1+%~bysge1p5t21k*oao&)*tXKP zNLiMrDc4In$*}WR4pg~ zE8LDpZPI&r(E#e1$C&qvnsrwOoGrf&GhzoNTsSn=VQ=6g9npfs#KY;14j8VxT9j9L zZ#i!R=HO|f{jAv}%)&AjitPI53juUrCl3Aa``-Ml5$$x#wB<$-GmWSGE{1LpAY!9% zLV;vj=0R1sI-A?bUP8oUI;2qXz!uCqlthj;3ab+9s_m5;8k34R`7Ba-L?sFx!Q~90 zu`opiuEPfX*C|3rD!5sVz92B<5Woyq$kWFFZ|XiZROw)Q@qpeuf4O~6NYk7$Rr%Lt zGxTT2@iIbU)bs~MYxVmA*}@WndgGI^KNmI7(ATyjP&I!PWeEf*$^Tp`mpFOwWdiy<(gm%!%Ta%JU8h9I_2>~C zhc#<1&LFfoNp7EJ>K~Q5W|h3Q)d<(-k7xmFSUX$tZDbiGKjpeD`~b4w3uurVZ*dzV z;TGaudVcl$A^V!m^Xt8xP}uft@H*H@=%$xXNJ-h7tz4xunI2@+BEKC?jfSB^JGuH` zFdcb1Ku>I$SG$3ZawTQ{Y}WaAI*ZL^tPdti<@QGSd44$u7*x#dSqH&j`+ujY2_?N#K4$l*Z)h~T^3Wzt7sudb1 zTE{h_c};AO#qJmg?)R=v=J0iS3?@s^Ip)vLjCz@1LN6Z1^z2IH2-x+^K3Y6P4@V?| z+4aBy(Dvd|!q~-^+gx%T2R3@(+jfVA3n>$s;_6srn%fM@`@$G6;8kO7Cw6?@EXdWa zfkyQY1i2db#h?S1#&HnCO#5*T`d72&Kx`YqidkecA9VgU{B)vV z&94xgL9SId`85Y}%nhbe_y?ph%nNbov*pGnqb*N#5##{o*Tn%QN^#xPZk(YX8$Ll9 zC47<{0Z}*5{@t`NSOTh*k!?Mr<}YeD^}&o*zMpESdt#ud4nnsEkU3?KejaJmaghXY zy-#8Dcb=fLL~Es z#~by6>2bzW_<_iER+%VTrQ6r99k^_088X%An;}dQ%MQ+L8Q%3G3!z!?hqhk4x|`?y z<9FVF01VjDXnad|>9FK{OI>Hfrri}jo8uMFnNlY5kKH%VY|Ap;H@-)uwre%b*YE5fD!H&3Zfa;U9u%pCpExrADZjQ zkk90l9~&|Ri~==tY2Ykvp0}{*_NMWq#IdC-k5XJywXPys27eiZ7q9MI4l5U^A3mq?U zv*i^&dxL*ciZRj$DI`2~P$(1$RMVxs{i=vZQz?)`nq^lvU9z7l-+6#1W(7YE1k(Af z8xH?y8pKz_PjL-8bB|WDl{=_tM&46UOa*F>>LXYA8lnt(^ZMOrr{HCn5`M8u+RaU7 z119S*Vo>{oQXBpgr8>kXe6hDn3h&`Y=Gy(K#FghwfjiSD2~O(=t&-Tf6bTrM!~!Kp z)=hp9PVi-X4VtXKaBsu#sdDqi?w727$?QP8c6`s&o^N)+?N=F;P*fB$-`ahVJot! zt{Mut(zP!&nfmwq?ENQNReG~IVYM(#zHw;ix`o4F7#0o}59w&A)q}rFU3rRy`U_-U zxuLknm`UwQopM3Lj}XhDWPhJADc@m!4@2o2*?t#pdvAXSX&{&VQL#rVe~t7#bgp;A zO4b7fOi&D*(Je?ek-btER^`m8wW*Q3uoRy1QoubgerBw>(prCol|V z&W2!8wnjAn)$^k4)$_9K^7|${+?K2aY%g_5((i0j<^lE?gHYyL#D2|o-n*v3&m$(i zZ}oO7B5%7QOm39yhpm{t6f?ZU2tq3+kolsok$z}yR_xl?VT<5?Cq4xNCrz6#ylzB! zt6|6_jv4m;f!03J$*yE{J-z~7;OLg1&u%$;Q>S5eVM*#dwK7_Y$ug|&v853sH)+D~ zNKQ@4#UBWxZ;aO=Lgpp|$c4^$6%BUsGVmzyOPlFla?`7a z_zY7-ucD{eMklopVJl9fg$8WLHA8A#oj&#_KQ;W;no(i~Xr-OmEWAoIVD2OL>nKi3 ztAy9Ab|X1SD3OdMIYEKOG+NdtBT&9ISaNZG$L@6MQe#lfkLWkvQRpVw&jen;QTr9J zKht~!a{${Q8WF{~V^;eDc zd!Obf0`7L|AyVEZt4*r4KtFpL>sV4nrj%800lJE)Vq))=?e~`@1!Yv}xqyVnt@c!Q zrODo1RemE3FEb+yo-cVfwcP2I<%P`M+7E9F#-F3NM%NOo&Lb}l0IoE(g69z$Y|$$< z57~X`lKfL}L1|FcNg3uoQ0S{FHm5+pOY9$<>n&alvvc4GAi9sW?;i%{h)6yqlAKW6 z036)UB2p3c$V;5>sw=AxICYXV@ZnEmC55EqaE_TfI_JLa&#EI>ubEr@=z{=C213`y zDZ>1Y-gNU$cpTl?F+T#-`)j^8t8YC`IKdp&gT72h=#9g{KazaT=+;kE3R}>w?l~bF zsZ6;NB{XfM7%GRJOp*6h*n!t6!{O~HgXY$A?oTjMxWztj1oxJ$ph0pxl9rkznwN%) z0b7*$&F_MtZM^jXjcNl3Lc%BFo7JbVmB+Zr5bOrX8Fy7A)6s3vS91Y>?y2C38R^=>&*1l&gv2oUZ&5d2m%HtoHOoqE0WiWJ`z^VSrPvbs2lrkXz2V}!==qzDgOytO)do=R_=bRi5FUQwGmHlUL ze$#mETcLsc2|7d}Q5(OoQ#_yiPvZ=KVhK>u0= z;N+!9`_D@>=|TYjj3fX6_y2VjV5R41s$*gO;C34Xgld>s}zcm-7w z&7`#Vf|UiiLG&KV9EMd1kBGDtLLyVEA-52#_ZJdCk{CcBl`Pe)C!_?;Th88-q>@G8 z=XV>ZXUQW$dB(f`(JTrO3u^{!Vho7{Qjz|Y28tIXTC-MFo&uH{;09x(iwKO1Fe@g9EQTCn@;^(LaBdgvYa~FhCUgs>7eIiXwF=ChJ4>{2H3FR{)gcf1$?dyA1qp_W z7atYs8rhW=O}7b7sWKcZh%+5v1rz~|E|8puCg?y+B1b(x0R%~RYxz>zKP4MIOjG7a zskuz5KN}rlQN2?+)=fWa9_M>!*7j+r=l8u7F@hPDE0O_4t&f5EzGbB!^~#LEn-Mn| zF4-_2P%WWiD{^C8@;zDHG(Fx|=@?{}TYy=!k~rZgm}<*yi zzFfs*8_7oOAYxg?<}7+Rs2=o1UsStee)h~r=!gEcq{`;-nIHS?<(`ZUCmzz{s^>Os ziecjq=1Vu;YfaUM;>U!PS~f2dmt4-Q%Ar1w2W8*3bh8Z9^g~+2%Vn)KKA__@H)NZ8 zuTWL^4fF6J_ieH%AUcy{nn+atbC)C13xX?Zwu1T%UUv91={Uo}&=o2lhz*Tc5E}^zSFqDdHRxt2Z zpdb8{-$;@?Y%OGv$}9wnp&yb*iDmvYmW-1Aoc>WuJvW~duW29J-7G@({rD#ct*fw_edw_xs!!SwDoEgJ&h*Zmlps!MVOBuGPx&3PU`T1nMt<%TOh}-I( z=vn9E-Qb6|VSb*W?c2qPIX!)5IOv*!D{SfL z&15;Br1Lj13TBnhdE=W3!#|g-#Hg-qAP`^sNexeX9kS<&6cyTpZM`%CTpW4XOLwS~ zmM24v78;cnokI8%+g3$`X;{D4se-h zTp*HyCVqaZd>Gowc5poC?PGnf6Ey5dwXP%?gKnksyw2_<@RLp@ib2U|4sw}@x|Aj{ zu_gMrci<}a7D}eJPCl>g?T=lp5u1|7KSEJn1Wol^|BP6)g1=XuSkT;PJ**<(UZ?opf94}JIY>r;Y;4!AfKl#SD>1Nnq*P4<)b-P%F*_WY%h zw|GzE-ijyjt6$qJUirCdELM|VACCo3UgNdx`dK3^SJHiUFwLM4T{CjRd(FXkPC_x& z{hy9Zopk2ar?{1QTv(Jrd0M7=&#kFr!)Si8AQdn!J(4#5%0pWYMPa zs5>6%pMYK>Tu=~~E!3T0Oqi9cRqr5jeJfYMQyh)c zg`Qt{Bae238oxVA#pT6F4!Fv~qe?+@W}ejKRHB3knIm~!3Wb|rW7>x6dH|Bfo* z!-(_Z-^yV`*FkO4JKm0QwS*08XCBpzxlW%nD@%-9d3M4Mp~7Xlwsh_lCv zpdvuz``x7-XW?>KI+4kyZ`AFH?#ePFL04x4v-g^6#AWtff>k|6-5&Cu?JYH~uBl5W z3Q~RDeD=OR4#uX;mqj9)JO}DRe4#;|9%cHJ>#Hi_zK$P@)VTG)o?^HidR=EN^UY zd#5iOf*4*)rh#`NCoFK%63DNt0PIlYPI#s|Q(bnz+0K|j(kYEMd+8YIBJIGzum_&k{qXJr6t2IeX5e ziF2OOd9)SGRpxi$xTXNK=`iy0IC@_N>!^7C5l}5$8ShTv8$y>%bor*M2J?X~D**(I z0{EYaZ(q~5uksgh?XU888r;7S|BeXvcgw#JME^kijdb=!S^F36-x1^fg%kfTxPQ{* z{%7=GM7h70zoNQ+$zSjJ_o#nI|BE*FPcQ$TIQN$qAOOKvUG*jZlb8PwjqYDR^xtoa ze|i2#5H`a9V~~HhMgE_|{JUBAzryg){2#;o4|8x?2{7=#yg__jBwq~4N7lcz{|9cj Bw2%M* literal 0 HcmV?d00001 From 01a11c94f8dc685df1b9e3bee8cff833e13428fd Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 23 Apr 2019 18:30:14 +0000 Subject: [PATCH 11/17] Add COMGR relative path for build machines [ROCm/hip commit: 920fe246d7403e8b297ffffc8a1dd8430eea0b7f] --- projects/hip/CMakeLists.txt | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt index a71d847af7..5073dc0905 100644 --- a/projects/hip/CMakeLists.txt +++ b/projects/hip/CMakeLists.txt @@ -200,7 +200,10 @@ endif() if(HIP_PLATFORM STREQUAL "hcc") find_package(amd_comgr REQUIRED CONFIG PATHS - /opt/rocm/lib/cmake/amd_comgr + /opt/rocm/ + PATH_SUFFIXES + cmake/amd_comgr + lib/cmake/amd_comgr ) MESSAGE(STATUS "Code Object Manager found at ${amd_comgr_DIR}.") endif() From c01236f6791ebea5f21b04a7b6975b336fd475f5 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 24 Apr 2019 04:20:03 -0700 Subject: [PATCH 12/17] Add hipMallocManaged default functional support (#1036) * Add hipMallocManaged default functional support * Fix build error * Add dtest [ROCm/hip commit: 2bc2c46d4d91970ac06f1ac825e96c9b216343c9] --- .../include/hip/hcc_detail/hip_runtime_api.h | 14 +++++ projects/hip/include/hip/hip_runtime_api.h | 6 +++ .../include/hip/nvcc_detail/hip_runtime_api.h | 7 +++ projects/hip/src/hip_memory.cpp | 23 ++++++-- .../runtimeApi/memory/hipMallocManaged.cpp | 52 +++++++++++++++++++ 5 files changed, 99 insertions(+), 3 deletions(-) create mode 100644 projects/hip/tests/src/runtimeApi/memory/hipMallocManaged.cpp diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index ccfe7aefc3..863e6c0e4a 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -177,6 +177,9 @@ enum hipLimit_t { 0x80000000 ///< Allocate non-coherent memory. Overrides HIP_COHERENT_HOST_ALLOC for specific ///< allocation. +#define hipMemAttachGlobal 0x0 +#define hipMemAttachHost 0x1 + #define hipDeviceMallocDefault 0x0 #define hipDeviceMallocFinegrained 0x1 ///< Memory is allocated in fine grained region of device. @@ -1103,6 +1106,17 @@ hipError_t hipMallocHost(void** ptr, size_t size); */ hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags); +/** + * @brief Allocates memory that will be automatically managed by the Unified Memory system. + * + * @param[out] ptr Pointer to the allocated managed memory + * @param[in] size Requested memory size + * @param[in] flags must be either hipMemAttachGlobal/hipMemAttachHost + * + * @return #hipSuccess, #hipErrorMemoryAllocation + */ +hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags __dparm(0)); + /** * @brief Allocate device accessible page locked host memory [Deprecated] * diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index 34363689e4..e7ecede8c1 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -337,6 +337,12 @@ static inline hipError_t hipHostMalloc(T** ptr, size_t size, unsigned int flags = hipHostMallocDefault) { return hipHostMalloc((void**)ptr, size, flags); } + +template +static inline hipError_t hipMallocManaged(T** devPtr, size_t size, + unsigned int flags = hipMemAttachGlobal) { + return hipMallocManaged((void**)devPtr, size, flags); +} #endif #endif diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index 4d84cfbc51..73ab91f25c 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -111,6 +111,9 @@ typedef enum hipChannelFormatKind { #define hipHostMallocCoherent 0x0 #define hipHostMallocNonCoherent 0x0 +#define hipMemAttachGlobal cudaMemAttachGlobal +#define hipMemAttachHost cudaMemAttachHost + #define hipHostRegisterDefault cudaHostRegisterDefault #define hipHostRegisterPortable cudaHostRegisterPortable #define hipHostRegisterMapped cudaHostRegisterMapped @@ -420,6 +423,10 @@ inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int fla return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags)); } +inline static hipError_t hipMallocManaged(void** ptr, size_t size, unsigned int flags) { + return hipCUDAErrorTohipError(cudaMallocManaged(ptr, size, flags)); +} + inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags __dparm(hipArrayDefault)) { diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index b6a1164a4a..b8632970d4 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -301,9 +301,7 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag return ihipLogStatus(hip_status); } -hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { - HIP_INIT_SPECIAL_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags); - HIP_SET_DEVICE(); +hipError_t ihipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { hipError_t hip_status = hipSuccess; if (HIP_SYNC_HOST_ALLOC) { @@ -368,6 +366,25 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { if (HIP_SYNC_HOST_ALLOC) { hipDeviceSynchronize(); } + return hip_status; +} + +hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { + HIP_INIT_SPECIAL_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags); + HIP_SET_DEVICE(); + hipError_t hip_status = hipSuccess; + hip_status = ihipHostMalloc(ptr, sizeBytes, flags); + return ihipLogStatus(hip_status); +} + +hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags) { + HIP_INIT_SPECIAL_API(hipMallocManaged, (TRACE_MEM), devPtr, size, flags); + HIP_SET_DEVICE(); + hipError_t hip_status = hipSuccess; + if(flags != hipMemAttachGlobal) + hip_status = hipErrorInvalidValue; + else + hip_status = ihipHostMalloc(devPtr, size, hipHostMallocDefault); return ihipLogStatus(hip_status); } diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMallocManaged.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMallocManaged.cpp new file mode 100644 index 0000000000..bf7662f6b7 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/memory/hipMallocManaged.cpp @@ -0,0 +1,52 @@ +#include +#include +#include "test_common.h" + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp + * RUN: %t -N 256M + * HIT_END + */ + +__global__ +void add(int n, float *x, float *y) +{ + int index = blockIdx.x * blockDim.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for (int i = index; i < n; i += stride) + y[i] = x[i] + y[i]; +} + +int main(int argc, char *argv[]) +{ + HipTest::parseStandardArguments(argc, argv, true); + int numElements = N; + bool testResult = true; + float *A, *B; + + hipMallocManaged(&A, numElements*sizeof(float)); + hipMallocManaged(&B, numElements*sizeof(float)); + + for (int i = 0; i < numElements; i++) { + A[i] = 1.0f; + B[i] = 2.0f; + } + + int blockSize = 256; + int numBlocks = (numElements + blockSize - 1) / blockSize; + dim3 dimGrid(numBlocks, 1, 1); + dim3 dimBlock(blockSize, 1, 1); + hipLaunchKernelGGL(add, dimGrid, dimBlock, 0, 0, numElements, A, B); + + hipDeviceSynchronize(); + + float maxError = 0.0f; + for (int i = 0; i < numElements; i++) + maxError = fmax(maxError, fabs(B[i]-3.0f)); + + hipFree(A); + hipFree(B); + if(maxError == 0.0f) + passed(); + failed("Output Mismatch\n"); +} From 43839c780ea8cebdbb979a77f74e6eb66fa1a15a Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 24 Apr 2019 17:40:35 +0300 Subject: [PATCH 13/17] [HIPIFY][doc] Update README.md + A few words about clang patches to work with CUDA 9.2 - 10.0 on Windows; + Fix cuDNN versions with correct values. [ROCm/hip commit: 4651dce3f0c182e792679942f1e4e4c60896fd13] --- projects/hip/hipify-clang/README.md | 16 ++++++++++------ 1 file changed, 10 insertions(+), 6 deletions(-) diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md index 37a1645c4b..0daaace2f0 100644 --- a/projects/hip/hipify-clang/README.md +++ b/projects/hip/hipify-clang/README.md @@ -32,10 +32,12 @@ ## Dependencies `hipify-clang` requires: -1. **LLVM+CLANG** of at least version 3.8.0, the latest stable and recommended release: **6.0.1 on Windows**, and **8.0.0 on Linux**. +1. **LLVM+CLANG** of at least version 3.8.0; the latest stable and recommended release: **6.0.1 on Windows**, and **8.0.0 on Linux**. 2. **CUDA** at least version 7.0, the latest supported version is **9.0 on Windows**, and **10.0 on Linux**. +If the target CUDA is 9.1, 9.2 or 10.0, to work on Windows you may apply patches*: [for LLVM 7.0.0](patches/patch_for_clang_7.0.0_bug_38811.zip), [for LLVM 7.0.1](patches/patch_for_clang_7.0.1_bug_38811.zip), [for LLVM 8.0.0](patches/patch_for_clang_8.0.0_bug_38811.zip). + | **LLVM release version** | **CUDA latest supported version** | **Windows** | **Linux** | |:------------------------:|:---------------------------------:|:------------:|:---------:| | 3.8.0 | 7.5 | + | + | @@ -176,9 +178,9 @@ To run it: On Linux the following configurations are tested: -Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-8.0 - cudnn-9.0 +Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.5.0.56 -Ubuntu 16-18: LLVM 8.0.0, CUDA 8.0 - 10.0, cudnn-8.0 - cudnn-10.0 +Ubuntu 16-18: LLVM 8.0.0, CUDA 8.0 - 10.0, cudnn-5.1.10 - cudnn-7.5.0.56 Build system for the above configurations: @@ -301,11 +303,13 @@ Testing Time: 2.51s ### Windows -On Windows the following configurations are tested: +On Windows 10 the following configurations are tested: -LLVM 5.0.0 - 5.0.2, CUDA 8.0, cudnn-8.0 +LLVM 5.0.0 - 5.0.2, CUDA 8.0, cudnn-5.1.10 - cudnn-7.1.4.18 -LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn-9.0 +LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn-7.0.5.15 - cudnn-7.5.0.56 + +LLVM 7.0.0 - 8.0.0 (with patch*), CUDA 7.5 - 10.0, cudnn-7.0.5.15 - cudnn-7.5.0.56 Build system for the above configurations: From d51679ed8e44a74ce033de6e138a2156dc006321 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 24 Apr 2019 20:35:52 +0300 Subject: [PATCH 14/17] [HIPIFY][fix][#204] Suppress warning message: #pragma once in main file [ROCm/hip commit: 6d3c4432345bdf36965d6d7f3af228227709b133] --- projects/hip/hipify-clang/src/LLVMCompat.cpp | 8 ++++++++ projects/hip/hipify-clang/src/LLVMCompat.h | 2 ++ projects/hip/hipify-clang/src/main.cpp | 3 +++ 3 files changed, 13 insertions(+) diff --git a/projects/hip/hipify-clang/src/LLVMCompat.cpp b/projects/hip/hipify-clang/src/LLVMCompat.cpp index 95c8ab7837..87dedc630e 100644 --- a/projects/hip/hipify-clang/src/LLVMCompat.cpp +++ b/projects/hip/hipify-clang/src/LLVMCompat.cpp @@ -110,4 +110,12 @@ std::error_code real_path(const Twine &path, SmallVectorImpl &output, #endif } +bool pragma_once_outside_header() { +#if LLVM_VERSION_MAJOR < 4 + return false; +#else + return true; +#endif +} + } // namespace llcompat diff --git a/projects/hip/hipify-clang/src/LLVMCompat.h b/projects/hip/hipify-clang/src/LLVMCompat.h index 98ff66fb41..069aeb91a3 100644 --- a/projects/hip/hipify-clang/src/LLVMCompat.h +++ b/projects/hip/hipify-clang/src/LLVMCompat.h @@ -81,4 +81,6 @@ void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, std::error_code real_path(const Twine &path, SmallVectorImpl &output, bool expand_tilde = false); +bool pragma_once_outside_header(); + } // namespace llcompat diff --git a/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp index 3b832deb8f..ae22f83ce0 100644 --- a/projects/hip/hipify-clang/src/main.cpp +++ b/projects/hip/hipify-clang/src/main.cpp @@ -222,6 +222,9 @@ int main(int argc, const char **argv) { #if defined(HIPIFY_CLANG_RES) Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES)); #endif + if (llcompat::pragma_once_outside_header()) { + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-Wno-pragma-once-outside-header", ct::ArgumentInsertPosition::BEGIN)); + } if (!MacroNames.empty()) { for (std::string s : MacroNames) { Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-D", ct::ArgumentInsertPosition::END)); From 86d68901da80198d4e1978c387ee1ceb44c21f41 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 25 Apr 2019 12:18:51 +0300 Subject: [PATCH 15/17] [HIPIFY][DNN] cudnnSetFilter4dDescriptor support [ROCm/hip commit: 3fee0f376582160c6cd3481b8cdb49671120c241] --- projects/hip/docs/markdown/CUDNN_API_supported_by_HIP.md | 2 +- projects/hip/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp | 2 +- .../unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/hip/docs/markdown/CUDNN_API_supported_by_HIP.md b/projects/hip/docs/markdown/CUDNN_API_supported_by_HIP.md index ba1eab75a9..43e7ca419e 100644 --- a/projects/hip/docs/markdown/CUDNN_API_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDNN_API_supported_by_HIP.md @@ -292,7 +292,7 @@ |`cudnnSetTensor` |`hipdnnSetTensor` | |`cudnnScaleTensor` |`hipdnnScaleTensor` | |`cudnnCreateFilterDescriptor` |`hipdnnCreateFilterDescriptor` | -|`cudnnSetFilter4dDescriptor` | | +|`cudnnSetFilter4dDescriptor` |`hipdnnSetFilter4dDescriptor` | |`cudnnGetFilter4dDescriptor` | | |`cudnnSetFilterNdDescriptor` |`hipdnnSetFilterNdDescriptor` | |`cudnnGetFilterNdDescriptor` |`hipdnnGetFilterNdDescriptor` | diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp index d99a201b8e..85e138a8fd 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp @@ -75,7 +75,7 @@ const std::map CUDA_DNN_FUNCTION_MAP{ // cuDNN Filter functions {"cudnnCreateFilterDescriptor", {"hipdnnCreateFilterDescriptor", "", CONV_LIB_FUNC, API_DNN}}, - {"cudnnSetFilter4dDescriptor", {"hipdnnSetFilter4dDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSetFilter4dDescriptor", {"hipdnnSetFilter4dDescriptor", "", CONV_LIB_FUNC, API_DNN}}, {"cudnnGetFilter4dDescriptor", {"hipdnnGetFilter4dDescriptor", "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnSetFilterNdDescriptor", {"hipdnnSetFilterNdDescriptor", "", CONV_LIB_FUNC, API_DNN}}, {"cudnnGetFilterNdDescriptor", {"hipdnnGetFilterNdDescriptor", "", CONV_LIB_FUNC, API_DNN}}, diff --git a/projects/hip/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu index 23880097ff..6aae90e403 100644 --- a/projects/hip/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu +++ b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu @@ -113,7 +113,7 @@ int main() { cudnnFilterDescriptor_t filt_desc; // CHECK: CUDNN_CALL(hipdnnCreateFilterDescriptor(&filt_desc)); CUDNN_CALL(cudnnCreateFilterDescriptor(&filt_desc)); - // CHECK-NOT: CUDNN_CALL(hipdnnSetFilter4dDescriptor( + // CHECK: CUDNN_CALL(hipdnnSetFilter4dDescriptor( CUDNN_CALL(cudnnSetFilter4dDescriptor( // CHECK: filt_desc, HIPDNN_DATA_FLOAT, HIPDNN_TENSOR_NCHW, filt_desc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, From d8f4eff1effd8b83c9fb1e8915f8dd7f31e95a40 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 25 Apr 2019 13:27:30 +0300 Subject: [PATCH 16/17] [HIPIFY][perl][fix][258] Memory fence device functions are supported now [ROCm/hip commit: abd1c53cf81ca89708172fdab90c4c0d89336ad7] --- projects/hip/bin/hipify-perl | 4 ---- 1 file changed, 4 deletions(-) diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index 3456f9fa41..3e39a5a11d 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -553,10 +553,6 @@ sub warnUnsupportedSpecialFunctions my $line_num = shift; my $m = 0; foreach $func ( - # memory fence: - "__threadfence_block", - "__threadfence", - "__threadfence_system", # Synchronization: "__syncthreads_count", "__syncthreads_and", From b9526364bb53a6335bfe0023d916c710050a8168 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 25 Apr 2019 11:47:48 +0530 Subject: [PATCH 17/17] Add dependency on amd_comgr in hip-config-*.cmake.in Change-Id: Iac1d851a8cfb99224e9c5926780273d9b9b08426 [ROCm/hip commit: ef99ffd9f4de3d2249807f041db751464644bce1] --- projects/hip/hip-config-clang.cmake.in | 1 + projects/hip/hip-config-hcc.cmake.in | 1 + 2 files changed, 2 insertions(+) diff --git a/projects/hip/hip-config-clang.cmake.in b/projects/hip/hip-config-clang.cmake.in index e4862f3d6c..9680b2ff8e 100644 --- a/projects/hip/hip-config-clang.cmake.in +++ b/projects/hip/hip-config-clang.cmake.in @@ -50,6 +50,7 @@ set_and_check( hip_BIN_INSTALL_DIR "@PACKAGE_BIN_INSTALL_DIR@" ) set_and_check(hip_HIPCC_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipcc") set_and_check(hip_HIPCONFIG_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipconfig") +find_dependency(amd_comgr) include( "${CMAKE_CURRENT_LIST_DIR}/hip-targets.cmake" ) set( hip_LIBRARIES hip::host hip::device) diff --git a/projects/hip/hip-config-hcc.cmake.in b/projects/hip/hip-config-hcc.cmake.in index 7e236ba0b4..c0ffc6e2af 100644 --- a/projects/hip/hip-config-hcc.cmake.in +++ b/projects/hip/hip-config-hcc.cmake.in @@ -51,6 +51,7 @@ set_and_check(hip_HIPCC_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipcc") set_and_check(hip_HIPCONFIG_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipconfig") find_dependency(hcc) +find_dependency(amd_comgr) include( "${CMAKE_CURRENT_LIST_DIR}/hip-targets.cmake" ) set( hip_LIBRARIES hip::host hip::device)