Merge 'master' into 'amd-master'
Change-Id: If65d05820af920addd3939c09cf76aafe99512aa
Этот коммит содержится в:
Jenkins
@@ -1,6 +1,11 @@
|
||||
cmake_minimum_required(VERSION 2.8.3)
|
||||
project(hip)
|
||||
|
||||
#############################
|
||||
# Options
|
||||
#############################
|
||||
option(BUILD_HIPIFY_CLANG "Enable building the CUDA->HIP converter" OFF)
|
||||
|
||||
#############################
|
||||
# Setup config generation
|
||||
#############################
|
||||
@@ -120,13 +125,6 @@ else()
|
||||
message(FATAL_ERROR "Don't know where to install HIP. Please specify absolute path using -DCMAKE_INSTALL_PREFIX")
|
||||
endif()
|
||||
|
||||
# Check if we need to build hipify-clang
|
||||
if(NOT DEFINED HIPIFY_CLANG_LLVM_DIR)
|
||||
if(DEFINED ENV{HIPIFY_CLANG_LLVM_DIR})
|
||||
set(HIPIFY_CLANG_LLVM_DIR $ENV{HIPIFY_CLANG_LLVM_DIR})
|
||||
endif()
|
||||
endif()
|
||||
|
||||
# Check if we need to enable ATP marker
|
||||
if(NOT DEFINED COMPILE_HIP_ATP_MARKER)
|
||||
if(NOT DEFINED ENV{COMPILE_HIP_ATP_MARKER})
|
||||
@@ -142,7 +140,9 @@ add_to_config(_buildInfo COMPILE_HIP_ATP_MARKER)
|
||||
# Build steps
|
||||
#############################
|
||||
# Build clang hipify if enabled
|
||||
add_subdirectory(hipify-clang)
|
||||
if (BUILD_HIPIFY_CLANG)
|
||||
add_subdirectory(hipify-clang)
|
||||
endif()
|
||||
|
||||
# Build hip_hcc if platform is hcc
|
||||
if(HIP_PLATFORM STREQUAL "hcc")
|
||||
@@ -287,23 +287,19 @@ set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_base)
|
||||
configure_file(packaging/hip_base.txt ${BUILD_DIR}/CMakeLists.txt @ONLY)
|
||||
configure_file(packaging/hip_base.postinst ${BUILD_DIR}/postinst @ONLY)
|
||||
configure_file(packaging/hip_base.prerm ${BUILD_DIR}/prerm @ONLY)
|
||||
if(NOT BUILD_HIPIFY_CLANG)
|
||||
add_custom_target(pkg_hip_base COMMAND ${CMAKE_COMMAND} .
|
||||
COMMAND rm -rf *.deb *.rpm *.tar.gz
|
||||
COMMAND make package
|
||||
COMMAND cp *.deb ${PROJECT_BINARY_DIR}
|
||||
COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
|
||||
COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
|
||||
WORKING_DIRECTORY ${BUILD_DIR})
|
||||
else()
|
||||
add_custom_target(pkg_hip_base COMMAND ${CMAKE_COMMAND} .
|
||||
COMMAND rm -rf *.deb *.rpm *.tar.gz
|
||||
COMMAND make package
|
||||
COMMAND cp *.deb ${PROJECT_BINARY_DIR}
|
||||
COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
|
||||
COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
|
||||
WORKING_DIRECTORY ${BUILD_DIR}
|
||||
DEPENDS hipify-clang)
|
||||
|
||||
add_custom_target(pkg_hip_base COMMAND ${CMAKE_COMMAND} .
|
||||
COMMAND rm -rf *.deb *.rpm *.tar.gz
|
||||
COMMAND make package
|
||||
COMMAND cp *.deb ${PROJECT_BINARY_DIR}
|
||||
COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
|
||||
COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
|
||||
WORKING_DIRECTORY ${BUILD_DIR}
|
||||
)
|
||||
|
||||
# Packaging needs to wait for hipify-clang to build if it's enabled...
|
||||
if (BUILD_HIPIFY_CLANG)
|
||||
add_dependencies(pkg_hip_base hipify-clang)
|
||||
endif()
|
||||
|
||||
# Package: hip_hcc
|
||||
|
||||
@@ -1,31 +1,17 @@
|
||||
cmake_minimum_required(VERSION 2.8.8)
|
||||
project(hipify-clang)
|
||||
|
||||
set(BUILD_HIPIFY_CLANG 0 CACHE INTERNAL "")
|
||||
find_package(LLVM REQUIRED)
|
||||
message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}:")
|
||||
message(STATUS " - CMake module path: ${LLVM_CMAKE_DIR}")
|
||||
message(STATUS " - Include path : ${LLVM_INCLUDE_DIRS}")
|
||||
message(STATUS " - Binary path : ${LLVM_TOOLS_BINARY_DIR}")
|
||||
|
||||
if (HIPIFY_CLANG_LLVM_DIR)
|
||||
find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED NO_DEFAULT_PATH)
|
||||
else()
|
||||
message(STATUS "hipify-clang will not be built. To build it please specify absolute path to LLVM 3.8 or higher using HIPIFY_CLANG_LLVM_DIR")
|
||||
return()
|
||||
endif()
|
||||
|
||||
option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" ON)
|
||||
|
||||
# Disable the tests if `lit` or `FileCheck` is not installed.
|
||||
find_program(LIT_COMMAND lit)
|
||||
find_program(FILECHECK_COMMAND FileCheck)
|
||||
find_program(SOCAT_COMMAND socat)
|
||||
if (NOT LIT_COMMAND OR NOT FILECHECK_COMMAND OR NOT SOCAT_COMMAND)
|
||||
set(HIPIFY_CLANG_TESTS OFF CACHE INTERNAL "")
|
||||
message(STATUS "hipify-clang's tests are not being built because `lit`,`FileCheck` or `socat` could not be found.")
|
||||
endif()
|
||||
option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" OFF)
|
||||
|
||||
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})
|
||||
@@ -81,6 +67,17 @@ install(TARGETS hipify-clang DESTINATION bin)
|
||||
if (HIPIFY_CLANG_TESTS)
|
||||
find_package(PythonInterp 2.7 REQUIRED EXACT)
|
||||
|
||||
function (require_program PROGRAM_NAME)
|
||||
find_program(FOUND_PROGRAM ${PROGRAM_NAME})
|
||||
if (NOT FOUND_PROGRAM)
|
||||
message(FATAL_ERROR "Can't find ${PROGRAM_NAME}. Either set HIPIFY_CLANG_TESTS to OFF to disable hipify tests, or install the missing program.")
|
||||
endif()
|
||||
endfunction()
|
||||
|
||||
require_program(lit)
|
||||
require_program(FileCheck)
|
||||
require_program(socat)
|
||||
|
||||
# Populates CUDA_TOOLKIT_ROOT_DIR, which is then applied to the lit config to give the
|
||||
# value of --cuda-path for the test runs.
|
||||
find_package(CUDA REQUIRED)
|
||||
@@ -102,5 +99,3 @@ if (HIPIFY_CLANG_TESTS)
|
||||
add_dependencies(test-hipify-clang test-hipify)
|
||||
set_target_properties(test-hipify-clang PROPERTIES FOLDER "Tests")
|
||||
endif()
|
||||
|
||||
set(BUILD_HIPIFY_CLANG 1 CACHE INTERNAL "")
|
||||
|
||||
@@ -1,3 +1,7 @@
|
||||
# hipify-clang
|
||||
|
||||
`hipify-clang` is a clang-based tool to automatically translate CUDA source code into portable HIP C++.
|
||||
|
||||
## Table of Contents
|
||||
|
||||
<!-- toc -->
|
||||
@@ -9,68 +13,65 @@
|
||||
|
||||
<!-- tocstop -->
|
||||
|
||||
## Using hipify-clang
|
||||
## Build and install
|
||||
|
||||
`hipify-clang` is a clang-based tool which can automate the translation of CUDA source code into portable HIP C++.
|
||||
The tool can automatically add extra HIP arguments (notably the "hipLaunchParm" required at the beginning of every HIP kernel call).
|
||||
`hipify-clang` has some additional dependencies explained below and can be built as a separate make step. The instructions below are specifically for **Ubuntu 14.04** and **Ubuntu 16.04**.
|
||||
### Dependencies
|
||||
|
||||
### Build and install
|
||||
`hipify-clang` requires clang+llvm of at least version 3.8.
|
||||
|
||||
- Download and unpack clang+llvm 3.8 binary package preqrequisite.
|
||||
In most cases, you can get a suitable version of clang+llvm with your package manager.
|
||||
|
||||
Failing that, you can [download a release archive](http://releases.llvm.org/), extract it somewhere, and set
|
||||
[CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.0/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it.
|
||||
|
||||
### Build
|
||||
|
||||
Assuming this repository is at `./HIP`:
|
||||
|
||||
**Ubuntu 14.04**:
|
||||
```shell
|
||||
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
|
||||
```
|
||||
**Ubuntu 16.04**:
|
||||
```shell
|
||||
wget http://llvm.org/releases/3.8.0/clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-16.04.tar.xz
|
||||
tar xvfJ clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-16.04.tar.xz
|
||||
```
|
||||
mkdir build inst
|
||||
|
||||
- Enable build of hipify-clang and specify path to LLVM.
|
||||
|
||||
Note HIPIFY_CLANG_LLVM_DIR must be a full absolute path to the location extracted above. Here's an example assuming we extract the clang 3.8 package into ~/HIP/clang+llvm-3.8.0/
|
||||
```shell
|
||||
cd HIP
|
||||
mkdir build
|
||||
cd build
|
||||
cmake -DHIPIFY_CLANG_LLVM_DIR=~/HIP/clang+llvm-3.8.0/ -DCMAKE_BUILD_TYPE=Release ..
|
||||
make
|
||||
make install
|
||||
cmake \
|
||||
-DCMAKE_INSTALL_PREFIX=../inst \
|
||||
-DCMAKE_BUILD_TYPE=Release \
|
||||
-DBUILD_HIPIFY_CLANG=ON \
|
||||
../HIP
|
||||
|
||||
make -j install
|
||||
```
|
||||
|
||||
### Running and using hipify-clang
|
||||
The binary can then be found at `./inst/bin/hipify-clang`.
|
||||
|
||||
`hipify-clang` performs an initial compile of the CUDA source code into a "symbol tree", and thus needs access to the appropriate header files.
|
||||
### Test
|
||||
|
||||
In the case when `hipify-clang` doesn't find cuda headers, it reports various errors about unknown keywords (e.g. '\__global\__'), API function names (e.g. 'cudaMalloc'), syntax (e.g. 'foo<<<1,n>>>(...)'), etc.
|
||||
`hipify-clang` has unit tests using LLVM [`lit`](https://llvm.org/docs/CommandGuide/lit.html)/[`FileCheck`](https://llvm.org/docs/CommandGuide/FileCheck.html).
|
||||
|
||||
To install CUDA headers, download the "deb(network)" variant of the target installer.
|
||||
To run it:
|
||||
1. Ensure `lit` and `FileCheck` are installed - these are distributed with LLVM.
|
||||
2. Ensure `socat` is installed - your distro almost certainly has a package for this.
|
||||
3. Build with the `HIPIFY_CLANG_TESTS` option turned on.
|
||||
4. `make test-hipify`
|
||||
|
||||
**Ubuntu 14.04**:
|
||||
```shell
|
||||
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
|
||||
```
|
||||
**Ubuntu 16.04**:
|
||||
```shell
|
||||
wget http://archive.ubuntu.com/ubuntu/pool/multiverse/n/nvidia-cuda-toolkit/nvidia-cuda-toolkit_7.5.18-0ubuntu1_amd64.deb
|
||||
sudo dpkg -i nvidia-cuda-toolkit_7.5.18-0ubuntu1_amd64.deb
|
||||
sudo apt-get update && sudo apt-get install cuda-minimal-build-7-5 cuda-curand-dev-7-5
|
||||
```
|
||||
To set additional options like Language Selection (only "-x cuda" is supported), Preprocessor Definition (-D), Include Path (-I), etc., options delimiter "--" should be used before them, for instance:
|
||||
## Running and using hipify-clang
|
||||
|
||||
To process a file, `hipify-clang` needs access to the same headers that would be needed to compile it with clang.
|
||||
|
||||
For example:
|
||||
|
||||
```shell
|
||||
./hipify-clang -print-stats sort_kernel.cu -- -x cuda -I/srv/git/HIP/include -I/usr/local/cuda-7.5/include -DX=1
|
||||
hipify-clang square.cu -- \
|
||||
-x cuda \
|
||||
--cuda-path=/opt/cuda \
|
||||
--cuda-gpu-arch=sm_30 \
|
||||
-isystem /opt/cuda/samples/common/inc
|
||||
```
|
||||
|
||||
Delimiter "--" is used to separate hipify-clang options (before the delimiter) from clang options (after the delimiter). It is strongly recommended to always specify the delimiter, even if there are no clang specific options at all, in order to avoid possible errors regarding compilation database; in such case delimeter should be the last option in hipify-clang's command line.
|
||||
`hipify-clang` arguments are given first, followed by a separator, and then the arguments you'd pass to `clang` if you
|
||||
were compiling the input file. The [Clang manual for compiling CUDA](https://llvm.org/docs/CompileCudaWithLLVM.html#compiling-cuda-code)
|
||||
may be useful.
|
||||
|
||||
Option "-x cuda" is also worth specifying in order to convert source CUDA files with extensions other than standard extensions (*.cu, *.cuh).
|
||||
For a list of `hipify-clang` options, run `hipify-clang --help`.
|
||||
|
||||
## Disclaimer
|
||||
|
||||
|
||||
@@ -1,7 +1,5 @@
|
||||
#include "CUDA2HipMap.h"
|
||||
|
||||
const std::set<llvm::StringRef> CUDA_EXCLUDES{"CHECK_CUDA_ERROR", "CUDA_SAFE_CALL"};
|
||||
|
||||
/// Maps the names of CUDA types to the corresponding hip types.
|
||||
const std::map<llvm::StringRef, hipCounter> CUDA_TYPE_NAME_MAP{
|
||||
// Error codes and return types
|
||||
|
||||
@@ -4,20 +4,9 @@
|
||||
#include <set>
|
||||
#include <map>
|
||||
|
||||
#include "Types.h"
|
||||
#include "Statistics.h"
|
||||
|
||||
// TODO: This shouldn't really be here. More restructuring needed...
|
||||
struct hipCounter {
|
||||
llvm::StringRef hipName;
|
||||
ConvTypes countType;
|
||||
ApiTypes countApiType;
|
||||
int unsupported;
|
||||
};
|
||||
|
||||
#define HIP_UNSUPPORTED -1
|
||||
|
||||
/// Macros to ignore.
|
||||
extern const std::set<llvm::StringRef> CUDA_EXCLUDES;
|
||||
#define HIP_UNSUPPORTED true
|
||||
|
||||
/// Maps cuda header names to hip header names.
|
||||
extern const std::map<llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP;
|
||||
|
||||
Разница между файлами не показана из-за своего большого размера
Загрузить разницу
@@ -0,0 +1,43 @@
|
||||
#include "LLVMCompat.h"
|
||||
|
||||
namespace llcompat {
|
||||
|
||||
void PrintStackTraceOnErrorSignal() {
|
||||
// The signature of PrintStackTraceOnErrorSignal changed in llvm 3.9. We don't support
|
||||
// anything older than 3.8, so let's specifically detect the one old version we support.
|
||||
#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8)
|
||||
llvm::sys::PrintStackTraceOnErrorSignal();
|
||||
#else
|
||||
llvm::sys::PrintStackTraceOnErrorSignal(clang::StringRef());
|
||||
#endif
|
||||
}
|
||||
|
||||
ct::Replacements& getReplacements(ct::RefactoringTool& Tool, clang::StringRef file) {
|
||||
#if LLVM_VERSION_MAJOR > 3
|
||||
// getReplacements() now returns a map from filename to Replacements - so create an entry
|
||||
// for this source file and return a reference to it.
|
||||
return Tool.getReplacements()[file];
|
||||
#else
|
||||
return Tool.getReplacements();
|
||||
#endif
|
||||
}
|
||||
|
||||
void insertReplacement(ct::Replacements& replacements, const ct::Replacement& rep) {
|
||||
#if LLVM_VERSION_MAJOR > 3
|
||||
// New clang added error checking to Replacements, and *insists* that you explicitly check it.
|
||||
llvm::Error e = replacements.add(rep);
|
||||
#else
|
||||
// In older versions, it's literally an std::set<Replacement>
|
||||
replacements.insert(rep);
|
||||
#endif
|
||||
}
|
||||
|
||||
void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token *start, size_t len, bool DisableMacroExpansion) {
|
||||
#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8)
|
||||
_pp.EnterTokenStream(start, len, false, DisableMacroExpansion);
|
||||
#else
|
||||
_pp.EnterTokenStream(clang::ArrayRef<clang::Token>{start, len}, DisableMacroExpansion);
|
||||
#endif
|
||||
}
|
||||
|
||||
} // namespace llcompat
|
||||
@@ -0,0 +1,49 @@
|
||||
#pragma once
|
||||
|
||||
#include <clang/Tooling/Core/Replacement.h>
|
||||
#include <clang/Tooling/Refactoring.h>
|
||||
#include <llvm/Support/Signals.h>
|
||||
#include <clang/Lex/Token.h>
|
||||
#include <clang/Lex/Preprocessor.h>
|
||||
|
||||
namespace ct = clang::tooling;
|
||||
|
||||
// Things for papering over the differences between different LLVM versions.
|
||||
|
||||
namespace llcompat {
|
||||
|
||||
|
||||
/**
|
||||
* The getNumArgs function on macros was rather unhelpfully renamed in clang 4.0. Its semantics
|
||||
* remain unchanged, so let's be slightly ugly about it here. :D
|
||||
*/
|
||||
#if LLVM_VERSION_MAJOR > 4
|
||||
#define GET_NUM_ARGS() getNumParams()
|
||||
#else
|
||||
#define GET_NUM_ARGS() getNumArgs()
|
||||
#endif
|
||||
|
||||
void PrintStackTraceOnErrorSignal();
|
||||
|
||||
/**
|
||||
* Get the replacement map for a given filename in a RefactoringTool.
|
||||
*
|
||||
* Older LLVM versions don't actually support multiple filenames, so everything all gets
|
||||
* smushed together. It is the caller's responsibility to cope with this.
|
||||
*/
|
||||
ct::Replacements& getReplacements(ct::RefactoringTool& Tool, clang::StringRef file);
|
||||
|
||||
/**
|
||||
* Add a Replacement to a Replacements.
|
||||
*/
|
||||
void insertReplacement(ct::Replacements& replacements, const ct::Replacement& rep);
|
||||
|
||||
/**
|
||||
* Version-agnostic version of Preprocessor::EnterTokenStream().
|
||||
*/
|
||||
void EnterPreprocessorTokenStream(clang::Preprocessor& _pp,
|
||||
const clang::Token *start,
|
||||
size_t len,
|
||||
bool DisableMacroExpansion);
|
||||
|
||||
} // namespace llcompat
|
||||
@@ -0,0 +1,227 @@
|
||||
#include "Statistics.h"
|
||||
#include <assert.h>
|
||||
#include <sstream>
|
||||
#include <iomanip>
|
||||
|
||||
|
||||
const char *counterNames[NUM_CONV_TYPES] = {
|
||||
"version", "init", "device", "mem", "kern", "coord_func", "math_func",
|
||||
"special_func", "stream", "event", "occupancy", "ctx", "peer", "module",
|
||||
"cache", "exec", "err", "def", "tex", "gl", "graphics",
|
||||
"surface", "jit", "d3d9", "d3d10", "d3d11", "vdpau", "egl",
|
||||
"thread", "other", "include", "include_cuda_main_header", "type", "literal",
|
||||
"numeric_literal"
|
||||
};
|
||||
|
||||
const char *apiNames[NUM_API_TYPES] = {
|
||||
"CUDA Driver API", "CUDA RT API", "CUBLAS API"
|
||||
};
|
||||
|
||||
namespace {
|
||||
|
||||
template<typename ST, typename ST2>
|
||||
void conditionalPrint(ST *stream1,
|
||||
ST2* stream2,
|
||||
const std::string& s1,
|
||||
const std::string& s2) {
|
||||
if (stream1) {
|
||||
*stream1 << s1;
|
||||
}
|
||||
|
||||
if (stream2) {
|
||||
*stream2 << s2;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Print a named stat value to both the terminal and the CSV file.
|
||||
*/
|
||||
template<typename T>
|
||||
void printStat(std::ostream *csv, llvm::raw_ostream* printOut, const std::string &name, T value) {
|
||||
if (printOut) {
|
||||
*printOut << " " << name << ": " << value << "\n";
|
||||
}
|
||||
|
||||
if (csv) {
|
||||
*csv << name << ";" << value << "\n";
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
} // Anonymous namespace
|
||||
|
||||
void StatCounter::incrementCounter(const hipCounter& counter, std::string name) {
|
||||
counters[name]++;
|
||||
apiCounters[(int) counter.countApiType]++;
|
||||
convTypeCounters[(int) counter.countType]++;
|
||||
}
|
||||
|
||||
void StatCounter::add(const StatCounter& other) {
|
||||
for (const auto& p : other.counters) {
|
||||
counters[p.first] += p.second;
|
||||
}
|
||||
|
||||
for (int i = 0; i < NUM_API_TYPES; i++) {
|
||||
apiCounters[i] += other.apiCounters[i];
|
||||
}
|
||||
|
||||
for (int i = 0; i < NUM_CONV_TYPES; i++) {
|
||||
convTypeCounters[i] += other.convTypeCounters[i];
|
||||
}
|
||||
}
|
||||
|
||||
int StatCounter::getConvSum() {
|
||||
int acc = 0;
|
||||
for (const int& i : convTypeCounters) {
|
||||
acc += i;
|
||||
}
|
||||
|
||||
return acc;
|
||||
}
|
||||
|
||||
void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, std::string prefix) {
|
||||
conditionalPrint(csv, printOut, "\nCUDA ref type;Count\n", "[HIPIFY] info: " + prefix + " refs by type:\n");
|
||||
for (int i = 0; i < NUM_CONV_TYPES; i++) {
|
||||
if (convTypeCounters[i] > 0) {
|
||||
printStat(csv, printOut, counterNames[i], convTypeCounters[i]);
|
||||
}
|
||||
}
|
||||
|
||||
conditionalPrint(csv, printOut, "\nCUDA API;Count\n", "[HIPIFY] info: " + prefix + " refs by API:\n");
|
||||
for (int i = 0; i < NUM_API_TYPES; i++) {
|
||||
printStat(csv, printOut, apiNames[i], apiCounters[i]);
|
||||
}
|
||||
|
||||
conditionalPrint(csv, printOut, "\nCUDA ref name;Count\n", "[HIPIFY] info: " + prefix + " refs by names:\n");
|
||||
for (const auto &it : counters) {
|
||||
printStat(csv, printOut, it.first, it.second);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
Statistics::Statistics(std::string name): fileName(name) {
|
||||
// Compute the total bytes/lines in the input file.
|
||||
std::ifstream src_file(name, std::ios::binary | std::ios::ate);
|
||||
src_file.clear();
|
||||
src_file.seekg(0);
|
||||
totalLines = (int) std::count(std::istreambuf_iterator<char>(src_file), std::istreambuf_iterator<char>(), '\n');
|
||||
totalBytes = (int) src_file.tellg();
|
||||
|
||||
// Mark the start time...
|
||||
startTime = chr::steady_clock::now();
|
||||
};
|
||||
|
||||
|
||||
///////// Counter update routines //////////
|
||||
|
||||
void Statistics::incrementCounter(const hipCounter &counter, std::string name) {
|
||||
if (counter.unsupported) {
|
||||
unsupported.incrementCounter(counter, name);
|
||||
} else {
|
||||
supported.incrementCounter(counter, name);
|
||||
}
|
||||
}
|
||||
|
||||
void Statistics::add(const Statistics &other) {
|
||||
supported.add(other.supported);
|
||||
unsupported.add(other.unsupported);
|
||||
totalBytes += other.totalBytes;
|
||||
totalLines += other.totalLines;
|
||||
touchedBytes += other.touchedBytes;
|
||||
}
|
||||
|
||||
void Statistics::lineTouched(int lineNumber) {
|
||||
touchedLines.insert(lineNumber);
|
||||
}
|
||||
void Statistics::bytesChanged(int bytes) {
|
||||
touchedBytes += bytes;
|
||||
}
|
||||
void Statistics::markCompletion() {
|
||||
completionTime = chr::steady_clock::now();
|
||||
}
|
||||
|
||||
|
||||
///////// Output functions //////////
|
||||
|
||||
void Statistics::print(std::ostream* csv, llvm::raw_ostream* printOut, bool skipHeader) {
|
||||
if (!skipHeader) {
|
||||
std::string str = "file \'" + fileName + "\' statistics:\n";
|
||||
conditionalPrint(csv, printOut, "\n" + str, "\n[HIPIFY] info: " + str);
|
||||
}
|
||||
|
||||
size_t changedLines = touchedLines.size();
|
||||
|
||||
// Total number of (un)supported refs that were converted.
|
||||
int supportedSum = supported.getConvSum();
|
||||
int unsupportedSum = unsupported.getConvSum();
|
||||
|
||||
printStat(csv, printOut, "CONVERTED refs count", supportedSum);
|
||||
printStat(csv, printOut, "UNCONVERTED refs count", unsupportedSum);
|
||||
printStat(csv, printOut, "CONVERSION %", 100 - std::lround(double(unsupportedSum * 100) / double(supportedSum + unsupportedSum)));
|
||||
printStat(csv, printOut, "REPLACED bytes", touchedBytes);
|
||||
printStat(csv, printOut, "TOTAL bytes", totalBytes);
|
||||
printStat(csv, printOut, "CHANGED lines of code", changedLines);
|
||||
printStat(csv, printOut, "TOTAL lines of code", totalLines);
|
||||
|
||||
if (totalBytes > 0) {
|
||||
printStat(csv, printOut, "CODE CHANGED (in bytes) %", std::lround(double(touchedBytes * 100) / double(totalBytes)));
|
||||
}
|
||||
|
||||
if (totalLines > 0) {
|
||||
printStat(csv, printOut, "CODE CHANGED (in lines) %", std::lround(double(changedLines * 100) / double(totalLines)));
|
||||
}
|
||||
|
||||
typedef std::chrono::duration<double, std::milli> duration;
|
||||
duration elapsed = completionTime - startTime;
|
||||
std::stringstream stream;
|
||||
stream << std::fixed << std::setprecision(2) << elapsed.count() / 1000;
|
||||
printStat(csv, printOut, "TIME ELAPSED s", stream.str());
|
||||
|
||||
supported.print(csv, printOut, "CONVERTED");
|
||||
unsupported.print(csv, printOut, "UNCONVERTED");
|
||||
}
|
||||
|
||||
void Statistics::printAggregate(std::ostream *csv, llvm::raw_ostream* printOut) {
|
||||
Statistics globalStats = getAggregate();
|
||||
|
||||
conditionalPrint(csv, printOut, "\nTOTAL statistics:\n", "\n[HIPIFY] info: TOTAL statistics:\n");
|
||||
|
||||
// A file is considered "converted" if we made any changes to it.
|
||||
int convertedFiles = 0;
|
||||
for (const auto& p : stats) {
|
||||
if (!p.second.touchedLines.empty()) {
|
||||
convertedFiles++;
|
||||
}
|
||||
}
|
||||
|
||||
printStat(csv, printOut, "CONVERTED files", convertedFiles);
|
||||
printStat(csv, printOut, "PROCESSED files", stats.size());
|
||||
|
||||
globalStats.print(csv, printOut);
|
||||
}
|
||||
|
||||
//// Static state management ////
|
||||
|
||||
Statistics Statistics::getAggregate() {
|
||||
Statistics globalStats("global");
|
||||
|
||||
for (const auto& p : stats) {
|
||||
globalStats.add(p.second);
|
||||
}
|
||||
|
||||
return globalStats;
|
||||
}
|
||||
|
||||
Statistics& Statistics::current() {
|
||||
assert(Statistics::currentStatistics);
|
||||
return *Statistics::currentStatistics;
|
||||
}
|
||||
|
||||
void Statistics::setActive(std::string name) {
|
||||
stats.emplace(std::make_pair(name, Statistics{name}));
|
||||
Statistics::currentStatistics = &stats.at(name);
|
||||
}
|
||||
|
||||
std::map<std::string, Statistics> Statistics::stats = {};
|
||||
Statistics* Statistics::currentStatistics = nullptr;
|
||||
@@ -0,0 +1,174 @@
|
||||
#pragma once
|
||||
|
||||
#include <chrono>
|
||||
#include <string>
|
||||
#include <fstream>
|
||||
#include <llvm/ADT/StringRef.h>
|
||||
#include <map>
|
||||
#include <set>
|
||||
#include <llvm/Support/raw_ostream.h>
|
||||
|
||||
namespace chr = std::chrono;
|
||||
|
||||
enum ConvTypes {
|
||||
CONV_VERSION = 0,
|
||||
CONV_INIT,
|
||||
CONV_DEVICE,
|
||||
CONV_MEM,
|
||||
CONV_KERN,
|
||||
CONV_COORD_FUNC,
|
||||
CONV_MATH_FUNC,
|
||||
CONV_SPECIAL_FUNC,
|
||||
CONV_STREAM,
|
||||
CONV_EVENT,
|
||||
CONV_OCCUPANCY,
|
||||
CONV_CONTEXT,
|
||||
CONV_PEER,
|
||||
CONV_MODULE,
|
||||
CONV_CACHE,
|
||||
CONV_EXEC,
|
||||
CONV_ERROR,
|
||||
CONV_DEF,
|
||||
CONV_TEX,
|
||||
CONV_GL,
|
||||
CONV_GRAPHICS,
|
||||
CONV_SURFACE,
|
||||
CONV_JIT,
|
||||
CONV_D3D9,
|
||||
CONV_D3D10,
|
||||
CONV_D3D11,
|
||||
CONV_VDPAU,
|
||||
CONV_EGL,
|
||||
CONV_THREAD,
|
||||
CONV_OTHER,
|
||||
CONV_INCLUDE,
|
||||
CONV_INCLUDE_CUDA_MAIN_H,
|
||||
CONV_TYPE,
|
||||
CONV_LITERAL,
|
||||
CONV_NUMERIC_LITERAL,
|
||||
CONV_LAST
|
||||
};
|
||||
constexpr int NUM_CONV_TYPES = (int) ConvTypes::CONV_LAST;
|
||||
|
||||
enum ApiTypes {
|
||||
API_DRIVER = 0,
|
||||
API_RUNTIME,
|
||||
API_BLAS,
|
||||
API_LAST
|
||||
};
|
||||
constexpr int NUM_API_TYPES = (int) ApiTypes::API_LAST;
|
||||
|
||||
// The names of various fields in in the statistics reports.
|
||||
extern const char *counterNames[NUM_CONV_TYPES];
|
||||
extern const char *apiNames[NUM_API_TYPES];
|
||||
|
||||
|
||||
struct hipCounter {
|
||||
llvm::StringRef hipName;
|
||||
ConvTypes countType;
|
||||
ApiTypes countApiType;
|
||||
bool unsupported;
|
||||
};
|
||||
|
||||
|
||||
/**
|
||||
* Tracks a set of named counters, as well as counters for each of the type enums defined above.
|
||||
*/
|
||||
class StatCounter {
|
||||
private:
|
||||
// Each thing we track is either "supported" or "unsupported"...
|
||||
std::map<std::string, int> counters;
|
||||
|
||||
int apiCounters[NUM_API_TYPES] = {};
|
||||
int convTypeCounters[NUM_CONV_TYPES] = {};
|
||||
|
||||
public:
|
||||
void incrementCounter(const hipCounter& counter, std::string name);
|
||||
|
||||
/**
|
||||
* Add the counters from `other` onto the counters of this object.
|
||||
*/
|
||||
void add(const StatCounter& other);
|
||||
|
||||
int getConvSum();
|
||||
|
||||
void print(std::ostream* csv, llvm::raw_ostream* printOut, std::string prefix);
|
||||
};
|
||||
|
||||
/**
|
||||
* Tracks the statistics for a single input file.
|
||||
*/
|
||||
class Statistics {
|
||||
StatCounter supported;
|
||||
StatCounter unsupported;
|
||||
|
||||
std::string fileName;
|
||||
|
||||
std::set<int> touchedLines = {};
|
||||
int touchedBytes = 0;
|
||||
|
||||
int totalLines = 0;
|
||||
int totalBytes = 0;
|
||||
|
||||
chr::steady_clock::time_point startTime;
|
||||
chr::steady_clock::time_point completionTime;
|
||||
|
||||
public:
|
||||
Statistics(std::string name);
|
||||
|
||||
void incrementCounter(const hipCounter &counter, std::string name);
|
||||
|
||||
/**
|
||||
* Add the counters from `other` onto the counters of this object.
|
||||
*/
|
||||
void add(const Statistics &other);
|
||||
|
||||
void lineTouched(int lineNumber);
|
||||
void bytesChanged(int bytes);
|
||||
|
||||
/**
|
||||
* Set the completion timestamp to now.
|
||||
*/
|
||||
void markCompletion();
|
||||
|
||||
/////// Output functions ///////
|
||||
|
||||
public:
|
||||
/**
|
||||
* Pretty-print the statistics stored in this object.
|
||||
*
|
||||
* @param csv Pointer to an output stream for the CSV to write. If null, no CSV is written
|
||||
* @param printOut Pointer to an output stream to print human-readable textual stats to. If null, no
|
||||
* such stats are produced.
|
||||
*/
|
||||
void print(std::ostream* csv, llvm::raw_ostream* printOut, bool skipHeader = false);
|
||||
|
||||
/// Print aggregated statistics for all registered counters.
|
||||
static void printAggregate(std::ostream *csv, llvm::raw_ostream* printOut);
|
||||
|
||||
/////// Static nonsense ///////
|
||||
|
||||
// The Statistics for each input file.
|
||||
static std::map<std::string, Statistics> stats;
|
||||
|
||||
// The Statistics objects for the currently-being-processed input file.
|
||||
static Statistics* currentStatistics;
|
||||
|
||||
/**
|
||||
* Aggregate statistics over all entries in `stats` and return the resulting Statistics object.
|
||||
*/
|
||||
static Statistics getAggregate();
|
||||
|
||||
/**
|
||||
* Convenient global entry point for updating the "active" Statistics. Since we operate single-threadedly
|
||||
* processing one file at a time, this allows us to simply expose the stats for the current file globally,
|
||||
* simplifying things.
|
||||
*/
|
||||
static Statistics& current();
|
||||
|
||||
/**
|
||||
* Set the active Statistics object to the named one, creating it if necessary, and write the completion
|
||||
* timestamp into the currently active one.
|
||||
*/
|
||||
static void setActive(std::string name);
|
||||
};
|
||||
@@ -0,0 +1,17 @@
|
||||
#include "StringUtils.h"
|
||||
|
||||
llvm::StringRef unquoteStr(llvm::StringRef s) {
|
||||
if (s.size() > 1 && s.front() == '"' && s.back() == '"') {
|
||||
return s.substr(1, s.size() - 2);
|
||||
}
|
||||
|
||||
return s;
|
||||
}
|
||||
|
||||
void removePrefixIfPresent(std::string &s, std::string prefix) {
|
||||
if (s.find(prefix) != 0) {
|
||||
return;
|
||||
}
|
||||
|
||||
s.erase(0, prefix.size());
|
||||
}
|
||||
@@ -0,0 +1,14 @@
|
||||
#pragma once
|
||||
|
||||
#include <string>
|
||||
#include "llvm/ADT/StringRef.h"
|
||||
|
||||
/**
|
||||
* Remove double-quotes from the start/end of a string, if present.
|
||||
*/
|
||||
llvm::StringRef unquoteStr(llvm::StringRef s);
|
||||
|
||||
/**
|
||||
* If `s` starts with `prefix`, remove it. Otherwise, does nothing.
|
||||
*/
|
||||
void removePrefixIfPresent(std::string &s, std::string prefix);
|
||||
@@ -1,47 +0,0 @@
|
||||
#pragma once
|
||||
|
||||
enum ConvTypes {
|
||||
CONV_VERSION = 0,
|
||||
CONV_INIT,
|
||||
CONV_DEVICE,
|
||||
CONV_MEM,
|
||||
CONV_KERN,
|
||||
CONV_COORD_FUNC,
|
||||
CONV_MATH_FUNC,
|
||||
CONV_SPECIAL_FUNC,
|
||||
CONV_STREAM,
|
||||
CONV_EVENT,
|
||||
CONV_OCCUPANCY,
|
||||
CONV_CONTEXT,
|
||||
CONV_PEER,
|
||||
CONV_MODULE,
|
||||
CONV_CACHE,
|
||||
CONV_EXEC,
|
||||
CONV_ERROR,
|
||||
CONV_DEF,
|
||||
CONV_TEX,
|
||||
CONV_GL,
|
||||
CONV_GRAPHICS,
|
||||
CONV_SURFACE,
|
||||
CONV_JIT,
|
||||
CONV_D3D9,
|
||||
CONV_D3D10,
|
||||
CONV_D3D11,
|
||||
CONV_VDPAU,
|
||||
CONV_EGL,
|
||||
CONV_THREAD,
|
||||
CONV_OTHER,
|
||||
CONV_INCLUDE,
|
||||
CONV_INCLUDE_CUDA_MAIN_H,
|
||||
CONV_TYPE,
|
||||
CONV_LITERAL,
|
||||
CONV_NUMERIC_LITERAL,
|
||||
CONV_LAST
|
||||
};
|
||||
|
||||
enum ApiTypes {
|
||||
API_DRIVER = 0,
|
||||
API_RUNTIME,
|
||||
API_BLAS,
|
||||
API_LAST
|
||||
};
|
||||
@@ -2,11 +2,23 @@
|
||||
|
||||
#include <iostream>
|
||||
|
||||
__global__ void axpy(float a, float* x, float* y) {
|
||||
|
||||
#define TOKEN_PASTE(X, Y) X ## Y
|
||||
#define ARG_LIST_AS_MACRO a, device_x, device_y
|
||||
#define KERNEL_CALL_AS_MACRO axpy<float><<<1, kDataLen>>>
|
||||
#define KERNEL_NAME_MACRO axpy<float>
|
||||
|
||||
// CHECK: #define COMPLETE_LAUNCH hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y)
|
||||
#define COMPLETE_LAUNCH axpy<<<1, kDataLen>>>(a, device_x, device_y)
|
||||
|
||||
|
||||
template<typename T>
|
||||
__global__ void axpy(T a, T *x, T *y) {
|
||||
// CHECK: y[hipThreadIdx_x] = a * x[hipThreadIdx_x];
|
||||
y[threadIdx.x] = a * x[threadIdx.x];
|
||||
}
|
||||
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
const int kDataLen = 4;
|
||||
|
||||
@@ -27,10 +39,29 @@ int main(int argc, char* argv[]) {
|
||||
// CHECK: hipMemcpy(device_x, host_x, kDataLen * sizeof(float), hipMemcpyHostToDevice);
|
||||
cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), cudaMemcpyHostToDevice);
|
||||
|
||||
// Launch the kernel.
|
||||
// Launch the kernel in numerous different strange ways to exercise the prerocessor.
|
||||
// CHECK: hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
|
||||
axpy<<<1, kDataLen>>>(a, device_x, device_y);
|
||||
|
||||
// CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
|
||||
axpy<float><<<1, kDataLen>>>(a, device_x, device_y);
|
||||
|
||||
// CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, a, TOKEN_PASTE(device, _x), device_y);
|
||||
axpy<float><<<1, kDataLen>>>(a, TOKEN_PASTE(device, _x), device_y);
|
||||
|
||||
// CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
|
||||
axpy<float><<<1, kDataLen>>>(ARG_LIST_AS_MACRO);
|
||||
|
||||
// CHECK: hipLaunchKernelGGL(KERNEL_NAME_MACRO, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
|
||||
KERNEL_NAME_MACRO<<<1, kDataLen>>>(ARG_LIST_AS_MACRO);
|
||||
|
||||
// CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
|
||||
KERNEL_CALL_AS_MACRO(ARG_LIST_AS_MACRO);
|
||||
|
||||
// CHECK: COMPLETE_LAUNCH;
|
||||
COMPLETE_LAUNCH;
|
||||
|
||||
|
||||
// Copy output data to host.
|
||||
// CHECK: hipDeviceSynchronize();
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
Ссылка в новой задаче
Block a user