diff --git a/bin/hipify-perl b/bin/hipify-perl index 2887402b4e..ef77fe0ea9 100755 --- a/bin/hipify-perl +++ b/bin/hipify-perl @@ -1,4 +1,5 @@ -#!/usr/bin/perl -w +#!/usr/bin/perl -w + ## # Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved. # @@ -20,60 +21,55 @@ # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN # THE SOFTWARE. ## -#usage hipify-perl [OPTIONS] INPUT_FILE -use Getopt::Long; -my $warn_whitelist =""; +#usage hipify-perl [OPTIONS] INPUT_FILE + +use Getopt::Long; +my $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. - , "warn-whitelist=s"=> \$warn_whitelist - , "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. + "examine" => \$examine # Combines -no-output and -print-stats options. + , "inplace" => \$inplace # Modify input file inplace, replacing input with hipified output, save backup in .prehip file. + , "no-output" => \$no_output # Don't write any translated output to stdout. + , "print-stats" => \$print_stats # Print translation statistics. + , "quiet-warnings" => \$quiet_warnings # Don't print warnings on unknown CUDA functions. + , "whitelist=s" => \$whitelist # TODO: test it beforehand ); -$print_stats = 1 if $n; -$no_output = 1 if $n; +$print_stats = 1 if $examine; +$no_output = 1 if $examine; -# These uses of cuda[A-Z] are commonly used in CUDA code but don't actually map to any CUDA API: -# TODO - use a hash lookup for these. -@warn_whitelist = ( - "cudaDevice" - ,"cudaDevice_t" - ,"cudaIDs" - ,"cudaGridDim" - ,"cudaDimGrid" - ,"cudaDimBlock" - ,"cudaDeviceId" - ,"cudaDevices", - ,"cudaGradOutput", - ,"cudaInput", - ,"cudaOutput", - ,"cudaGradInput", - ,"cudaIndices", - ,"cudaGaugeField" - ,"cudaMom" - ,"cudaGauge" - ,"cudaInGauge" - ,"cudaColorSpinorField" - ,"cudaSiteLink" - ,"cudaFatLink" - ,"cudaStaple" - ,"cudaCloverField" - ,"cudaParam" - ); -#print "WW=@warn_whitelist\n"; +# Whitelist of cuda[A-Z] identifiers, which are commonly used in CUDA sources but don't map to any CUDA API: +@whitelist = ( + "cudaCloverField" + , "cudaColorSpinorField" + , "cudaDevice" + , "cudaDeviceId" + , "cudaDevice_t" + , "cudaDevices" + , "cudaDimBlock" + , "cudaDimGrid" + , "cudaFatLink" + , "cudaGauge" + , "cudaGaugeField" + , "cudaGradInput" + , "cudaGradOutput" + , "cudaGridDim" + , "cudaIDs" + , "cudaInGauge" + , "cudaIndices" + , "cudaInput" + , "cudaMom" + , "cudaOutput" + , "cudaParam" + , "cudaSiteLink" + , "cudaStaple" +); -# Allow users to add their own functions. -push (@warn_whitelist, split(',',$warn_whitelist)); +push(@whitelist, split(',', $whitelist)); -#Stats tracking code: -@statNames = ("error", "init", "version", "device", "context", "module", "memory", "addressing", "stream", "event", "external_resource_interop", "stream_memory", "execution", "graph", "occupancy", "texture", "surface", "peer", "graphics", "profiler", "openGL", "D3D9", "D3D10", "D3D11", "VDPAU", "EGL", "thread", "complex", "library", "device_library", "include", "include_cuda_main_header", "type", "literal", "numeric_literal", "define", "kernel_func", "extern_shared", "kern_launch"); +@statNames = ("error", "init", "version", "device", "context", "module", "memory", "addressing", "stream", "event", "external_resource_interop", "stream_memory", "execution", "graph", "occupancy", "texture", "surface", "peer", "graphics", "profiler", "openGL", "D3D9", "D3D10", "D3D11", "VDPAU", "EGL", "thread", "complex", "library", "device_library", "device_function", "include", "include_cuda_main_header", "type", "literal", "numeric_literal", "define", "extern_shared", "kernel_launch"); -#Compute total of all individual counts: sub totalStats { my %count = %{ shift() }; my $total = 0; @@ -90,32 +86,30 @@ sub printStats { my $warnings = shift(); my $loc = shift(); my $total = totalStats(\%counts); - printf STDERR "%s %d CUDA->HIP refs( ", $label, $total; + 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; -} + printf STDERR ")\n 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() }; + 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 $dest_ref = shift(); my @statNames = @{ shift() }; - foreach $stat (@statNames) { - $dest_ref->{$stat} = 0; + foreach $stat(@statNames) { + $dest_ref->{$stat} = 0; } -} +}; -# count of transforms in all files: +# Count of transforms in all files my %tt; clearStats(\%tt, \@statNames); $Twarnings = 0; @@ -146,22 +140,19 @@ while (@ARGV) { 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: + # Note : \b is used in perl to indicate the start of a word + # Count of transforms in this file my %ft; clearStats(\%ft, \@statNames); my $countIncludes = 0; - my $countKeywords = 0; # keywords like __global__, __shared__ - not converted by hipify-perl, but counted here. - my $warnings = 0; - my $warningsCublas = 0; - my $warningsCurand = 0; - my %warningTags; # hash with counts of particular unknown keywords. + my $countKeywords = 0; + my $warnings = 0; + my %warningTags; my $lineCount = 0; - undef $/; # Read whole file at once, so we can match newlines. + undef $/; + # Read whole file at once, so we can match newlines while () { - # chomp; - # next if /^(\s*(#.*)?)?$/; $ft{'error'} += s/\bcudaGetErrorName\b/hipGetErrorName/g; $ft{'error'} += s/\bcudaGetErrorString\b/hipGetErrorString/g; $ft{'error'} += s/\bcudaGetLastError\b/hipGetLastError/g; @@ -1651,13 +1642,11 @@ while (@ARGV) { # 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. + # 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: + # Match patterns for the below regular expression: #'extern __shared__ double foo[];' #'extern __shared__ unsigned int foo[];' #'extern volatile __shared__ double foo[];' @@ -1668,12 +1657,12 @@ while (@ARGV) { #'extern __shared__ blah::type s[];' #'extern __shared__ typename mapper::type s_data[];' #'extern __attribute__((used)) __shared__ typename mapper::type s_data[];' + $k += s/extern\s+([\w\(\)]+)?\s*__shared__\s+([\w:<>\s]+)\s+(\w+)\s*\[\s*\]\s*;/HIP_DYNAMIC_SHARED($1 $2, $3)/g; $ft{'extern_shared'} += $k; } - # CUDA Launch Syntax. Note these only work if launch is on a single line. { - # match uses ? for <.*> which will be unitialized if this is not present in launch syntax. + # Match uses ? for <.*> which will be unitialized if this is not present in launch syntax no warnings qw/uninitialized/; my $k = 0; @@ -1708,35 +1697,32 @@ while (@ARGV) { $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0, /g; if ($k) { - $ft{'kern_launch'} += $k; + $ft{'kernel_launch'} += $k; $Tkernels{$1} ++; } } - if ($count_conversions) { + if ($print_stats) { 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{'kernel_func'}; + my $hasDeviceCode = $countKeywords + $ft{'device_function'}; unless ($quiet_warnings) { - #print STDERR "Check WARNINGs\n"; - # copy into array of lines, process line-by-line to show warnings: - if ($hasDeviceCode or (/\bcuda/) or (/<<<.*>>>/) or (/(\bcublas[A-Z]\w+)/) or (/(\bcurand[A-Z]\w+)/) ) { + # Copy into array of lines, process line-by-line to show warnings + if ($hasDeviceCode or (/\bcu/) or (/\bCU_/) or (/\bCUDA_/) or (/<<<.*>>>/)) { my @lines = split /\n/, $_; - my $tmp = $_; # copies the whole file, could be a little smarter here... + # Copy the whole file + my $tmp = $_; my $line_num = 0; foreach (@lines) { $line_num ++; - # remove any whitelisted words: - foreach $w (@warn_whitelist) { + # Remove any whitelisted words + foreach $w (@whitelist) { s/\b$w\b/ZAP/ } 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 + # Flag any remaining code that look like cuda API calls: may want to add these to hipify $tag = (defined $1) ? $1 : "Launch"; } if (defined $tag) { @@ -1751,26 +1737,18 @@ while (@ARGV) { $_ = $tmp; } } - - # To limit bogus translations, try to make sure we are in a kernel: if ($hasDeviceCode > 0) { - $ft{'kernel_func'} += countSupportedDeviceFunctions(); + $ft{'device_function'} += countSupportedDeviceFunctions(); } - transformHostFunctions(); - - # Print it! - # TODO - would like to move this code outside loop but it uses $_ which contains the whole file. + # TODO: would like to move this code outside loop but it uses $_ which contains the whole file unless ($no_output) { my $apiCalls = $ft{'error'} + $ft{'init'} + $ft{'version'} + $ft{'device'} + $ft{'context'} + $ft{'module'} + $ft{'memory'} + $ft{'addressing'} + $ft{'stream'} + $ft{'event'} + $ft{'external_resource_interop'} + $ft{'stream_memory'} + $ft{'execution'} + $ft{'graph'} + $ft{'occupancy'} + $ft{'texture'} + $ft{'surface'} + $ft{'peer'} + $ft{'graphics'} + $ft{'profiler'} + $ft{'openGL'} + $ft{'D3D9'} + $ft{'D3D10'} + $ft{'D3D11'} + $ft{'VDPAU'} + $ft{'EGL'} + $ft{'thread'} + $ft{'complex'} + $ft{'library'} + $ft{'device_library'} + $ft{'include'} + $ft{'include_cuda_main_header'} + $ft{'type'} + $ft{'literal'} + $ft{'numeric_literal'} + $ft{'define'}; - my $kernStuff = $hasDeviceCode + $ft{'kern_launch'} + $ft{'kernel_func'}; + my $kernStuff = $hasDeviceCode + $ft{'kernel_launch'} + $ft{'device_function'}; my $totalCalls = $apiCalls + $kernStuff; $is_dos = m/\r\n$/; if ($totalCalls and ($countIncludes == 0) and ($kernStuff != 0)) { # TODO: implement hipify-clang's logic with header files AMAP - # 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. print $OUTFILE '#include "hip/hip_runtime.h"' . ($is_dos ? "\r\n" : "\n"); } print $OUTFILE "$_"; @@ -1778,12 +1756,11 @@ while (@ARGV) { $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); + printStats(" info: converted", \@statNames, \%ft, $warnings, $lineCount); print STDERR " in '$fileName'\n"; } - # Update totals for all files: + # Update totals for all files addStats(\%tt, \%ft); $Twarnings += $warnings; $TlineCount += $lineCount; @@ -1791,10 +1768,10 @@ while (@ARGV) { $TwarningTags{$key} += $warningTags{$key}; } } -#-- Print total stats for all files processed: +# Print total stats for all files processed: if ($print_stats and ($fileCount > 1)) { print STDERR "\n"; - printStats("info: TOTAL-converted", \@statNames, \%tt, $Twarnings, $TlineCount); + 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}; @@ -1807,7 +1784,7 @@ if ($print_stats and ($fileCount > 1)) { print STDERR "\n"; print STDERR "\n"; } -if ($count_conversions) { +if ($print_stats) { foreach my $key (sort { $convertedTags{$b} <=> $convertedTags{$a} } keys %convertedTags) { printf STDERR " %s %d\n", $key, $convertedTags{$key}; } diff --git a/docs/markdown/hip_porting_driver_api.md b/docs/markdown/hip_porting_driver_api.md index 7d02666af3..8e66780add 100644 --- a/docs/markdown/hip_porting_driver_api.md +++ b/docs/markdown/hip_porting_driver_api.md @@ -105,11 +105,13 @@ hip-clang emits a global variable `__hip_gpubin_handle` of void** type with link #### Kernel Launching hip-clang supports kernel launching by CUDA `<<<>>>` syntax, hipLaunchKernel, and hipLaunchKernelGGL. The latter two are macros which expand to CUDA `<<<>>>` syntax. -In host code, hip-clang emits a stub function with the same name and arguments as the kernel. In the body of this function, hipSetupArgument is called for each kernel argument, then hipLaunchByPtr is called with a function pointer to the stub function. - When the executable or shared library is loaded by the dynamic linker, the initilization functions are called. In the initialization functions, when `__hipRegisterFatBinary` is called, the code objects containing all kernels are loaded; when `__hipRegisterFunction` is called, the stub functions are associated with the corresponding kernels in code objects. -In the host code, for the `<<<>>>` statement, hip-clang first emits call of hipConfigureCall to set up the threads and grids, then emits call of the stub function with the given arguments. In the stub function, when the runtime host API function hipLaunchByPtr is called, the real kernel associated with the stub function is launched. +hip-clang implements two sets of kernel launching APIs. + +By default, in the host code, for the `<<<>>>` statement, hip-clang first emits call of hipConfigureCall to set up the threads and grids, then emits call of the stub function with the given arguments. In the stub function, hipSetupArgument is called for each kernel argument, then hipLaunchByPtr is called with a function pointer to the stub function. In hipLaunchByPtr, the real kernel associated with the stub function is launched. + +If HIP program is compiled with -fhip-new-launch-api, in the host code, for the `<<<>>>` statement, hip-clang first emits call of `__hipPushCallConfiguration` to save the grid dimension, block dimension, shared memory usage and stream to a stack, then emits call of the stub function with the given arguments. In the stub function, `__hipPopCallConfiguration` is called to get the saved grid dimension, block dimension, shared memory usage and stream, then hipLaunchKernel is called with a function pointer to the stub function. In hipLaunchKernel, the real kernel associated with the stub function is launched. ### NVCC Implementation Notes diff --git a/docs/markdown/hip_profiling.md b/docs/markdown/hip_profiling.md index a659216044..8a44368680 100644 --- a/docs/markdown/hip_profiling.md +++ b/docs/markdown/hip_profiling.md @@ -130,7 +130,7 @@ void FunctionFoo(...) The HIP marker API is only supported on ROCm platform. The marker macros are defined on CUDA platforms and will compile, but are silently ignored at runtime. -This [HIP sample](samples/2_Cookbook/2_Profiler/) shows the profiler marker API used in a small application. +This [HIP sample](https://github.com/ROCm-Developer-Tools/HIP/tree/master/samples/2_Cookbook/2_Profiler) shows the profiler marker API used in a small application. More information on the marker API can be found in the profiler header file and PDF in a ROCm installation: - /opt/rocm/profiler/CXLActivityLogger/include/CXLActivityLogger.h @@ -236,7 +236,7 @@ Here's an example for one API followed by a description for the sections of the - The second line shows the completion of the API, including the numeric return value (`ret= 0`) as well as an string representation for the error code (`hipSuccess`). If the returned error code is non-zero, then the csecond line message is shown in red (unless HIP_TRACE_API_COLOR is "none" - see below). -Heres a specific example showing the output of the [square](samples/0_Intro/square) program running on HIP: +Heres a specific example showing the output of the [square](https://github.com/ROCm-Developer-Tools/HIP/tree/master/samples/0_Intro/square) program running on HIP: ``` $ HIP_TRACE_API=1 ./square.hip.out diff --git a/hipify-clang/README.md b/hipify-clang/README.md index 4631b0de40..c7524ad0f1 100644 --- a/hipify-clang/README.md +++ b/hipify-clang/README.md @@ -144,9 +144,9 @@ To run it: * Path to cuDNN should be specified by the `CUDA_DNN_ROOT_DIR` option: - - Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.3.30` + - Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.4.38` - - Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.3.30` + - Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.4.38` 5. Ensure [`python`](https://www.python.org/downloads) of minimum required version 2.7 is installed. @@ -180,9 +180,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-5.0.5 - cudnn-7.6.3.30 +Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.6.4.38 -Ubuntu 16-18: LLVM 8.0.0 - 9.0.0, CUDA 8.0 - 10.1, cudnn-5.1.10 - cudnn-7.6.3.30 +Ubuntu 16-18: LLVM 8.0.0 - 9.0.0, CUDA 8.0 - 10.1, cudnn-5.1.10 - cudnn-7.6.4.38 Build system for the above configurations: @@ -197,7 +197,7 @@ cmake -DCMAKE_INSTALL_PREFIX=../dist \ -DCMAKE_PREFIX_PATH=/srv/git/LLVM/9.0.0/dist \ -DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-10.1 \ - -DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.3.30 \ + -DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.4.38 \ -DLLVM_EXTERNAL_LIT=/srv/git/LLVM/9.0.0/build/bin/llvm-lit \ .. ``` @@ -325,9 +325,9 @@ On Windows 10 the following configurations are tested: 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-7.0.5.15 - cudnn-7.6.3.30 +LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn-7.0.5.15 - cudnn-7.6.4.38 -LLVM 7.0.0 - 9.0.0, CUDA 7.5 - 10.1, cudnn-7.0.5.15 - cudnn-7.6.3.30 +LLVM 7.0.0 - 9.0.0, CUDA 7.5 - 10.1, cudnn-7.0.5.15 - cudnn-7.6.4.38 Build system for the above configurations: @@ -344,7 +344,7 @@ cmake -DCMAKE_PREFIX_PATH=f:/LLVM/9.0.0/dist \ -DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v10.1" \ -DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v10.1" \ - -DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.3.30 \ + -DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.4.38 \ -DLLVM_EXTERNAL_LIT=f:/LLVM/9.0.0/build/Release/bin/llvm-lit.py \ -Thost=x64 .. diff --git a/hipify-clang/src/CUDA2HIP_Perl.cpp b/hipify-clang/src/CUDA2HIP_Perl.cpp index 5407cae35c..db9831b315 100644 --- a/hipify-clang/src/CUDA2HIP_Perl.cpp +++ b/hipify-clang/src/CUDA2HIP_Perl.cpp @@ -36,6 +36,29 @@ using namespace llvm; namespace perl { + const std::string sCopyright = + "##\n" + "# Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved.\n" + "#\n" + "# Permission is hereby granted, free of charge, to any person obtaining a copy\n" + "# of this software and associated documentation files (the \"Software\"), to deal\n" + "# in the Software without restriction, including without limitation the rights\n" + "# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell\n" + "# copies of the Software, and to permit persons to whom the Software is\n" + "# furnished to do so, subject to the following conditions:\n" + "#\n" + "# The above copyright notice and this permission notice shall be included in\n" + "# all copies or substantial portions of the Software.\n" + "#\n" + "# THE SOFTWARE IS PROVIDED \"AS IS\", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR\n" + "# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,\n" + "# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE\n" + "# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER\n" + "# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,\n" + "# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN\n" + "# THE SOFTWARE.\n" + "##\n"; + const std::string tab = " "; const std::string double_tab = tab + tab; const std::string triple_tab = double_tab + tab; @@ -44,46 +67,130 @@ namespace perl { const std::string sReturn_m = "return $m;\n"; const std::string sForeach = "foreach $func (\n"; const std::string sMy = "my $m = 0;\n"; + const std::string sCudaDevice = "cudaDevice"; + const std::string sCudaDeviceId = "cudaDeviceId"; + const std::string sCudaDevices = "cudaDevices"; + const std::string sCudaDevice_t = "cudaDevice_t"; + const std::string sCudaIDs = "cudaIDs"; + const std::string sCudaGridDim = "cudaGridDim"; + const std::string sCudaDimGrid = "cudaDimGrid"; + const std::string sCudaDimBlock = "cudaDimBlock"; + const std::string sCudaGradInput = "cudaGradInput"; + const std::string sCudaGradOutput = "cudaGradOutput"; + const std::string sCudaInput = "cudaInput"; + const std::string sCudaOutput = "cudaOutput"; + const std::string sCudaIndices = "cudaIndices"; + const std::string sCudaGaugeField = "cudaGaugeField"; + const std::string sCudaMom = "cudaMom"; + const std::string sCudaGauge = "cudaGauge"; + const std::string sCudaInGauge = "cudaInGauge"; + const std::string sCudaColorSpinorField = "cudaColorSpinorField"; + const std::string sCudaSiteLink = "cudaSiteLink"; + const std::string sCudaFatLink = "cudaFatLink"; + const std::string sCudaStaple = "cudaStaple"; + const std::string sCudaCloverField = "cudaCloverField"; + const std::string sCudaParam = "cudaParam"; - void generateHostFunctions(std::unique_ptr& perlStreamPtr) { - *perlStreamPtr.get() << "\n" << sSub << " transformHostFunctions\n" << "{\n" << tab << sMy; + const std::set Whitelist{ + {sCudaDevice}, {sCudaDevice_t}, {sCudaIDs}, {sCudaGridDim}, {sCudaDimGrid}, {sCudaDimBlock}, {sCudaDeviceId}, {sCudaDevices}, + {sCudaGradInput}, {sCudaGradOutput}, {sCudaInput}, {sCudaOutput}, {sCudaIndices}, {sCudaGaugeField}, {sCudaMom}, {sCudaGauge}, + {sCudaInGauge}, {sCudaColorSpinorField}, {sCudaSiteLink}, {sCudaFatLink}, {sCudaStaple}, {sCudaCloverField}, {sCudaParam} + }; + + void generateHeader(std::unique_ptr& streamPtr) { + *streamPtr.get() << "#!/usr/bin/perl -w" << std::endl << std::endl; + *streamPtr.get() << sCopyright << std::endl; + *streamPtr.get() << "#usage hipify-perl [OPTIONS] INPUT_FILE" << std::endl << std::endl; + *streamPtr.get() << "use Getopt::Long;" << std::endl; + *streamPtr.get() << "my $whitelist = \"\";" << std::endl << std::endl; + *streamPtr.get() << "GetOptions(" << std::endl; + *streamPtr.get() << tab << " \"examine\" => \\$examine # Combines -no-output and -print-stats options." << std::endl; + *streamPtr.get() << tab << ", \"inplace\" => \\$inplace # Modify input file inplace, replacing input with hipified output, save backup in .prehip file." << std::endl; + *streamPtr.get() << tab << ", \"no-output\" => \\$no_output # Don't write any translated output to stdout." << std::endl; + *streamPtr.get() << tab << ", \"print-stats\" => \\$print_stats # Print translation statistics." << std::endl; + *streamPtr.get() << tab << ", \"quiet-warnings\" => \\$quiet_warnings # Don't print warnings on unknown CUDA functions." << std::endl; + *streamPtr.get() << tab << ", \"whitelist=s\" => \\$whitelist # TODO: test it beforehand" << std::endl; + *streamPtr.get() << ");" << std::endl << std::endl; + *streamPtr.get() << "$print_stats = 1 if $examine;" << std::endl; + *streamPtr.get() << "$no_output = 1 if $examine;" << std::endl << std::endl; + *streamPtr.get() << "# Whitelist of cuda[A-Z] identifiers, which are commonly used in CUDA sources but don't map to any CUDA API:" << std::endl; + *streamPtr.get() << "@whitelist = ("; + unsigned int num = 0; + for (const std::string &m : Whitelist) { + *streamPtr.get() << std::endl << tab << (num ? ", " : " ") << "\"" << m << "\""; + ++num; + } + *streamPtr.get() << std::endl << ");" << std::endl << std::endl; + *streamPtr.get() << "push(@whitelist, split(',', $whitelist));" << std::endl << std::endl; + } + + void generateStatFunctions(std::unique_ptr& streamPtr) { + *streamPtr.get() << std::endl << sSub << " totalStats" << " {" << std::endl; + *streamPtr.get() << tab << "my %count = %{ shift() };" << std::endl; + *streamPtr.get() << tab << "my $total = 0;" << std::endl; + *streamPtr.get() << tab << "foreach $key (keys %count) {" << std::endl; + *streamPtr.get() << double_tab << "$total += $count{$key};" << std::endl << tab << "}" << std::endl; + *streamPtr.get() << tab << "return $total;" << std::endl << "};" << std::endl; + *streamPtr.get() << std::endl << sSub << " printStats" << " {" << std::endl; + *streamPtr.get() << tab << "my $label = shift();" << std::endl; + *streamPtr.get() << tab << "my @statNames = @{ shift() };" << std::endl; + *streamPtr.get() << tab << "my %counts = %{ shift() };" << std::endl; + *streamPtr.get() << tab << "my $warnings = shift();" << std::endl; + *streamPtr.get() << tab << "my $loc = shift();" << std::endl; + *streamPtr.get() << tab << "my $total = totalStats(\\%counts);" << std::endl; + *streamPtr.get() << tab << "printf STDERR \"%s %d CUDA->HIP refs ( \", $label, $total;" << std::endl; + *streamPtr.get() << tab << "foreach $stat (@statNames) {" << std::endl; + *streamPtr.get() << double_tab << "printf STDERR \"%s:%d \", $stat, $counts{$stat};" << std::endl; + *streamPtr.get() << tab << "}" << std::endl; + *streamPtr.get() << tab << "printf STDERR \")\\n warn:%d LOC:%d\", $warnings, $loc;" << std::endl << "};" << std::endl; + for (int i = 0; i < 2; ++i) { + *streamPtr.get() << std::endl << sSub << " " << (i ? "clearStats" : "addStats") << " {" << std::endl; + *streamPtr.get() << tab << "my $dest_ref = shift();" << std::endl; + *streamPtr.get() << tab << (i ? "my @statNames = @{ shift() };" : "my %adder = %{ shift() };") << std::endl; + *streamPtr.get() << tab << "foreach " << (i ? "$stat(@statNames)" : "$key (keys %adder)") << " {" << std::endl; + *streamPtr.get() << double_tab << "$dest_ref->" << (i ? "{$stat} = 0;" : "{$key} += $adder{$key};") << std::endl << tab << "}" << std::endl << "};" << std::endl; + } + } + + void generateHostFunctions(std::unique_ptr& streamPtr) { + *streamPtr.get() << std::endl << sSub << " transformHostFunctions" << "{" << std::endl << tab << sMy; std::set &funcSet = DeviceSymbolFunctions0; const std::string s0 = "$m += s/(?second.hipName.str() << "\""; + *streamPtr.get() << (count ? ",\n" : "") << double_tab << "\"" << found->second.hipName.str() << "\""; count++; } } - *perlStreamPtr.get() << "\n" << tab << ")\n" << tab << "{\n" << double_tab; + *streamPtr.get() << std::endl << tab << ")" << std::endl << tab << "{" << std::endl << double_tab; switch (i) { - case 0: - default: - *perlStreamPtr.get() << s0 << sHIP_SYMBOL << "\\($2\\),/g\n"; break; - case 1: - *perlStreamPtr.get() << s1 << sHIP_SYMBOL << "\\($3\\)$4/g;\n"; break; - case 2: - *perlStreamPtr.get() << s0 << s_reinterpret_cast << "\\($2\\),/g\n"; break; - case 3: - *perlStreamPtr.get() << s1 << s_reinterpret_cast << "\\($3\\)$4/g;\n"; break; + case 0: + default: + *streamPtr.get() << s0 << sHIP_SYMBOL << "\\($2\\),/g" << std::endl; break; + case 1: + *streamPtr.get() << s1 << sHIP_SYMBOL << "\\($3\\)$4/g;" << std::endl; break; + case 2: + *streamPtr.get() << s0 << s_reinterpret_cast << "\\($2\\),/g" << std::endl; break; + case 3: + *streamPtr.get() << s1 << s_reinterpret_cast << "\\($3\\)$4/g;" << std::endl; break; } - *perlStreamPtr.get() << tab << "}\n"; + *streamPtr.get() << tab << "}" << std::endl; } - *perlStreamPtr.get() << tab << sReturn_m << "}\n"; + *streamPtr.get() << tab << sReturn_m << "}" << std::endl; } - void generateDeviceFunctions(std::unique_ptr& perlStreamPtr) { + void generateDeviceFunctions(std::unique_ptr& streamPtr) { unsigned int countUnsupported = 0; unsigned int countSupported = 0; std::stringstream sSupported; @@ -98,39 +205,39 @@ namespace perl { std::stringstream subWarnUnsupported; std::stringstream subCommon; std::string sCommon = tab + sMy + tab + sForeach; - subCountSupported << "\n" << sSub << " countSupportedDeviceFunctions\n" << "{\n" << (countSupported ? sCommon : tab + sReturn_0); - subWarnUnsupported << "\n" << sSub << " warnUnsupportedDeviceFunctions\n" << "{\n" << (countUnsupported ? tab + "my $line_num = shift;\n" + sCommon : tab + sReturn_0); + subCountSupported << std::endl << sSub << " countSupportedDeviceFunctions" << " {" << std::endl << (countSupported ? sCommon : tab + sReturn_0); + subWarnUnsupported << std::endl << sSub << " warnUnsupportedDeviceFunctions" << " {" << std::endl << (countUnsupported ? tab + "my $line_num = shift;\n" + sCommon : tab + sReturn_0); if (countSupported) { - subCountSupported << sSupported.str() << "\n" << tab << ")\n"; + subCountSupported << sSupported.str() << std::endl << tab << ")" << std::endl; } if (countUnsupported) { - subWarnUnsupported << sUnsupported.str() << "\n" << tab << ")\n"; + subWarnUnsupported << sUnsupported.str() << std::endl << tab << ")" << std::endl; } if (countSupported || countUnsupported) { - subCommon << tab << "{\n"; - subCommon << double_tab << "# match device function from the list, except those, which have a namespace prefix (aka somenamespace::umin(...));\n"; - subCommon << double_tab << "# function with only global namespace qualifier '::' (aka ::umin(...)) should be treated as a device function (and warned as well as without such qualifier);\n"; - subCommon << double_tab << "my $mt_namespace = m/(\\w+)::($func)\\s*\\(\\s*.*\\s*\\)/g;\n"; - subCommon << double_tab << "my $mt = m/($func)\\s*\\(\\s*.*\\s*\\)/g;\n"; - subCommon << double_tab << "if ($mt && !$mt_namespace) {\n"; - subCommon << triple_tab << "$m += $mt;\n"; + subCommon << tab << "{" << std::endl; + subCommon << double_tab << "# match device function from the list, except those, which have a namespace prefix (aka somenamespace::umin(...));" << std::endl; + subCommon << double_tab << "# function with only global namespace qualifier '::' (aka ::umin(...)) should be treated as a device function (and warned as well as without such qualifier);" << std::endl; + subCommon << double_tab << "my $mt_namespace = m/(\\w+)::($func)\\s*\\(\\s*.*\\s*\\)/g;" << std::endl; + subCommon << double_tab << "my $mt = m/($func)\\s*\\(\\s*.*\\s*\\)/g;" << std::endl; + subCommon << double_tab << "if ($mt && !$mt_namespace) {" << std::endl; + subCommon << triple_tab << "$m += $mt;" << std::endl; } if (countSupported) { subCountSupported << subCommon.str(); } if (countUnsupported) { subWarnUnsupported << subCommon.str(); - subWarnUnsupported << triple_tab << "print STDERR \" warning: $fileName:$line_num: unsupported device function \\\"$func\\\": $_\\n\";\n"; + subWarnUnsupported << triple_tab << "print STDERR \" warning: $fileName:$line_num: unsupported device function \\\"$func\\\": $_\\n\";" << std::endl; } if (countSupported || countUnsupported) { sCommon = double_tab + "}\n" + tab + "}\n" + tab + sReturn_m; } if (countSupported) subCountSupported << sCommon; if (countUnsupported) subWarnUnsupported << sCommon; - subCountSupported << "}\n"; - subWarnUnsupported << "}\n"; - *perlStreamPtr.get() << subCountSupported.str(); - *perlStreamPtr.get() << subWarnUnsupported.str(); + subCountSupported << "}" << std::endl; + subWarnUnsupported << "}" << std::endl; + *streamPtr.get() << subCountSupported.str(); + *streamPtr.get() << subWarnUnsupported.str(); } bool generate(bool Generate) { @@ -150,15 +257,17 @@ namespace perl { llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << tmpFile << "\n"; return false; } - std::unique_ptr perlStreamPtr = std::unique_ptr(new std::ofstream(tmpFile.c_str(), std::ios_base::trunc)); + std::unique_ptr streamPtr = std::unique_ptr(new std::ofstream(tmpFile.c_str(), std::ios_base::trunc)); + generateHeader(streamPtr); std::string sConv = "my $conversions = "; - *perlStreamPtr.get() << "@statNames = ("; + *streamPtr.get() << "@statNames = ("; for (int i = 0; i < NUM_CONV_TYPES - 1; ++i) { - *perlStreamPtr.get() << "\"" << counterNames[i] << "\", "; + *streamPtr.get() << "\"" << counterNames[i] << "\", "; sConv += "$ft{'" + std::string(counterNames[i]) + "'} + "; } - *perlStreamPtr.get() << "\"" << counterNames[NUM_CONV_TYPES - 1] << "\");\n\n"; - *perlStreamPtr.get() << sConv << "$ft{'" << counterNames[NUM_CONV_TYPES - 1] << "'};\n\n"; + *streamPtr.get() << "\"" << counterNames[NUM_CONV_TYPES - 1] << "\");" << std::endl; + generateStatFunctions(streamPtr); + *streamPtr.get() << std::endl << sConv << "$ft{'" << counterNames[NUM_CONV_TYPES - 1] << "'};" << std::endl << std::endl; for (int i = 0; i < NUM_CONV_TYPES; ++i) { if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) { for (auto& ma : CUDA_INCLUDE_MAP) { @@ -168,21 +277,22 @@ namespace perl { std::string sHIP = ma.second.hipName.str(); sCUDA = std::regex_replace(sCUDA, std::regex("/"), "\\/"); sHIP = std::regex_replace(sHIP, std::regex("/"), "\\/"); - *perlStreamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << sCUDA << "\\b/" << sHIP << "/g;\n"; + *streamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << sCUDA << "\\b/" << sHIP << "/g;" << std::endl; } } - } else { + } + else { for (auto& ma : CUDA_RENAMES_MAP()) { if (Statistics::isUnsupported(ma.second)) continue; if (i == ma.second.type) { - *perlStreamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << ma.first.str() << "\\b/" << ma.second.hipName.str() << "/g;\n"; + *streamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << ma.first.str() << "\\b/" << ma.second.hipName.str() << "/g;" << std::endl; } } } } - generateHostFunctions(perlStreamPtr); - generateDeviceFunctions(perlStreamPtr); - perlStreamPtr.get()->flush(); + generateHostFunctions(streamPtr); + generateDeviceFunctions(streamPtr); + streamPtr.get()->flush(); bool ret = true; EC = sys::fs::copy_file(tmpFile, dstPerlMap); if (EC) { diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index e331cfbec5..cb0a5eedc9 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/hipify-clang/src/HipifyAction.cpp @@ -380,7 +380,7 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc ct::Replacement Rep(*SM, launchStart, length, OS.str()); clang::FullSourceLoc fullSL(launchStart, *SM); insertReplacement(Rep, fullSL); - hipCounter counter = {"hipLaunchKernelGGL", "", ConvTypes::CONV_EXECUTION, ApiTypes::API_RUNTIME}; + hipCounter counter = {"hipLaunchKernelGGL", "", ConvTypes::CONV_KERNEL_LAUNCH, ApiTypes::API_RUNTIME}; Statistics::current().incrementCounter(counter, refName.str()); return true; } @@ -423,7 +423,7 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match ct::Replacement Rep(*SM, slStart, repLength, repName); clang::FullSourceLoc fullSL(slStart, *SM); insertReplacement(Rep, fullSL); - hipCounter counter = {sHIP_DYNAMIC_SHARED, "", ConvTypes::CONV_MEMORY, ApiTypes::API_RUNTIME}; + hipCounter counter = {sHIP_DYNAMIC_SHARED, "", ConvTypes::CONV_EXTERN_SHARED, ApiTypes::API_RUNTIME}; Statistics::current().incrementCounter(counter, refName.str()); return true; } diff --git a/hipify-clang/src/Statistics.cpp b/hipify-clang/src/Statistics.cpp index d3efd4a5d5..9751763be3 100644 --- a/hipify-clang/src/Statistics.cpp +++ b/hipify-clang/src/Statistics.cpp @@ -63,7 +63,9 @@ const char *counterNames[NUM_CONV_TYPES] = { "type", // CONV_TYPE "literal", // CONV_LITERAL "numeric_literal", // CONV_NUMERIC_LITERAL - "define" // CONV_DEFINE + "define", // CONV_DEFINE + "extern_shared", // CONV_EXTERN_SHARED + "kernel_launch" // CONV_KERNEL_LAUNCH }; const char *counterTypes[NUM_CONV_TYPES] = { @@ -102,7 +104,9 @@ const char *counterTypes[NUM_CONV_TYPES] = { "CONV_TYPE", "CONV_LITERAL", "CONV_NUMERIC_LITERAL", - "CONV_DEFINE" + "CONV_DEFINE", + "CONV_EXTERN_SHARED", + "CONV_KERNEL_LAUNCH" }; const char *apiNames[NUM_API_TYPES] = { diff --git a/hipify-clang/src/Statistics.h b/hipify-clang/src/Statistics.h index 91f493f4a5..051f680fb1 100644 --- a/hipify-clang/src/Statistics.h +++ b/hipify-clang/src/Statistics.h @@ -119,6 +119,8 @@ enum ConvTypes { CONV_LITERAL, CONV_NUMERIC_LITERAL, CONV_DEFINE, + CONV_EXTERN_SHARED, + CONV_KERNEL_LAUNCH, CONV_LAST }; constexpr int NUM_CONV_TYPES = (int) ConvTypes::CONV_LAST; diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 1199f7b663..a6f9cc0826 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -690,12 +690,7 @@ extern "C" uint64_t __clock_u64() __HC__; __device__ inline __attribute((always_inline)) long long int __clock64() { -// ToDo: Unify HCC and HIP implementation. -#if __HCC__ - return (long long int) __clock_u64(); -#else - return (long long int) __builtin_amdgcn_s_memrealtime(); -#endif +return (long long int) __builtin_readcyclecounter(); } __device__ diff --git a/include/hip/hcc_detail/hip_fp16.h b/include/hip/hcc_detail/hip_fp16.h index 11e1b6be11..18ad566121 100644 --- a/include/hip/hcc_detail/hip_fp16.h +++ b/include/hip/hcc_detail/hip_fp16.h @@ -223,6 +223,7 @@ THE SOFTWARE. Enable_if_t< std::is_floating_point{} && !std::is_same{}>* = nullptr> + __host__ __device__ operator T() const { return data; } #endif __host__ __device__ diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index a614e3599d..0b70570352 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -37,6 +37,7 @@ THE SOFTWARE. //#include #if __cplusplus #include +#include #else #include #include @@ -198,35 +199,93 @@ __device__ int __hip_move_dpp_N(int src); #if defined __HCC__ -template < - typename std::common_type::type f> -class Coordinates { - using R = decltype(f(0)); +namespace hip_impl { + struct GroupId { + using R = decltype(hc_get_group_id(0)); - struct X { - __device__ operator R() const { return f(0); } - __device__ uint32_t operator=(R _) { return f(0); } - }; - struct Y { - __device__ operator R() const { return f(1); } - __device__ uint32_t operator=(R _) { return f(1); } - }; - struct Z { - __device__ operator R() const { return f(2); } - __device__ uint32_t operator=(R _) { return f(2); } - }; + __device__ + R operator()(std::uint32_t x) const noexcept { return hc_get_group_id(x); } + }; + struct GroupSize { + using R = decltype(hc_get_group_size(0)); - public: - static constexpr X x{}; - static constexpr Y y{}; - static constexpr Z z{}; + __device__ + R operator()(std::uint32_t x) const noexcept { + return hc_get_group_size(x); + } + }; + struct NumGroups { + using R = decltype(hc_get_num_groups(0)); + + __device__ + R operator()(std::uint32_t x) const noexcept { + return hc_get_num_groups(x); + } + }; + struct WorkitemId { + using R = decltype(hc_get_workitem_id(0)); + + __device__ + R operator()(std::uint32_t x) const noexcept { + return hc_get_workitem_id(x); + } + }; +} // Namespace hip_impl. + +template +struct Coordinates { + using R = decltype(F{}(0)); + + struct X { __device__ operator R() const noexcept { return F{}(0); } }; + struct Y { __device__ operator R() const noexcept { return F{}(1); } }; + struct Z { __device__ operator R() const noexcept { return F{}(2); } }; + + static constexpr X x{}; + static constexpr Y y{}; + static constexpr Z z{}; }; -static constexpr Coordinates blockDim; -static constexpr Coordinates blockIdx; -static constexpr Coordinates gridDim; -static constexpr Coordinates threadIdx; +inline +__device__ +std::uint32_t operator*(Coordinates::X, + Coordinates::X) noexcept { + return hc_get_grid_size(0); +} +inline +__device__ +std::uint32_t operator*(Coordinates::X, + Coordinates::X) noexcept { + return hc_get_grid_size(0); +} +inline +__device__ +std::uint32_t operator*(Coordinates::Y, + Coordinates::Y) noexcept { + return hc_get_grid_size(1); +} +inline +__device__ +std::uint32_t operator*(Coordinates::Y, + Coordinates::Y) noexcept { + return hc_get_grid_size(1); +} +inline +__device__ +std::uint32_t operator*(Coordinates::Z, + Coordinates::Z) noexcept { + return hc_get_grid_size(2); +} +inline +__device__ +std::uint32_t operator*(Coordinates::Z, + Coordinates::Z) noexcept { + return hc_get_grid_size(2); +} + +static constexpr Coordinates blockDim{}; +static constexpr Coordinates blockIdx{}; +static constexpr Coordinates gridDim{}; +static constexpr Coordinates threadIdx{}; #define hipThreadIdx_x (hc_get_workitem_id(0)) #define hipThreadIdx_y (hc_get_workitem_id(1)) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 35b08be2ff..edd63b96f0 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -1519,34 +1519,6 @@ hipError_t hipMemcpyToSymbol(void*, const void*, size_t, size_t, hipMemcpyKind, } // Namespace hip_impl. #endif -#ifdef __cplusplus -extern "C" { -#endif - -/** - * @brief C compliant kernel launch API - * - * @param [in] function_address - kernel function pointer. - * @param [in] numBlocks - number of blocks - * @param [in] dimBlocks - dimension of a block - * @param [in] args - kernel arguments - * @param [in] sharedMemBytes - Amount of dynamic shared memory to allocate for this kernel. The - * Kernel can access this with HIP_DYNAMIC_SHARED. - * @param [in] stream - Stream where the kernel should be dispatched. May be 0, in which case th - * default stream is used with associated synchronization rules. - * - * @returns #hipSuccess, #hipErrorInvalidValue, hipInvalidDevice - * - */ - -hipError_t hipLaunchKernel(const void* function_address, - dim3 numBlocks, dim3 dimBlocks, void** args, - size_t sharedMemBytes, hipStream_t stream); - -#ifdef __cplusplus -} -#endif - #if defined(__cplusplus) extern "C" { #endif @@ -3055,6 +3027,65 @@ hipError_t hipSetupArgument(const void* arg, size_t size, size_t offset); hipError_t hipLaunchByPtr(const void* func); +/** + * @brief Push configuration of a kernel launch. + * + * @param [in] gridDim grid dimension specified as multiple of blockDim. + * @param [in] blockDim block dimensions specified in work-items + * @param [in] sharedMem Amount of dynamic shared memory to allocate for this kernel. The + * kernel can access this with HIP_DYNAMIC_SHARED. + * @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case the + * default stream is used with associated synchronization rules. + * + * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue + * + */ + +hipError_t __hipPushCallConfiguration(dim3 gridDim, + dim3 blockDim, + size_t sharedMem __dparm(0), + hipStream_t stream __dparm(0)); + +/** + * @brief Pop configuration of a kernel launch. + * + * @param [out] gridDim grid dimension specified as multiple of blockDim. + * @param [out] blockDim block dimensions specified in work-items + * @param [out] sharedMem Amount of dynamic shared memory to allocate for this kernel. The + * kernel can access this with HIP_DYNAMIC_SHARED. + * @param [out] stream Stream where the kernel should be dispatched. May be 0, in which case the + * default stream is used with associated synchronization rules. + * + * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue + * + */ +hipError_t __hipPopCallConfiguration(dim3 *gridDim, + dim3 *blockDim, + size_t *sharedMem, + hipStream_t *stream); + +/** + * @brief C compliant kernel launch API + * + * @param [in] function_address - kernel stub function pointer. + * @param [in] numBlocks - number of blocks + * @param [in] dimBlocks - dimension of a block + * @param [in] args - kernel arguments + * @param [in] sharedMemBytes - Amount of dynamic shared memory to allocate for this kernel. The + * Kernel can access this with HIP_DYNAMIC_SHARED. + * @param [in] stream - Stream where the kernel should be dispatched. May be 0, in which case th + * default stream is used with associated synchronization rules. + * + * @returns #hipSuccess, #hipErrorInvalidValue, hipInvalidDevice + * + */ + +hipError_t hipLaunchKernel(const void* function_address, + dim3 numBlocks, + dim3 dimBlocks, + void** args, + size_t sharedMemBytes __dparm(0), + hipStream_t stream __dparm(0)); /** * @} @@ -3065,7 +3096,9 @@ hipError_t hipLaunchByPtr(const void* func); } /* extern "c" */ #endif +#if USE_PROF_API #include +#endif #ifdef __cplusplus extern "C" { diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 4c07b9777b..5fe6ce1996 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -132,20 +132,6 @@ extern hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, int device); return ihipLogStatus(hipStatus); \ } -hipError_t hipModuleUnload(hipModule_t hmod) { - HIP_INIT_API(hipModuleUnload, hmod); - - // TODO - improve this synchronization so it is thread-safe. - // Currently we want for all inflight activity to complete, but don't prevent another - // thread from launching new kernels before we finish this operation. - ihipSynchronize(tls); - - delete hmod; // The ihipModule_t dtor will clean everything up. - hmod = nullptr; - - return ihipLogStatus(hipSuccess); -} - hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, uint32_t localWorkSizeX, uint32_t localWorkSizeY, @@ -821,6 +807,18 @@ inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t ag return HSA_STATUS_SUCCESS; } +inline hsa_status_t remove_agent_global_variables(hsa_executable_t, hsa_agent_t agent, + hsa_executable_symbol_t x, void* unused) { + hsa_symbol_kind_t t = {}; + hsa_executable_symbol_get_info(x, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &t); + + if (t == HSA_SYMBOL_KIND_VARIABLE) { + hc::am_memtracker_remove(hip_impl::address(x)); + } + + return HSA_STATUS_SUCCESS; +} + hsa_executable_symbol_t find_kernel_by_name(hsa_executable_t executable, const char* kname, hsa_agent_t* agent = nullptr) { using namespace hip_impl; @@ -898,8 +896,28 @@ namespace hip_impl { return r; } + void remove_agent_globals_from_tracker(hsa_agent_t agent, hsa_executable_t executable) { + hsa_executable_iterate_agent_symbols(executable, agent, remove_agent_global_variables, NULL); + } } // Namespace hip_impl. +hipError_t hipModuleUnload(hipModule_t hmod) { + HIP_INIT_API(hipModuleUnload, hmod); + + // TODO - improve this synchronization so it is thread-safe. + // Currently we want for all inflight activity to complete, but don't prevent another + // thread from launching new kernels before we finish this operation. + ihipSynchronize(tls); + + // deleting ihipModule_t does not remove agent globals from hc_am memtracker + hip_impl::remove_agent_globals_from_tracker(hip_impl::this_agent(), hip_impl::executable_for(hmod)); + + delete hmod; // The ihipModule_t dtor will clean everything up. + hmod = nullptr; + + return ihipLogStatus(hipSuccess); +} + hipError_t ihipModuleGetFunction(TlsData *tls, hipFunction_t* func, hipModule_t hmod, const char* name, hsa_agent_t *agent = nullptr) { using namespace hip_impl; diff --git a/tests/src/deviceLib/hipTestHalf.cpp b/tests/src/deviceLib/hipTestHalf.cpp index d48ce9b4f6..2056cffa91 100644 --- a/tests/src/deviceLib/hipTestHalf.cpp +++ b/tests/src/deviceLib/hipTestHalf.cpp @@ -30,6 +30,11 @@ THE SOFTWARE. #if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__ || __HIP_ARCH_GFX908__ +__device__ void test_convert() { + __half x; + float y = (float)x; +} + __global__ void __halfMath(bool* result, __half a) { result[0] = __heq(__hadd(a, __half{1}), __half{2}); diff --git a/tests/src/runtimeApi/module/vcpy_kernel.code b/tests/src/runtimeApi/module/vcpy_kernel.code index 40894f972a..2aa9f253b3 100644 Binary files a/tests/src/runtimeApi/module/vcpy_kernel.code and b/tests/src/runtimeApi/module/vcpy_kernel.code differ