AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Merge 'master' into 'amd-master'

Просмотр исходного кода 
Change-Id: Ieef32efec384d27df2a0fbe262be139b5e734b43
Этот коммит содержится в:
Jenkins
2019-12-10 11:06:12 +00:00
родитель 15aa9de37f 280a8ebadf
Коммит b79041c97e
3 изменённых файлов: 221 добавлений и 68 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
hipamd/hipify-clang/src/main.cpp
+122 -51
Просмотреть файл
@@ -34,6 +34,14 @@ THE SOFTWARE.
#include "ArgParse.h"
#include "StringUtils.h"
#include "llvm/Support/Debug.h"
#include "clang/Basic/Diagnostic.h"
#include "clang/Basic/DiagnosticIDs.h"
#include "clang/Basic/DiagnosticOptions.h"
#include "clang/Driver/Driver.h"
#include "clang/Driver/Compilation.h"
#include "clang/Driver/Tool.h"
#include "clang/Frontend/TextDiagnosticPrinter.h"
#if LLVM_VERSION_MAJOR < 8
#include "llvm/Support/Path.h"
#endif
@@ -42,6 +50,114 @@ constexpr auto DEBUG_TYPE = "cuda2hip";
namespace ct = clang::tooling;
void cleanupHipifyOptions(std::vector<const char*> &args) {
std::vector<std::string> hipifyOptions = {"-perl", "-python", "-roc", "-inplace",
"-no-backup", "-no-output", "-print-stats",
"-print-stats-csv", "-examine", "-save-temps",
"-skip-excluded-preprocessor-conditional-blocks"};
for (const auto &a : hipifyOptions) {
args.erase(std::remove(args.begin(), args.end(), a), args.end());
args.erase(std::remove(args.begin(), args.end(), "-" + a), args.end());
}
std::vector<std::string> hipifyDirOptions = {"-o-dir", "-o-hipify-perl-dir", "-o-stats",
"-o-python-map-dir", "-temp-dir"};
for (const auto &a : hipifyDirOptions) {
// remove all pairs of arguments "-option value"
auto it = args.erase(std::remove(args.begin(), args.end(), a), args.end());
if (it != args.end()) {
args.erase(it);
}
// remove all pairs of arguments "--option value"
it = args.erase(std::remove(args.begin(), args.end(), "-" + a), args.end());
if (it != args.end()) {
args.erase(it);
}
// remove all "-option=value" and "--option=value"
args.erase(
std::remove_if(args.begin(), args.end(),
[a](const std::string &s) { return s.find(a + "=") == 0 || s.find("-" + a + "=") == 0; }
),
args.end()
);
}
}
void sortInputFiles(int argc, const char **argv, std::vector<std::string> &files) {
if (files.size() < 2) return;
IntrusiveRefCntPtr<clang::DiagnosticOptions> diagOpts(new clang::DiagnosticOptions());
clang::TextDiagnosticPrinter diagClient(llvm::errs(), &*diagOpts);
clang::DiagnosticsEngine Diagnostics(IntrusiveRefCntPtr<clang::DiagnosticIDs>(new clang::DiagnosticIDs()), &*diagOpts, &diagClient, false);
std::unique_ptr<clang::driver::Driver> driver(new clang::driver::Driver("", "nvptx64-nvidia-cuda", Diagnostics));
std::vector<const char*> Args(argv, argv + argc);
cleanupHipifyOptions(Args);
std::unique_ptr<clang::driver::Compilation> C(driver->BuildCompilation(Args));
std::vector<std::string> sortedFiles;
for (const auto &J : C->getJobs()) {
if (std::string(J.getCreator().getName()) != "clang") continue;
const auto &JA = J.getArguments();
for (size_t i = 0; i < JA.size(); ++i) {
const auto &A = std::string(JA[i]);
if (std::find(files.begin(), files.end(), A) != files.end() &&
i > 0 && std::string(JA[i - 1]) == "-main-file-name") {
sortedFiles.push_back(A);
}
}
}
if (sortedFiles.empty()) return;
std::reverse(sortedFiles.begin(), sortedFiles.end());
files.assign(sortedFiles.begin(), sortedFiles.end());
}
void appendArgumentsAdjusters(ct::RefactoringTool &Tool, const std::string &sSourceAbsPath) {
if (!IncludeDirs.empty()) {
for (std::string s : IncludeDirs) {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-I", ct::ArgumentInsertPosition::BEGIN));
}
}
if (!MacroNames.empty()) {
for (std::string s : MacroNames) {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-D", ct::ArgumentInsertPosition::BEGIN));
}
}
// Includes for clang's CUDA wrappers for using by packaged hipify-clang
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include", ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-isystem", ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include/cuda_wrappers", ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-isystem", ct::ArgumentInsertPosition::BEGIN));
// Ensure at least c++11 is used.
std::string stdCpp = "-std=c++11";
#if defined(_MSC_VER)
stdCpp = "-std=c++14";
#endif
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(stdCpp.c_str(), ct::ArgumentInsertPosition::BEGIN));
std::string sInclude = "-I" + sys::path::parent_path(sSourceAbsPath).str();
#if defined(HIPIFY_CLANG_RES)
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES, ct::ArgumentInsertPosition::BEGIN));
#endif
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sInclude.c_str(), ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-fno-delayed-template-parsing", ct::ArgumentInsertPosition::BEGIN));
if (llcompat::pragma_once_outside_header()) {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-Wno-pragma-once-outside-header", ct::ArgumentInsertPosition::BEGIN));
}
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN));
if (!CudaGpuArch.empty()) {
std::string sCudaGpuArch = "--cuda-gpu-arch=" + CudaGpuArch;
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaGpuArch.c_str(), ct::ArgumentInsertPosition::BEGIN));
}
if (!CudaPath.empty()) {
std::string sCudaPath = "--cuda-path=" + CudaPath;
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaPath.c_str(), ct::ArgumentInsertPosition::BEGIN));
}
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("cuda", ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-x", ct::ArgumentInsertPosition::BEGIN));
if (Verbose) {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-v", ct::ArgumentInsertPosition::END));
}
Tool.appendArgumentsAdjuster(ct::getClangSyntaxOnlyAdjuster());
}
int main(int argc, const char **argv) {
std::vector<const char*> new_argv(argv, argv + argc);
if (std::find(new_argv.begin(), new_argv.end(), std::string("--")) == new_argv.end()) {
@@ -119,7 +235,7 @@ int main(int argc, const char **argv) {
}
// Arguments for the Statistics print routines.
std::unique_ptr<std::ostream> csv = nullptr;
llvm::raw_ostream* statPrint = nullptr;
llvm::raw_ostream *statPrint = nullptr;
bool create_csv = false;
if (!OutputStatsFilename.empty()) {
PrintStatsCSV = true;
@@ -139,7 +255,8 @@ int main(int argc, const char **argv) {
if (PrintStats) {
statPrint = &llvm::errs();
}
for (const auto & src : fileSources) {
sortInputFiles(argc, argv, fileSources);
for (const auto &src : fileSources) {
// Create a copy of the file to work on. When we're done, we'll move this onto the
// output (which may mean overwriting the input, if we're in-place).
// Should we fail for some reason, we'll just leak this file and not corrupt the input.
@@ -192,56 +309,10 @@ int main(int argc, const char **argv) {
// because that'll break relative includes, and we don't want to overwrite the input file.
// So what we do is operate on a copy, which we then move to the output.
ct::RefactoringTool Tool(OptionsParser.getCompilations(), std::string(tmpFile.c_str()));
ct::Replacements& replacementsToUse = llcompat::getReplacements(Tool, tmpFile.c_str());
ct::Replacements &replacementsToUse = llcompat::getReplacements(Tool, tmpFile.c_str());
ReplacementsFrontendActionFactory<HipifyAction> actionFactory(&replacementsToUse);
if (!IncludeDirs.empty()) {
for (std::string s : IncludeDirs) {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-I", ct::ArgumentInsertPosition::BEGIN));
}
}
if (!MacroNames.empty()) {
for (std::string s : MacroNames) {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-D", ct::ArgumentInsertPosition::BEGIN));
}
}
// Includes for clang's CUDA wrappers for using by packaged hipify-clang
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include", ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-isystem", ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include/cuda_wrappers", ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-isystem", ct::ArgumentInsertPosition::BEGIN));
// Ensure at least c++11 is used.
std::string stdCpp = "-std=c++11";
#if defined(_MSC_VER)
stdCpp = "-std=c++14";
#endif
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(stdCpp.c_str(), ct::ArgumentInsertPosition::BEGIN));
std::string sInclude = "-I" + sys::path::parent_path(sSourceAbsPath).str();
#if defined(HIPIFY_CLANG_RES)
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES, ct::ArgumentInsertPosition::BEGIN));
#endif
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sInclude.c_str(), ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-fno-delayed-template-parsing", ct::ArgumentInsertPosition::BEGIN));
if (llcompat::pragma_once_outside_header()) {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-Wno-pragma-once-outside-header", ct::ArgumentInsertPosition::BEGIN));
}
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN));
if (!CudaGpuArch.empty()) {
std::string sCudaGpuArch = "--cuda-gpu-arch=" + CudaGpuArch;
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaGpuArch.c_str(), ct::ArgumentInsertPosition::BEGIN));
}
if (!CudaPath.empty()) {
std::string sCudaPath = "--cuda-path=" + CudaPath;
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaPath.c_str(), ct::ArgumentInsertPosition::BEGIN));
}
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("cuda", ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-x", ct::ArgumentInsertPosition::BEGIN));
if (Verbose) {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-v", ct::ArgumentInsertPosition::END));
}
Tool.appendArgumentsAdjuster(ct::getClangSyntaxOnlyAdjuster());
Statistics& currentStat = Statistics::current();
appendArgumentsAdjusters(Tool, sSourceAbsPath);
Statistics &currentStat = Statistics::current();
// Hipify _all_ the things!
if (Tool.runAndSave(&actionFactory)) {
currentStat.hasErrors = true;
hipamd/include/hip/hcc_detail/hip_runtime_api.h
+1 -1
Просмотреть файл
@@ -2956,7 +2956,7 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk, unsigned int flags);
#if defined(__clang__) && defined(__HIP__)
#if __HIP_VDI__
/**
* @brief Launches kernels on multiple devices and guarantees all specified kernels are dispatched
* on respective streams before enqueuing any other work on the specified streams from any other threads
hipamd/include/hip/hcc_detail/hip_vector_types.h
+98 -16
Просмотреть файл
@@ -34,15 +34,31 @@ THE SOFTWARE.
#include "hip/hcc_detail/host_defines.h"
#if !defined(_MSC_VER) || __clang__
#if defined(__clang__)
#define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n)))
#if !defined(_MSC_VER)
#if __has_attribute(ext_vector_type)
#define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n)))
#else
#define __NATIVE_VECTOR__(n, ...) [n]
#endif
#if defined(__cplusplus) && defined(__clang__)
#if defined(__cplusplus)
#include <array>
#include <iosfwd>
#include <type_traits>
namespace hip_impl {
template<typename, typename, unsigned int> struct Scalar_accessor;
} // Namespace hip_impl.
namespace std {
template<typename T, typename U, unsigned int n>
struct is_integral<hip_impl::Scalar_accessor<T, U, n>>
: is_integral<T> {};
template<typename T, typename U, unsigned int n>
struct is_floating_point<hip_impl::Scalar_accessor<T, U, n>>
: is_floating_point<T> {};
} // Namespace std.
namespace hip_impl {
template<typename T, typename Vector, unsigned int idx>
struct Scalar_accessor {
@@ -93,6 +109,27 @@ THE SOFTWARE.
operator T() const noexcept { return data[idx]; }
__host__ __device__
operator T() const volatile noexcept { return data[idx]; }
// The conversions to enum are fairly ghastly, but unfortunately used in
// some pre-existing, difficult to modify, code.
template<
typename U,
typename std::enable_if<
!std::is_same<U, T>{} &&
std::is_enum<U>{} &&
std::is_convertible<
T, typename std::enable_if<std::is_enum<U>::value, std::underlying_type<U>>::type::type>{}>::type* = nullptr>
__host__ __device__
operator U() const noexcept { return static_cast<U>(data[idx]); }
template<
typename U,
typename std::enable_if<
!std::is_same<U, T>{} &&
std::is_enum<U>{} &&
std::is_convertible<
T, typename std::enable_if<std::is_enum<U>::value, std::underlying_type<U>>::type::type>{}>::type* = nullptr>
__host__ __device__
operator U() const volatile noexcept { return static_cast<U>(data[idx]); }
__host__ __device__
operator T&() noexcept {
@@ -107,7 +144,13 @@ THE SOFTWARE.
__host__ __device__
Address operator&() const noexcept { return Address{this}; }
__host__ __device__
Scalar_accessor& operator=(const Scalar_accessor& x) noexcept {
data[idx] = x.data[idx];
return *this;
}
__host__ __device__
Scalar_accessor& operator=(T x) noexcept {
data[idx] = x;
@@ -179,7 +222,7 @@ THE SOFTWARE.
typename std::enable_if<
std::is_convertible<U, T>{}>::type* = nullptr>
__host__ __device__
Scalar_accessor& operator/=(T x) noexcept {
Scalar_accessor& operator/=(U x) noexcept {
data[idx] /= x;
return *this;
}
@@ -245,7 +288,7 @@ THE SOFTWARE.
template<typename T>
struct HIP_vector_base<T, 1> {
typedef T Native_vec_ __NATIVE_VECTOR__(1, T);
using Native_vec_ = T __NATIVE_VECTOR__(1, T);
union {
Native_vec_ data;
@@ -253,11 +296,22 @@ THE SOFTWARE.
};
using value_type = T;
__host__ __device__
HIP_vector_base& operator=(const HIP_vector_base& x) noexcept {
#if __has_attribute(ext_vector_type)
data = x.data;
#else
data[0] = x.data[0];
#endif
return *this;
}
};
template<typename T>
struct HIP_vector_base<T, 2> {
typedef T Native_vec_ __NATIVE_VECTOR__(2, T);
using Native_vec_ = T __NATIVE_VECTOR__(2, T);
union {
Native_vec_ data;
@@ -266,6 +320,18 @@ THE SOFTWARE.
};
using value_type = T;
__host__ __device__
HIP_vector_base& operator=(const HIP_vector_base& x) noexcept {
#if __has_attribute(ext_vector_type)
data = x.data;
#else
data[0] = x.data[0];
data[1] = x.data[1];
#endif
return *this;
}
};
template<typename T>
@@ -404,15 +470,11 @@ THE SOFTWARE.
return *this;
}
using Vec3_cmp = int __NATIVE_VECTOR__(3, int);
using Vec3_cmp = int __attribute__((vector_size(4 * sizeof(int))));
__host__ __device__
Vec3_cmp operator==(const Native_vec_& x) const noexcept
{
Vec3_cmp r;
r[0] = d[0] == x.d[0];
r[1] = d[1] == x.d[1];
r[2] = d[2] == x.d[2];
return r;
return Vec3_cmp{d[0] == x.d[0], d[1] == x.d[1], d[2] == x.d[2]};
}
};
@@ -430,7 +492,7 @@ THE SOFTWARE.
template<typename T>
struct HIP_vector_base<T, 4> {
typedef T Native_vec_ __NATIVE_VECTOR__(4, T);
using Native_vec_ = T __NATIVE_VECTOR__(4, T);
union {
Native_vec_ data;
@@ -441,6 +503,20 @@ THE SOFTWARE.
};
using value_type = T;
__host__ __device__
HIP_vector_base& operator=(const HIP_vector_base& x) noexcept {
#if __has_attribute(ext_vector_type)
data = x.data;
#else
data[0] = x.data[0];
data[1] = x.data[1];
data[2] = x.data[2];
data[3] = x.data[3];
#endif
return *this;
}
};
template<typename T, unsigned int rank>
@@ -464,7 +540,14 @@ THE SOFTWARE.
typename std::enable_if<
(rank > 1) && sizeof...(Us) == rank>::type* = nullptr>
inline __host__ __device__
HIP_vector_type(Us... xs) noexcept { data = Native_vec_{static_cast<T>(xs)...}; }
HIP_vector_type(Us... xs) noexcept
{
#if __has_attribute(ext_vector_type)
new (&data) Native_vec_{static_cast<T>(xs)...};
#else
new (&data) std::array<T, rank>{static_cast<T>(xs)...};
#endif
}
inline __host__ __device__
HIP_vector_type(const HIP_vector_type&) = default;
inline __host__ __device__
@@ -651,7 +734,6 @@ THE SOFTWARE.
}
};
template<typename T, unsigned int n>
inline __host__ __device__
HIP_vector_type<T, n> operator+(