From 1d1e50fa771ef1c4752744a73154e195ca4e21b4 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 7 Jun 2017 00:15:05 -0500 Subject: [PATCH 01/12] Enable HCC_OPT_FLUSH=1. Requires appropriate HCC with this support : commit 38e392b517a46a09a3b1c8f388e6a0db3741c510 [ROCm/hip commit: ac634bf33428d8fe844f5e8009e706cfd1612a39] --- projects/hip/src/hip_hcc.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 08a2cdbfcf..d826a0cec3 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -98,7 +98,7 @@ int HIP_SYNC_NULL_STREAM = 0; // HIP needs to change some behavior based on HCC_OPT_FLUSH : // TODO - set this to 1 -int HCC_OPT_FLUSH = 0; +int HCC_OPT_FLUSH = 1; From 7213c7a1ad532594fc3ced8180b0db6d13bd63a1 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 7 Jun 2017 15:23:37 +0530 Subject: [PATCH 02/12] p2p_copy_coherency test: gracefully handle single gpu case Change-Id: I216663f67ef58c673136332635dab8b57079b909 [ROCm/hip commit: 1efb6ce994e50015f0b1f5352c846259f6314bf2] --- .../tests/src/runtimeApi/memory/p2p_copy_coherency.cpp | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp b/projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp index 459c0054c9..6bc6235454 100644 --- a/projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp @@ -156,7 +156,13 @@ int main(int argc, char *argv[]) int dev0 = 0; int dev1 = 1; - // TODO - only works on multi-GPU system: + int numDevices; + HIPCHECK(hipGetDeviceCount(&numDevices)); + if (numDevices == 1) { + printf("warning : test requires atleast two gpus\n"); + passed(); + } + if (enablePeers(dev0,dev1) == -1) { printf ("warning : could not find peer gpus\n"); return -1; From ee92e19e7a90a7e32d8a3bb7830a2d0ecf6675db Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 7 Jun 2017 15:24:44 +0530 Subject: [PATCH 03/12] hipDeviceMemcpy test: make it functional on nvcc path Change-Id: Id10c79b48747ed701adbd0a233c53cd60cfa743b [ROCm/hip commit: e6cafbf34207ce8f3fb7d6eb60459dae44b01f5c] --- projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp b/projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp index 3843c07bb9..527df9bab1 100644 --- a/projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp +++ b/projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp @@ -4,7 +4,7 @@ #include "../test_common.h" -#define LEN 1030 +#define LEN 1024 #define SIZE LEN << 2 /* HIT_START @@ -17,13 +17,13 @@ __global__ void cpy(hipLaunchParm lp, uint32_t *Out, uint32_t *In) { int tx = hipThreadIdx_x; - memcpy(Out + tx, In + tx, SIZE/LEN); + memcpy(Out + tx, In + tx, sizeof(uint32_t)); } __global__ void set(hipLaunchParm lp, uint32_t *ptr, uint8_t val, size_t size) { int tx = hipThreadIdx_x; - memset(ptr + tx, val, size); + memset(ptr + tx, val, (sizeof(uint32_t)*(size/LEN))); } int main() From 59a792d892c4e6f8e61a743329d0c4e13b919d08 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 7 Jun 2017 15:25:54 +0530 Subject: [PATCH 04/12] hipMemcpy-size test: reduce max size to make it work correctly on nvcc path Change-Id: I9ce9f5a9e141ffd8ddf961269010b33358e02771 [ROCm/hip commit: 1c93d8592e7b548cd56ab2b2026721d55423cc50] --- projects/hip/tests/src/runtimeApi/memory/hipMemcpy.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpy.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpy.cpp index 749ec0de77..e8e803e44c 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemcpy.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpy.cpp @@ -304,7 +304,7 @@ void memcpytest2_sizes(size_t maxElem=0) HIPCHECK(hipMemGetInfo(&free, &total)); if (maxElem == 0) { - maxElem = free/sizeof(T)/5; + maxElem = free/sizeof(T)/8; } printf (" device#%d: hipMemGetInfo: free=%zu (%4.2fMB) total=%zu (%4.2fMB) maxSize=%6.1fMB\n", From cd3beb8c66d2e6e0cb3a18a5b2a77535565567e3 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 7 Jun 2017 15:50:28 +0530 Subject: [PATCH 05/12] hip_hcc package: add libstdc++-static as a rpm dependency Change-Id: I83a79353492a6be3d788b7c0ce4a8f3aa740d9d9 [ROCm/hip commit: 6b768c2f0abb95ab31e1a758c2c341eb7547a4ab] --- projects/hip/packaging/hip_hcc.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/packaging/hip_hcc.txt b/projects/hip/packaging/hip_hcc.txt index b0808aa0bc..284d97e2e5 100644 --- a/projects/hip/packaging/hip_hcc.txt +++ b/projects/hip/packaging/hip_hcc.txt @@ -42,9 +42,9 @@ set(CPACK_RPM_POST_INSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/postinst") set(CPACK_RPM_PRE_UNINSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/prerm") set(CPACK_RPM_PACKAGE_AUTOREQPROV " no") if(@COMPILE_HIP_ATP_MARKER@) - set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${CPACK_PACKAGE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, rocm-profiler") + set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${CPACK_PACKAGE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, rocm-profiler, libstdc++-static") else() - set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${CPACK_PACKAGE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@") + set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${CPACK_PACKAGE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, libstdc++-static") endif() set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") set(CPACK_SOURCE_GENERATOR "TGZ") From ef7c9be9c6dbd3d04beb7fae13fd4ce2ce67d96a Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 7 Jun 2017 09:05:30 -0500 Subject: [PATCH 06/12] Use amHostCoherentFlag. Requires new HCC version. [ROCm/hip commit: 99e9c7cca50a01745f71eeb3869ce2e51a23992f] --- projects/hip/src/hip_memory.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 3ab7713afa..c04c2611c3 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -243,6 +243,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) } + hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { HIP_INIT_SPECIAL_API((TRACE_MEM), ptr, sizeBytes, flags); @@ -289,10 +290,10 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) if (flags & hipHostMallocCoherent) { amFlags = amHostCoherent; } else if (flags & hipHostMallocNonCoherent) { - amFlags = amHostPinned; + amFlags = amHostNonCoherent; } else { // depends on env variables: - amFlags = HIP_COHERENT_HOST_ALLOC ? amHostCoherent : amHostPinned; + amFlags = HIP_COHERENT_HOST_ALLOC ? amHostCoherent : amHostNonCoherent; } From 8b145076508e9d0da3832256f86061d721974713 Mon Sep 17 00:00:00 2001 From: "Sun, Peng" Date: Thu, 8 Jun 2017 19:20:10 -0500 Subject: [PATCH 07/12] Add clang version guard so the hip_fp16.h header won't be picked up by gcc Change-Id: Ia21335a455bc93210901b44bc8c76a7f4a385b55 [ROCm/hip commit: 06816fb68b446c31300d85758302c9d0bb7dd425] --- projects/hip/include/hip/hcc_detail/hip_fp16.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_fp16.h b/projects/hip/include/hip/hcc_detail/hip_fp16.h index b1ecc61cb0..4d90ec82b2 100644 --- a/projects/hip/include/hip/hcc_detail/hip_fp16.h +++ b/projects/hip/include/hip/hcc_detail/hip_fp16.h @@ -24,7 +24,7 @@ THE SOFTWARE. #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H #include "hip/hcc_detail/hip_vector_types.h" - +#if ( __clang_major__ > 3) typedef __fp16 __half; typedef __fp16 __half1 __attribute__((ext_vector_type(1))); typedef __fp16 __half2 __attribute__((ext_vector_type(2))); @@ -454,6 +454,6 @@ __device__ static inline __half2 h2trunc(const __half2 h) { a.xy = __hip_hc_ir_h2trunc_int(h.xy); return a; } - +#endif //clang_major > 3 #endif From 47921e39d24933fbde5cbbb686549d9d72b64d15 Mon Sep 17 00:00:00 2001 From: "Sun, Peng" Date: Thu, 8 Jun 2017 19:24:22 -0500 Subject: [PATCH 08/12] Fix error related to undefined reference of __get_dynamicgroupbaseptr(). Change-Id: I14951e1725e35dd5f5e53805f81cdb58661f59f2 [ROCm/hip commit: 43df5ba6604eac6c19d37139c3d8761406fb745d] --- projects/hip/include/hip/hcc_detail/hip_runtime.h | 8 ++++---- projects/hip/src/device_util.cpp | 6 ++++-- 2 files changed, 8 insertions(+), 6 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h index 129020d9cd..95826f9b60 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h @@ -305,7 +305,7 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask __host__ __device__ int min(int arg1, int arg2); __host__ __device__ int max(int arg1, int arg2); -__device__ ADDRESS_SPACE_3 void* __get_dynamicgroupbaseptr(); +__device__ void* __get_dynamicgroupbaseptr(); /** @@ -464,10 +464,10 @@ do {\ // Macro to replace extern __shared__ declarations // to local variable definitions #define HIP_DYNAMIC_SHARED(type, var) \ - ADDRESS_SPACE_3 type* var = \ - (ADDRESS_SPACE_3 type*)__get_dynamicgroupbaseptr(); \ + type* var = \ + (type*)__get_dynamicgroupbaseptr(); \ -#define HIP_DYNAMIC_SHARED_ATTRIBUTE ADDRESS_SPACE_3 +#define HIP_DYNAMIC_SHARED_ATTRIBUTE diff --git a/projects/hip/src/device_util.cpp b/projects/hip/src/device_util.cpp index e59a44e5ba..062372f0f4 100644 --- a/projects/hip/src/device_util.cpp +++ b/projects/hip/src/device_util.cpp @@ -1101,11 +1101,13 @@ __host__ __device__ int max(int arg1, int arg2) return (int)(hc::precise_math::fmax((float)arg1, (float)arg2)); } -__device__ ADDRESS_SPACE_3 void* __get_dynamicgroupbaseptr() -{ +__device__ void* __get_dynamicgroupbaseptr() { return hc::get_dynamic_group_segment_base_pointer(); } +__host__ void* __get_dynamicgroupbaseptr() { + return nullptr; +} // Precise Math Functions __device__ float __hip_precise_cosf(float x) { From ce0ecfd06724470d7e746faac85a2faa8d874a25 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 12 Jun 2017 09:57:17 +0530 Subject: [PATCH 09/12] Initial implementation of hipify-cmakefile Change-Id: Id365da9f887b5c3409639f000b430d093fd4f6b3 [ROCm/hip commit: 5339320485d79b676f8b4b65a71f3995c2ba4530] --- projects/hip/bin/hipify-cmakefile | 279 ++++++++++++++++++++++++++++++ 1 file changed, 279 insertions(+) create mode 100755 projects/hip/bin/hipify-cmakefile diff --git a/projects/hip/bin/hipify-cmakefile b/projects/hip/bin/hipify-cmakefile new file mode 100755 index 0000000000..b11de4adc1 --- /dev/null +++ b/projects/hip/bin/hipify-cmakefile @@ -0,0 +1,279 @@ +#!/usr/bin/perl -w +## +# Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. +## +#usage hipify-cmakefile [OPTIONS] INPUT_FILE +use Getopt::Long; + +GetOptions( + "print-stats" => \$print_stats # print the command-line, like a header. + , "quiet-warnings" => \$quiet_warnings # don't print warnings on unknown CUDA functions. + , "no-output" => \$no_output # don't write any translated output to stdout. + , "inplace" => \$inplace # modify input file inplace, save backup in ".prehip" file. + , "n" => \$n # combination of print_stats + no-output. +); + +$print_stats = 1 if $n; +$no_output = 1 if $n; + +@warn_whitelist = (); + +#--- +#Stats tracking code: +@statNames = ( "macro", "include", "option", "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 @statNames = @{ shift() }; + my %counts = %{ shift() }; + my $warnings = shift(); + my $loc = shift(); + + 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: +sub addStats { + my $dest_ref = shift(); + my %adder = %{ shift() }; + + foreach $key ( keys %adder ) { + $dest_ref->{$key} += $adder{$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 ); + +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; + } + 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; + } + + # count of transforms in this file, init to 0 here: + my %ft; + clearStats( \%ft, \@statNames ); + + my $lineCount = 0; + + undef $/; # Read whole file at once, so we can match newlines. + while () { + + # Replace find_package(CUDA) with find_package(HIP) + $ft{'include'} += s/\bfind_package[ ]*\([ ]*CUDA[ ]*[0-9.]*/find_package(HIP/ig; + + # Replace macros + $ft{'macro'} += s/\bCUDA_ADD_EXECUTABLE/HIP_ADD_EXECUTABLE/ig; + $ft{'macro'} += s/\bCUDA_ADD_LIBRARY/HIP_ADD_LIBRARY/ig; + $ft{'macro'} += s/\bCUDA_INCLUDE_DIRECTORIES/HIP_INCLUDE_DIRECTORIES/ig; + + # Replace options + $ft{'option'} += s/\bCUDA_NVCC_FLAGS/HIP_NVCC_FLAGS/ig; + $ft{'option'} += s/\bCUDA_HOST_COMPILATION_CPP/HIP_HOST_COMPILATION_CPP/ig; + $ft{'option'} += s/\bCUDA_SOURCE_PROPERTY_FORMAT/HIP_SOURCE_PROPERTY_FORMAT/ig; + + # Replace variables + $ft{'other'} += s/\bCUDA_FOUND/HIP_FOUND/ig; + $ft{'other'} += s/\bCUDA_VERSION/HIP_VERSION/ig; + $ft{'other'} += s/\bCUDA_TOOLKIT_ROOT_DIR/HIP_ROOT_DIR/ig; + + unless ($quiet_warnings) { + + #print STDERR "Check WARNINGs\n"; + # copy into array of lines, process line-by-line to show warnings: + 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/; + } + + $s = warnUnsupportedSpecialFunctions($line_num); + $warnings += $s; + } + + $_ = $tmp; + } + + #-------- + # Print it! + unless ($no_output) { + print $OUTFILE "$_"; + } + $lineCount = $_ =~ tr/\n//; + } + + my $totalConverted = totalStats( \%ft ); + + if ( ( $totalConverted + $warnings ) and $print_stats ) { + printStats( "info: converted", \@statNames, \%ft, $warnings, $lineCount ); + print STDERR " in '$fileName'\n"; + print STDERR "You may need to hand-edit '$fileName' to add steps to build correctly on HCC path\n"; + } + + # Update totals for all files: + addStats( \%tt, \%ft ); + $Twarnings += $warnings; + $TlineCount += $lineCount; +} + +#-- 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"; +} + +#--- +sub warnUnsupportedSpecialFunctions { + my $line_num = shift; + my $m = 0; + + foreach $func ( + # macros: + "CUDA_ADD_CUFFT_TO_TARGET", + "CUDA_ADD_CUBLAS_TO_TARGET", + #"CUDA_ADD_EXECUTABLE", + #"CUDA_ADD_LIBRARY", + "CUDA_BUILD_CLEAN_TARGET", + "CUDA_COMPILE", + "CUDA_COMPILE_PTX", + "CUDA_COMPILE_FATBIN", + "CUDA_COMPILE_CUBIN", + "CUDA_COMPUTE_SEPARABLE_COMPILATION_OBJECT_FILE_NAME", + #"CUDA_INCLUDE_DIRECTORIES", + "CUDA_LINK_SEPARABLE_COMPILATION_OBJECTS", + "CUDA_SELECT_NVCC_ARCH_FLAGS", + "CUDA_WRAP_SRCS", + + # options: + "CUDA_64_BIT_DEVICE_CODE", + "CUDA_ATTACH_VS_BUILD_RULE_TO_CUDA_FILE", + "CUDA_BUILD_CUBIN", + "CUDA_BUILD_EMULATION", + "CUDA_LINK_LIBRARIES_KEYWORD", + "CUDA_GENERATED_OUTPUT_DIR", + #"CUDA_HOST_COMPILATION_CPP", + "CUDA_HOST_COMPILER", + #"CUDA_NVCC_FLAGS", + #"CUDA_NVCC_FLAGS_", + "CUDA_PROPAGATE_HOST_FLAGS", + "CUDA_SEPARABLE_COMPILATION", + #"CUDA_SOURCE_PROPERTY_FORMAT", + "CUDA_USE_STATIC_CUDA_RUNTIME", + "CUDA_VERBOSE_BUILD", + + # others: + #"CUDA_VERSION_MAJOR", + #"CUDA_VERSION_MINOR", + #"CUDA_VERSION", + #"CUDA_VERSION_STRING", + "CUDA_HAS_FP16", + #"CUDA_TOOLKIT_ROOT_DIR", + "CUDA_SDK_ROOT_DIR", + "CUDA_INCLUDE_DIRS", + "CUDA_LIBRARIES", + "CUDA_CUFFT_LIBRARIES", + "CUDA_CUBLAS_LIBRARIES", + "CUDA_cudart_static_LIBRARY", + "CUDA_cudadevrt_LIBRARY", + "CUDA_cupti_LIBRARY", + "CUDA_curand_LIBRARY", + "CUDA_cusolver_LIBRARY", + "CUDA_cusparse_LIBRARY", + "CUDA_npp_LIBRARY", + "CUDA_nppc_LIBRARY", + "CUDA_nppi_LIBRARY", + "CUDA_npps_LIBRARY", + "CUDA_nvcuvenc_LIBRARY", + "CUDA_nvcuvid_LIBRARY" + ) + { + my $mt = m/\b($func)/g; + if ($mt) { + $m += $mt; + print STDERR " warning: $fileName:#$line_num : unsupported macro/option : $_\n"; + } + } + + return $m; +} From 3b5222dc4284c56c16e8e3238f65c70bb19757e0 Mon Sep 17 00:00:00 2001 From: Patrick Flick Date: Sun, 4 Jun 2017 10:24:00 -0400 Subject: [PATCH 10/12] fix typo [ROCm/hip commit: 5dfe207eb94bba66f62a6fba587c14defba07774] --- projects/hip/README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/README.md b/projects/hip/README.md index d04d63714f..565fd6a36d 100644 --- a/projects/hip/README.md +++ b/projects/hip/README.md @@ -134,7 +134,7 @@ The README with the procedures and tips the team used during this porting effort * **bin**: Tools and scripts to help with hip porting * **hipify** : Tool to convert CUDA code to portable CPP. Converts CUDA APIs and kernel builtins. - * **hipcc** : Compiler driver that can be used to replace nvcc in existing CUDA code. hipcc ill call nvcc or hcc depending on platform, and include appropriate platform-specific headers and libraries. + * **hipcc** : Compiler driver that can be used to replace nvcc in existing CUDA code. hipcc will call nvcc or hcc depending on platform, and include appropriate platform-specific headers and libraries. * **hipconfig** : Print HIP configuration (HIP_PATH, HIP_PLATFORM, CXX config flags, etc) * **hipexamine.sh** : Script to scan directory, find all code, and report statistics on how much can be ported with HIP (and identify likely features not yet supported) From 5a2d48117dfd6cb73a0983f45cf78cf0367fe2fa Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 12 Jun 2017 11:19:55 +0530 Subject: [PATCH 11/12] Update directed tests README.md Change-Id: I395245454d376508f04e5a4a62c8933895cb3867 [ROCm/hip commit: b850a08d99184610f930e98d2370bc7a4c89f253] --- projects/hip/tests/README.md | 86 +++++++++++++++++++++++++++--------- 1 file changed, 64 insertions(+), 22 deletions(-) diff --git a/projects/hip/tests/README.md b/projects/hip/tests/README.md index cb41cc10cd..27cde7c534 100644 --- a/projects/hip/tests/README.md +++ b/projects/hip/tests/README.md @@ -1,39 +1,78 @@ # HIP testing environment. -This document explains how to use the HIP CMAKE testing environment. +This document explains how to use the HIP CMAKE testing environment. +We make use of the HIT Integrated Tester (HIT) framework to automatically find and add test cases to the CMAKE testing environment. ### Quickstart -Usage : + +HIP unit tests are integrated into the top-level cmake project. The tests depend upon the installed version of HIP. +Typical usage (paths relative to top of the HIP repo): ``` $ mkdir build $ cd build -$ cmake ../src +$ cmake .. -DCMAKE_INSTALL_PREFIX=$PWD/install $ make +$ make install +$ make build_tests $ make test ``` ### How to add a new test -The tests/src/runtimeApi/memory/hipMemtest.cpp file contains a simple unit test and is a good starting point for other tests. -Copy this to a new test name and modify tests/src/CMakefiles.txt to add the test to the build environment. - -Recent versions of the test infrastructure use a hierarchy of folders. Each folder contains src and CMakefiles.txt file. -See the CMakefiles.txt files for description of the intended purpose for each sub-directory. +The test infrastructure use a hierarchy of folders. So add the new test to the appropriate folder. +The tests/src/runtimeApi/memory/hipMemset.cpp file contains a simple unit test and is a good starting point for other tests. +Copy this to a new test name and modify it. -#### Edit CMakefiles.txt: -// Example: +### HIP Integrated Tester (HIT) + +The HIT framework sutomatically finds and adds test cases to the CMAKE testing environment. It achives this by parsing all files in the tests/src folder. +The parser looks for a code block similar to the one below. ``` -# Build the test executable: -build_hip_executable (hipMemset hipMemset.cpp) - - -# This runs the tests with the specified command-line testing. -# Multiple make_test may be specified. -make_test(hipMemset " ") +/* HIT_START + * BUILD: %t %s ../../test_common.cpp + * RUN: %t + * //Small copy + * RUN: %t -N 10 --memsetval 0x42 + * // Oddball size + * RUN: %t -N 10013 --memsetval 0x5a + * // Big copy + * RUN: %t -N 256M --memsetval 0xa6 + * HIT_END + */ ``` +In the above, BUILD commands provide instructions on how to build the test case while RUN commands provide instructions on how to execute the test case. -It is recommended to place the build and run steps adjacent in the CMakefiles.txt. +#### BUILD command + +The supported syntax for the BUILD command is: +``` +BUILD: %t %s HIPCC_OPTIONS HCC_OPTIONS NVCC_OPTIONS EXCLUDE_HIP_PLATFORM +``` +%s: refers to current source file name. Additional source files needed for the test can be specified by name (including relative path). +%t: refers to target executable named derived by removing the extension from the current source file. Alternatively a target executable name can be specified. +HIPCC_OPTIONS: All options specified after this delimiter are passed to hipcc on both HCC and NVCC platforms. +HCC_OPTIONS: All options specified after this delimiter are passed to hipcc on HCC platform only. +NVCC_OPTIONS: All options specified after this delimiter are passed to hipcc on NVCC platform only. +EXCLUDE_HIP_PLATFORM: This can be used to exclude a test case from HCC, NVCC or both platforms. + + +#### RUN command + +The supported syntax for the RUN command is: +``` +RUN: %t EXCLUDE_HIP_PLATFORM +``` +%t: refers to target executable named derived by removing the extension from the current source file. Alternatively a target executable name can be specified. +EXCLUDE_HIP_PLATFORM: This can be used to exclude a test case from HCC, NVCC or both platforms. Note that if the test has been excluded for a specific platform in the BUILD command, it is automatically excluded from the RUN command as well for the same platform. + + +#### RUN_NAMED command + +When using the RUN command, HIT will squash and append the arguments specified to the test executable name to generate the CMAKE test name. Sometimes we might want to specify a more descriptive name. The RUN_NAMED command is used for that. The supported syntax for the RUN_NAMED command is: +``` +RUN: %t CMAKE_TEST_NAME EXCLUDE_HIP_PLATFORM +``` ### Running tests: @@ -43,11 +82,14 @@ ctest ### Run subsets of all tests: ``` -# Run one test on the commandline (obtain commandline parms from CMakefiles.tst) -./hipMemset +# Run one test on the commandline +./directed_tests/runtime/memory/hipMemset -# Run all the memory tests: +# Run all the hipMemcpy tests: ctest -R Memcpy + +# Run all tests in a specific folder: +ctest -R memory ``` @@ -55,7 +97,7 @@ ctest -R Memcpy Find the test and commandline that fail: -(From the test build directory, perhaps hip/tests/build) +(From the build directory, perhaps hip/build) grep -IR hipMemcpy-modes -IR ../tests/ ../tests/src/runtimeApi/memory/hipMemcpy.cpp: * RUN_NAMED: %t hipMemcpy-modes --tests 0x1 From e69bd819d58b7345a1f7d08dd48033fe693d1068 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 12 Jun 2017 11:20:28 +0530 Subject: [PATCH 12/12] Updated RELEASE.md Change-Id: Ic451612555c66f3ed7131514fc97fcc41091370a [ROCm/hip commit: ad33c9406ae7af93cffcf0a943a927b9795ace77] --- projects/hip/RELEASE.md | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/projects/hip/RELEASE.md b/projects/hip/RELEASE.md index 21fd8da7bb..5787c59881 100644 --- a/projects/hip/RELEASE.md +++ b/projects/hip/RELEASE.md @@ -13,6 +13,15 @@ Upcoming: ## Revision History: +=================================================================================================== +- new APIs: hipMemcpy2DAsync, hipMallocPitch, hipHostMallocCoherent, hipHostMallocNonCoherent +- added support for building hipify-clang using clang 3.9 +- hipify-clang updates for CUDA 8.0 runtime+driver support +- renamed hipify to hipify-perl +- initial implementation of hipify-cmakefile +- several documentation updates & bug fixes + + =================================================================================================== Release: 1.0.17102 Date: 2017.03.07