AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

CUDA names in string literals replacment added

Browse Source 
[ROCm/hip commit: 012afc743f]
This commit is contained in:
atimofee
2016-02-24 19:23:23 +03:00
parent 91e750e1ce 28d550b858
commit 1582663239
7 changed files with 194 additions and 39 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/CMakeLists.txt
+7 -2
View File
@@ -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
projects/hip/README.md
+12 -1
View File
@@ -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
```
projects/hip/src/Cuda2Hip.cpp
+36 -36
View File
@@ -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) ?
projects/hip/test/CMakeLists.txt
+27
View File
@@ -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")
projects/hip/test/axpy.cu
+43
View File
@@ -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;
}
projects/hip/test/lit.cfg
+54
View File
@@ -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))
projects/hip/test/lit.site.cfg.in
+15
View File
@@ -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))