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) 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 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; +} 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 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/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") 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) { 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; 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; } 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 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() 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", 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;