Merge branch 'amd-develop' into amd-master
Change-Id: I52830df409da1f021c32ea569d4ae671aeb57b03
[ROCm/clr commit: 9b8079ad28]
This commit is contained in:
Maneesh Gupta
@@ -177,7 +177,8 @@ if(HIP_PLATFORM STREQUAL "hcc")
|
||||
set(SOURCE_FILES_DEVICE
|
||||
src/device_util.cpp
|
||||
src/hip_ldg.cpp
|
||||
src/hip_fp16.cpp)
|
||||
src/hip_fp16.cpp
|
||||
src/device_functions.cpp)
|
||||
|
||||
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L${HCC_HOME}/lib -lmcwamp -Wl,-Bsymbolic")
|
||||
add_library(hip_hcc SHARED ${SOURCE_FILES_RUNTIME})
|
||||
|
||||
@@ -93,7 +93,7 @@ Differences or limitations of HIP APIs as compared to CUDA APIs should be clearl
|
||||
## Coding Guidelines (in brief)
|
||||
- Code Indentation:
|
||||
- Tabs should be expanded to spaces.
|
||||
- Use 4 spaces indendation.
|
||||
- Use 4 spaces indentation.
|
||||
- Capitalization and Naming
|
||||
- Prefer camelCase for HIP interfaces and internal symbols. Note HCC uses _ for separator.
|
||||
This guideline is not yet consistently followed in HIP code - eventual compliance is aspirational.
|
||||
@@ -120,7 +120,7 @@ Differences or limitations of HIP APIs as compared to CUDA APIs should be clearl
|
||||
|
||||
- HIP_INIT_API() should be placed at the start of each top-level HIP API. This function will make sure the HIP runtime is initialized,
|
||||
and also constructs an appropriate API string for tracing and CodeXL marker tracing. The arguments to HIP_INIT_API should match
|
||||
those of the parent fucntion.
|
||||
those of the parent function.
|
||||
- ihipLogStatus should only be called from top-level HIP APIs,and should be called to log and return the error code. The error code
|
||||
is used by the GetLastError and PeekLastError functions - if a HIP API simply returns, then the error will not be logged correctly.
|
||||
|
||||
@@ -161,4 +161,4 @@ doxygen bug list.
|
||||
|
||||
## Other Tips:
|
||||
### Markdown Editing
|
||||
Recommended to use an offline Markown viewer to review documentation, such as Markdown Preview Plus extension in Chrome browser, or Remarkable.
|
||||
Recommended to use an offline Markdown viewer to review documentation, such as Markdown Preview Plus extension in Chrome browser, or Remarkable.
|
||||
|
||||
@@ -88,8 +88,8 @@ The HIP Runtime API code and compute kernel definition can exist in the same sou
|
||||
|
||||
## HIP Portability and Compiler Technology
|
||||
HIP C++ code can be compiled with either :
|
||||
- On the Nvidia CUDA platform, HIP provides header file which translate from the HIP runtime APIs to CUDA runtime APIs. The header file contains mostly inlined
|
||||
functions and thus has very low overhead - developers coding in HIP should expect the same perforamnce as coding in native CUDA. The code is then
|
||||
- On the NVIDIA CUDA platform, HIP provides header file which translate from the HIP runtime APIs to CUDA runtime APIs. The header file contains mostly inlined
|
||||
functions and thus has very low overhead - developers coding in HIP should expect the same performance as coding in native CUDA. The code is then
|
||||
compiled with nvcc, the standard C++ compiler provided with the CUDA SDK. Developers can use any tools supported by the CUDA SDK including the CUDA
|
||||
profiler and debugger.
|
||||
- On the AMD ROCm platform, HIP provides a header and runtime library built on top of hcc compiler. The HIP runtime implements HIP streams, events, and memory APIs,
|
||||
@@ -120,7 +120,7 @@ make
|
||||
|
||||
|
||||
## More Examples
|
||||
The GitHub repot [HIP-Examples](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP-Examples.git) contains a hipified vesion of the popular Rodinia benchmark suite.
|
||||
The GitHub repository [HIP-Examples](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP-Examples.git) contains a hipified version of the popular Rodinia benchmark suite.
|
||||
The README with the procedures and tips the team used during this porting effort is here: [Rodinia Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP-Examples/blob/master/rodinia_3.0/hip/README.hip_porting)
|
||||
|
||||
## Tour of the HIP Directories
|
||||
|
||||
@@ -3,7 +3,7 @@
|
||||
We have attempted to document known bugs and limitations - in particular the [HIP Kernel Language](docs/markdown/hip_kernel_language.md) document uses the phrase "Under Development", and the [HIP Runtime API bug list](http://gpuopen-professionalcompute-tools.github.io/HIP/bug.html) lists known bugs.
|
||||
|
||||
Upcoming:
|
||||
- Stability: Enforce perioidic host synchronization to reclaim resources if the application has launched a large
|
||||
- Stability: Enforce periodic host synchronization to reclaim resources if the application has launched a large
|
||||
number of commands (>1K) without synchronizing.
|
||||
- Register keyword now silently ignored on HCC (previously would emit warning).
|
||||
- Doc updates: Add some more frequently asked questions to FAQ, fix TOC in some files, review.
|
||||
@@ -73,7 +73,7 @@ Date: 2016.04.25
|
||||
- Create static library and link.
|
||||
- Set HIP_PATH to install.
|
||||
- Make hipDevice and hipStream thread-safe.
|
||||
- Prefered hipStream usage is still to create new streams for each new thread, but it works even if you don;t.
|
||||
- Preferred hipStream usage is still to create new streams for each new thread, but it works even if you don;t.
|
||||
- Improve automated platform detection: If AMD GPU is installed and detected by driver, default HIP_PLATFORM to hcc.
|
||||
- HIP_TRACE_API now prints arguments to the HIP function (in addition to name of function).
|
||||
- Deprecate hipDeviceGetProp (Replace with hipGetDeviceProp)
|
||||
|
||||
@@ -10,7 +10,7 @@
|
||||
- [What specific version of CUDA does HIP support?](#what-specific-version-of-cuda-does-hip-support)
|
||||
- [What libraries does HIP support?](#what-libraries-does-hip-support)
|
||||
- [How does HIP compare with OpenCL?](#how-does-hip-compare-with-opencl)
|
||||
- [How does porting CUDA to HIP compare to porting CUDA to OpenCL?]
|
||||
- [How does porting CUDA to HIP compare to porting CUDA to OpenCL?](#how-does-porting-cuda-to-hip-compare-to-porting-cuda-to-opencl)
|
||||
- [What hardware does HIP support?](#what-hardware-does-hip-support)
|
||||
- [Does Hipify automatically convert all source code?](#does-hipify-automatically-convert-all-source-code)
|
||||
- [What is NVCC?](#what-is-nvcc)
|
||||
@@ -53,7 +53,6 @@ At a high-level, the following features are not supported:
|
||||
- Graphics interoperation with OpenGL or Direct3D
|
||||
- CUDA Driver API (Under Development)
|
||||
- CUDA IPC Functions (Under Development)
|
||||
|
||||
- CUDA array, mipmappedArray and pitched memory
|
||||
- MemcpyToSymbol functions
|
||||
- Queue priority controls
|
||||
@@ -107,7 +106,7 @@ However, we can provide a rough summary of the features included in each CUDA SD
|
||||
|
||||
### What libraries does HIP support?
|
||||
HIP includes growing support for the 4 key math libraries using hcBlas, hcFft, hcrng, and hcsparse).
|
||||
These offer pointer-based memory interfaces (as opposed to opaque buffers) and can be easily interfaces with other HCC code. Developers should use conditional compliation if portability to nvcc systems is desired - using calls to cu* routines on one path and hc* routines on the other.
|
||||
These offer pointer-based memory interfaces (as opposed to opaque buffers) and can be easily interfaces with other HCC code. Developers should use conditional compilation if portability to nvcc systems is desired - using calls to cu* routines on one path and hc* routines on the other.
|
||||
|
||||
- [hcblas](https://bitbucket.org/multicoreware/hcblas)
|
||||
- [hcfft](https://bitbucket.org/multicoreware/hcfft)
|
||||
@@ -132,7 +131,7 @@ HIP offers several benefits over OpenCL:
|
||||
Both HIP and CUDA are dialects of C++, and thus porting between them is relatively straightforward.
|
||||
Both dialects support templates, classes, lambdas, and other C++ constructs.
|
||||
As one example, the hipify tool was originally a perl script that used simple text conversions from CUDA to HIP.
|
||||
HIP and CUDA provide similar math library calls as well. In summary, the HIP philospohy was to make the HIP language close enough to CUDA that the porting effort is relatively simple.
|
||||
HIP and CUDA provide similar math library calls as well. In summary, the HIP philosophy was to make the HIP language close enough to CUDA that the porting effort is relatively simple.
|
||||
This reduces the potential for error, and also makes it easy to automate the translation. HIP's goal is to quickly get the ported program running on both platforms with little manual intervention,
|
||||
so that the programmer can focus on performance optimizations.
|
||||
|
||||
@@ -196,8 +195,8 @@ HIP will set the platform to HCC if it sees that the AMD graphics driver is inst
|
||||
Sometimes this isn't what you want - you can force HIP to recognize the platform by setting HIP_PLATFORM to hcc (or nvcc)
|
||||
```
|
||||
export HIP_PLATFORM=hcc
|
||||
|
||||
```
|
||||
|
||||
One symptom of this problem is the message "error: 'unknown error'(11) at square.hipref.cpp:56". This can occur if you have a CUDA installation on an AMD platform, and HIP incorrectly detects the platform as nvcc. HIP may be able to compile the application using the nvcc tool-chain, but will generate this error at runtime since the platform does not have a CUDA device. The fix is to set HIP_PLATFORM=hcc and rebuild the issue.
|
||||
|
||||
If you see issues related to incorrect platform detection, please file an issue with the GitHub issue tracker so we can improve HIP's platform detection logic.
|
||||
|
||||
@@ -105,7 +105,7 @@ HIP parses the `__noinline__` and `__forceinline__` keywords and converts them t
|
||||
|
||||
|
||||
```
|
||||
// Example psuedocode introducing hipLaunchKernel:
|
||||
// Example pseudo code introducing hipLaunchKernel:
|
||||
__global__ MyKernel(hipLaunchParm lp, float *A, float *B, float *C, size_t N)
|
||||
{
|
||||
...
|
||||
|
||||
@@ -31,8 +31,8 @@ accelerator language front-end. Here, NVCC is not used. Instead, the environm
|
||||
different kernel language or different compilation flow.
|
||||
Other environments have many kernels and do not want them to be all loaded automatically.
|
||||
The Module functions can be used to load the generated code objects and launch kernels.
|
||||
As we will see below, HIP defines a Module API which provides similar explict control over code
|
||||
object managemenet.
|
||||
As we will see below, HIP defines a Module API which provides similar explicit control over code
|
||||
object management.
|
||||
|
||||
### cuCtx API
|
||||
The Driver API defines "Context" and "Devices" as separate entities.
|
||||
@@ -144,7 +144,7 @@ table shows the type equivalence to enable this interaction.
|
||||
The hipModule interface does not support the hipModuleLoadEx function, which is used to control PTX compilaton options.
|
||||
HCC does not use PTX and does not support the same compilation options.
|
||||
In fact, HCC code objects always contain fully compiled ISA and do not require additional compilation as part of the load step.
|
||||
Code which requires this functionaly should use platform-specific coding, calling `cuModuleLoadEx`
|
||||
Code which requires this functionally should use platform-specific coding, calling `cuModuleLoadEx`
|
||||
on the NVCC path and hipModuleLoad on the hcc path. For example:
|
||||
|
||||
```
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
# HIP Porting Guide
|
||||
In addition to providing a portable C++ programmming environement for GPUs, HIP is designed to ease
|
||||
In addition to providing a portable C++ programming environment for GPUs, HIP is designed to ease
|
||||
the porting of existing CUDA code into the HIP environment. This section describes the available tools
|
||||
and provides practical suggestions on how to port CUDA code and work through common issues.
|
||||
|
||||
|
||||
@@ -13,7 +13,7 @@ The default device can be set with hipSetDevice.
|
||||
- hcc = Heterogeneous Compute Compiler (https://bitbucket.org/multicoreware/hcc/wiki/Home).
|
||||
|
||||
- hipify - tool to convert CUDA(R) code to portable C++ code.
|
||||
- hipconfig - tool to report various confoguration properties of the target platform.
|
||||
- hipconfig - tool to report various configuration properties of the target platform.
|
||||
|
||||
- nvcc = nvcc compiler, do not capitalize.
|
||||
- hcc = heterogeneous compute compiler, do not capitalize.
|
||||
|
||||
@@ -85,6 +85,7 @@ const char *counterNames[CONV_LAST] = {
|
||||
"special_func", "stream", "event", "ctx", "module", "cache",
|
||||
"err", "def", "tex", "other", "include", "include_cuda_main_header",
|
||||
"type", "literal", "numeric_literal"};
|
||||
|
||||
enum ApiTypes {
|
||||
API_DRIVER = 0,
|
||||
API_RUNTIME,
|
||||
@@ -97,6 +98,15 @@ const char *apiNames[API_LAST] = {
|
||||
|
||||
namespace {
|
||||
|
||||
int64_t countRepsTotal[CONV_LAST] = { 0 };
|
||||
int64_t countApiRepsTotal[API_LAST] = { 0 };
|
||||
|
||||
struct hipCounter {
|
||||
StringRef hipName;
|
||||
ConvTypes countType;
|
||||
ApiTypes countApiType;
|
||||
};
|
||||
|
||||
struct cuda2hipMap {
|
||||
cuda2hipMap() {
|
||||
|
||||
@@ -1317,13 +1327,7 @@ struct cuda2hipMap {
|
||||
//cuda2hipRename["cublasDrotmg_v2"] = {"hipblasDrotmg", CONV_MATH_FUNC, API_BLAS};
|
||||
}
|
||||
|
||||
struct HipNames {
|
||||
StringRef hipName;
|
||||
ConvTypes countType;
|
||||
ApiTypes countApiType;
|
||||
};
|
||||
|
||||
SmallDenseMap<StringRef, HipNames> cuda2hipRename;
|
||||
SmallDenseMap<StringRef, hipCounter> cuda2hipRename;
|
||||
std::set<StringRef> cudaExcludes;
|
||||
};
|
||||
|
||||
@@ -1333,36 +1337,52 @@ StringRef unquoteStr(StringRef s) {
|
||||
return s;
|
||||
}
|
||||
|
||||
static void processString(StringRef s, const cuda2hipMap &map,
|
||||
Replacements *Replace, SourceManager &SM,
|
||||
SourceLocation start,
|
||||
int64_t countReps[CONV_LAST],
|
||||
int64_t countApiReps[API_LAST]) {
|
||||
size_t begin = 0;
|
||||
while ((begin = s.find("cu", 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()) {
|
||||
StringRef repName = found->second.hipName;
|
||||
countReps[CONV_LITERAL]++;
|
||||
countApiReps[API_RUNTIME]++;
|
||||
SourceLocation sl = start.getLocWithOffset(begin + 1);
|
||||
Replacement Rep(SM, sl, name.size(), repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
if (end == StringRef::npos)
|
||||
break;
|
||||
begin = end + 1;
|
||||
class Cuda2Hip {
|
||||
public:
|
||||
Cuda2Hip(Replacements *R): Replace(R) {}
|
||||
|
||||
int64_t countReps[CONV_LAST] = { 0 };
|
||||
int64_t countApiReps[API_LAST] = { 0 };
|
||||
|
||||
protected:
|
||||
struct cuda2hipMap N;
|
||||
Replacements *Replace;
|
||||
|
||||
virtual void updateCounters(const hipCounter & counter) {
|
||||
countReps[counter.countType]++;
|
||||
countRepsTotal[counter.countType]++;
|
||||
countApiReps[counter.countApiType]++;
|
||||
countApiRepsTotal[counter.countApiType]++;
|
||||
}
|
||||
}
|
||||
|
||||
void processString(StringRef s, SourceManager &SM, SourceLocation start) {
|
||||
size_t begin = 0;
|
||||
while ((begin = s.find("cu", begin)) != StringRef::npos) {
|
||||
const size_t end = s.find_first_of(" ", begin + 4);
|
||||
StringRef name = s.slice(begin, end);
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
StringRef repName = found->second.hipName;
|
||||
hipCounter counter = { "", CONV_LITERAL, API_RUNTIME };
|
||||
updateCounters(counter);
|
||||
SourceLocation sl = start.getLocWithOffset(begin + 1);
|
||||
Replacement Rep(SM, sl, name.size(), repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
if (end == StringRef::npos) {
|
||||
break;
|
||||
}
|
||||
begin = end + 1;
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
class Cuda2HipCallback;
|
||||
|
||||
class HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks {
|
||||
class HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks, public Cuda2Hip {
|
||||
public:
|
||||
HipifyPPCallbacks(Replacements *R)
|
||||
: SeenEnd(false), _sm(nullptr), _pp(nullptr), Replace(R) {}
|
||||
: Cuda2Hip(R), SeenEnd(false), _sm(nullptr), _pp(nullptr) {}
|
||||
|
||||
virtual bool handleBeginSource(CompilerInstance &CI,
|
||||
StringRef Filename) override {
|
||||
@@ -1389,8 +1409,7 @@ public:
|
||||
const auto found = N.cuda2hipRename.find(file_name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
StringRef repName = found->second.hipName;
|
||||
countReps[found->second.countType]++;
|
||||
countApiReps[found->second.countApiType]++;
|
||||
updateCounters(found->second);
|
||||
DEBUG(dbgs() << "Include file found: " << file_name << "\n"
|
||||
<< "SourceLocation:"
|
||||
<< filename_range.getBegin().printToString(*_sm) << "\n"
|
||||
@@ -1418,8 +1437,7 @@ public:
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
StringRef repName = found->second.hipName;
|
||||
countReps[found->second.countType]++;
|
||||
countApiReps[found->second.countApiType]++;
|
||||
updateCounters(found->second);
|
||||
SourceLocation sl = T.getLocation();
|
||||
DEBUG(dbgs() << "Identifier " << name
|
||||
<< " found in definition of macro "
|
||||
@@ -1465,8 +1483,7 @@ public:
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
StringRef repName = found->second.hipName;
|
||||
countReps[found->second.countType]++;
|
||||
countApiReps[found->second.countApiType]++;
|
||||
updateCounters(found->second);
|
||||
DEBUG(dbgs()
|
||||
<< "Identifier " << name
|
||||
<< " found as an actual argument in expansion of macro "
|
||||
@@ -1488,16 +1505,14 @@ public:
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
sl = sl_macro;
|
||||
StringRef repName = found->second.hipName;
|
||||
countReps[found->second.countType]++;
|
||||
countApiReps[found->second.countApiType]++;
|
||||
updateCounters(found->second);
|
||||
Replacement Rep(*_sm, sl, length, repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
} else {
|
||||
if (tok.is(tok::string_literal)) {
|
||||
StringRef s(tok.getLiteralData(), tok.getLength());
|
||||
processString(unquoteStr(s), N, Replace, *_sm, tok.getLocation(),
|
||||
countReps, countApiReps);
|
||||
processString(unquoteStr(s), *_sm, tok.getLocation());
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -1513,18 +1528,14 @@ public:
|
||||
void setSourceManager(SourceManager *sm) { _sm = sm; }
|
||||
void setPreprocessor(Preprocessor *pp) { _pp = pp; }
|
||||
void setMatch(Cuda2HipCallback *match) { Match = match; }
|
||||
int64_t countReps[CONV_LAST] = { 0 };
|
||||
int64_t countApiReps[API_LAST] = { 0 };
|
||||
|
||||
private:
|
||||
SourceManager *_sm;
|
||||
Preprocessor *_pp;
|
||||
Cuda2HipCallback *Match;
|
||||
Replacements *Replace;
|
||||
struct cuda2hipMap N;
|
||||
};
|
||||
|
||||
class Cuda2HipCallback : public MatchFinder::MatchCallback {
|
||||
class Cuda2HipCallback : public MatchFinder::MatchCallback, public Cuda2Hip {
|
||||
private:
|
||||
void convertKernelDecl(const FunctionDecl *kernelDecl, const MatchFinder::MatchResult &Result) {
|
||||
SourceManager *SM = Result.SourceManager;
|
||||
@@ -1579,8 +1590,7 @@ private:
|
||||
}
|
||||
}
|
||||
if (bReplace) {
|
||||
countReps[found->second.countType]++;
|
||||
countApiReps[found->second.countApiType]++;
|
||||
updateCounters(found->second);
|
||||
Replacement Rep(*SM, sl, length, repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
@@ -1653,8 +1663,8 @@ private:
|
||||
SM->getCharacterData(launchKernel->getLocStart());
|
||||
Replacement Rep(*SM, launchKernel->getLocStart(), length, OS.str());
|
||||
Replace->insert(Rep);
|
||||
countReps[CONV_KERN]++;
|
||||
countApiReps[API_RUNTIME]++;
|
||||
hipCounter counter = { "", CONV_KERN, API_RUNTIME };
|
||||
updateCounters(counter);
|
||||
return true;
|
||||
}
|
||||
return false;
|
||||
@@ -1675,8 +1685,7 @@ private:
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
StringRef repName = found->second.hipName;
|
||||
countReps[found->second.countType]++;
|
||||
countApiReps[found->second.countApiType]++;
|
||||
updateCounters(found->second);
|
||||
SourceLocation sl = threadIdx->getLocStart();
|
||||
SourceManager *SM = Result.SourceManager;
|
||||
Replacement Rep(*SM, sl, name.size(), repName);
|
||||
@@ -1695,8 +1704,7 @@ private:
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
StringRef repName = found->second.hipName;
|
||||
countReps[found->second.countType]++;
|
||||
countApiReps[found->second.countApiType]++;
|
||||
updateCounters(found->second);
|
||||
SourceLocation sl = enumConstantRef->getLocStart();
|
||||
SourceManager *SM = Result.SourceManager;
|
||||
Replacement Rep(*SM, sl, name.size(), repName);
|
||||
@@ -1719,8 +1727,7 @@ private:
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
StringRef repName = found->second.hipName;
|
||||
countReps[found->second.countType]++;
|
||||
countApiReps[found->second.countApiType]++;
|
||||
updateCounters(found->second);
|
||||
SourceLocation sl = enumConstantDecl->getLocStart();
|
||||
SourceManager *SM = Result.SourceManager;
|
||||
Replacement Rep(*SM, sl, name.size(), repName);
|
||||
@@ -1742,8 +1749,7 @@ private:
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
StringRef repName = found->second.hipName;
|
||||
countReps[found->second.countType]++;
|
||||
countApiReps[found->second.countApiType]++;
|
||||
updateCounters(found->second);
|
||||
SourceLocation sl = typedefVar->getLocStart();
|
||||
SourceManager *SM = Result.SourceManager;
|
||||
Replacement Rep(*SM, sl, name.size(), repName);
|
||||
@@ -1763,8 +1769,7 @@ private:
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
StringRef repName = found->second.hipName;
|
||||
countReps[found->second.countType]++;
|
||||
countApiReps[found->second.countApiType]++;
|
||||
updateCounters(found->second);
|
||||
TypeLoc TL = structVar->getTypeSourceInfo()->getTypeLoc();
|
||||
SourceLocation sl = TL.getUnqualifiedLoc().getLocStart();
|
||||
SourceManager *SM = Result.SourceManager;
|
||||
@@ -1784,8 +1789,7 @@ private:
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
StringRef repName = found->second.hipName;
|
||||
countReps[found->second.countType]++;
|
||||
countApiReps[found->second.countApiType]++;
|
||||
updateCounters(found->second);
|
||||
TypeLoc TL = structVarPtr->getTypeSourceInfo()->getTypeLoc();
|
||||
SourceLocation sl = TL.getUnqualifiedLoc().getLocStart();
|
||||
SourceManager *SM = Result.SourceManager;
|
||||
@@ -1807,8 +1811,7 @@ private:
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
StringRef repName = found->second.hipName;
|
||||
countReps[found->second.countType]++;
|
||||
countApiReps[found->second.countApiType]++;
|
||||
updateCounters(found->second);
|
||||
TypeLoc TL = typeInfo->getTypeLoc();
|
||||
SourceLocation sl = TL.getUnqualifiedLoc().getLocStart();
|
||||
SourceManager *SM = Result.SourceManager;
|
||||
@@ -1852,8 +1855,8 @@ private:
|
||||
StringRef repName = Twine("HIP_DYNAMIC_SHARED(" + typeName + ", " + varName + ")").toStringRef(tmpData);
|
||||
Replacement Rep(*SM, slStart, repLength, repName);
|
||||
Replace->insert(Rep);
|
||||
countReps[CONV_MEM]++;
|
||||
countApiReps[API_RUNTIME]++;
|
||||
hipCounter counter = { "", CONV_MEM, API_RUNTIME };
|
||||
updateCounters(counter);
|
||||
}
|
||||
}
|
||||
return true;
|
||||
@@ -1872,8 +1875,7 @@ private:
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
StringRef repName = found->second.hipName;
|
||||
countReps[found->second.countType]++;
|
||||
countApiReps[found->second.countApiType]++;
|
||||
updateCounters(found->second);
|
||||
TypeLoc TL = paramDecl->getTypeSourceInfo()->getTypeLoc();
|
||||
SourceLocation sl = TL.getUnqualifiedLoc().getLocStart();
|
||||
SourceManager *SM = Result.SourceManager;
|
||||
@@ -1897,8 +1899,7 @@ private:
|
||||
const auto found = N.cuda2hipRename.find(name);
|
||||
if (found != N.cuda2hipRename.end()) {
|
||||
StringRef repName = found->second.hipName;
|
||||
countReps[found->second.countType]++;
|
||||
countApiReps[found->second.countApiType]++;
|
||||
updateCounters(found->second);
|
||||
TypeLoc TL = paramDeclPtr->getTypeSourceInfo()->getTypeLoc();
|
||||
SourceLocation sl = TL.getUnqualifiedLoc().getLocStart();
|
||||
SourceManager *SM = Result.SourceManager;
|
||||
@@ -1925,7 +1926,7 @@ private:
|
||||
if (sLiteral->getCharByteWidth() == 1) {
|
||||
StringRef s = sLiteral->getString();
|
||||
SourceManager *SM = Result.SourceManager;
|
||||
processString(s, N, Replace, *SM, sLiteral->getLocStart(), countReps, countApiReps);
|
||||
processString(s, *SM, sLiteral->getLocStart());
|
||||
}
|
||||
return true;
|
||||
}
|
||||
@@ -1934,7 +1935,7 @@ private:
|
||||
|
||||
public:
|
||||
Cuda2HipCallback(Replacements *Replace, ast_matchers::MatchFinder *parent, HipifyPPCallbacks *PPCallbacks)
|
||||
: Replace(Replace), owner(parent), PP(PPCallbacks) {
|
||||
: Cuda2Hip(Replace), owner(parent), PP(PPCallbacks) {
|
||||
PP->setMatch(this);
|
||||
}
|
||||
|
||||
@@ -1962,19 +1963,14 @@ public:
|
||||
SourceManager *SM = Result.SourceManager;
|
||||
Replacement Rep(*SM, SM->getLocForStartOfFile(SM->getMainFileID()), 0, repName);
|
||||
Replace->insert(Rep);
|
||||
countReps[CONV_INCLUDE_CUDA_MAIN_H]++;
|
||||
countApiReps[API_RUNTIME]++;
|
||||
hipCounter counter = { "", CONV_INCLUDE_CUDA_MAIN_H, API_RUNTIME };
|
||||
updateCounters(counter);
|
||||
}
|
||||
}
|
||||
|
||||
int64_t countReps[CONV_LAST] = { 0 };
|
||||
int64_t countApiReps[API_LAST] = { 0 };
|
||||
|
||||
private:
|
||||
Replacements *Replace;
|
||||
ast_matchers::MatchFinder *owner;
|
||||
HipifyPPCallbacks *PP;
|
||||
struct cuda2hipMap N;
|
||||
};
|
||||
|
||||
void HipifyPPCallbacks::handleEndSource() {
|
||||
@@ -1983,8 +1979,8 @@ void HipifyPPCallbacks::handleEndSource() {
|
||||
StringRef repName = "#include <hip/hip_runtime.h>\n";
|
||||
Replacement Rep(*_sm, _sm->getLocForStartOfFile(_sm->getMainFileID()), 0, repName);
|
||||
Replace->insert(Rep);
|
||||
countReps[CONV_INCLUDE_CUDA_MAIN_H]++;
|
||||
countApiReps[API_RUNTIME]++;
|
||||
hipCounter counter = { "", CONV_INCLUDE_CUDA_MAIN_H, API_RUNTIME };
|
||||
updateCounters(counter);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -2019,7 +2015,7 @@ static cl::opt<bool> PrintStats("print-stats",
|
||||
cl::value_desc("print-stats"),
|
||||
cl::cat(ToolTemplateCategory));
|
||||
|
||||
static cl::opt<bool> N("n",
|
||||
static cl::opt<bool> Examine("examine",
|
||||
cl::desc("Combines -no-output and -print-stats options"),
|
||||
cl::value_desc("n"),
|
||||
cl::cat(ToolTemplateCategory));
|
||||
@@ -2088,12 +2084,12 @@ void addAllMatchers(ast_matchers::MatchFinder &Finder, Cuda2HipCallback *Callbac
|
||||
Callback);
|
||||
}
|
||||
|
||||
void printStats(std::string fileSource, HipifyPPCallbacks &PPCallbacks, Cuda2HipCallback &Callback) {
|
||||
int64_t printStats(std::string fileSource, HipifyPPCallbacks &PPCallbacks, Cuda2HipCallback &Callback) {
|
||||
int64_t sum = 0;
|
||||
for (int i = 0; i < CONV_LAST; i++) {
|
||||
sum += Callback.countReps[i] + PPCallbacks.countReps[i];
|
||||
}
|
||||
llvm::outs() << "info: converted " << sum << " CUDA->HIP refs ( ";
|
||||
llvm::outs() << "Info: converted " << sum << " CUDA->HIP refs ( ";
|
||||
for (int i = 0; i < CONV_LAST; i++) {
|
||||
llvm::outs() << counterNames[i] << ':' << Callback.countReps[i] + PPCallbacks.countReps[i] << ' ';
|
||||
}
|
||||
@@ -2102,15 +2098,33 @@ void printStats(std::string fileSource, HipifyPPCallbacks &PPCallbacks, Cuda2Hip
|
||||
llvm::outs() << apiNames[i] << ':' << Callback.countApiReps[i] + PPCallbacks.countApiReps[i] << ' ';
|
||||
}
|
||||
llvm::outs() << ") in \'" << fileSource << "\'\n";
|
||||
return sum;
|
||||
}
|
||||
|
||||
void printAllStats(int64_t totalFiles, int64_t convertedFiles) {
|
||||
int64_t sum = 0;
|
||||
for (int i = 0; i < CONV_LAST; i++) {
|
||||
sum += countRepsTotal[i];
|
||||
}
|
||||
llvm::outs() << "Info: totally converted " << sum << " CUDA->HIP refs ( ";
|
||||
for (int i = 0; i < CONV_LAST; i++) {
|
||||
llvm::outs() << counterNames[i] << ':' << countRepsTotal[i] << ' ';
|
||||
}
|
||||
llvm::outs() << "), by APIs ( ";
|
||||
for (int i = 0; i < API_LAST; i++) {
|
||||
llvm::outs() << apiNames[i] << ':' << countApiRepsTotal[i] << ' ';
|
||||
}
|
||||
llvm::outs() << ") in " << convertedFiles << " files of " << totalFiles << " processed files.\n";
|
||||
}
|
||||
|
||||
int main(int argc, const char **argv) {
|
||||
llvm::sys::PrintStackTraceOnErrorSignal();
|
||||
CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::Required);
|
||||
CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::OneOrMore);
|
||||
std::vector<std::string> fileSources = OptionsParser.getSourcePathList();
|
||||
std::string dst = OutputFilename;
|
||||
if (N) {
|
||||
NoOutput = PrintStats = true;
|
||||
if (!dst.empty() && fileSources.size() > 1) {
|
||||
llvm::errs() << "Conflict: -o and multiple source files are specified.\n";
|
||||
return 1;
|
||||
}
|
||||
if (NoOutput) {
|
||||
if (Inplace) {
|
||||
@@ -2122,84 +2136,98 @@ int main(int argc, const char **argv) {
|
||||
return 1;
|
||||
}
|
||||
}
|
||||
if (dst.empty()) {
|
||||
dst = fileSources[0];
|
||||
if (!Inplace) {
|
||||
size_t pos = dst.rfind(".");
|
||||
if (pos != std::string::npos && pos+1 < dst.size()) {
|
||||
dst = dst.substr(0, pos) + ".hip." + dst.substr(pos+1, dst.size()-pos-1);
|
||||
} else {
|
||||
dst += ".hip.cu";
|
||||
if (Examine) {
|
||||
NoOutput = PrintStats = true;
|
||||
}
|
||||
int Result = 0;
|
||||
size_t filesTransleted = fileSources.size();
|
||||
for (const auto & src : fileSources) {
|
||||
if (dst.empty()) {
|
||||
dst = src;
|
||||
if (!Inplace) {
|
||||
size_t pos = dst.rfind(".");
|
||||
if (pos != std::string::npos && pos + 1 < dst.size()) {
|
||||
dst = dst.substr(0, pos) + ".hip." + dst.substr(pos + 1, dst.size() - pos - 1);
|
||||
}
|
||||
else {
|
||||
dst += ".hip.cu";
|
||||
}
|
||||
}
|
||||
}
|
||||
} else {
|
||||
if (Inplace) {
|
||||
llvm::errs() << "Conflict: both -o and -inplace options are specified.\n";
|
||||
return 1;
|
||||
else {
|
||||
if (Inplace) {
|
||||
llvm::errs() << "Conflict: both -o and -inplace options are specified.\n";
|
||||
return 1;
|
||||
}
|
||||
dst += ".hip";
|
||||
}
|
||||
// backup source file since tooling may change "inplace"
|
||||
if (!NoBackup || !Inplace) {
|
||||
std::ifstream source(src, std::ios::binary);
|
||||
std::ofstream dest(Inplace ? dst + ".prehip" : dst, std::ios::binary);
|
||||
dest << source.rdbuf();
|
||||
source.close();
|
||||
dest.close();
|
||||
}
|
||||
dst += ".hip";
|
||||
}
|
||||
// backup source file since tooling may change "inplace"
|
||||
if (!NoBackup || !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;
|
||||
HipifyPPCallbacks PPCallbacks(&Tool.getReplacements());
|
||||
Cuda2HipCallback Callback(&Tool.getReplacements(), &Finder, &PPCallbacks);
|
||||
RefactoringTool Tool(OptionsParser.getCompilations(), dst);
|
||||
ast_matchers::MatchFinder Finder;
|
||||
HipifyPPCallbacks PPCallbacks(&Tool.getReplacements());
|
||||
Cuda2HipCallback Callback(&Tool.getReplacements(), &Finder, &PPCallbacks);
|
||||
|
||||
addAllMatchers(Finder, &Callback);
|
||||
addAllMatchers(Finder, &Callback);
|
||||
|
||||
auto action = newFrontendActionFactory(&Finder, &PPCallbacks);
|
||||
std::vector<const char*> compilationStages;
|
||||
compilationStages.push_back("--cuda-host-only");
|
||||
Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster(compilationStages[0], ArgumentInsertPosition::BEGIN));
|
||||
Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("-std=c++11"));
|
||||
auto action = newFrontendActionFactory(&Finder, &PPCallbacks);
|
||||
std::vector<const char*> compilationStages;
|
||||
compilationStages.push_back("--cuda-host-only");
|
||||
Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster(compilationStages[0], ArgumentInsertPosition::BEGIN));
|
||||
Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("-std=c++11"));
|
||||
#if defined(HIPIFY_CLANG_RES)
|
||||
Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES));
|
||||
Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES));
|
||||
#endif
|
||||
Tool.appendArgumentsAdjuster(getClangSyntaxOnlyAdjuster());
|
||||
int Result = Tool.run(action.get());
|
||||
Tool.clearArgumentsAdjusters();
|
||||
Tool.appendArgumentsAdjuster(getClangSyntaxOnlyAdjuster());
|
||||
Result += Tool.run(action.get());
|
||||
Tool.clearArgumentsAdjusters();
|
||||
|
||||
LangOptions DefaultLangOptions;
|
||||
IntrusiveRefCntPtr<DiagnosticOptions> DiagOpts = new DiagnosticOptions();
|
||||
TextDiagnosticPrinter DiagnosticPrinter(llvm::errs(), &*DiagOpts);
|
||||
DiagnosticsEngine Diagnostics(
|
||||
LangOptions DefaultLangOptions;
|
||||
IntrusiveRefCntPtr<DiagnosticOptions> DiagOpts = new DiagnosticOptions();
|
||||
TextDiagnosticPrinter DiagnosticPrinter(llvm::errs(), &*DiagOpts);
|
||||
DiagnosticsEngine Diagnostics(
|
||||
IntrusiveRefCntPtr<DiagnosticIDs>(new DiagnosticIDs()), &*DiagOpts,
|
||||
&DiagnosticPrinter, false);
|
||||
|
||||
DEBUG(dbgs() << "Replacements collected by the tool:\n");
|
||||
for (const auto &r : Tool.getReplacements()) {
|
||||
DEBUG(dbgs() << r.toString() << "\n");
|
||||
}
|
||||
DEBUG(dbgs() << "Replacements collected by the tool:\n");
|
||||
for (const auto &r : Tool.getReplacements()) {
|
||||
DEBUG(dbgs() << r.toString() << "\n");
|
||||
}
|
||||
|
||||
SourceManager Sources(Diagnostics, Tool.getFiles());
|
||||
Rewriter Rewrite(Sources, DefaultLangOptions);
|
||||
SourceManager Sources(Diagnostics, Tool.getFiles());
|
||||
Rewriter Rewrite(Sources, DefaultLangOptions);
|
||||
|
||||
if (!Tool.applyAllReplacements(Rewrite)) {
|
||||
DEBUG(dbgs() << "Skipped some replacements.\n");
|
||||
}
|
||||
if (!NoOutput) {
|
||||
Result = Rewrite.overwriteChangedFiles();
|
||||
}
|
||||
if (!Inplace && !NoOutput) {
|
||||
size_t pos = dst.rfind(".");
|
||||
if (pos != std::string::npos) {
|
||||
rename(dst.c_str(), dst.substr(0, pos).c_str());
|
||||
if (!Tool.applyAllReplacements(Rewrite)) {
|
||||
DEBUG(dbgs() << "Skipped some replacements.\n");
|
||||
}
|
||||
if (!NoOutput) {
|
||||
Result += Rewrite.overwriteChangedFiles();
|
||||
}
|
||||
if (!Inplace && !NoOutput) {
|
||||
size_t pos = dst.rfind(".");
|
||||
if (pos != std::string::npos) {
|
||||
rename(dst.c_str(), dst.substr(0, pos).c_str());
|
||||
}
|
||||
}
|
||||
if (NoOutput) {
|
||||
remove(dst.c_str());
|
||||
}
|
||||
dst.clear();
|
||||
if (PrintStats) {
|
||||
if (0 == printStats(src, PPCallbacks, Callback)) {
|
||||
filesTransleted--;
|
||||
}
|
||||
}
|
||||
}
|
||||
if (NoOutput) {
|
||||
remove(dst.c_str());
|
||||
if (PrintStats && fileSources.size() > 1) {
|
||||
printAllStats(fileSources.size(), filesTransleted);
|
||||
}
|
||||
if (PrintStats) {
|
||||
printStats(fileSources[0], PPCallbacks, Callback);
|
||||
}
|
||||
|
||||
return Result;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,33 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#ifndef HIP_DEVICE_FUNCTIONS_H
|
||||
#define HIP_DEVICE_FUNCTIONS_H
|
||||
|
||||
#include<hip/hip_common.h>
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
|
||||
#include <hip/hcc_detail/device_functions.h>
|
||||
#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
|
||||
#include <device_functions.h>
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,32 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#ifndef HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H
|
||||
#define HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H
|
||||
|
||||
#include "hip_runtime.h"
|
||||
|
||||
__device__ float __int_as_float (int x);
|
||||
|
||||
__device__ double __hiloint2double (int hi, int lo);
|
||||
|
||||
extern __HIP_DEVICE__ double __longlong_as_double(long long int x);
|
||||
extern __HIP_DEVICE__ long long int __double_as_longlong(double x);
|
||||
|
||||
#endif
|
||||
@@ -25,6 +25,67 @@ THE SOFTWARE.
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
#define __CLANG_VERSION__ __clang_major__ * 10 + __clang_minor__
|
||||
|
||||
#ifdef HIP_HALF_HW_SUPPORT
|
||||
|
||||
typedef __fp16 __half;
|
||||
|
||||
typedef struct __attribute__((aligned(4))){
|
||||
int a;
|
||||
} __half2;
|
||||
|
||||
extern "C" __half __hip_hadd_gfx803(__half a, __half b);
|
||||
extern "C" __half __hip_hfma_gfx803(__half a, __half b);
|
||||
extern "C" __half __hip_hmul_gfx803(__half a, __half b);
|
||||
extern "C" __half __hip_hsub_gfx803(__half a, __half b);
|
||||
|
||||
extern "C" int __hip_hadd2_gfx803(int a, int b);
|
||||
extern "C" int __hip_hfma2_gfx803(int a, int b);
|
||||
extern "C" int __hip_hmul2_gfx803(int a, int b);
|
||||
extern "C" int __hip_hsub2_gfx803(int a, int b);
|
||||
|
||||
__device__ inline __half __hadd(__half a, __half b) {
|
||||
return __hip_hadd_gfx803(a, b);
|
||||
}
|
||||
|
||||
__device__ inline __half __hadd_sat(__half a, __half b) {
|
||||
return __hip_hadd_gfx803(a, b);
|
||||
}
|
||||
|
||||
__device__ inline __half __hfma(__half a, __half b) {
|
||||
return __hip_hfma_gfx803(a, b);
|
||||
}
|
||||
|
||||
__device__ inline __half __hfma_sat(__half a, __half b) {
|
||||
return __hip_hfma_gfx803(a, b);
|
||||
}
|
||||
|
||||
__device__ inline __half __hmul(__half a, __half b) {
|
||||
return __hip_hmul_gfx803(a, b);
|
||||
}
|
||||
|
||||
__device__ inline __half __hmul_sat(__half a, __half b) {
|
||||
return __hip_hmul_gfx803(a, b);
|
||||
}
|
||||
|
||||
__device__ inline __half __hsub(__half a, __half b) {
|
||||
return __hip_hsub_gfx803(a, b);
|
||||
}
|
||||
|
||||
__device__ inline __half __hsub_sat(__half a, __half b) {
|
||||
return __hip_hsub_gfx803(a, b);
|
||||
}
|
||||
|
||||
|
||||
__device__ inline __half2 __hadd2(__half2 a, __half2 b) {
|
||||
__half2 ret;
|
||||
ret.a = __hip_hadd2_gfx803(a.a, b.a);
|
||||
return ret;
|
||||
}
|
||||
|
||||
#else
|
||||
|
||||
typedef struct{
|
||||
unsigned x: 16;
|
||||
} __half;
|
||||
@@ -175,3 +236,4 @@ __device__ __half2 __lowhigh2highlow(const __half2 a);
|
||||
__device__ __half2 __low2half2(const __half2 a, const __half2 b);
|
||||
|
||||
#endif
|
||||
#endif
|
||||
|
||||
@@ -343,9 +343,6 @@ __HIP_DEVICE__ double2 make_double2(double, double );
|
||||
__HIP_DEVICE__ double3 make_double3(double, double, double );
|
||||
__HIP_DEVICE__ double4 make_double4(double, double, double, double );
|
||||
|
||||
extern __HIP_DEVICE__ double __longlong_as_double(long long int x);
|
||||
extern __HIP_DEVICE__ long long int __double_as_longlong(double x);
|
||||
|
||||
|
||||
/*
|
||||
///---
|
||||
|
||||
@@ -1,6 +1,6 @@
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
HIPCC=$(HIP_PATH)/bin/hipcc
|
||||
@@ -22,14 +22,15 @@ CXX=$(HIPCC)
|
||||
|
||||
|
||||
$(EXECUTABLE): $(OBJECTS)
|
||||
$(HIPCC) $(OBJECTS) -o $@
|
||||
$(HIPCC) $(OBJECTS) -o $@
|
||||
|
||||
|
||||
test: $(EXECUTABLE)
|
||||
$(EXECUTABLE)
|
||||
$(EXECUTABLE)
|
||||
|
||||
|
||||
clean:
|
||||
rm -f $(EXECUTABLE)
|
||||
rm -f $(OBJECTS)
|
||||
rm -f $(HIP_PATH)/src/*.o
|
||||
rm -f $(EXECUTABLE)
|
||||
rm -f $(OBJECTS)
|
||||
rm -f $(HIP_PATH)/src/*.o
|
||||
|
||||
|
||||
@@ -12,7 +12,7 @@ endif
|
||||
ifeq (${HIP_PLATFORM}, hcc)
|
||||
HCBLAS_ROOT?= $(wildcard /opt/rocm/hcblas)
|
||||
HIPCC_FLAGS += -stdlib=libc++ -I$(HCBLAS_ROOT)/include
|
||||
LIBS = -L$(HCBLAS_ROOT)/lib -lhcblas
|
||||
LIBS = -L$(HCBLAS_ROOT)/lib -lhipblas -rpath $(HIP_PATH)/lib
|
||||
endif
|
||||
|
||||
|
||||
|
||||
@@ -0,0 +1,58 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#include <hip/device_functions.h>
|
||||
|
||||
extern "C" float __hip_int_as_float(int);
|
||||
|
||||
typedef struct {
|
||||
signed int hi;
|
||||
signed int lo;
|
||||
} __hip_signed_2;
|
||||
|
||||
typedef struct {
|
||||
union {
|
||||
double d;
|
||||
long long int lli;
|
||||
__hip_signed_2 s2;
|
||||
} ;
|
||||
} __hip_64bit_struct;
|
||||
|
||||
typedef struct {
|
||||
union {
|
||||
float f;
|
||||
unsigned int ui;
|
||||
signed int si;
|
||||
};
|
||||
} __hip_32bit_struct;
|
||||
|
||||
__device__ float __int_as_float (int x) {
|
||||
__hip_32bit_struct s;
|
||||
s.si = x;
|
||||
return s.f;
|
||||
}
|
||||
|
||||
__device__ double __hiloint2double (int hi, int lo) {
|
||||
__hip_64bit_struct s;
|
||||
s.s2.hi = hi;
|
||||
s.s2.lo = lo;
|
||||
return s.d;
|
||||
}
|
||||
|
||||
|
||||
@@ -22,6 +22,8 @@ THE SOFTWARE.
|
||||
|
||||
#include"hip/hip_fp16.h"
|
||||
|
||||
#if __CLANG_VERSION__ == 35
|
||||
|
||||
static const unsigned sign_val = 0x8000;
|
||||
static const __half __half_value_one_float = {0x3C00};
|
||||
static const __half __half_value_zero_float = {0x0};
|
||||
@@ -372,3 +374,4 @@ __device__ __half2 __lowhigh2highlow(const __half2 a){
|
||||
__device__ __half2 __low2half2(const __half2 a, const __half2 b){
|
||||
return {a.q, b.q};
|
||||
}
|
||||
#endif
|
||||
|
||||
@@ -34,6 +34,7 @@ THE SOFTWARE.
|
||||
#endif
|
||||
|
||||
#define USE_DISPATCH_HSA_KERNEL 1
|
||||
#define USE_IPC 0
|
||||
//
|
||||
|
||||
|
||||
@@ -375,7 +376,9 @@ struct LockedBase {
|
||||
class ihipIpcMemHandle_t
|
||||
{
|
||||
public:
|
||||
#if USE_IPC
|
||||
hsa_amd_ipc_memory_t ipc_handle; ///< ipc memory handle on ROCr
|
||||
#endif
|
||||
char reserved[HIP_IPC_HANDLE_SIZE];
|
||||
size_t psize;
|
||||
};
|
||||
|
||||
@@ -2,12 +2,12 @@ target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-i64:
|
||||
target triple = "amdgcn--amdhsa"
|
||||
|
||||
|
||||
define void @__threadfence() #1 {
|
||||
define linkonce_odr spir_func void @__threadfence() #1 {
|
||||
fence syncscope(2) seq_cst
|
||||
ret void
|
||||
}
|
||||
|
||||
define void @__threadfence_block() #1 {
|
||||
define linkonce_odr spir_func void @__threadfence_block() #1 {
|
||||
fence syncscope(3) seq_cst
|
||||
ret void
|
||||
}
|
||||
@@ -34,4 +34,29 @@ define linkonce_odr spir_func i32 @__rocm_hadd(i32 %in1, i32 %in2) {
|
||||
ret i32 %val
|
||||
}
|
||||
|
||||
define linkonce_odr spir_func half @__hip_hadd_gfx803(half %a, half %b) #1 {
|
||||
%val = tail call half asm "v_add_f16 $0, $1, $2","=v,v,v"(half %a, half %b)
|
||||
ret half %val
|
||||
}
|
||||
|
||||
define linkonce_odr spir_func half @__hip_hfma_gfx803(half %a, half %b, half %c) #1 {
|
||||
%val = tail call half asm "v_fma_f16 $0, $1, $2, $3","=v,v,v,v"(half %a, half %b, half %c)
|
||||
ret half %val
|
||||
}
|
||||
|
||||
define linkonce_odr spir_func half @__hip_hmul_gfx803(half %a, half %b) #1 {
|
||||
%val = tail call half asm "v_mul_f16 $0, $1, $2","=v,v,v"(half %a, half %b)
|
||||
ret half %val
|
||||
}
|
||||
|
||||
define linkonce_odr spir_func half @__hip_hsub_gfx803(half %a, half %b) #1 {
|
||||
%val = tail call half asm "v_sub_f16 $0, $1, $2","=v,v,v"(half %a, half %b)
|
||||
ret half %val
|
||||
}
|
||||
|
||||
define linkonce_odr spir_func i32 @__hip_hadd2_gfx803(i32 %a, i32 %b) #1 {
|
||||
%val = tail call i32 asm "v_add_f16_sdwa $0, $1, $2 dst_sel:WORD_0 dst_unused:UNUSED_PRESERVE src0_sel:WORD_0 src1_sel:WORD_0","=v,v,v"(i32 %a, i32 %b)
|
||||
ret i32 %val
|
||||
}
|
||||
|
||||
attributes #1 = { alwaysinline nounwind }
|
||||
|
||||
@@ -1056,11 +1056,15 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr){
|
||||
// Save the size of the pointer to hipIpcMemHandle
|
||||
(*handle)->psize = psize;
|
||||
|
||||
#if USE_IPC
|
||||
// Create HSA ipc memory
|
||||
hsa_status_t hsa_status =
|
||||
hsa_amd_ipc_memory_create(devPtr, psize, &(*handle)->ipc_handle);
|
||||
if(hsa_status!= HSA_STATUS_SUCCESS)
|
||||
hipStatus = hipErrorMemoryAllocation;
|
||||
#else
|
||||
hipStatus = hipErrorRuntimeOther;
|
||||
#endif
|
||||
|
||||
return hipStatus;
|
||||
}
|
||||
@@ -1069,6 +1073,7 @@ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned
|
||||
// HIP_INIT_API ( devPtr, handle.handle , flags);
|
||||
hipError_t hipStatus = hipSuccess;
|
||||
|
||||
#if USE_IPC
|
||||
// Get the current device agent.
|
||||
hc::accelerator acc;
|
||||
hsa_agent_t *agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
|
||||
@@ -1080,7 +1085,9 @@ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned
|
||||
hsa_amd_ipc_memory_attach(&handle->ipc_handle, handle->psize, 1, agent, devPtr);
|
||||
if(hsa_status != HSA_STATUS_SUCCESS)
|
||||
hipStatus = hipErrorMapBufferObjectFailed;
|
||||
|
||||
#else
|
||||
hipStatus = hipErrorRuntimeOther;
|
||||
#endif
|
||||
return hipStatus;
|
||||
}
|
||||
|
||||
@@ -1088,10 +1095,14 @@ hipError_t hipIpcCloseMemHandle(void *devPtr){
|
||||
HIP_INIT_API ( devPtr );
|
||||
hipError_t hipStatus = hipSuccess;
|
||||
|
||||
#if USE_IPC
|
||||
hsa_status_t hsa_status =
|
||||
hsa_amd_ipc_memory_detach(devPtr);
|
||||
if(hsa_status != HSA_STATUS_SUCCESS)
|
||||
return hipErrorInvalidResourceHandle;
|
||||
#else
|
||||
hipStatus = hipErrorRuntimeOther;
|
||||
#endif
|
||||
return hipStatus;
|
||||
}
|
||||
|
||||
|
||||
مرجع در شماره جدید
Block a user