Merge branch 'privatestaging' into grid_launch
[ROCm/clr commit: 20043d602e]
Этот коммит содержится в:
Ben Sander
@@ -8,6 +8,15 @@ set(HIP_VERSION_MAJOR "0")
|
||||
set(HIP_VERSION_MINOR "84")
|
||||
set(HIP_VERSION_PATCH "0")
|
||||
|
||||
|
||||
if(NOT DEFINED BUILD_CLANG_HIPIFY)
|
||||
set(BUILD_CLANG_HIPIFY 0)
|
||||
endif()
|
||||
|
||||
if(BUILD_CLANG_HIPIFY)
|
||||
add_subdirectory(clang-hipify)
|
||||
endif()
|
||||
|
||||
#############################
|
||||
# Configure variables
|
||||
#############################
|
||||
|
||||
@@ -0,0 +1,189 @@
|
||||
<!-- START doctoc generated TOC please keep comment here to allow auto update -->
|
||||
<!-- DON'T EDIT THIS SECTION, INSTEAD RE-RUN doctoc TO UPDATE -->
|
||||
**Installation**
|
||||
|
||||
- [Installing pre-built packages:](#installing-pre-built-packages)
|
||||
- [Prerequisites](#prerequisites)
|
||||
- [AMD (hcc)](#amd-hcc)
|
||||
- [NVIDIA (nvcc)](#nvidia-nvcc)
|
||||
- [Verify your installation](#verify-your-installation)
|
||||
- [Building HIP from source](#building-hip-from-source)
|
||||
- [HCC Options](#hcc-options)
|
||||
- [Using HIP with the AMD Native-GCN compiler.](#using-hip-with-the-amd-native-gcn-compiler)
|
||||
- [Compiling CodeXL markers for HIP Functions](#compiling-codexl-markers-for-hip-functions)
|
||||
- [Using clang-hipify](#using-clang-hipify)
|
||||
- [Building](#building)
|
||||
- [Running and using clang-hipify](#running-and-using-clang-hipify)
|
||||
|
||||
<!-- END doctoc generated TOC please keep comment here to allow auto update -->
|
||||
|
||||
# Installing pre-built packages
|
||||
|
||||
HIP can be easily installed using pre-built binary packages using the package manager for your platform.
|
||||
|
||||
## Prerequisites
|
||||
HIP code can be developed either on AMD ROCm platform using hcc compiler, or a CUDA platform with nvcc installed:
|
||||
|
||||
## AMD-hcc
|
||||
|
||||
* Install the [rocm](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/) packages. Rocm will install all of the necessary components, including the kernel driver, runtime software, HCC compiler, and HIP.
|
||||
|
||||
* Default paths and environment variables:
|
||||
|
||||
* By default HIP looks for hcc in /opt/rocm/hcc (can be overridden by setting HCC_HOME environment variable)
|
||||
* By default HIP looks for HSA in /opt/rocm/hsa (can be overridden by setting HSA_PATH environment variable)
|
||||
* By default HIP is installed into /opt/rocm/hip (can be overridden by setting HIP_PATH environment variable).
|
||||
* Optionally, consider adding /opt/rocm/bin to your path to make it easier to use the tools.
|
||||
|
||||
|
||||
## NVIDIA-nvcc
|
||||
* Configure the additional package server as described [here](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/).
|
||||
* Install the "hip_nvcc" package. This will install CUDA SDK and the HIP porting layer.
|
||||
```
|
||||
apt-get install hip_nvcc
|
||||
```
|
||||
|
||||
* Default paths and environment variables:
|
||||
* By default HIP looks for CUDA SDK in /usr/local/cuda (can be overriden by setting CUDA_PATH env variable)
|
||||
* By default HIP is installed into /opt/rocm/hip (can be overridden by setting HIP_PATH environment variable).
|
||||
* Optionally, consider adding /opt/rocm/bin to your path to make it easier to use the tools.
|
||||
|
||||
|
||||
## Verify your installation
|
||||
Run hipconfig (instructions below assume default installation path) :
|
||||
```shell
|
||||
/opt/rocm/bin/hipconfig --full
|
||||
```
|
||||
|
||||
Compile and run the [square sample](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/tree/master/samples/0_Intro/square).
|
||||
|
||||
|
||||
# Building HIP from source
|
||||
HIP source code is available and the project can be built from source on the HCC platform.
|
||||
|
||||
1. Follow the above steps to install and validate the binary packages.
|
||||
2. Download HIP source code (from the [GitHub repot](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP).)
|
||||
3. Build and install HIP (This is the simple version assuming default paths ; see below for additional options.)
|
||||
```
|
||||
cd HIP-privatestaging
|
||||
mkdir build
|
||||
cd build
|
||||
cmake ..
|
||||
make
|
||||
make install
|
||||
```
|
||||
|
||||
* Default paths:
|
||||
* By default cmake looks for hcc in /opt/rocm/hcc (can be overridden by setting ```-DHCC_HOME=/path/to/hcc``` in the cmake step).*
|
||||
* By default cmake looks for HSA in /opt/rocm/hsa (can be overridden by setting ```-DHSA_PATH=/path/to/hsa``` in the cmake step).*
|
||||
* By default cmake installs HIP to /opt/rocm/hip (can be overridden by setting ```-DCMAKE_INSTALL_PREFIX=/where/to/install/hip``` in the cmake step).*
|
||||
|
||||
Here's a richer command-line that overrides the default paths:
|
||||
|
||||
```shell
|
||||
cd HIP-privatestaging
|
||||
mkdir build
|
||||
cd build
|
||||
cmake -DHSA_PATH=/path/to/hsa -DHCC_HOME=/path/to/hcc -DCMAKE_INSTALL_PREFIX=/where/to/install/hip -DCMAKE_BUILD_TYPE=Release ..
|
||||
make
|
||||
make install
|
||||
```
|
||||
|
||||
* After installation, make sure HIP_PATH is pointed to `/where/to/install/hip`.
|
||||
|
||||
## HCC Options
|
||||
|
||||
### Using HIP with the AMD Native-GCN compiler.
|
||||
AMD recently released a direct-to-GCN-ISA target. This compiler generates GCN ISA directly from LLVM, without going through an intermediate compiler
|
||||
IR such as HSAIL or PTX.
|
||||
The native GCN target is included with upstream LLVM, and has also been integrated with HCC compiler and can be used to compiler HIP programs for AMD.
|
||||
Binary packages for the direct-to-isa package are included with the [rocm](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/) package.
|
||||
Alternatively, this sections describes how to build it from source:
|
||||
|
||||
1. Install the rocm packages as described above.
|
||||
2. Follow the instructions [here](https://github.com/RadeonOpenCompute/HCC-Native-GCN-ISA/wiki)
|
||||
* In the make step for HCC, we recommend setting -DCMAKE_INSTALL_PREFIX.
|
||||
* Set HCC_HOME environment variable before compiling HIP program to point to the native compiler:
|
||||
```shell
|
||||
export HCC_HOME=/path/to/native/hcc
|
||||
```
|
||||
|
||||
|
||||
### Compiling CodeXL markers for HIP Functions
|
||||
HIP can generate markers at function begin/end which are displayed on the CodeXL timeline view. To do this, you need to install CodeXL, tell HIP
|
||||
where the CodeXL install directory lives, and enable HIP to generate the markers:
|
||||
|
||||
1. Install CodeXL
|
||||
See [CodeXL Download](http://developer.amd.com/tools-and-sdks/opencl-zone/codexl/?webSyncID=9d9c2cb9-3d73-5e65-268a-c7b06428e5e0&sessionGUID=29beacd0-d654-ddc6-a3e2-b9e6c0b0cc77) for the installation file.
|
||||
Also this [blog](http://gpuopen.com/getting-up-to-speed-with-the-codexl-gpu-profiler-and-radeon-open-compute/) provides more information and tips for using CodeXL. In addition to installing the CodeXL profiling
|
||||
and visualization tools, CodeXL also comes with an SDK that allow applications to add markers to the timeline viewer. We'll be linking HIP against this library.
|
||||
|
||||
2. Set CODEXL_PATH
|
||||
```shell
|
||||
# set to your code-xl installation location:
|
||||
export CODEXL_PATH=/opt/AMD/CodeXL
|
||||
```
|
||||
|
||||
3. Enable in source code.
|
||||
In src/hip_hcc.cpp, enable the define
|
||||
```c
|
||||
#define COMPILE_TRACE_MARKER 1
|
||||
```
|
||||
|
||||
|
||||
Then recompile the target application, run with profiler enabled to generate ATP file or trace log.
|
||||
```shell
|
||||
# Use profiler to generate timeline view:
|
||||
$CODEXL_PATH/CodeXLGpuProfiler -A -o ./myHipApp
|
||||
...
|
||||
Session output path: /home/me/HIP-privatestaging/tests/b1/mytrace.atp
|
||||
```
|
||||
|
||||
You can also print the HIP function strings to stderr using HIP_TRACE_API environment variable. This can be useful for tracing application flow. Also can be combined with the more detailed debug information provided
|
||||
by the HIP_DB switch. For example:
|
||||
```shell
|
||||
# Trace to stderr showing begin/end of each function (with arguments) + intermediate debug trace during the execution of each function.
|
||||
HIP_TRACE_API=1 HIP_DB=0x2 ./myHipApp
|
||||
```
|
||||
|
||||
Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors.
|
||||
|
||||
|
||||
### Using clang-hipify
|
||||
|
||||
Clang-hipify is a clang-based tool which can automate the translation of CUDA source code into portable HIP C++.
|
||||
The clang-hipify tool can automatically add extra HIP arguments (notably the "hipLaunchParm" required at the
|
||||
beginning of every HIP kernel call). Clang-hipify has some additional dependencies explained below and
|
||||
can be built as a separate make step.
|
||||
|
||||
|
||||
#### Building
|
||||
|
||||
1. Download and unpack clang+llvm 3.8 binary package preqrequisite:
|
||||
```
|
||||
wget http://llvm.org/releases/3.8.0/clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-14.04.tar.xz
|
||||
tar xvfJ clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-14.04.tar.xz
|
||||
```
|
||||
|
||||
2. Enable build of clang-hipify and specify path to LLVM:
|
||||
Note LLVM_DIR must be a full absolute path (not relative) to the location extracted above. Here's an example assuming we
|
||||
extract the clang 3.8 package into ~/HIP-privatestaging/clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-14.04/.
|
||||
```
|
||||
cd HIP-privatestaging
|
||||
mkdir build.clang-hipify
|
||||
cd build.clang-hipify
|
||||
cmake -DBUILD_CLANG_HIPIFY=1 -DLLVM_DIR=~/HIP-privatestaging/clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-14.04/ -DCMAKE_BUILD_TYPE=Release ..
|
||||
make
|
||||
make install
|
||||
```
|
||||
|
||||
#### Running and using clang-hipify
|
||||
clang-hipify performs an initial compile of the CUDA source code into a "symbol tree", and thus needs access to the appropriate header files:
|
||||
1. Download "deb(network)" variant of target installer from https://developer.nvidia.com/cuda-downloads. The commands below show how to download and install a recent version from the http://developer.download.nvidia.com/compute/cuda/repos/ubuntu1404/x86_64/cuda-repo-ubuntu1404_7.5-18_amd64.deb.
|
||||
|
||||
```
|
||||
wget http://developer.download.nvidia.com/compute/cuda/repos/ubuntu1404/x86_64/cuda-repo-ubuntu1404_7.5-18_amd64.deb
|
||||
sudo dpkg -i cuda-repo-ubuntu1404_7.5-18_amd64.deb
|
||||
sudo apt-get update && sudo apt-get install cuda-minimal-build-7-5 cuda-curand-dev-7-5
|
||||
```
|
||||
|
||||
@@ -10,24 +10,9 @@ Key features include:
|
||||
* Developers can specialize for the platform (CUDA or hcc) to tune for performance or handle tricky cases
|
||||
|
||||
New projects can be developed directly in the portable HIP C++ language and can run on either NVIDIA or AMD platforms. Additionally, HIP provides porting tools which make it easy to port existing CUDA codes to the HIP layer, with no loss of performance as compared to the original CUDA application. HIP is not intended to be a drop-in replacement for CUDA, and developers should expect to do some manual coding and performance tuning work to complete the port.
|
||||
## Installation
|
||||
```shell
|
||||
cd HIP-privatestaging
|
||||
mkdir build
|
||||
cd build
|
||||
cmake ..
|
||||
make
|
||||
make install
|
||||
```
|
||||
*By default cmake looks for hcc in /opt/rocm/hcc (can be overridden by setting ```-DHCC_HOME=/path/to/hcc``` in the cmake step).*
|
||||
|
||||
*By default cmake looks for HSA in /opt/rocm/hsa (can be overridden by setting ```-DHSA_PATH=/path/to/hsa``` in the cmake step).*
|
||||
|
||||
*By default cmake installs HIP to /opt/rocm/hip (can be overridden by setting ```-DCMAKE_INSTALL_PREFIX=/where/to/install/hip``` in the cmake step).*
|
||||
|
||||
*Make sure HIP_PATH is pointed to `/where/to/install/hip` and PATH includes `$HIP_PATH/bin`. This requirement is optional, but required to run any HIP test infrastructure.*
|
||||
|
||||
## More Info:
|
||||
- [Installation](INSTALL.md) and [clang-hipify](INSTALL.md#use-clang-hipify.md)
|
||||
- [HIP FAQ](docs/markdown/hip_faq.md)
|
||||
- [HIP Kernel Language](docs/markdown/hip_kernel_language.md)
|
||||
- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP)
|
||||
@@ -39,92 +24,8 @@ make install
|
||||
|
||||
## How do I get set up?
|
||||
|
||||
### Prerequisites - Choose Your Platform
|
||||
HIP code can be developed either on AMD ROCm platform using hcc compiler, or a CUDA platform with nvcc installed:
|
||||
See the [Installation](INSTALL.md) notes.
|
||||
|
||||
#### AMD (hcc):
|
||||
|
||||
* Install [hcc](https://bitbucket.org/multicoreware/hcc/wiki/Home) including supporting HSA kernel and runtime driver stack
|
||||
* By default HIP looks for hcc in /opt/rocm/hcc (can be overridden by setting HCC_HOME environment variable)
|
||||
* By default HIP looks for HSA in /opt/rocm/hsa (can be overridden by setting HSA_PATH environment variable)
|
||||
* Ensure that ROCR runtime is installed and added to LD_LIBRARY_PATH
|
||||
* Install HIP (from this GitHub repot). By default HIP is installed into /opt/rocm/hip (can be overridden by setting HIP_PATH environment variable).
|
||||
|
||||
* Optionally, consider adding /opt/rocm/bin to your path to make it easier to use the tools.
|
||||
|
||||
#### NVIDIA (nvcc)
|
||||
* Install CUDA SDK from manufacturer website
|
||||
* By default HIP looks for CUDA SDK in /usr/local/cuda (can be overriden by setting CUDA_PATH env variable)
|
||||
* Install HIP (from this GitHub repot). By default HIP is installed into /opt/rocm/hip (can be overridden by setting HIP_PATH environment variable).
|
||||
|
||||
* Optionally, consider adding /opt/rocm/bin to your path to make it easier to use the tools.
|
||||
|
||||
#### Verify your installation
|
||||
Run hipconfig (instructions below assume default installation path) :
|
||||
```shell
|
||||
/opt/rocm/bin/hipconfig --full
|
||||
```
|
||||
|
||||
Compile and run the [square sample](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/tree/master/samples/0_Intro/square).
|
||||
|
||||
|
||||
|
||||
### HCC Options
|
||||
|
||||
#### Compiling CodeXL markers for HIP Functions
|
||||
HIP can generate markers at function begin/end which are displayed on the CodeXL timeline view. To do this, you need to install CodeXL, tell HIP
|
||||
where the CodeXL install directory lives, and enable HIP to generate the markers:
|
||||
|
||||
1. Install CodeXL
|
||||
See [CodeXL Download](http://developer.amd.com/tools-and-sdks/opencl-zone/codexl/?webSyncID=9d9c2cb9-3d73-5e65-268a-c7b06428e5e0&sessionGUID=29beacd0-d654-ddc6-a3e2-b9e6c0b0cc77) for the installation file.
|
||||
Also this [blog](http://gpuopen.com/getting-up-to-speed-with-the-codexl-gpu-profiler-and-radeon-open-compute/) provides more information and tips for using CodeXL. In addition to installing the CodeXL profiling
|
||||
and visualization tools, CodeXL also comes with an SDK that allow applications to add markers to the timeline viewer. We'll be linking HIP against this library.
|
||||
|
||||
2. Set CODEXL_PATH
|
||||
```shell
|
||||
# set to your code-xl installation location:
|
||||
export CODEXL_PATH=/opt/AMD/CodeXL
|
||||
```
|
||||
|
||||
3. Enable in source code.
|
||||
In src/hip_hcc.cpp, enable the define
|
||||
```c
|
||||
#define COMPILE_TRACE_MARKER 1
|
||||
```
|
||||
|
||||
|
||||
Then recompile the target application, run with profiler enabled to generate ATP file or trace log.
|
||||
```shell
|
||||
# Use profiler to generate timeline view:
|
||||
$CODEXL_PATH/CodeXLGpuProfiler -A -o ./myHipApp
|
||||
...
|
||||
Session output path: /home/me/HIP-privatestaging/tests/b1/mytrace.atp
|
||||
```
|
||||
|
||||
You can also print the HIP function strings to stderr using HIP_TRACE_API environment variable. This can be useful for tracing application flow. Also can be combined with the more detailed debug information provided
|
||||
by the HIP_DB switch. For example:
|
||||
```shell
|
||||
# Trace to stderr showing begin/end of each function (with arguments) + intermediate debug trace during the execution of each function.
|
||||
HIP_TRACE_API=1 HIP_DB=0x2 ./myHipApp
|
||||
```
|
||||
|
||||
Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors.
|
||||
|
||||
|
||||
#### Using HIP with the AMD Native-GCN compiler.
|
||||
AMD recently released a direct-to-GCN-ISA target. This compiler generates GCN ISA directly from LLVM, without going through an intermediate compiler
|
||||
IR such as HSAIL or PTX.
|
||||
The native GCN target is included with upstream LLVM, and has also been integrated with HCC compiler and can be used to compiler HIP programs for AMD.
|
||||
Here's how to use it with HIP:
|
||||
|
||||
- Follow the instructions here to compile the HCC and native LLVM compiler:
|
||||
> https://github.com/RadeonOpenCompute/HCC-Native-GCN-ISA/wiki
|
||||
> (In the make step for HCC, we recommend setting -DCMAKE_INSTALL_PREFIX=/opt/hcc-native)
|
||||
|
||||
Set HCC_HOME environment variable before compiling HIP program to point to the native compiler:
|
||||
```shell
|
||||
export HCC_HOME=/opt/hcc-native
|
||||
```
|
||||
|
||||
## Examples and Getting Started:
|
||||
|
||||
|
||||
@@ -33,8 +33,6 @@ $HIP_PATH=dirname (dirname $0) unless defined $HIP_PATH; # use parent dir
|
||||
|
||||
#print "HIP_PATH=$HIP_PATH\n";
|
||||
|
||||
$CUDA_PATH=$ENV{'CUDA_PATH'};
|
||||
$CUDA_PATH='/usr/local/cuda' unless defined $CUDA_PATH;
|
||||
|
||||
$CODEXL_PATH = $ENV{'CODEXL_PATH'};
|
||||
$CODEXL_PATH = "/opt/AMD/CodeXL" unless defined $CODEXL_PATH;
|
||||
@@ -42,17 +40,14 @@ $marker_path = "$CODEXL_PATH/SDK/AMDTActivityLogger";
|
||||
|
||||
#---
|
||||
#HIP_PLATFORM controls whether to use NVCC or HCC for compilation:
|
||||
$HIP_PLATFORM=$ENV{'HIP_PLATFORM'};
|
||||
if (not defined $HIP_PLATFORM and (-e "$CUDA_PATH/bin/nvcc")) {
|
||||
$HIP_PLATFORM="nvcc";
|
||||
}
|
||||
$HIP_PLATFORM= `$HIP_PATH/bin/hipconfig --platform`;
|
||||
|
||||
$HIP_PLATFORM="hcc" unless defined $HIP_PLATFORM;
|
||||
|
||||
|
||||
if ($verbose & 0x2) {
|
||||
print ("HIP_PATH=$HIP_PATH\n");
|
||||
print ("HIP_PLATFORM=$HIP_PLATFORM\n");
|
||||
print ("CUDA_PATH=$CUDA_PATH\n");
|
||||
}
|
||||
|
||||
# set if user explicitly requests -stdlib=libc++. (else we default to libstdc++ for better interop with g++):
|
||||
@@ -109,6 +104,11 @@ if ($HIP_PLATFORM eq "hcc") {
|
||||
}
|
||||
|
||||
} elsif ($HIP_PLATFORM eq "nvcc") {
|
||||
if ($verbose & 0x2) {
|
||||
print ("CUDA_PATH=$CUDA_PATH\n");
|
||||
}
|
||||
$CUDA_PATH=$ENV{'CUDA_PATH'};
|
||||
$CUDA_PATH='/usr/local/cuda' unless defined $CUDA_PATH;
|
||||
|
||||
$HIPCC="$CUDA_PATH/bin/nvcc";
|
||||
$HIPCXXFLAGS .= " -I$CUDA_PATH/include";
|
||||
|
||||
@@ -0,0 +1,74 @@
|
||||
cmake_minimum_required(VERSION 2.8.8)
|
||||
|
||||
project(hipify-clang)
|
||||
|
||||
find_package(LLVM 3.8 REQUIRED PATHS ${LLVM_DIR} NO_DEFAULT_PATH)
|
||||
|
||||
list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR})
|
||||
include(AddLLVM)
|
||||
|
||||
message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}")
|
||||
|
||||
include_directories(${LLVM_INCLUDE_DIRS})
|
||||
link_directories(${LLVM_LIBRARY_DIRS})
|
||||
add_definitions(${LLVM_DEFINITIONS})
|
||||
add_llvm_executable(hipify-clang src/Cuda2Hip.cpp )
|
||||
find_program(LIT_COMMAND lit)
|
||||
|
||||
# Link against LLVM and CLANG tools libraries
|
||||
target_link_libraries(hipify-clang
|
||||
clangASTMatchers
|
||||
clangFrontend
|
||||
clangTooling
|
||||
clangParse
|
||||
clangSerialization
|
||||
clangSema
|
||||
clangEdit
|
||||
clangLex
|
||||
clangAnalysis
|
||||
clangDriver
|
||||
clangAST
|
||||
clangToolingCore
|
||||
clangRewrite
|
||||
clangBasic
|
||||
LLVMProfileData
|
||||
LLVMSupport
|
||||
LLVMMCParser
|
||||
LLVMMC
|
||||
LLVMBitReader
|
||||
LLVMOption
|
||||
LLVMCore)
|
||||
|
||||
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}")
|
||||
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}")
|
||||
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden")
|
||||
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DIRS}/clang/${LLVM_VERSION_MAJOR}.${LLVM_VERSION_MINOR}.${LLVM_VERSION_PATCH}\\\"")
|
||||
|
||||
install(TARGETS hipify-clang
|
||||
DESTINATION bin)
|
||||
|
||||
# tests
|
||||
set(Python_ADDITIONAL_VERSIONS 2.7)
|
||||
include(FindPythonInterp)
|
||||
if( NOT PYTHONINTERP_FOUND )
|
||||
message(FATAL_ERROR
|
||||
"Unable to find Python interpreter, required for builds and testing\n\n"
|
||||
"Please install Python or specify the PYTHON_EXECUTABLE CMake variable.")
|
||||
endif()
|
||||
|
||||
set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR} )
|
||||
|
||||
configure_file(
|
||||
${CMAKE_SOURCE_DIR}/tests/clang-hipify/lit.site.cfg.in
|
||||
${CMAKE_CURRENT_BINARY_DIR}/tests/clang-hipify/lit.site.cfg
|
||||
@ONLY)
|
||||
|
||||
add_lit_testsuite(test-hipify "Running HIPify regression tests"
|
||||
${CMAKE_SOURCE_DIR}/tests/clang-hipify
|
||||
PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/clang-hipify/lit.site.cfg
|
||||
DEPENDS hipify-clang lit
|
||||
)
|
||||
|
||||
add_custom_target(test-clang-hipify)
|
||||
add_dependencies(test-clang-hipify test-hipify)
|
||||
set_target_properties(test-clang-hipify PROPERTIES FOLDER "Tests")
|
||||
@@ -0,0 +1,937 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
/**
|
||||
* @file Cuda2Hip.cpp
|
||||
*
|
||||
* This file is compiled and linked into clang based hipify tool.
|
||||
*/
|
||||
#include "clang/ASTMatchers/ASTMatchFinder.h"
|
||||
#include "clang/ASTMatchers/ASTMatchers.h"
|
||||
#include "clang/Basic/SourceManager.h"
|
||||
#include "clang/Frontend/CompilerInstance.h"
|
||||
#include "clang/Frontend/FrontendActions.h"
|
||||
#include "clang/Frontend/TextDiagnosticPrinter.h"
|
||||
#include "clang/Lex/Lexer.h"
|
||||
#include "clang/Lex/MacroArgs.h"
|
||||
#include "clang/Lex/MacroInfo.h"
|
||||
#include "clang/Lex/PPCallbacks.h"
|
||||
#include "clang/Lex/Preprocessor.h"
|
||||
#include "clang/Rewrite/Core/Rewriter.h"
|
||||
#include "clang/Tooling/CommonOptionsParser.h"
|
||||
#include "clang/Tooling/Refactoring.h"
|
||||
#include "clang/Tooling/Tooling.h"
|
||||
#include "llvm/Support/CommandLine.h"
|
||||
#include "llvm/Support/Debug.h"
|
||||
#include "llvm/Support/MemoryBuffer.h"
|
||||
#include "llvm/Support/Signals.h"
|
||||
|
||||
#include <cstdio>
|
||||
#include <fstream>
|
||||
|
||||
using namespace clang;
|
||||
using namespace clang::ast_matchers;
|
||||
using namespace clang::tooling;
|
||||
using namespace llvm;
|
||||
|
||||
#define DEBUG_TYPE "cuda2hip"
|
||||
|
||||
enum ConvTypes {
|
||||
CONV_DEV = 0,
|
||||
CONV_MEM,
|
||||
CONV_KERN,
|
||||
CONV_COORD_FUNC,
|
||||
CONV_MATH_FUNC,
|
||||
CONV_SPECIAL_FUNC,
|
||||
CONV_STREAM,
|
||||
CONV_EVENT,
|
||||
CONV_ERR,
|
||||
CONV_DEF,
|
||||
CONV_TEX,
|
||||
CONV_OTHER,
|
||||
CONV_INCLUDE,
|
||||
CONV_LITERAL,
|
||||
CONV_LAST
|
||||
};
|
||||
|
||||
const char *counterNames[ConvTypes::CONV_LAST] = {
|
||||
"dev", "mem", "kern", "coord_func", "math_func",
|
||||
"special_func", "stream", "event", "err", "def",
|
||||
"tex", "other", "include", "literal"};
|
||||
|
||||
namespace {
|
||||
|
||||
struct cuda2hipMap {
|
||||
cuda2hipMap() {
|
||||
// defines
|
||||
cuda2hipRename["__CUDACC__"] = {"__HIPCC__", CONV_DEF};
|
||||
|
||||
// includes
|
||||
cuda2hipRename["cuda_runtime.h"] = {"hip_runtime.h", CONV_INCLUDE};
|
||||
cuda2hipRename["cuda_runtime_api.h"] = {"hip_runtime_api.h", CONV_INCLUDE};
|
||||
|
||||
// Error codes and return types:
|
||||
cuda2hipRename["cudaError_t"] = {"hipError_t", CONV_ERR};
|
||||
cuda2hipRename["cudaError"] = {"hipError", CONV_ERR};
|
||||
cuda2hipRename["cudaSuccess"] = {"hipSuccess", CONV_ERR};
|
||||
|
||||
cuda2hipRename["cudaErrorUnknown"] = {"hipErrorUnknown", CONV_ERR};
|
||||
cuda2hipRename["cudaErrorMemoryAllocation"] = {"hipErrorMemoryAllocation",
|
||||
CONV_ERR};
|
||||
cuda2hipRename["cudaErrorMemoryFree"] = {"hipErrorMemoryFree", CONV_ERR};
|
||||
cuda2hipRename["cudaErrorUnknownSymbol"] = {"hipErrorUnknownSymbol",
|
||||
CONV_ERR};
|
||||
cuda2hipRename["cudaErrorOutOfResources"] = {"hipErrorOutOfResources",
|
||||
CONV_ERR};
|
||||
cuda2hipRename["cudaErrorInvalidValue"] = {"hipErrorInvalidValue",
|
||||
CONV_ERR};
|
||||
cuda2hipRename["cudaErrorInvalidResourceHandle"] = {
|
||||
"hipErrorInvalidResourceHandle", CONV_ERR};
|
||||
cuda2hipRename["cudaErrorInvalidDevice"] = {"hipErrorInvalidDevice",
|
||||
CONV_ERR};
|
||||
cuda2hipRename["cudaErrorNoDevice"] = {"hipErrorNoDevice", CONV_ERR};
|
||||
cuda2hipRename["cudaErrorNotReady"] = {"hipErrorNotReady", CONV_ERR};
|
||||
cuda2hipRename["cudaErrorUnknown"] = {"hipErrorUnknown", CONV_ERR};
|
||||
|
||||
// error APIs:
|
||||
cuda2hipRename["cudaGetLastError"] = {"hipGetLastError", CONV_ERR};
|
||||
cuda2hipRename["cudaPeekAtLastError"] = {"hipPeekAtLastError", CONV_ERR};
|
||||
cuda2hipRename["cudaGetErrorName"] = {"hipGetErrorName", CONV_ERR};
|
||||
cuda2hipRename["cudaGetErrorString"] = {"hipGetErrorString", CONV_ERR};
|
||||
|
||||
// Memcpy
|
||||
cuda2hipRename["cudaMemcpy"] = {"hipMemcpy", CONV_MEM};
|
||||
cuda2hipRename["cudaMemcpyHostToHost"] = {"hipMemcpyHostToHost", CONV_MEM};
|
||||
cuda2hipRename["cudaMemcpyHostToDevice"] = {"hipMemcpyHostToDevice",
|
||||
CONV_MEM};
|
||||
cuda2hipRename["cudaMemcpyDeviceToHost"] = {"hipMemcpyDeviceToHost",
|
||||
CONV_MEM};
|
||||
cuda2hipRename["cudaMemcpyDeviceToDevice"] = {"hipMemcpyDeviceToDevice",
|
||||
CONV_MEM};
|
||||
cuda2hipRename["cudaMemcpyDefault"] = {"hipMemcpyDefault", CONV_MEM};
|
||||
cuda2hipRename["cudaMemcpyToSymbol"] = {"hipMemcpyToSymbol", CONV_MEM};
|
||||
cuda2hipRename["cudaMemset"] = {"hipMemset", CONV_MEM};
|
||||
cuda2hipRename["cudaMemsetAsync"] = {"hipMemsetAsync", CONV_MEM};
|
||||
cuda2hipRename["cudaMemcpyAsync"] = {"hipMemcpyAsync", CONV_MEM};
|
||||
cuda2hipRename["cudaMemGetInfo"] = {"hipMemGetInfo", CONV_MEM};
|
||||
cuda2hipRename["cudaMemcpyKind"] = {"hipMemcpyKind", CONV_MEM};
|
||||
|
||||
// Memory management :
|
||||
cuda2hipRename["cudaMalloc"] = {"hipMalloc", CONV_MEM};
|
||||
cuda2hipRename["cudaMallocHost"] = {"hipHostAlloc", CONV_MEM};
|
||||
cuda2hipRename["cudaFree"] = {"hipFree", CONV_MEM};
|
||||
cuda2hipRename["cudaFreeHost"] = {"hipHostFree", CONV_MEM};
|
||||
|
||||
// Coordinate Indexing and Dimensions:
|
||||
cuda2hipRename["threadIdx.x"] = {"hipThreadIdx_x", CONV_COORD_FUNC};
|
||||
cuda2hipRename["threadIdx.y"] = {"hipThreadIdx_y", CONV_COORD_FUNC};
|
||||
cuda2hipRename["threadIdx.z"] = {"hipThreadIdx_z", CONV_COORD_FUNC};
|
||||
|
||||
cuda2hipRename["blockIdx.x"] = {"hipBlockIdx_x", CONV_COORD_FUNC};
|
||||
cuda2hipRename["blockIdx.y"] = {"hipBlockIdx_y", CONV_COORD_FUNC};
|
||||
cuda2hipRename["blockIdx.z"] = {"hipBlockIdx_z", CONV_COORD_FUNC};
|
||||
|
||||
cuda2hipRename["blockDim.x"] = {"hipBlockDim_x", CONV_COORD_FUNC};
|
||||
cuda2hipRename["blockDim.y"] = {"hipBlockDim_y", CONV_COORD_FUNC};
|
||||
cuda2hipRename["blockDim.z"] = {"hipBlockDim_z", CONV_COORD_FUNC};
|
||||
|
||||
cuda2hipRename["gridDim.x"] = {"hipGridDim_x", CONV_COORD_FUNC};
|
||||
cuda2hipRename["gridDim.y"] = {"hipGridDim_y", CONV_COORD_FUNC};
|
||||
cuda2hipRename["gridDim.z"] = {"hipGridDim_z", CONV_COORD_FUNC};
|
||||
|
||||
cuda2hipRename["blockIdx.x"] = {"hipBlockIdx_x", CONV_COORD_FUNC};
|
||||
cuda2hipRename["blockIdx.y"] = {"hipBlockIdx_y", CONV_COORD_FUNC};
|
||||
cuda2hipRename["blockIdx.z"] = {"hipBlockIdx_z", CONV_COORD_FUNC};
|
||||
|
||||
cuda2hipRename["blockDim.x"] = {"hipBlockDim_x", CONV_COORD_FUNC};
|
||||
cuda2hipRename["blockDim.y"] = {"hipBlockDim_y", CONV_COORD_FUNC};
|
||||
cuda2hipRename["blockDim.z"] = {"hipBlockDim_z", CONV_COORD_FUNC};
|
||||
|
||||
cuda2hipRename["gridDim.x"] = {"hipGridDim_x", CONV_COORD_FUNC};
|
||||
cuda2hipRename["gridDim.y"] = {"hipGridDim_y", CONV_COORD_FUNC};
|
||||
cuda2hipRename["gridDim.z"] = {"hipGridDim_z", CONV_COORD_FUNC};
|
||||
|
||||
cuda2hipRename["warpSize"] = {"hipWarpSize", CONV_SPECIAL_FUNC};
|
||||
|
||||
// Events
|
||||
cuda2hipRename["cudaEvent_t"] = {"hipEvent_t", CONV_EVENT};
|
||||
cuda2hipRename["cudaEventCreate"] = {"hipEventCreate", CONV_EVENT};
|
||||
cuda2hipRename["cudaEventCreateWithFlags"] = {"hipEventCreateWithFlags",
|
||||
CONV_EVENT};
|
||||
cuda2hipRename["cudaEventDestroy"] = {"hipEventDestroy", CONV_EVENT};
|
||||
cuda2hipRename["cudaEventRecord"] = {"hipEventRecord", CONV_EVENT};
|
||||
cuda2hipRename["cudaEventElapsedTime"] = {"hipEventElapsedTime",
|
||||
CONV_EVENT};
|
||||
cuda2hipRename["cudaEventSynchronize"] = {"hipEventSynchronize",
|
||||
CONV_EVENT};
|
||||
|
||||
// Streams
|
||||
cuda2hipRename["cudaStream_t"] = {"hipStream_t", CONV_STREAM};
|
||||
cuda2hipRename["cudaStreamCreate"] = {"hipStreamCreate", CONV_STREAM};
|
||||
cuda2hipRename["cudaStreamCreateWithFlags"] = {"hipStreamCreateWithFlags",
|
||||
CONV_STREAM};
|
||||
cuda2hipRename["cudaStreamDestroy"] = {"hipStreamDestroy", CONV_STREAM};
|
||||
cuda2hipRename["cudaStreamWaitEvent"] = {"hipStreamWaitEven", CONV_STREAM};
|
||||
cuda2hipRename["cudaStreamSynchronize"] = {"hipStreamSynchronize",
|
||||
CONV_STREAM};
|
||||
cuda2hipRename["cudaStreamDefault"] = {"hipStreamDefault", CONV_STREAM};
|
||||
cuda2hipRename["cudaStreamNonBlocking"] = {"hipStreamNonBlocking",
|
||||
CONV_STREAM};
|
||||
|
||||
// Other synchronization
|
||||
cuda2hipRename["cudaDeviceSynchronize"] = {"hipDeviceSynchronize",
|
||||
CONV_DEV};
|
||||
cuda2hipRename["cudaThreadSynchronize"] = {
|
||||
"hipDeviceSynchronize",
|
||||
CONV_DEV}; // translate deprecated cudaThreadSynchronize
|
||||
cuda2hipRename["cudaDeviceReset"] = {"hipDeviceReset", CONV_DEV};
|
||||
cuda2hipRename["cudaThreadExit"] = {
|
||||
"hipDeviceReset", CONV_DEV}; // translate deprecated cudaThreadExit
|
||||
cuda2hipRename["cudaSetDevice"] = {"hipSetDevice", CONV_DEV};
|
||||
cuda2hipRename["cudaGetDevice"] = {"hipGetDevice", CONV_DEV};
|
||||
|
||||
// Attribute
|
||||
cuda2hipRename["bcudaDeviceAttr"] = {"hipDeviceAttribute_t", CONV_DEV};
|
||||
cuda2hipRename["bcudaDeviceGetAttribute"] = {"hipDeviceGetAttribute",
|
||||
CONV_DEV};
|
||||
|
||||
// Device
|
||||
cuda2hipRename["cudaDeviceProp"] = {"hipDeviceProp_t", CONV_DEV};
|
||||
cuda2hipRename["cudaGetDeviceProperties"] = {"hipGetDeviceProperties",
|
||||
CONV_DEV};
|
||||
|
||||
// Cache config
|
||||
cuda2hipRename["cudaDeviceSetCacheConfig"] = {"hipDeviceSetCacheConfig",
|
||||
CONV_DEV};
|
||||
cuda2hipRename["cudaThreadSetCacheConfig"] = {
|
||||
"hipDeviceSetCacheConfig", CONV_DEV}; // translate deprecated
|
||||
cuda2hipRename["cudaDeviceGetCacheConfig"] = {"hipDeviceGetCacheConfig",
|
||||
CONV_DEV};
|
||||
cuda2hipRename["cudaThreadGetCacheConfig"] = {
|
||||
"hipDeviceGetCacheConfig", CONV_DEV}; // translate deprecated
|
||||
cuda2hipRename["cudaFuncCache"] = {"hipFuncCache", CONV_DEV};
|
||||
cuda2hipRename["cudaFuncCachePreferNone"] = {"hipFuncCachePreferNone",
|
||||
CONV_DEV};
|
||||
cuda2hipRename["cudaFuncCachePreferShared"] = {"hipFuncCachePreferShared",
|
||||
CONV_DEV};
|
||||
cuda2hipRename["cudaFuncCachePreferL1"] = {"hipFuncCachePreferL1",
|
||||
CONV_DEV};
|
||||
cuda2hipRename["cudaFuncCachePreferEqual"] = {"hipFuncCachePreferEqual",
|
||||
CONV_DEV};
|
||||
// function
|
||||
cuda2hipRename["cudaFuncSetCacheConfig"] = {"hipFuncSetCacheConfig",
|
||||
CONV_DEV};
|
||||
|
||||
cuda2hipRename["cudaDriverGetVersion"] = {"hipDriverGetVersion", CONV_DEV};
|
||||
// cuda2hipRename["cudaRuntimeGetVersion"] = {"hipRuntimeGetVersion",
|
||||
// CONV_DEV};
|
||||
|
||||
// Peer2Peer
|
||||
cuda2hipRename["cudaDeviceCanAccessPeer"] = {"hipDeviceCanAccessPeer",
|
||||
CONV_DEV};
|
||||
cuda2hipRename["cudaDeviceDisablePeerAccess"] = {
|
||||
"hipDeviceDisablePeerAccess", CONV_DEV};
|
||||
cuda2hipRename["cudaDeviceEnablePeerAccess"] = {"hipDeviceEnablePeerAccess",
|
||||
CONV_DEV};
|
||||
cuda2hipRename["cudaMemcpyPeerAsync"] = {"hipMemcpyPeerAsync", CONV_MEM};
|
||||
cuda2hipRename["cudaMemcpyPeer"] = {"hipMemcpyPeer", CONV_MEM};
|
||||
|
||||
// Shared mem:
|
||||
cuda2hipRename["cudaDeviceSetSharedMemConfig"] = {
|
||||
"hipDeviceSetSharedMemConfig", CONV_DEV};
|
||||
cuda2hipRename["cudaThreadSetSharedMemConfig"] = {
|
||||
"hipDeviceSetSharedMemConfig", CONV_DEV}; // translate deprecated
|
||||
cuda2hipRename["cudaDeviceGetSharedMemConfig"] = {
|
||||
"hipDeviceGetSharedMemConfig", CONV_DEV};
|
||||
cuda2hipRename["cudaThreadGetSharedMemConfig"] = {
|
||||
"hipDeviceGetSharedMemConfig", CONV_DEV}; // translate deprecated
|
||||
cuda2hipRename["cudaSharedMemConfig"] = {"hipSharedMemConfig", CONV_DEV};
|
||||
cuda2hipRename["cudaSharedMemBankSizeDefault"] = {
|
||||
"hipSharedMemBankSizeDefault", CONV_DEV};
|
||||
cuda2hipRename["cudaSharedMemBankSizeFourByte"] = {
|
||||
"hipSharedMemBankSizeFourByte", CONV_DEV};
|
||||
cuda2hipRename["cudaSharedMemBankSizeEightByte"] = {
|
||||
"hipSharedMemBankSizeEightByte", CONV_DEV};
|
||||
|
||||
cuda2hipRename["cudaGetDeviceCount"] = {"hipGetDeviceCount", CONV_DEV};
|
||||
|
||||
// Profiler
|
||||
// cuda2hipRename["cudaProfilerInitialize"] = "hipProfilerInitialize"; //
|
||||
// see if these are called anywhere.
|
||||
cuda2hipRename["cudaProfilerStart"] = {"hipProfilerStart", CONV_OTHER};
|
||||
cuda2hipRename["cudaProfilerStop"] = {"hipProfilerStop", CONV_OTHER};
|
||||
|
||||
cuda2hipRename["cudaChannelFormatDesc"] = {"hipChannelFormatDesc",
|
||||
CONV_TEX};
|
||||
cuda2hipRename["cudaFilterModePoint"] = {"hipFilterModePoint", CONV_TEX};
|
||||
cuda2hipRename["cudaReadModeElementType"] = {"hipReadModeElementType",
|
||||
CONV_TEX};
|
||||
|
||||
cuda2hipRename["cudaCreateChannelDesc"] = {"hipCreateChannelDesc",
|
||||
CONV_TEX};
|
||||
cuda2hipRename["cudaBindTexture"] = {"hipBindTexture", CONV_TEX};
|
||||
cuda2hipRename["cudaUnbindTexture"] = {"hipUnbindTexture", CONV_TEX};
|
||||
}
|
||||
|
||||
struct HipNames {
|
||||
StringRef hipName;
|
||||
ConvTypes countType;
|
||||
};
|
||||
|
||||
SmallDenseMap<StringRef, HipNames> cuda2hipRename;
|
||||
};
|
||||
|
||||
StringRef unquoteStr(StringRef s) {
|
||||
if (s.size() > 1 && s.front() == '"' && s.back() == '"')
|
||||
return s.substr(1, s.size() - 2);
|
||||
return s;
|
||||
}
|
||||
|
||||
static void processString(StringRef s, const cuda2hipMap &map,
|
||||
Replacements *Replace, SourceManager &SM,
|
||||
SourceLocation start,
|
||||
int64_t countReps[ConvTypes::CONV_LAST]) {
|
||||
size_t begin = 0;
|
||||
while ((begin = s.find("cuda", begin)) != StringRef::npos) {
|
||||
const size_t end = s.find_first_of(" ", begin + 4);
|
||||
StringRef name = s.slice(begin, end);
|
||||
const auto found = map.cuda2hipRename.find(name);
|
||||
if (found != map.cuda2hipRename.end()) {
|
||||
countReps[CONV_LITERAL]++;
|
||||
StringRef repName = found->second.hipName;
|
||||
SourceLocation sl = start.getLocWithOffset(begin + 1);
|
||||
Replacement Rep(SM, sl, name.size(), repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
if (end == StringRef::npos)
|
||||
break;
|
||||
begin = end + 1;
|
||||
}
|
||||
}
|
||||
|
||||
struct HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks {
|
||||
HipifyPPCallbacks(Replacements *R)
|
||||
: SeenEnd(false), _sm(nullptr), _pp(nullptr), Replace(R) {}
|
||||
|
||||
virtual bool handleBeginSource(CompilerInstance &CI,
|
||||
StringRef Filename) override {
|
||||
Preprocessor &PP = CI.getPreprocessor();
|
||||
SourceManager &SM = CI.getSourceManager();
|
||||
setSourceManager(&SM);
|
||||
PP.addPPCallbacks(std::unique_ptr<HipifyPPCallbacks>(this));
|
||||
PP.Retain();
|
||||
setPreprocessor(&PP);
|
||||
return true;
|
||||
}
|
||||
|
||||
virtual void InclusionDirective(SourceLocation hash_loc,
|
||||
const Token &include_token,
|
||||
StringRef file_name, bool is_angled,
|
||||
CharSourceRange filename_range,
|
||||
const FileEntry *file, StringRef search_path,
|
||||
StringRef relative_path,
|
||||
const clang::Module *imported) override {
|
||||
if (_sm->isWrittenInMainFile(hash_loc)) {
|
||||
if (is_angled) {
|
||||
const auto found = N.cuda2hipRename.find(file_name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
countReps[found->second.countType]++;
|
||||
StringRef repName = found->second.hipName;
|
||||
DEBUG(dbgs() << "Include file found: " << file_name << "\n"
|
||||
<< "SourceLocation:"
|
||||
<< filename_range.getBegin().printToString(*_sm) << "\n"
|
||||
<< "Will be replaced with " << repName << "\n");
|
||||
SourceLocation sl = filename_range.getBegin();
|
||||
SourceLocation sle = filename_range.getEnd();
|
||||
const char *B = _sm->getCharacterData(sl);
|
||||
const char *E = _sm->getCharacterData(sle);
|
||||
SmallString<128> tmpData;
|
||||
Replacement Rep(*_sm, sl, E - B,
|
||||
Twine("<" + repName + ">").toStringRef(tmpData));
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
virtual void MacroDefined(const Token &MacroNameTok,
|
||||
const MacroDirective *MD) override {
|
||||
if (_sm->isWrittenInMainFile(MD->getLocation()) &&
|
||||
MD->getKind() == MacroDirective::MD_Define) {
|
||||
for (auto T : MD->getMacroInfo()->tokens()) {
|
||||
if (T.isAnyIdentifier()) {
|
||||
StringRef name = T.getIdentifierInfo()->getName();
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
countReps[found->second.countType]++;
|
||||
StringRef repName = found->second.hipName;
|
||||
SourceLocation sl = T.getLocation();
|
||||
DEBUG(dbgs() << "Identifier " << name
|
||||
<< " found in definition of macro "
|
||||
<< MacroNameTok.getIdentifierInfo()->getName() << "\n"
|
||||
<< "will be replaced with: " << repName << "\n"
|
||||
<< "SourceLocation: " << sl.printToString(*_sm)
|
||||
<< "\n");
|
||||
Replacement Rep(*_sm, sl, name.size(), repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
virtual void MacroExpands(const Token &MacroNameTok,
|
||||
const MacroDefinition &MD, SourceRange Range,
|
||||
const MacroArgs *Args) override {
|
||||
if (_sm->isWrittenInMainFile(MacroNameTok.getLocation())) {
|
||||
for (unsigned int i = 0; Args && i < MD.getMacroInfo()->getNumArgs();
|
||||
i++) {
|
||||
StringRef macroName = MacroNameTok.getIdentifierInfo()->getName();
|
||||
std::vector<Token> toks;
|
||||
// Code below is a kind of stolen from 'MacroArgs::getPreExpArgument'
|
||||
// to workaround the 'const' MacroArgs passed into this hook.
|
||||
const Token *start = Args->getUnexpArgument(i);
|
||||
size_t len = Args->getArgLength(start) + 1;
|
||||
#if (LLVM_VERSION_MAJOR >= 3) && (LLVM_VERSION_MINOR >= 9)
|
||||
_pp->EnterTokenStream(ArrayRef<Token>(start, len), false);
|
||||
#else
|
||||
_pp->EnterTokenStream(start, len, false, false);
|
||||
#endif
|
||||
do {
|
||||
toks.push_back(Token());
|
||||
Token &tk = toks.back();
|
||||
_pp->Lex(tk);
|
||||
} while (toks.back().isNot(tok::eof));
|
||||
_pp->RemoveTopOfLexerStack();
|
||||
// end of stolen code
|
||||
for (auto tok : toks) {
|
||||
if (tok.isAnyIdentifier()) {
|
||||
StringRef name = tok.getIdentifierInfo()->getName();
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
countReps[found->second.countType]++;
|
||||
StringRef repName = found->second.hipName;
|
||||
DEBUG(dbgs()
|
||||
<< "Identifier " << name
|
||||
<< " found as an actual argument in expansion of macro "
|
||||
<< macroName << "\n"
|
||||
<< "will be replaced with: " << repName << "\n");
|
||||
SourceLocation sl = tok.getLocation();
|
||||
Replacement Rep(*_sm, sl, name.size(), repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
}
|
||||
if (tok.is(tok::string_literal)) {
|
||||
StringRef s(tok.getLiteralData(), tok.getLength());
|
||||
processString(unquoteStr(s), N, Replace, *_sm, tok.getLocation(),
|
||||
countReps);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
void EndOfMainFile() override {}
|
||||
|
||||
bool SeenEnd;
|
||||
void setSourceManager(SourceManager *sm) { _sm = sm; }
|
||||
void setPreprocessor(Preprocessor *pp) { _pp = pp; }
|
||||
|
||||
int64_t countReps[ConvTypes::CONV_LAST] = {0};
|
||||
|
||||
private:
|
||||
SourceManager *_sm;
|
||||
Preprocessor *_pp;
|
||||
|
||||
Replacements *Replace;
|
||||
struct cuda2hipMap N;
|
||||
};
|
||||
|
||||
class Cuda2HipCallback : public MatchFinder::MatchCallback {
|
||||
public:
|
||||
Cuda2HipCallback(Replacements *Replace, ast_matchers::MatchFinder *parent)
|
||||
: Replace(Replace), owner(parent) {}
|
||||
|
||||
void convertKernelDecl(const FunctionDecl *kernelDecl,
|
||||
const MatchFinder::MatchResult &Result) {
|
||||
SourceManager *SM = Result.SourceManager;
|
||||
LangOptions DefaultLangOptions;
|
||||
|
||||
SmallString<40> XStr;
|
||||
raw_svector_ostream OS(XStr);
|
||||
StringRef initialParamList;
|
||||
OS << "hipLaunchParm lp";
|
||||
size_t replacementLength = OS.str().size();
|
||||
SourceLocation sl = kernelDecl->getNameInfo().getEndLoc();
|
||||
SourceLocation kernelArgListStart = clang::Lexer::findLocationAfterToken(
|
||||
sl, clang::tok::l_paren, *SM, DefaultLangOptions, true);
|
||||
DEBUG(dbgs() << kernelArgListStart.printToString(*SM));
|
||||
if (kernelDecl->getNumParams() > 0) {
|
||||
const ParmVarDecl *pvdFirst = kernelDecl->getParamDecl(0);
|
||||
const ParmVarDecl *pvdLast =
|
||||
kernelDecl->getParamDecl(kernelDecl->getNumParams() - 1);
|
||||
SourceLocation kernelArgListStart(pvdFirst->getLocStart());
|
||||
SourceLocation kernelArgListEnd(pvdLast->getLocEnd());
|
||||
SourceLocation stop = clang::Lexer::getLocForEndOfToken(
|
||||
kernelArgListEnd, 0, *SM, DefaultLangOptions);
|
||||
replacementLength +=
|
||||
SM->getCharacterData(stop) - SM->getCharacterData(kernelArgListStart);
|
||||
initialParamList = StringRef(SM->getCharacterData(kernelArgListStart),
|
||||
replacementLength);
|
||||
OS << ", " << initialParamList;
|
||||
}
|
||||
DEBUG(dbgs() << "initial paramlist: " << initialParamList << "\n"
|
||||
<< "new paramlist: " << OS.str() << "\n");
|
||||
Replacement Rep0(*(Result.SourceManager), kernelArgListStart,
|
||||
replacementLength, OS.str());
|
||||
Replace->insert(Rep0);
|
||||
}
|
||||
|
||||
void run(const MatchFinder::MatchResult &Result) override {
|
||||
SourceManager *SM = Result.SourceManager;
|
||||
LangOptions DefaultLangOptions;
|
||||
|
||||
if (const CallExpr *call =
|
||||
Result.Nodes.getNodeAs<clang::CallExpr>("cudaCall")) {
|
||||
const FunctionDecl *funcDcl = call->getDirectCallee();
|
||||
StringRef name = funcDcl->getDeclName().getAsString();
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
countReps[found->second.countType]++;
|
||||
StringRef repName = found->second.hipName;
|
||||
SourceLocation sl = call->getLocStart();
|
||||
Replacement Rep(*SM, SM->isMacroArgExpansion(sl)
|
||||
? SM->getImmediateSpellingLoc(sl)
|
||||
: sl,
|
||||
name.size(), repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
}
|
||||
|
||||
if (const CUDAKernelCallExpr *launchKernel =
|
||||
Result.Nodes.getNodeAs<clang::CUDAKernelCallExpr>(
|
||||
"cudaLaunchKernel")) {
|
||||
SmallString<40> XStr;
|
||||
raw_svector_ostream OS(XStr);
|
||||
StringRef calleeName;
|
||||
const FunctionDecl *kernelDecl = launchKernel->getDirectCallee();
|
||||
if (kernelDecl) {
|
||||
calleeName = kernelDecl->getName();
|
||||
convertKernelDecl(kernelDecl, Result);
|
||||
} else {
|
||||
const Expr *e = launchKernel->getCallee();
|
||||
if (const UnresolvedLookupExpr *ule =
|
||||
dyn_cast<UnresolvedLookupExpr>(e)) {
|
||||
calleeName = ule->getName().getAsIdentifierInfo()->getName();
|
||||
owner->addMatcher(functionTemplateDecl(hasName(calleeName))
|
||||
.bind("unresolvedTemplateName"),
|
||||
this);
|
||||
}
|
||||
}
|
||||
|
||||
XStr.clear();
|
||||
OS << "hipLaunchKernel(HIP_KERNEL_NAME(" << calleeName << "),";
|
||||
|
||||
const CallExpr *config = launchKernel->getConfig();
|
||||
DEBUG(dbgs() << "Kernel config arguments:"
|
||||
<< "\n");
|
||||
for (unsigned argno = 0; argno < config->getNumArgs(); argno++) {
|
||||
const Expr *arg = config->getArg(argno);
|
||||
if (!isa<CXXDefaultArgExpr>(arg)) {
|
||||
const ParmVarDecl *pvd =
|
||||
config->getDirectCallee()->getParamDecl(argno);
|
||||
|
||||
SourceLocation sl(arg->getLocStart());
|
||||
SourceLocation el(arg->getLocEnd());
|
||||
SourceLocation stop =
|
||||
clang::Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions);
|
||||
StringRef outs(SM->getCharacterData(sl),
|
||||
SM->getCharacterData(stop) - SM->getCharacterData(sl));
|
||||
DEBUG(dbgs() << "args[ " << argno << "]" << outs << " <"
|
||||
<< pvd->getType().getAsString() << ">"
|
||||
<< "\n");
|
||||
if (pvd->getType().getAsString().compare("dim3") == 0)
|
||||
OS << " dim3(" << outs << "),";
|
||||
else
|
||||
OS << " " << outs << ",";
|
||||
} else
|
||||
OS << " 0,";
|
||||
}
|
||||
|
||||
for (unsigned argno = 0; argno < launchKernel->getNumArgs(); argno++) {
|
||||
const Expr *arg = launchKernel->getArg(argno);
|
||||
SourceLocation sl(arg->getLocStart());
|
||||
SourceLocation el(arg->getLocEnd());
|
||||
SourceLocation stop =
|
||||
clang::Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions);
|
||||
std::string outs(SM->getCharacterData(sl),
|
||||
SM->getCharacterData(stop) - SM->getCharacterData(sl));
|
||||
DEBUG(dbgs() << outs << "\n");
|
||||
OS << " " << outs << ",";
|
||||
}
|
||||
XStr.pop_back();
|
||||
OS << ")";
|
||||
size_t length =
|
||||
SM->getCharacterData(clang::Lexer::getLocForEndOfToken(
|
||||
launchKernel->getLocEnd(), 0, *SM, DefaultLangOptions)) -
|
||||
SM->getCharacterData(launchKernel->getLocStart());
|
||||
Replacement Rep(*SM, launchKernel->getLocStart(), length, OS.str());
|
||||
Replace->insert(Rep);
|
||||
countReps[ConvTypes::CONV_KERN]++;
|
||||
}
|
||||
|
||||
if (const FunctionTemplateDecl *templateDecl =
|
||||
Result.Nodes.getNodeAs<clang::FunctionTemplateDecl>(
|
||||
"unresolvedTemplateName")) {
|
||||
FunctionDecl *kernelDecl = templateDecl->getTemplatedDecl();
|
||||
convertKernelDecl(kernelDecl, Result);
|
||||
}
|
||||
|
||||
if (const MemberExpr *threadIdx =
|
||||
Result.Nodes.getNodeAs<clang::MemberExpr>("cudaBuiltin")) {
|
||||
if (const OpaqueValueExpr *refBase =
|
||||
dyn_cast<OpaqueValueExpr>(threadIdx->getBase())) {
|
||||
if (const DeclRefExpr *declRef =
|
||||
dyn_cast<DeclRefExpr>(refBase->getSourceExpr())) {
|
||||
StringRef name = declRef->getDecl()->getName();
|
||||
StringRef memberName = threadIdx->getMemberDecl()->getName();
|
||||
size_t pos = memberName.find_first_not_of("__fetch_builtin_");
|
||||
memberName = memberName.slice(pos, memberName.size());
|
||||
SmallString<128> tmpData;
|
||||
name = Twine(name + "." + memberName).toStringRef(tmpData);
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
countReps[found->second.countType]++;
|
||||
StringRef repName = found->second.hipName;
|
||||
SourceLocation sl = threadIdx->getLocStart();
|
||||
Replacement Rep(*SM, sl, name.size(), repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
if (const DeclRefExpr *cudaEnumConstantRef =
|
||||
Result.Nodes.getNodeAs<clang::DeclRefExpr>("cudaEnumConstantRef")) {
|
||||
StringRef name = cudaEnumConstantRef->getDecl()->getNameAsString();
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
countReps[found->second.countType]++;
|
||||
StringRef repName = found->second.hipName;
|
||||
SourceLocation sl = cudaEnumConstantRef->getLocStart();
|
||||
Replacement Rep(*SM, sl, name.size(), repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
}
|
||||
|
||||
if (const VarDecl *cudaEnumConstantDecl =
|
||||
Result.Nodes.getNodeAs<clang::VarDecl>("cudaEnumConstantDecl")) {
|
||||
StringRef name =
|
||||
cudaEnumConstantDecl->getType()->getAsTagDecl()->getNameAsString();
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
countReps[found->second.countType]++;
|
||||
StringRef repName = found->second.hipName;
|
||||
SourceLocation sl = cudaEnumConstantDecl->getLocStart();
|
||||
Replacement Rep(*SM, sl, name.size(), repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
}
|
||||
|
||||
if (const VarDecl *cudaStructVar =
|
||||
Result.Nodes.getNodeAs<clang::VarDecl>("cudaStructVar")) {
|
||||
StringRef name = cudaStructVar->getType()
|
||||
->getAsStructureType()
|
||||
->getDecl()
|
||||
->getNameAsString();
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
countReps[found->second.countType]++;
|
||||
StringRef repName = found->second.hipName;
|
||||
TypeLoc TL = cudaStructVar->getTypeSourceInfo()->getTypeLoc();
|
||||
SourceLocation sl = TL.getUnqualifiedLoc().getLocStart();
|
||||
Replacement Rep(*SM, sl, name.size(), repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
}
|
||||
|
||||
if (const VarDecl *cudaStructVarPtr =
|
||||
Result.Nodes.getNodeAs<clang::VarDecl>("cudaStructVarPtr")) {
|
||||
const Type *t = cudaStructVarPtr->getType().getTypePtrOrNull();
|
||||
if (t) {
|
||||
StringRef name = t->getPointeeCXXRecordDecl()->getName();
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
countReps[found->second.countType]++;
|
||||
StringRef repName = found->second.hipName;
|
||||
TypeLoc TL = cudaStructVarPtr->getTypeSourceInfo()->getTypeLoc();
|
||||
SourceLocation sl = TL.getUnqualifiedLoc().getLocStart();
|
||||
Replacement Rep(*SM, sl, name.size(), repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
if (const ParmVarDecl *cudaParamDecl =
|
||||
Result.Nodes.getNodeAs<clang::ParmVarDecl>("cudaParamDecl")) {
|
||||
QualType QT = cudaParamDecl->getOriginalType().getUnqualifiedType();
|
||||
StringRef name = QT.getAsString();
|
||||
const Type *t = QT.getTypePtr();
|
||||
if (t->isStructureOrClassType()) {
|
||||
name = t->getAsCXXRecordDecl()->getName();
|
||||
}
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
countReps[found->second.countType]++;
|
||||
StringRef repName = found->second.hipName;
|
||||
TypeLoc TL = cudaParamDecl->getTypeSourceInfo()->getTypeLoc();
|
||||
SourceLocation sl = TL.getUnqualifiedLoc().getLocStart();
|
||||
Replacement Rep(*SM, sl, name.size(), repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
}
|
||||
|
||||
if (const ParmVarDecl *cudaParamDeclPtr =
|
||||
Result.Nodes.getNodeAs<clang::ParmVarDecl>("cudaParamDeclPtr")) {
|
||||
const Type *pt = cudaParamDeclPtr->getType().getTypePtrOrNull();
|
||||
if (pt) {
|
||||
QualType QT = pt->getPointeeType();
|
||||
const Type *t = QT.getTypePtr();
|
||||
StringRef name = t->isStructureOrClassType()
|
||||
? t->getAsCXXRecordDecl()->getName()
|
||||
: StringRef(QT.getAsString());
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
countReps[found->second.countType]++;
|
||||
StringRef repName = found->second.hipName;
|
||||
TypeLoc TL = cudaParamDeclPtr->getTypeSourceInfo()->getTypeLoc();
|
||||
SourceLocation sl = TL.getUnqualifiedLoc().getLocStart();
|
||||
Replacement Rep(*SM, sl, name.size(), repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
if (const StringLiteral *stringLiteral =
|
||||
Result.Nodes.getNodeAs<clang::StringLiteral>("stringLiteral")) {
|
||||
if (stringLiteral->getCharByteWidth() == 1) {
|
||||
StringRef s = stringLiteral->getString();
|
||||
processString(s, N, Replace, *SM, stringLiteral->getLocStart(),
|
||||
countReps);
|
||||
}
|
||||
}
|
||||
|
||||
if (const UnaryExprOrTypeTraitExpr *expr =
|
||||
Result.Nodes.getNodeAs<clang::UnaryExprOrTypeTraitExpr>(
|
||||
"cudaStructSizeOf")) {
|
||||
TypeSourceInfo *typeInfo = expr->getArgumentTypeInfo();
|
||||
QualType QT = typeInfo->getType().getUnqualifiedType();
|
||||
const Type *type = QT.getTypePtr();
|
||||
StringRef name = type->getAsCXXRecordDecl()->getName();
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
countReps[found->second.countType]++;
|
||||
StringRef repName = found->second.hipName;
|
||||
TypeLoc TL = typeInfo->getTypeLoc();
|
||||
SourceLocation sl = TL.getUnqualifiedLoc().getLocStart();
|
||||
Replacement Rep(*SM, sl, name.size(), repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
int64_t countReps[ConvTypes::CONV_LAST] = {0};
|
||||
|
||||
private:
|
||||
Replacements *Replace;
|
||||
ast_matchers::MatchFinder *owner;
|
||||
struct cuda2hipMap N;
|
||||
};
|
||||
|
||||
} // end anonymous namespace
|
||||
|
||||
// Set up the command line options
|
||||
static cl::OptionCategory
|
||||
ToolTemplateCategory("CUDA to HIP source translator options");
|
||||
static cl::extrahelp MoreHelp("<source0> specify the path of source file\n\n");
|
||||
|
||||
static cl::opt<std::string> OutputFilename("o", cl::desc("Output filename"),
|
||||
cl::value_desc("filename"),
|
||||
cl::cat(ToolTemplateCategory));
|
||||
|
||||
static cl::opt<bool>
|
||||
Inplace("inplace",
|
||||
cl::desc("Modify input file inplace, replacing input with hipified "
|
||||
"output, save backup in .prehip file. "),
|
||||
cl::value_desc("inplace"), cl::cat(ToolTemplateCategory));
|
||||
|
||||
static cl::opt<bool>
|
||||
NoOutput("no-output",
|
||||
cl::desc("don't write any translated output to stdout"),
|
||||
cl::value_desc("no-output"), cl::cat(ToolTemplateCategory));
|
||||
static cl::opt<bool>
|
||||
PrintStats("print-stats", cl::desc("print the command-line, like a header"),
|
||||
cl::value_desc("print-stats"), cl::cat(ToolTemplateCategory));
|
||||
|
||||
int main(int argc, const char **argv) {
|
||||
|
||||
llvm::sys::PrintStackTraceOnErrorSignal();
|
||||
|
||||
int Result;
|
||||
|
||||
CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory,
|
||||
llvm::cl::Required);
|
||||
std::string dst = OutputFilename;
|
||||
std::vector<std::string> fileSources = OptionsParser.getSourcePathList();
|
||||
if (dst.empty()) {
|
||||
dst = fileSources[0];
|
||||
if (!Inplace) {
|
||||
size_t pos = dst.rfind(".cu");
|
||||
if (pos != std::string::npos) {
|
||||
dst = dst.substr(0, pos) + ".hip.cu";
|
||||
} else {
|
||||
llvm::errs() << "Input .cu file was not specified.\n";
|
||||
return 1;
|
||||
}
|
||||
}
|
||||
} else {
|
||||
if (Inplace) {
|
||||
llvm::errs() << "Conflict: both -o and -inplace options are specified.";
|
||||
}
|
||||
dst += ".cu";
|
||||
}
|
||||
|
||||
// copy source file since tooling makes changes "inplace"
|
||||
std::ifstream source(fileSources[0], std::ios::binary);
|
||||
std::ofstream dest(Inplace ? dst + ".prehip" : dst, std::ios::binary);
|
||||
dest << source.rdbuf();
|
||||
source.close();
|
||||
dest.close();
|
||||
|
||||
RefactoringTool Tool(OptionsParser.getCompilations(), dst);
|
||||
ast_matchers::MatchFinder Finder;
|
||||
Cuda2HipCallback Callback(&Tool.getReplacements(), &Finder);
|
||||
HipifyPPCallbacks PPCallbacks(&Tool.getReplacements());
|
||||
Finder.addMatcher(callExpr(isExpansionInMainFile(),
|
||||
callee(functionDecl(matchesName("cuda.*"))))
|
||||
.bind("cudaCall"),
|
||||
&Callback);
|
||||
Finder.addMatcher(cudaKernelCallExpr().bind("cudaLaunchKernel"), &Callback);
|
||||
Finder.addMatcher(memberExpr(isExpansionInMainFile(),
|
||||
hasObjectExpression(hasType(cxxRecordDecl(
|
||||
matchesName("__cuda_builtin_")))))
|
||||
.bind("cudaBuiltin"),
|
||||
&Callback);
|
||||
Finder.addMatcher(declRefExpr(isExpansionInMainFile(),
|
||||
to(enumConstantDecl(matchesName("cuda.*"))))
|
||||
.bind("cudaEnumConstantRef"),
|
||||
&Callback);
|
||||
Finder.addMatcher(
|
||||
varDecl(isExpansionInMainFile(), hasType(enumDecl(matchesName("cuda.*"))))
|
||||
.bind("cudaEnumConstantDecl"),
|
||||
&Callback);
|
||||
Finder.addMatcher(varDecl(isExpansionInMainFile(),
|
||||
hasType(cxxRecordDecl(matchesName("cuda.*"))))
|
||||
.bind("cudaStructVar"),
|
||||
&Callback);
|
||||
Finder.addMatcher(
|
||||
varDecl(isExpansionInMainFile(),
|
||||
hasType(pointsTo(cxxRecordDecl(matchesName("cuda.*")))))
|
||||
.bind("cudaStructVarPtr"),
|
||||
&Callback);
|
||||
Finder.addMatcher(parmVarDecl(isExpansionInMainFile(),
|
||||
hasType(namedDecl(matchesName("cuda.*"))))
|
||||
.bind("cudaParamDecl"),
|
||||
&Callback);
|
||||
Finder.addMatcher(
|
||||
parmVarDecl(isExpansionInMainFile(),
|
||||
hasType(pointsTo(namedDecl(matchesName("cuda.*")))))
|
||||
.bind("cudaParamDeclPtr"),
|
||||
&Callback);
|
||||
Finder.addMatcher(expr(isExpansionInMainFile(),
|
||||
sizeOfExpr(hasArgumentOfType(recordType(hasDeclaration(
|
||||
cxxRecordDecl(matchesName("cuda.*")))))))
|
||||
.bind("cudaStructSizeOf"),
|
||||
&Callback);
|
||||
Finder.addMatcher(
|
||||
stringLiteral(isExpansionInMainFile()).bind("stringLiteral"), &Callback);
|
||||
|
||||
auto action = newFrontendActionFactory(&Finder, &PPCallbacks);
|
||||
|
||||
std::vector<const char *> compilationStages;
|
||||
compilationStages.push_back("--cuda-host-only");
|
||||
//compilationStages.push_back("--cuda-device-only");
|
||||
|
||||
for (auto Stage : compilationStages) {
|
||||
Tool.appendArgumentsAdjuster(
|
||||
getInsertArgumentAdjuster(Stage, ArgumentInsertPosition::BEGIN));
|
||||
Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("-std=c++11"));
|
||||
#if defined(HIPIFY_CLANG_RES)
|
||||
Tool.appendArgumentsAdjuster(
|
||||
getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES));
|
||||
#endif
|
||||
Tool.appendArgumentsAdjuster(getClangSyntaxOnlyAdjuster());
|
||||
Result = Tool.run(action.get());
|
||||
|
||||
Tool.clearArgumentsAdjusters();
|
||||
}
|
||||
|
||||
LangOptions DefaultLangOptions;
|
||||
IntrusiveRefCntPtr<DiagnosticOptions> DiagOpts = new DiagnosticOptions();
|
||||
TextDiagnosticPrinter DiagnosticPrinter(llvm::errs(), &*DiagOpts);
|
||||
DiagnosticsEngine Diagnostics(
|
||||
IntrusiveRefCntPtr<DiagnosticIDs>(new DiagnosticIDs()), &*DiagOpts,
|
||||
&DiagnosticPrinter, false);
|
||||
SourceManager Sources(Diagnostics, Tool.getFiles());
|
||||
|
||||
DEBUG(dbgs() << "Replacements collected by the tool:\n");
|
||||
for (const auto &r : Tool.getReplacements()) {
|
||||
DEBUG(dbgs() << r.toString() << "\n");
|
||||
}
|
||||
|
||||
Rewriter Rewrite(Sources, DefaultLangOptions);
|
||||
|
||||
if (!Tool.applyAllReplacements(Rewrite)) {
|
||||
DEBUG(dbgs() << "Skipped some replacements.\n");
|
||||
}
|
||||
|
||||
Result = Rewrite.overwriteChangedFiles();
|
||||
|
||||
if (!Inplace) {
|
||||
size_t pos = dst.rfind(".cu");
|
||||
if (pos != std::string::npos) {
|
||||
rename(dst.c_str(), dst.substr(0, pos).c_str());
|
||||
}
|
||||
}
|
||||
if (PrintStats) {
|
||||
int64_t sum = 0;
|
||||
for (int i = 0; i < ConvTypes::CONV_LAST; i++) {
|
||||
sum += Callback.countReps[i] + PPCallbacks.countReps[i];
|
||||
}
|
||||
llvm::outs() << "info: converted " << sum << " CUDA->HIP refs ( ";
|
||||
for (int i = 0; i < ConvTypes::CONV_LAST; i++) {
|
||||
llvm::outs() << counterNames[i] << ':'
|
||||
<< Callback.countReps[i] + PPCallbacks.countReps[i] << ' ';
|
||||
}
|
||||
llvm::outs() << ") in \'" << fileSources[0] << "\'\n";
|
||||
}
|
||||
return Result;
|
||||
}
|
||||
@@ -25,7 +25,7 @@ THE SOFTWARE.
|
||||
#include "hip/hcc_detail/staging_buffer.h"
|
||||
|
||||
|
||||
#if defined(__HCC__) && (__hcc_workweek__ < 1502)
|
||||
#if defined(__HCC__) && (__hcc_workweek__ < 16155)
|
||||
#error("This version of HIP requires a newer version of HCC.");
|
||||
#endif
|
||||
|
||||
@@ -41,7 +41,7 @@ THE SOFTWARE.
|
||||
|
||||
|
||||
// Use new am_memory_host_lock APIs:
|
||||
#define USE_HCC_LOCK_API 0
|
||||
#define USE_HCC_LOCK_API 1
|
||||
|
||||
|
||||
//---
|
||||
@@ -440,9 +440,13 @@ public:
|
||||
hc::accelerator_view _av;
|
||||
unsigned _flags;
|
||||
|
||||
private: // Critical Data. THis MUST be accessed through LockedAccessor_StreamCrit_t
|
||||
private:
|
||||
// Critical Data. THis MUST be accessed through LockedAccessor_StreamCrit_t
|
||||
ihipStreamCritical_t _criticalData;
|
||||
|
||||
// Array of dependency completion_future.
|
||||
std::vector<hc::completion_future> _depFutures;
|
||||
|
||||
private:
|
||||
void enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal);
|
||||
void waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal);
|
||||
|
||||
@@ -181,6 +181,8 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty
|
||||
// Reset the stream to "empty" - next command will not set up an inpute dependency on any older signal.
|
||||
crit->_last_command_type = ihipCommandCopyH2D;
|
||||
crit->_last_copy_signal = NULL;
|
||||
|
||||
_depFutures.clear();
|
||||
}
|
||||
|
||||
|
||||
@@ -428,6 +430,9 @@ int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t
|
||||
needSync = 1;
|
||||
hsa_signal_t *hsaSignal = (static_cast<hsa_signal_t*> (crit->_last_kernel_future.get_native_handle()));
|
||||
if (hsaSignal) {
|
||||
// Keep reference to the kernel future in order to keep the
|
||||
// dependent signal alive.
|
||||
_depFutures.push_back(crit->_last_kernel_future);
|
||||
*waitSignal = * hsaSignal;
|
||||
} else {
|
||||
assert(0); // if NULL signal, and we return 1, hsa_amd_memory_copy_async will fail. Confirm this never happens.
|
||||
|
||||
@@ -427,10 +427,53 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s
|
||||
|
||||
hipError_t hipMemset(void* dst, int value, size_t sizeBytes )
|
||||
{
|
||||
HIP_INIT_API(dst, value, sizeBytes);
|
||||
|
||||
hipStream_t stream = hipStreamNull;
|
||||
// TODO - call an ihip memset so HIP_TRACE is correct.
|
||||
return hipMemsetAsync(dst, value, sizeBytes, hipStreamNull);
|
||||
HIP_INIT_API(dst, value, sizeBytes, stream);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
stream = ihipSyncAndResolveStream(stream);
|
||||
|
||||
if (stream) {
|
||||
stream->lockopen_preKernelCommand();
|
||||
|
||||
hc::completion_future cf ;
|
||||
|
||||
if ((sizeBytes & 0x3) == 0) {
|
||||
// use a faster dword-per-workitem copy:
|
||||
try {
|
||||
value = value & 0xff;
|
||||
unsigned value32 = (value << 24) | (value << 16) | (value << 8) | (value) ;
|
||||
cf = ihipMemsetKernel<unsigned> (stream, static_cast<unsigned*> (dst), value32, sizeBytes/sizeof(unsigned));
|
||||
}
|
||||
catch (std::exception &ex) {
|
||||
e = hipErrorInvalidValue;
|
||||
}
|
||||
} else {
|
||||
// use a slow byte-per-workitem copy:
|
||||
try {
|
||||
cf = ihipMemsetKernel<char> (stream, static_cast<char*> (dst), value, sizeBytes);
|
||||
}
|
||||
catch (std::exception &ex) {
|
||||
e = hipErrorInvalidValue;
|
||||
}
|
||||
}
|
||||
cf.wait();
|
||||
|
||||
stream->lockclose_postKernelCommand(cf);
|
||||
|
||||
|
||||
if (HIP_LAUNCH_BLOCKING) {
|
||||
tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING wait for memset [stream:%p].\n", __func__, (void*)stream);
|
||||
cf.wait();
|
||||
tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING memset completed [stream:%p].\n", __func__, (void*)stream);
|
||||
}
|
||||
} else {
|
||||
e = hipErrorInvalidValue;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
|
||||
|
||||
@@ -24,7 +24,7 @@ THE SOFTWARE.
|
||||
#include "hcc_detail/trace_helper.h"
|
||||
|
||||
/**
|
||||
* @warning HCC returns 0 in *canAccessPeer ; Need to update this function when RT supports P2P
|
||||
* HCC returns 0 in *canAccessPeer ; Need to update this function when RT supports P2P
|
||||
*/
|
||||
//---
|
||||
hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDeviceId)
|
||||
@@ -145,7 +145,7 @@ hipError_t hipMemcpyPeer (void* dst, int dstDevice, const void* src, int srcDe
|
||||
|
||||
|
||||
/**
|
||||
* @bug This function uses a synchronous copy
|
||||
* This function uses a synchronous copy
|
||||
*/
|
||||
//---
|
||||
hipError_t hipMemcpyPeerAsync (void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream)
|
||||
|
||||
@@ -0,0 +1,45 @@
|
||||
// RUN: hipify "%s" -o=%t --
|
||||
|
||||
#include <iostream>
|
||||
|
||||
__global__ void axpy(float a, float* x, float* y) {
|
||||
// RUN: sh -c "test `grep -c -F 'y[hipThreadIdx_x] = a * x[hipThreadIdx_x];' %t` -eq 2"
|
||||
y[threadIdx.x] = a * x[threadIdx.x];
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
const int kDataLen = 4;
|
||||
|
||||
float a = 2.0f;
|
||||
float host_x[kDataLen] = {1.0f, 2.0f, 3.0f, 4.0f};
|
||||
float host_y[kDataLen];
|
||||
|
||||
// Copy input data to device.
|
||||
float* device_x;
|
||||
float* device_y;
|
||||
// RUN: sh -c "test `grep -c -F 'hipMalloc(&device_x, kDataLen * sizeof(float));' %t` -eq 2"
|
||||
cudaMalloc(&device_x, kDataLen * sizeof(float));
|
||||
// RUN: sh -c "test `grep -c -F 'hipMalloc(&device_y, kDataLen * sizeof(float));' %t` -eq 2"
|
||||
cudaMalloc(&device_y, kDataLen * sizeof(float));
|
||||
// RUN: sh -c "test `grep -c -F 'hipMemcpy(device_x, host_x, kDataLen * sizeof(float), hipMemcpyHostToDevice);' %t` -eq 2"
|
||||
cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), cudaMemcpyHostToDevice);
|
||||
|
||||
// Launch the kernel.
|
||||
// RUN: sh -c "test `grep -c -F 'hipLaunchKernel(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);' %t` -eq 2"
|
||||
axpy<<<1, kDataLen>>>(a, device_x, device_y);
|
||||
|
||||
// Copy output data to host.
|
||||
// RUN: sh -c "test `grep -c -F 'hipDeviceSynchronize();' %t` -eq 2"
|
||||
cudaDeviceSynchronize();
|
||||
// RUN: sh -c "test `grep -c -F 'hipMemcpy(host_y, device_y, kDataLen * sizeof(float), hipMemcpyDeviceToHost);' %t` -eq 2"
|
||||
cudaMemcpy(host_y, device_y, kDataLen * sizeof(float), cudaMemcpyDeviceToHost);
|
||||
|
||||
// Print the results.
|
||||
for (int i = 0; i < kDataLen; ++i) {
|
||||
std::cout << "y[" << i << "] = " << host_y[i] << "\n";
|
||||
}
|
||||
|
||||
// RUN: sh -c "test `grep -c -F 'hipDeviceReset();' %t` -eq 2"
|
||||
cudaDeviceReset();
|
||||
return 0;
|
||||
}
|
||||
@@ -0,0 +1,48 @@
|
||||
# -*- Python -*-
|
||||
import os
|
||||
import platform
|
||||
import re
|
||||
import subprocess
|
||||
|
||||
import lit.formats
|
||||
import lit.util
|
||||
|
||||
# Configuration file for the 'lit' test runner.
|
||||
|
||||
# name: The name of this test suite.
|
||||
config.name = 'hipify'
|
||||
|
||||
# suffixes: CUDA source is only supported
|
||||
config.suffixes = ['.cu']
|
||||
|
||||
# testFormat: The test format to use to interpret tests.
|
||||
config.test_format = lit.formats.ShTest()
|
||||
|
||||
# test_source_root: The root path where tests are located.
|
||||
config.test_source_root = os.path.dirname(__file__)
|
||||
|
||||
# test_exec_root: The path where tests are located (default is the test suite root).
|
||||
#config.test_exec_root = config.test_source_root
|
||||
|
||||
# target_triple: Used by ShTest and TclTest formats for XFAIL checks.
|
||||
config.target_triple = '(unused)'
|
||||
|
||||
# available_features: Used by ShTest and TclTest formats for REQUIRES checks.
|
||||
config.available_features = []
|
||||
|
||||
site_cfg = lit_config.params.get('site_config', None)
|
||||
lit_config.load_config(config, site_cfg)
|
||||
|
||||
obj_root = getattr(config, 'obj_root', None)
|
||||
if obj_root is not None:
|
||||
config.test_exec_root = obj_root
|
||||
|
||||
if obj_root is not None:
|
||||
llvm_tools_dir = getattr(config, 'llvm_tools_dir', None)
|
||||
if not llvm_tools_dir:
|
||||
lit_config.fatal('No LLVM tools dir set!')
|
||||
path = os.path.pathsep.join((llvm_tools_dir, config.environment['PATH']))
|
||||
config.environment['PATH'] = path
|
||||
|
||||
config.substitutions.append(("hipify", obj_root+"/hipify-clang"))
|
||||
|
||||
@@ -0,0 +1,15 @@
|
||||
import sys
|
||||
|
||||
config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@"
|
||||
config.obj_root = "@BINARY_DIR@"
|
||||
|
||||
# Support substitution of the tools and libs dirs with user parameters. This is
|
||||
# used when we can't determine the tool dir at configuration time.
|
||||
try:
|
||||
config.llvm_tools_dir = config.llvm_tools_dir % lit_config.params
|
||||
config.obj_root = config.obj_root % lit_config.params
|
||||
except KeyError:
|
||||
e = sys.exc_info()[1]
|
||||
key, = e.args
|
||||
lit_config.fatal("unable to find %r parameter, use '--param=%s=VALUE'" % (key,key))
|
||||
|
||||
@@ -167,7 +167,6 @@ make_hip_executable (hipMultiThreadStreams2 hipMultiThreadStreams2.cpp)
|
||||
make_hip_executable (hipHostAlloc hipHostAlloc.cpp)
|
||||
make_hip_executable (hipStreamL5 hipStreamL5.cpp)
|
||||
make_hip_executable (hipHostGetFlags hipHostGetFlags.cpp)
|
||||
#TODO - re-enable. This requires working hipHostRegister call, waiting on HCC feature.
|
||||
make_hip_executable (hipHostRegister hipHostRegister.cpp)
|
||||
make_hip_executable (hipRandomMemcpyAsync hipRandomMemcpyAsync.cpp)
|
||||
make_hip_executable (hipMemoryAllocate hipMemoryAllocate.cpp)
|
||||
@@ -213,7 +212,7 @@ make_test(hipHostAlloc " ")
|
||||
# BS- comment out since test appears broken - asks for device pointer but pointer was never allocated.
|
||||
#make_test(hipHostGetFlags " ")
|
||||
make_test(hipHcc " " )
|
||||
#make_test(hipHostRegister " ")
|
||||
make_test(hipHostRegister " ")
|
||||
make_test(hipStreamL5 " ")
|
||||
make_test(hipRandomMemcpyAsync " ")
|
||||
#make_test(hipAPIStreamEnable " ")
|
||||
|
||||
@@ -20,6 +20,8 @@ THE USE OR OTHER DEALINGS IN THE SOFTWARE. */
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <assert.h>
|
||||
#include <string>
|
||||
|
||||
using namespace std;
|
||||
|
||||
//./hipEnvVar -c -d 0 -h
|
||||
|
||||
@@ -40,6 +40,8 @@ int num;
|
||||
hipGetDeviceCount(&num);
|
||||
if(num < 2)
|
||||
{
|
||||
printf ("warning: Not enough GPUs to run the test, exiting without running.\n");
|
||||
passed();
|
||||
return 0;
|
||||
}
|
||||
|
||||
|
||||
Ссылка в новой задаче
Block a user