Skip to content

Commit

Permalink
Merge branch 'amd-develop' into amd-master
Browse files Browse the repository at this point in the history
Change-Id: I7f2fba7875ed1c89dfc768f7415ed6fb0d1c6407
  • Loading branch information
mangupta committed Jun 12, 2017
2 parents 75f8840 + ad33c94 commit 080eb12
Show file tree
Hide file tree
Showing 13 changed files with 377 additions and 38 deletions.
2 changes: 1 addition & 1 deletion README.md
Original file line number Diff line number Diff line change
Expand Up @@ -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)

Expand Down
9 changes: 9 additions & 0 deletions RELEASE.md
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
279 changes: 279 additions & 0 deletions bin/hipify-cmakefile
Original file line number Diff line number Diff line change
@@ -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 (<INFILE>) {

# 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_<CONFIG>",
"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;
}
4 changes: 2 additions & 2 deletions include/hip/hcc_detail/hip_fp16.h
Original file line number Diff line number Diff line change
Expand Up @@ -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)));
Expand Down Expand Up @@ -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
8 changes: 4 additions & 4 deletions include/hip/hcc_detail/hip_runtime.h
Original file line number Diff line number Diff line change
Expand Up @@ -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();


/**
Expand Down Expand Up @@ -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



Expand Down
4 changes: 2 additions & 2 deletions packaging/hip_hcc.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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")
Expand Down
6 changes: 4 additions & 2 deletions src/device_util.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand Down
2 changes: 1 addition & 1 deletion src/hip_hcc.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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;



Expand Down
Loading

0 comments on commit 080eb12

Please sign in to comment.