AI/rocm-systems
2
0
Fork 0
程式碼 問題 合併請求 Actions 套件 專案 版本發布 Wiki 動態

Merge 'master' into 'amd-master'

瀏覽代碼 
Change-Id: I00292478e5ccfbb1007e449673e09c11d717782d


[ROCm/hip commit: 1dd9292d8d]
此提交包含在：
Jenkins
2019-04-25 17:18:46 -07:00
父節點 fa6b0ab49c b9526364bb
當前提交 f3bfdab643
共有 26 個檔案被更改，包括 256 行新增 和 244 行删除
下載 Patch 檔 下載差異檔 展開所有檔案 折疊所有檔案
projects/hip/CMakeLists.txt
+4 -1
查看文件
@@ -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()
projects/hip/bin/hipify-perl
+62 -207
查看文件
@@ -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 (<INFILE>)
{
#--------
# 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<hip\/hip_runtime.h>/;
$countIncludes += s/(\s*#\s*include\s+)[<"]cuda_runtime_api\.h[>"]/$1<hip\/hip_runtime_api.h>/;
$countIncludes += s/(\s*#\s*include\s+)[<"]cuda_fp16\.h[>"]/$1<hip\/hip_fp16.h>/;
#--------
# 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<T>::type s[];'
#'extern __shared__ typename mapper<Float>::type s_data[];'
#'extern __attribute__((used)) __shared__ typename mapper<Float>::type s_data[];'
$ft{'extern_shared'} += $k;
}
#--------
# CUDA Launch Syntax
# Note these only work if launch is on a single line.
# Handle the <<numBlocks, blockDim>>> 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 <<numBlocks, blockDim, sharedSize, stream>>> 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 <<numBlocks, blockDim, sharedSize>>> 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 <<numBlocks, blockDim>>> 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,35 @@ 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 +585,6 @@ sub warnUnsupportedSpecialFunctions
$m += $mt;
print STDERR " warning: $fileName:#$line_num : unsupported device function : $_\n";
}
}
return $m;
}
projects/hip/docs/markdown/CUDNN_API_supported_by_HIP.md
+1 -1
查看文件
@@ -292,7 +292,7 @@
|`cudnnSetTensor` |`hipdnnSetTensor` |
|`cudnnScaleTensor` |`hipdnnScaleTensor` |
|`cudnnCreateFilterDescriptor` |`hipdnnCreateFilterDescriptor` |
|`cudnnSetFilter4dDescriptor` | |
|`cudnnSetFilter4dDescriptor` |`hipdnnSetFilter4dDescriptor` |
|`cudnnGetFilter4dDescriptor` | |
|`cudnnSetFilterNdDescriptor` |`hipdnnSetFilterNdDescriptor` |
|`cudnnGetFilterNdDescriptor` |`hipdnnGetFilterNdDescriptor` |
projects/hip/hip-config-clang.cmake.in
+1
查看文件
@@ -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)
projects/hip/hip-config-hcc.cmake.in
+1
查看文件
@@ -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)
projects/hip/hipify-clang/README.md
+15 -9
查看文件
@@ -32,10 +32,12 @@
## <a name="dependencies"></a> 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 | + | + |
@@ -49,12 +51,14 @@
| 5.0.2 | 8.0 | + | + |
| 6.0.0 | 9.0 | + | + |
| **6.0.1** | **9.0** | + <br/> **LATEST STABLE RELEASE** | + |
| 7.0.0 | 9.2 | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
| 7.0.1 | 9.2 | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
| 7.0.0 | 9.2 | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.0.0_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
| 7.0.1 | 9.2 | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.0.1_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
| 7.1.0 | 9.2 (?) | - <br/> LLVM 7.1.0 <br/> is not yet released | - <br/> LLVM 7.1.0 <br/> is not yet released |
| **8.0.0** | **10.0** | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | + <br/> **LATEST STABLE RELEASE** |
| **8.0.0** | **10.0** | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_8.0.0_bug_38811.zip)*</br> | + <br/> **LATEST STABLE RELEASE** |
| 8.0.1 | 10.1 (?) | - <br/> LLVM 8.0.1 <br/> is not yet released | - <br/> LLVM 8.0.1 <br/> 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
@@ -174,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:
@@ -299,11 +303,13 @@ Testing Time: 2.51s
### <a name="windows"></a >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:
projects/hip/hipify-clang/patches/patch_for_clang_7.0.0_bug_38811.zip
二進制
查看文件
未顯示二進位檔案。
projects/hip/hipify-clang/patches/patch_for_clang_7.0.1_bug_38811.zip
二進制
查看文件
未顯示二進位檔案。
projects/hip/hipify-clang/patches/patch_for_clang_8.0.0_bug_38811.zip
二進制
查看文件
未顯示二進位檔案。
projects/hip/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp
+1 -1
查看文件
@@ -75,7 +75,7 @@ const std::map<llvm::StringRef, hipCounter> 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}},
projects/hip/hipify-clang/src/LLVMCompat.cpp
+8
查看文件
@@ -110,4 +110,12 @@ std::error_code real_path(const Twine &path, SmallVectorImpl<char> &output,
#endif
}
bool pragma_once_outside_header() {
#if LLVM_VERSION_MAJOR < 4
return false;
#else
return true;
#endif
}
} // namespace llcompat
projects/hip/hipify-clang/src/LLVMCompat.h
+2
查看文件
@@ -81,4 +81,6 @@ void EnterPreprocessorTokenStream(clang::Preprocessor& _pp,
std::error_code real_path(const Twine &path, SmallVectorImpl<char> &output,
bool expand_tilde = false);
bool pragma_once_outside_header();
} // namespace llcompat
projects/hip/hipify-clang/src/main.cpp
+3
查看文件
@@ -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));
projects/hip/include/hip/hcc_detail/device_functions.h
+3 -3
查看文件
@@ -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) {
@@ -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) {
projects/hip/include/hip/hcc_detail/hip_common.h
+40
查看文件
@@ -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
projects/hip/include/hip/hcc_detail/hip_fp16.h
+7
查看文件
@@ -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 <hip/hcc_detail/hip_common.h>
#include "hip/hcc_detail/host_defines.h"
#include <assert.h>
#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
projects/hip/include/hip/hcc_detail/hip_ldg.h
+2 -2
查看文件
@@ -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__)
projects/hip/include/hip/hcc_detail/hip_runtime.h
+1 -13
查看文件
@@ -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 <hip/hcc_detail/hip_common.h>
//---
// Top part of file can be compiled with any compiler
projects/hip/include/hip/hcc_detail/hip_runtime_api.h
+16
查看文件
@@ -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]
*
@@ -2966,7 +2980,9 @@ hipError_t hipLaunchByPtr(const void* func);
} /* extern "c" */
#endif
#ifdef __cplusplus
#include <hip/hcc_detail/hip_prof_api.h>
#endif
#ifdef __cplusplus
extern "C" {
projects/hip/include/hip/hip_runtime_api.h
+6
查看文件
@@ -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 <class T>
static inline hipError_t hipMallocManaged(T** devPtr, size_t size,
unsigned int flags = hipMemAttachGlobal) {
return hipMallocManaged((void**)devPtr, size, flags);
}
#endif
#endif
projects/hip/include/hip/nvcc_detail/hip_runtime_api.h
+8 -1
查看文件
@@ -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)) {
@@ -866,7 +873,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) {
projects/hip/src/hip_memory.cpp
+20 -3
查看文件
@@ -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);
}
projects/hip/src/hip_module.cpp
+1 -1
查看文件
@@ -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));
}
projects/hip/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu
+1 -1
查看文件
@@ -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,
projects/hip/tests/src/runtimeApi/memory/hipMallocManaged.cpp
+52
查看文件
@@ -0,0 +1,52 @@
#include <hip/hip_runtime.h>
#include <math.h>
#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");
}
projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp
+1 -1
查看文件
@@ -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));