CUDA names in string literals replacment added
This commit is contained in:
atimofee
@@ -1,17 +1,21 @@
|
||||
if(${HIPIFY_STANDLONE})
|
||||
cmake_minimum_required(VERSION 2.8.8)
|
||||
|
||||
|
||||
project(hipify)
|
||||
|
||||
find_package(LLVM REQUIRED PATHS ${LLVM_DIR} NO_DEFAULT_PATH)
|
||||
|
||||
list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR})
|
||||
include(AddLLVM)
|
||||
|
||||
message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}")
|
||||
message(STATUS "Using LLVMConfig.cmake in: ${LLVM_DIR}")
|
||||
|
||||
include_directories(${LLVM_INCLUDE_DIRS})
|
||||
link_directories(${LLVM_LIBRARY_DIRS})
|
||||
add_definitions(${LLVM_DEFINITIONS})
|
||||
add_executable( hipify src/Cuda2Hip.cpp )
|
||||
add_llvm_executable( hipify src/Cuda2Hip.cpp )
|
||||
|
||||
else()
|
||||
set(LLVM_LINK_COMPONENTS
|
||||
Option
|
||||
@@ -48,6 +52,7 @@ if(${HIPIFY_STANDLONE})
|
||||
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} -Wall -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden")
|
||||
add_subdirectory(test)
|
||||
else()
|
||||
install(TARGETS hipify
|
||||
RUNTIME DESTINATION bin
|
||||
|
||||
+12
-1
@@ -9,7 +9,7 @@ git clone http://llvm.org/git/llvm.git llvm
|
||||
git clone http://llvm.org/git/clang.git llvm/tools/clang
|
||||
|
||||
mkdir llvm_build && cd llvm_build
|
||||
cmake -DCMAKE_INSTALL_PREFIX="LLVM_INSTALL_PATH" -DLLVM_TARGETS_TO_BUILD="X86;NVPTX;AMDGPU" ../llvm
|
||||
cmake -DCMAKE_INSTALL_PREFIX="LLVM_INSTALL_PATH" -DLLVM_INSTALL_UTILS=ON -DLLVM_TARGETS_TO_BUILD="X86;NVPTX;AMDGPU" ../llvm
|
||||
make && make install
|
||||
|
||||
git clone https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP-hipify.git path_to_hipify_src
|
||||
@@ -34,3 +34,14 @@ cmake -DLLVM_TARGETS_TO_BUILD="X86;NVPTX;AMDGPU" ../llvm
|
||||
make -C tools/clang/tools/extra/hipify
|
||||
```
|
||||
|
||||
# How to run tests (for *standalone* tool only)
|
||||
- install Python and add python-setuptools
|
||||
- install lit python script
|
||||
- make sure that FileCheck util is installed to **LLVM_INSTALL_PATH/bin/FileCheck**
|
||||
- run tests from path_to_hipify_src/build
|
||||
```
|
||||
sudo apt-get install python python-setuptools
|
||||
sudo easy_install lit
|
||||
make -C path_to_hipify_src/build test
|
||||
```
|
||||
|
||||
|
||||
+36
-36
@@ -229,22 +229,22 @@ namespace {
|
||||
DenseMap<StringRef, StringRef> cuda2hipRename;
|
||||
};
|
||||
|
||||
struct HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks {
|
||||
HipifyPPCallbacks(Replacements * R)
|
||||
: SeenEnd(false), _sm(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();
|
||||
return true;
|
||||
}
|
||||
|
||||
struct HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks {
|
||||
HipifyPPCallbacks(Replacements * R)
|
||||
: SeenEnd(false), _sm(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();
|
||||
return true;
|
||||
}
|
||||
|
||||
virtual void InclusionDirective(
|
||||
SourceLocation hash_loc,
|
||||
const Token &include_token,
|
||||
@@ -265,8 +265,8 @@ namespace {
|
||||
llvm::errs() << "\nWill 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);
|
||||
const char* B = _sm->getCharacterData(sl);
|
||||
const char* E = _sm->getCharacterData(sle);
|
||||
repName = "<" + repName + ">";
|
||||
Replacement Rep(*_sm, sl, E - B, repName);
|
||||
Replace->insert(Rep);
|
||||
@@ -300,22 +300,22 @@ namespace {
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
void EndOfMainFile() override
|
||||
{
|
||||
|
||||
}
|
||||
|
||||
bool SeenEnd;
|
||||
void setSourceManager(SourceManager * sm) { _sm = sm; }
|
||||
|
||||
private:
|
||||
|
||||
SourceManager * _sm;
|
||||
Replacements * Replace;
|
||||
struct hipName N;
|
||||
};
|
||||
|
||||
|
||||
void EndOfMainFile() override
|
||||
{
|
||||
|
||||
}
|
||||
|
||||
bool SeenEnd;
|
||||
void setSourceManager(SourceManager * sm) { _sm = sm; }
|
||||
|
||||
private:
|
||||
|
||||
SourceManager * _sm;
|
||||
Replacements * Replace;
|
||||
struct hipName N;
|
||||
};
|
||||
|
||||
class Cuda2HipCallback : public MatchFinder::MatchCallback {
|
||||
public:
|
||||
Cuda2HipCallback(Replacements *Replace) : Replace(Replace) {}
|
||||
@@ -452,8 +452,8 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback {
|
||||
|
||||
if (const ParmVarDecl * cudaParamDecl = Result.Nodes.getNodeAs<clang::ParmVarDecl>("cudaParamDecl"))
|
||||
{
|
||||
QualType QT = cudaParamDecl->getOriginalType();
|
||||
std::string name = QT.getAsString();
|
||||
QualType QT = cudaParamDecl->getOriginalType();
|
||||
std::string name = QT.getAsString();
|
||||
std::string repName = N.cuda2hipRename[name];
|
||||
SourceLocation sl = cudaParamDecl->getLocStart();
|
||||
Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ?
|
||||
|
||||
@@ -0,0 +1,27 @@
|
||||
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_CURRENT_SOURCE_DIR}/lit.site.cfg.in
|
||||
${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg
|
||||
@ONLY)
|
||||
|
||||
add_lit_testsuite(check-hipify "Running HIPify regression tests"
|
||||
${CMAKE_CURRENT_SOURCE_DIR}
|
||||
PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg
|
||||
DEPENDS hipify
|
||||
)
|
||||
|
||||
add_custom_target(check)
|
||||
add_dependencies(check check-hipify)
|
||||
add_custom_target(test)
|
||||
add_dependencies(test check-hipify)
|
||||
set_target_properties(check PROPERTIES FOLDER "Tests")
|
||||
|
||||
@@ -0,0 +1,43 @@
|
||||
// RUN: hipify "%s" 2>&1 | FileCheck %s
|
||||
|
||||
#include <helper_cuda.h> // for checkCudaErrors
|
||||
|
||||
#include <iostream>
|
||||
|
||||
__global__ void axpy(float a, float* x, float* y) {
|
||||
// CHECK: hipThreadIdx_x
|
||||
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;
|
||||
checkCudaErrors(cudaMalloc(&device_x, kDataLen * sizeof(float)));
|
||||
checkCudaErrors(cudaMalloc(&device_y, kDataLen * sizeof(float)));
|
||||
checkCudaErrors(cudaMemcpy(device_x, host_x, kDataLen * sizeof(float),
|
||||
cudaMemcpyHostToDevice));
|
||||
|
||||
// Launch the kernel.
|
||||
// CHECK: hipLaunchKernel(HIP_KERNEL_NAME
|
||||
axpy<<<1, kDataLen>>>(a, device_x, device_y);
|
||||
|
||||
// Copy output data to host.
|
||||
checkCudaErrors(cudaDeviceSynchronize());
|
||||
checkCudaErrors(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";
|
||||
}
|
||||
|
||||
checkCudaErrors(cudaDeviceReset());
|
||||
return 0;
|
||||
}
|
||||
@@ -0,0 +1,54 @@
|
||||
# -*- 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
|
||||
|
||||
tool_name = "FileCheck"
|
||||
tool_path = lit.util.which(tool_name, llvm_tools_dir)
|
||||
if not tool_path:
|
||||
# Warn, but still provide a substitution.
|
||||
lit_config.note('Did not find ' + tool_name + ' in ' + llvm_tools_dir)
|
||||
tool_path = llvm_tools_dir + '/' + tool_name
|
||||
config.substitutions.append((tool_name, tool_path))
|
||||
|
||||
@@ -0,0 +1,15 @@
|
||||
import sys
|
||||
|
||||
config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@"
|
||||
config.obj_root = "@CMAKE_CURRENT_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))
|
||||
|
||||
Verwijs in nieuw issue
Block a user