Merge 'master' into 'amd-master'
Change-Id: I83daa8e77a448c4ae80ad9007eac5f786f8f4954
Dieser Commit ist enthalten in:
Jenkins
+321
-6
@@ -160,6 +160,8 @@ while (@ARGV) {
|
||||
undef $/; # Read whole file at once, so we can match newlines.
|
||||
while (<INFILE>)
|
||||
{
|
||||
# chomp;
|
||||
# next if /^(\s*(#.*)?)?$/;
|
||||
$ft{'error'} += s/\bcudaGetErrorName\b/hipGetErrorName/g;
|
||||
$ft{'error'} += s/\bcudaGetErrorString\b/hipGetErrorString/g;
|
||||
$ft{'error'} += s/\bcudaGetLastError\b/hipGetLastError/g;
|
||||
@@ -363,6 +365,7 @@ while (@ARGV) {
|
||||
$ft{'execution'} += s/\bcuFuncGetAttribute\b/hipFuncGetAttribute/g;
|
||||
$ft{'execution'} += s/\bcuLaunchKernel\b/hipModuleLaunchKernel/g;
|
||||
$ft{'execution'} += s/\bcudaConfigureCall\b/hipConfigureCall/g;
|
||||
$ft{'execution'} += s/\bcudaFuncGetAttributes\b/hipFuncGetAttributes/g;
|
||||
$ft{'execution'} += s/\bcudaLaunch\b/hipLaunchByPtr/g;
|
||||
$ft{'execution'} += s/\bcudaLaunchCooperativeKernel\b/hipLaunchCooperativeKernel/g;
|
||||
$ft{'execution'} += s/\bcudaLaunchCooperativeKernelMultiDevice\b/hipLaunchCooperativeKernelMultiDevice/g;
|
||||
@@ -1294,6 +1297,7 @@ while (@ARGV) {
|
||||
$ft{'numeric_literal'} += s/\bCU_CTX_SCHED_MASK\b/hipDeviceScheduleMask/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_CTX_SCHED_SPIN\b/hipDeviceScheduleSpin/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_CTX_SCHED_YIELD\b/hipDeviceScheduleYield/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY\b/hipDeviceAttributeCanMapHostMemory/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_CLOCK_RATE\b/hipDeviceAttributeClockRate/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR\b/hipDeviceAttributeComputeCapabilityMajor/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR\b/hipDeviceAttributeComputeCapabilityMinor/g;
|
||||
@@ -1301,8 +1305,10 @@ while (@ARGV) {
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_CONCURRENT_KERNELS\b/hipDeviceAttributeConcurrentKernels/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_COOPERATIVE_LAUNCH\b/hipDeviceAttributeCooperativeLaunch/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_COOPERATIVE_MULTI_DEVICE_LAUNCH\b/hipDeviceAttributeCooperativeMultiDeviceLaunch/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_ECC_ENABLED\b/hipDeviceAttributeEccEnabled/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH\b/hipDeviceAttributeMemoryBusWidth/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_INTEGRATED\b/hipDeviceAttributeIntegrated/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT\b/hipDeviceAttributeKernelExecTimeout/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_L2_CACHE_SIZE\b/hipDeviceAttributeL2CacheSize/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH\b/hipDeviceAttributeMaxTexture1DWidth/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT\b/hipDeviceAttributeMaxTexture2DHeight/g;
|
||||
@@ -1316,6 +1322,7 @@ while (@ARGV) {
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X\b/hipDeviceAttributeMaxGridDimX/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Y\b/hipDeviceAttributeMaxGridDimY/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Z\b/hipDeviceAttributeMaxGridDimZ/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAX_PITCH\b/hipDeviceAttributeMaxPitch/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK\b/hipDeviceAttributeMaxRegistersPerBlock/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK\b/hipDeviceAttributeMaxSharedMemoryPerBlock/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_MULTIPROCESSOR\b/hipDeviceAttributeMaxSharedMemoryPerMultiprocessor/g;
|
||||
@@ -1328,6 +1335,7 @@ while (@ARGV) {
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_PCI_DEVICE_ID\b/hipDeviceAttributePciDeviceId/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_REGISTERS_PER_BLOCK\b/hipDeviceAttributeMaxRegistersPerBlock/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_SHARED_MEMORY_PER_BLOCK\b/hipDeviceAttributeMaxSharedMemoryPerBlock/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT\b/hipDeviceAttributeTextureAlignment/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY\b/hipDeviceAttributeTotalConstantMemory/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_WARP_SIZE\b/hipDeviceAttributeWarpSize/g;
|
||||
$ft{'numeric_literal'} += s/\bCU_EVENT_BLOCKING_SYNC\b/hipEventBlockingSync/g;
|
||||
@@ -1755,6 +1763,8 @@ while (@ARGV) {
|
||||
$ft{'kernel_func'} += countSupportedDeviceFunctions();
|
||||
}
|
||||
|
||||
$ft{'memory'} += transformSymbolFunctions();
|
||||
|
||||
# Print it!
|
||||
# TODO - would like to move this code outside loop but it uses $_ which contains the whole file.
|
||||
unless ($no_output) {
|
||||
@@ -1809,18 +1819,323 @@ if ($count_conversions) {
|
||||
}
|
||||
}
|
||||
|
||||
sub transformSymbolFunctions
|
||||
{
|
||||
my $m = 0;
|
||||
foreach $func (
|
||||
"hipMemcpyToSymbol",
|
||||
"hipMemcpyToSymbolAsync"
|
||||
)
|
||||
{
|
||||
$m += s/(?<!\/\/ CHECK: )($func)\s*\(\s*([^,]+)\s*,/$func\(HIP_SYMBOL\($2\),/g
|
||||
}
|
||||
foreach $func (
|
||||
"hipGetSymbolAddress",
|
||||
"hipGetSymbolSize",
|
||||
"hipMemcpyFromSymbol",
|
||||
"hipMemcpyFromSymbolAsync"
|
||||
)
|
||||
{
|
||||
$m += s/(?<!\/\/ CHECK: )($func)\s*\(\s*([^,]+)\s*,\s*([^,\)]+)\s*(,\s*|\))\s*/$func\($2, HIP_SYMBOL\($3\)$4/g;
|
||||
}
|
||||
return $m;
|
||||
}
|
||||
|
||||
sub countSupportedDeviceFunctions
|
||||
{
|
||||
my $m = 0;
|
||||
# TODO: list all of the supported functions
|
||||
# TODO: split the list on math, device, and maybe fp16
|
||||
foreach $func (
|
||||
# Synchronization:
|
||||
"__brev",
|
||||
"__brevll",
|
||||
"__byte_perm",
|
||||
"__clz",
|
||||
"__clzll",
|
||||
"__cosf",
|
||||
"__double2int_rz",
|
||||
"__double2ll_rz",
|
||||
"__double2uint_rz",
|
||||
"__double2ull_rz",
|
||||
"__exp10f",
|
||||
"__expf",
|
||||
"__fadd_rd",
|
||||
"__fadd_rn",
|
||||
"__fadd_ru",
|
||||
"__fadd_rz",
|
||||
"__fdiv_rd",
|
||||
"__fdiv_rn",
|
||||
"__fdiv_ru",
|
||||
"__fdiv_rz",
|
||||
"__fdividef",
|
||||
"__ffs",
|
||||
"__ffsll",
|
||||
"__float2int_rd",
|
||||
"__float2int_rn",
|
||||
"__float2int_ru",
|
||||
"__float2int_rz",
|
||||
"__float2ll_rd",
|
||||
"__float2ll_rn",
|
||||
"__float2ll_ru",
|
||||
"__float2ll_rz",
|
||||
"__float2uint_rd",
|
||||
"__float2uint_rn",
|
||||
"__float2uint_ru",
|
||||
"__float2uint_rz",
|
||||
"__float2ull_rd",
|
||||
"__float2ull_rn",
|
||||
"__float2ull_ru",
|
||||
"__float2ull_rz",
|
||||
"__float_as_int",
|
||||
"__float_as_uint",
|
||||
"__fmaf_rd",
|
||||
"__fmaf_rn",
|
||||
"__fmaf_ru",
|
||||
"__fmaf_rz",
|
||||
"__fmul_rd",
|
||||
"__fmul_rn",
|
||||
"__fmul_ru",
|
||||
"__fmul_rz",
|
||||
"__frcp_rd",
|
||||
"__frcp_rn",
|
||||
"__frcp_ru",
|
||||
"__frcp_rz",
|
||||
"__frsqrt_rn",
|
||||
"__fsqrt_rd",
|
||||
"__fsqrt_rn",
|
||||
"__fsqrt_ru",
|
||||
"__fsqrt_rz",
|
||||
"__fsub_rd",
|
||||
"__fsub_rn",
|
||||
"__fsub_ru",
|
||||
"__fsub_rz",
|
||||
"__hadd",
|
||||
"__int2float_rd",
|
||||
"__int2float_rn",
|
||||
"__int2float_ru",
|
||||
"__int2float_rz",
|
||||
"__int_as_float",
|
||||
"__ll2float_rd",
|
||||
"__ll2float_rn",
|
||||
"__ll2float_ru",
|
||||
"__ll2float_rz",
|
||||
"__log10f",
|
||||
"__log2f",
|
||||
"__logf",
|
||||
"__mul24",
|
||||
"__mul64hi",
|
||||
"__mulhi",
|
||||
"__popc",
|
||||
"__popcll",
|
||||
"__powf",
|
||||
"__rhadd",
|
||||
"__sad",
|
||||
"__saturatef",
|
||||
"__sincosf",
|
||||
"__sinf",
|
||||
"__syncthreads",
|
||||
"__tanf",
|
||||
"__threadfence",
|
||||
"__threadfence_block",
|
||||
"__uhadd",
|
||||
"__uint2float_rd",
|
||||
"__uint2float_rn",
|
||||
"__uint2float_ru",
|
||||
"__uint2float_rz",
|
||||
"__uint_as_float",
|
||||
"__ull2float_rd",
|
||||
"__ull2float_rn",
|
||||
"__ull2float_ru",
|
||||
"__ull2float_rz",
|
||||
"__umul24",
|
||||
"__umul64hi",
|
||||
"__umulhi",
|
||||
"__urhadd",
|
||||
"__usad",
|
||||
"abs",
|
||||
"acos",
|
||||
"acosf",
|
||||
"acosh",
|
||||
"acoshf",
|
||||
"asin",
|
||||
"asinf",
|
||||
"asinh",
|
||||
"asinhf",
|
||||
"atan",
|
||||
"atan2",
|
||||
"atan2f",
|
||||
"atanf",
|
||||
"atanh",
|
||||
"atanhf",
|
||||
"cbrt",
|
||||
"cbrtf",
|
||||
"ceil",
|
||||
"ceilf",
|
||||
"copysign",
|
||||
"copysignf",
|
||||
"cos",
|
||||
"cosf",
|
||||
"cosh",
|
||||
"coshf",
|
||||
"cospi",
|
||||
"cospif",
|
||||
"cyl_bessel_i0",
|
||||
"cyl_bessel_i0f",
|
||||
"cyl_bessel_i1",
|
||||
"cyl_bessel_i1f",
|
||||
"erf",
|
||||
"erfc",
|
||||
"erfcf",
|
||||
"erfcinv",
|
||||
"erfcinvf",
|
||||
"erfcx",
|
||||
"erfcxf",
|
||||
"erff",
|
||||
"erfinv",
|
||||
"erfinvf",
|
||||
"exp",
|
||||
"exp10",
|
||||
"exp10f",
|
||||
"exp2",
|
||||
"exp2f",
|
||||
"expf",
|
||||
"expm1",
|
||||
"expm1f",
|
||||
"fabs",
|
||||
"fabsf",
|
||||
"fdim",
|
||||
"fdimf",
|
||||
"fdivide",
|
||||
"fdividef",
|
||||
"floor",
|
||||
"floorf",
|
||||
"fma",
|
||||
"fmaf",
|
||||
"fmax",
|
||||
"fmaxf",
|
||||
"fmin",
|
||||
"fminf",
|
||||
"fmod",
|
||||
"fmodf",
|
||||
"frexp",
|
||||
"frexpf",
|
||||
"hypot",
|
||||
"hypotf",
|
||||
"ilogb",
|
||||
"ilogbf",
|
||||
"isfinite",
|
||||
"isinf",
|
||||
"isnan",
|
||||
"j0",
|
||||
"j0f",
|
||||
"j1",
|
||||
"j1f",
|
||||
"jn",
|
||||
"jnf",
|
||||
"labs",
|
||||
"ldexp",
|
||||
"ldexpf",
|
||||
"lgamma",
|
||||
"lgammaf",
|
||||
"llabs",
|
||||
"llrint",
|
||||
"llrintf",
|
||||
"llround",
|
||||
"llroundf",
|
||||
"log",
|
||||
"log10",
|
||||
"log10f",
|
||||
"log1p",
|
||||
"log1pf",
|
||||
"log2",
|
||||
"log2f",
|
||||
"logb",
|
||||
"logbf",
|
||||
"logf",
|
||||
"lrint",
|
||||
"lrintf",
|
||||
"lround",
|
||||
"lroundf",
|
||||
"max",
|
||||
"min",
|
||||
"modf",
|
||||
"modff",
|
||||
"nan",
|
||||
"nanf",
|
||||
"nearbyint",
|
||||
"nearbyintf",
|
||||
"nextafter",
|
||||
"nextafterf",
|
||||
"norm",
|
||||
"norm3d",
|
||||
"norm3df",
|
||||
"norm4d",
|
||||
"norm4df",
|
||||
"normcdf",
|
||||
"normcdff",
|
||||
"normcdfinv",
|
||||
"normcdfinvf",
|
||||
"normf",
|
||||
"pow",
|
||||
"powf",
|
||||
"rcbrt",
|
||||
"rcbrtf",
|
||||
"remainder",
|
||||
"remainderf",
|
||||
"remquo",
|
||||
"remquof",
|
||||
"rhypot",
|
||||
"rhypotf",
|
||||
"rint",
|
||||
"rintf",
|
||||
"rnorm",
|
||||
"rnorm3d",
|
||||
"rnorm3df",
|
||||
"rnorm4d",
|
||||
"rnorm4df",
|
||||
"rnormf",
|
||||
"round",
|
||||
"roundf",
|
||||
"rsqrt",
|
||||
"rsqrtf",
|
||||
"scalbln",
|
||||
"scalblnf",
|
||||
"scalbn",
|
||||
"scalbnf",
|
||||
"signbit",
|
||||
"sin",
|
||||
"sincos",
|
||||
"sincosf",
|
||||
"sincospi",
|
||||
"sincospif",
|
||||
"sinf",
|
||||
"sinh",
|
||||
"sinhf",
|
||||
"sinpi",
|
||||
"sinpif",
|
||||
"sqrt",
|
||||
"sqrtf",
|
||||
"tan",
|
||||
"tanf",
|
||||
"tanh",
|
||||
"tanhf",
|
||||
"tgamma",
|
||||
"tgammaf",
|
||||
"trunc",
|
||||
"truncf",
|
||||
"y0",
|
||||
"y0f",
|
||||
"y1",
|
||||
"y1f",
|
||||
"yn",
|
||||
"ynf"
|
||||
)
|
||||
{
|
||||
# match device func at the beginning of a word, but not if it already has a namespace qualifier ('::') :
|
||||
$m += m/[:]?[:]?\b($func)\b(\w*\()/g;
|
||||
# match device function from the list, except those, which have a namespace prefix (aka somenamespace::umin(...));
|
||||
# function with only global namespace qualifier '::' (aka ::umin(...)) should be treated as a device function (and warned as well as without such qualifier);
|
||||
my $mt_namespace = m/(\w+)::($func)\s*\(\s*.*\s*\)/g;
|
||||
my $mt = m/($func)\s*\(\s*.*\s*\)/g;
|
||||
if ($mt && !$mt_namespace) {
|
||||
$m += $mt;
|
||||
}
|
||||
}
|
||||
return $m;
|
||||
}
|
||||
@@ -1956,7 +2271,7 @@ sub warnUnsupportedDeviceFunctions
|
||||
"umul24"
|
||||
)
|
||||
{
|
||||
# match device function from the list of unsupported, except those, which have a namespace prefix (aka somenamespace::umin(...));
|
||||
# match device function from the list, except those, which have a namespace prefix (aka somenamespace::umin(...));
|
||||
# function with only global namespace qualifier '::' (aka ::umin(...)) should be treated as a device function (and warned as well as without such qualifier);
|
||||
my $mt_namespace = m/(\w+)::($func)\s*\(\s*.*\s*\)/g;
|
||||
my $mt = m/($func)\s*\(\s*.*\s*\)/g;
|
||||
|
||||
@@ -101,7 +101,7 @@
|
||||
|
||||
| **CUDA** | **HIP** |**CUDA version\***|
|
||||
|-----------------------------------------------------------|---------------------------------------|:----------------:|
|
||||
| `cudaFuncGetAttributes` | |
|
||||
| `cudaFuncGetAttributes` |`hipFuncGetAttributes` |
|
||||
| `cudaFuncSetAttribute` | | 9.0 |
|
||||
| `cudaFuncSetCacheConfig` |`hipFuncSetCacheConfig` |
|
||||
| `cudaFuncSetSharedMemConfig` | |
|
||||
|
||||
@@ -32,11 +32,10 @@
|
||||
## <a name="dependencies"></a> Dependencies
|
||||
|
||||
`hipify-clang` requires:
|
||||
1. [**LLVM+CLANG**](http://releases.llvm.org) of at least version [3.8.0](http://releases.llvm.org/download.html#3.8.0); the latest stable and recommended release: [**6.0.1**](http://releases.llvm.org/download.html#6.0.1) on **Windows**, and [**8.0.1**](http://releases.llvm.org/download.html#8.0.1) on **Linux**.
|
||||
|
||||
2. **CUDA** at least version [7.0](https://developer.nvidia.com/cuda-toolkit-70), the latest supported version is [**9.0**](https://developer.nvidia.com/cuda-90-download-archive) on **Windows**, and [**10.0**](https://developer.nvidia.com/cuda-10.0-download-archive) on **Linux**.
|
||||
1. [**LLVM+CLANG**](http://releases.llvm.org) of at least version [3.8.0](http://releases.llvm.org/download.html#3.8.0); the latest stable and recommended release: [**9.0.0**](http://releases.llvm.org/download.html#9.0.0).
|
||||
|
||||
If the target CUDA is [9.1](https://developer.nvidia.com/cuda-91-download-archive), [9.2](https://developer.nvidia.com/cuda-92-download-archive) or [10.0](https://developer.nvidia.com/cuda-10.0-download-archive), to work on Windows you may apply patches* for LLVM: [7.0.0](patches/patch_for_clang_7.0.0_bug_38811.zip), [7.0.1](patches/patch_for_clang_7.0.1_bug_38811.zip), [7.1.0](patches/patch_for_clang_7.1.0_bug_38811.zip), [8.0.0](patches/patch_for_clang_8.0.0_bug_38811.zip), [8.0.1](patches/patch_for_clang_8.0.1_bug_38811.zip).
|
||||
2. [**CUDA**](https://developer.nvidia.com/cuda-downloads) at least version [7.0](https://developer.nvidia.com/cuda-toolkit-70), the latest supported version is [**10.1 Update 2**](https://developer.nvidia.com/cuda-downloads).
|
||||
|
||||
| **LLVM release version** | **CUDA latest supported version** | **Windows** | **Linux** |
|
||||
|:----------------------------------------------------------:|:-------------------------------------------------------------------:|:------------:|:---------:|
|
||||
@@ -50,20 +49,20 @@ If the target CUDA is [9.1](https://developer.nvidia.com/cuda-91-download-archiv
|
||||
| [5.0.1](http://releases.llvm.org/download.html#5.0.1) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
|
||||
| [5.0.2](http://releases.llvm.org/download.html#5.0.2) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
|
||||
| [6.0.0](http://releases.llvm.org/download.html#6.0.0) | [9.0](https://developer.nvidia.com/cuda-90-download-archive) | + | + |
|
||||
| [**6.0.1**](http://releases.llvm.org/download.html#6.0.1) | [**9.0**](https://developer.nvidia.com/cuda-90-download-archive) | + <br/> **LATEST STABLE RELEASE** | + |
|
||||
| [6.0.1](http://releases.llvm.org/download.html#6.0.1) | [9.0](https://developer.nvidia.com/cuda-90-download-archive) | + | + |
|
||||
| [7.0.0](http://releases.llvm.org/download.html#7.0.0) | [9.2](https://developer.nvidia.com/cuda-92-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.0.0_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
|
||||
| [7.0.1](http://releases.llvm.org/download.html#7.0.1) | [9.2](https://developer.nvidia.com/cuda-92-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.0.1_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
|
||||
| [7.1.0](http://releases.llvm.org/download.html#7.1.0) | [9.2](https://developer.nvidia.com/cuda-92-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.1.0_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
|
||||
| [8.0.0](http://releases.llvm.org/download.html#8.0.0) | [10.0](https://developer.nvidia.com/cuda-10.0-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_8.0.0_bug_38811.zip)*</br> | + |
|
||||
| [**8.0.1**](http://releases.llvm.org/download.html#8.0.1) | [**10.0**](https://developer.nvidia.com/cuda-10.0-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_8.0.1_bug_38811.zip)*</br> | + <br/> **LATEST STABLE RELEASE** |
|
||||
| 9.0.0 | 10.1 | <br/> LLVM 9.0.0 <br/> is not yet released | <br/> LLVM 9.0.0 <br/> is not yet released |
|
||||
| [8.0.0](http://releases.llvm.org/download.html#8.0.0) | [10.0](https://developer.nvidia.com/cuda-10.0-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_8.0.0_bug_38811.zip)*</br> | + |
|
||||
| [8.0.1](http://releases.llvm.org/download.html#8.0.1) | [10.0](https://developer.nvidia.com/cuda-10.0-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_8.0.1_bug_38811.zip)*</br> | + |
|
||||
| [**9.0.0**](http://releases.llvm.org/download.html#9.0.0) | [**10.1**](https://developer.nvidia.com/cuda-downloads) | + <br/> **LATEST STABLE RELEASE** | + <br/> **LATEST STABLE RELEASE** |
|
||||
|
||||
`*` Download the patch and unpack it into your LLVM distributive directory; a few header files will be overwritten; rebuilding of LLVM is not needed.
|
||||
|
||||
In most cases, you can get a suitable version of LLVM+CLANG with your package manager.
|
||||
|
||||
Failing that or having multiple versions of LLVM, you can [download a release archive](http://releases.llvm.org/), build or install it, and set
|
||||
[CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.12/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it; for instance: `-DCMAKE_PREFIX_PATH=f:\LLVM\6.0.1\dist`
|
||||
[CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.12/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it; for instance: `-DCMAKE_PREFIX_PATH=f:\LLVM\9.0.0\dist`
|
||||
|
||||
## <a name="build-and-install"></a> Build and install
|
||||
|
||||
@@ -83,10 +82,10 @@ cmake \
|
||||
|
||||
make -j install
|
||||
```
|
||||
On Windows, the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"`; the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make.`
|
||||
On Windows, the following option should be specified for `cmake` at first place: `-G "Visual Studio 16 2019 Win64"`; the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make.`
|
||||
|
||||
Debug build type `-DCMAKE_BUILD_TYPE=Debug` is also supported and tested; `LLVM+CLANG` should be built in `Debug` mode as well.
|
||||
64-bit build mode `-Thost=x64` is supported as well; `LLVM+CLANG` should be built in 64-bit mode as well.
|
||||
64-bit build mode (`-Thost=x64` on Windows) is also supported; `LLVM+CLANG` should be built in 64-bit mode as well.
|
||||
|
||||
The binary can then be found at `./dist/bin/hipify-clang`.
|
||||
|
||||
@@ -97,8 +96,8 @@ The binary can then be found at `./dist/bin/hipify-clang`.
|
||||
**LLVM+CLANG should be built from sources, pre-built binaries are not exhaustive for testing.**
|
||||
|
||||
To run it:
|
||||
1. Download [`LLVM`](http://releases.llvm.org/6.0.1/llvm-6.0.1.src.tar.xz)+[`CLANG`](http://releases.llvm.org/6.0.1/cfe-6.0.1.src.tar.xz) sources.
|
||||
2. Build [`LLVM+CLANG`](http://llvm.org/docs/CMake.html):
|
||||
1. download [`LLVM`](http://releases.llvm.org/9.0.0/llvm-9.0.0.src.tar.xz)+[`CLANG`](http://releases.llvm.org/9.0.0/cfe-9.0.0.src.tar.xz) sources;
|
||||
2. build [`LLVM+CLANG`](http://llvm.org/docs/CMake.html):
|
||||
```shell
|
||||
cd llvm
|
||||
mkdir build dist
|
||||
@@ -119,7 +118,7 @@ To run it:
|
||||
- **Windows**:
|
||||
```shell
|
||||
cmake \
|
||||
-G "Visual Studio 15 2017 Win64" \
|
||||
-G "Visual Studio 16 2019 Win64" \
|
||||
-DCMAKE_INSTALL_PREFIX=../dist \
|
||||
-DLLVM_SOURCE_DIR=../llvm \
|
||||
-DLLVM_TARGETS_TO_BUILD="NVPTX" \
|
||||
@@ -128,26 +127,26 @@ To run it:
|
||||
../llvm
|
||||
```
|
||||
|
||||
Run `Visual Studio 15 2017`, open the generated `LLVM.sln`, build all, build project `INSTALL`.
|
||||
Run `Visual Studio 16 2019`, open the generated `LLVM.sln`, build all, build project `INSTALL`.
|
||||
|
||||
|
||||
3. Ensure [`CUDA`](https://developer.nvidia.com/cuda-toolkit-archive) of minimum version 7.0 is installed.
|
||||
|
||||
* Having multiple CUDA installations to choose a particular version the `DCUDA_TOOLKIT_ROOT_DIR` option should be specified:
|
||||
|
||||
- Linux: `-DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-10.0`
|
||||
- Linux: `-DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-10.1`
|
||||
|
||||
- Windows: `-DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0"`
|
||||
- Windows: `-DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v10.1"`
|
||||
|
||||
`-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0"`
|
||||
`-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v10.1"`
|
||||
|
||||
4. Ensure [`cuDNN`](https://developer.nvidia.com/rdp/cudnn-archive) of the version corresponding to CUDA's version is installed.
|
||||
|
||||
* Path to cuDNN should be specified by the `CUDA_DNN_ROOT_DIR` option:
|
||||
|
||||
- Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.0-v7.6.3.30`
|
||||
- Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.3.30`
|
||||
|
||||
- Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.6.3.30`
|
||||
- Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.3.30`
|
||||
|
||||
5. Ensure [`python`](https://www.python.org/downloads) of minimum required version 2.7 is installed.
|
||||
|
||||
@@ -155,15 +154,15 @@ To run it:
|
||||
|
||||
* Install `lit` into `python`:
|
||||
|
||||
- Linux: `python /srv/git/LLVM/8.0.1/llvm/utils/lit/setup.py install`
|
||||
- Linux: `python /srv/git/LLVM/9.0.0/llvm/utils/lit/setup.py install`
|
||||
|
||||
- Windows: `python f:/LLVM/6.0.1/llvm/utils/lit/setup.py install`
|
||||
- Windows: `python f:/LLVM/9.0.0/llvm/utils/lit/setup.py install`
|
||||
|
||||
* Starting with LLVM 6.0.1 path to `llvm-lit` python script should be specified by the `LLVM_EXTERNAL_LIT` option:
|
||||
|
||||
- Linux: `-DLLVM_EXTERNAL_LIT=/srv/git/LLVM/8.0.1/build/bin/llvm-lit`
|
||||
- Linux: `-DLLVM_EXTERNAL_LIT=/srv/git/LLVM/9.0.0/build/bin/llvm-lit`
|
||||
|
||||
- Windows: `-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.1/build/Release/bin/llvm-lit.py`
|
||||
- Windows: `-DLLVM_EXTERNAL_LIT=f:/LLVM/9.0.0/build/Release/bin/llvm-lit.py`
|
||||
|
||||
7. Set `HIPIFY_CLANG_TESTS` option turned on: `-DHIPIFY_CLANG_TESTS=1`.
|
||||
|
||||
@@ -175,7 +174,7 @@ To run it:
|
||||
|
||||
- Linux: `make test-hipify`.
|
||||
|
||||
- Windows: run `Visual Studio 15 2017`, open the generated `hipify-clang.sln`, build project `test-hipify`.
|
||||
- Windows: run `Visual Studio 16 2019`, open the generated `hipify-clang.sln`, build project `test-hipify`.
|
||||
|
||||
### <a name="linux"></a >Linux
|
||||
|
||||
@@ -183,7 +182,7 @@ On Linux the following configurations are tested:
|
||||
|
||||
Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.6.3.30
|
||||
|
||||
Ubuntu 16-18: LLVM 8.0.0 - 8.0.1, CUDA 8.0 - 10.0, cudnn-5.1.10 - cudnn-7.6.3.30
|
||||
Ubuntu 16-18: LLVM 8.0.0 - 9.0.0, CUDA 8.0 - 10.1, cudnn-5.1.10 - cudnn-7.6.3.30
|
||||
|
||||
Build system for the above configurations:
|
||||
|
||||
@@ -196,16 +195,16 @@ cmake
|
||||
-DHIPIFY_CLANG_TESTS=1 \
|
||||
-DCMAKE_BUILD_TYPE=Release \
|
||||
-DCMAKE_INSTALL_PREFIX=../dist \
|
||||
-DCMAKE_PREFIX_PATH=/srv/git/LLVM/8.0.1/dist \
|
||||
-DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-10.0 \
|
||||
-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.0-v7.6.3.30 \
|
||||
-DLLVM_EXTERNAL_LIT=/srv/git/LLVM/8.0.1/build/bin/llvm-lit \
|
||||
-DCMAKE_PREFIX_PATH=/srv/git/LLVM/9.0.0/dist \
|
||||
-DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-10.1 \
|
||||
-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.3.30 \
|
||||
-DLLVM_EXTERNAL_LIT=/srv/git/LLVM/9.0.0/build/bin/llvm-lit \
|
||||
..
|
||||
```
|
||||
*A corresponding successful output:*
|
||||
```shell
|
||||
-- The C compiler identification is GNU 5.4.0
|
||||
-- The CXX compiler identification is GNU 5.4.0
|
||||
-- The C compiler identification is GNU 7.4.0
|
||||
-- The CXX compiler identification is GNU 7.4.0
|
||||
-- Check for working C compiler: /usr/bin/cc
|
||||
-- Check for working C compiler: /usr/bin/cc -- works
|
||||
-- Detecting C compiler ABI info
|
||||
@@ -218,14 +217,14 @@ cmake
|
||||
-- Detecting CXX compiler ABI info - done
|
||||
-- Detecting CXX compile features
|
||||
-- Detecting CXX compile features - done
|
||||
-- Found LLVM 8.0.1:
|
||||
-- - CMake module path: /srv/git/LLVM/8.0.1/dist/lib/cmake/llvm
|
||||
-- - Include path : /srv/git/LLVM/8.0.1/dist/include
|
||||
-- - Binary path : /srv/git/LLVM/8.0.1/dist/bin
|
||||
-- Found LLVM 9.0.0:
|
||||
-- - CMake module path: /srv/git/LLVM/9.0.0/dist/lib/cmake/llvm
|
||||
-- - Include path : /srv/git/LLVM/9.0.0/dist/include
|
||||
-- - Binary path : /srv/git/LLVM/9.0.0/dist/bin
|
||||
-- Linker detection: GNU ld
|
||||
-- Found PythonInterp: /usr/bin/python2.7 (found suitable version "2.7.12", minimum required is "2.7")
|
||||
-- Found lit: /usr/local/bin/lit
|
||||
-- Found FileCheck: /srv/git/LLVM/8.0.1/dist/bin/FileCheck
|
||||
-- Found FileCheck: /srv/git/LLVM/9.0.0/dist/bin/FileCheck
|
||||
-- Looking for pthread.h
|
||||
-- Looking for pthread.h - found
|
||||
-- Looking for pthread_create
|
||||
@@ -235,7 +234,7 @@ cmake
|
||||
-- Looking for pthread_create in pthread
|
||||
-- Looking for pthread_create in pthread - found
|
||||
-- Found Threads: TRUE
|
||||
-- Found CUDA: /usr/local/cuda-10.0 (found version "10.0")
|
||||
-- Found CUDA: /usr/local/cuda-10.1 (found version "10.1")
|
||||
-- Configuring done
|
||||
-- Generating done
|
||||
-- Build files have been written to: /srv/git/HIP/hipify-clang/build
|
||||
@@ -247,8 +246,8 @@ make test-hipify
|
||||
```shell
|
||||
Running HIPify regression tests
|
||||
========================================
|
||||
CUDA 10.0 - will be used for testing
|
||||
LLVM 8.0.1 - will be used for testing
|
||||
CUDA 10.1 - will be used for testing
|
||||
LLVM 9.0.0 - will be used for testing
|
||||
x86_64 - Platform architecture
|
||||
Linux 5.2.0 - Platform OS
|
||||
64 - hipify-clang binary bitness
|
||||
@@ -326,38 +325,38 @@ LLVM 5.0.0 - 5.0.2, CUDA 8.0, cudnn-5.1.10 - cudnn-7.1.4.18
|
||||
|
||||
LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn-7.0.5.15 - cudnn-7.6.3.30
|
||||
|
||||
LLVM 7.0.0 - 8.0.1 (with patch*), CUDA 7.5 - 10.0, cudnn-7.0.5.15 - cudnn-7.6.3.30
|
||||
LLVM 7.0.0 - 9.0.0, CUDA 7.5 - 10.1, cudnn-7.0.5.15 - cudnn-7.6.3.30
|
||||
|
||||
Build system for the above configurations:
|
||||
|
||||
Python 3.6 (min), cmake 3.12.3 (min), Visual Studio 2017 (15.5.2) - 2019 (16.2.5).
|
||||
|
||||
Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 15 2017`:
|
||||
Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 16 2019`:
|
||||
|
||||
```shell
|
||||
cmake
|
||||
-G "Visual Studio 15 2017 Win64" \
|
||||
-G "Visual Studio 16 2019 Win64" \
|
||||
-DHIPIFY_CLANG_TESTS=1 \
|
||||
-DCMAKE_BUILD_TYPE=Release \
|
||||
-DCMAKE_INSTALL_PREFIX=../dist \
|
||||
-DCMAKE_PREFIX_PATH=f:/LLVM/6.0.1/dist \
|
||||
-DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0" \
|
||||
-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0" \
|
||||
-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.6.3.30 \
|
||||
-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.1/build/Release/bin/llvm-lit.py \
|
||||
-DCMAKE_PREFIX_PATH=f:/LLVM/9.0.0/dist \
|
||||
-DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v10.1" \
|
||||
-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v10.1" \
|
||||
-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.3.30 \
|
||||
-DLLVM_EXTERNAL_LIT=f:/LLVM/9.0.0/build/Release/bin/llvm-lit.py \
|
||||
-Thost=x64
|
||||
..
|
||||
```
|
||||
*A corresponding successful output:*
|
||||
```shell
|
||||
-- Found LLVM 6.0.1:
|
||||
-- - CMake module path: F:/LLVM/6.0.1/dist/lib/cmake/llvm
|
||||
-- - Include path : F:/LLVM/6.0.1/dist/include
|
||||
-- - Binary path : F:/LLVM/6.0.1/dist/bin
|
||||
-- Found LLVM 9.0.0:
|
||||
-- - CMake module path: F:/LLVM/9.0.0/dist/lib/cmake/llvm
|
||||
-- - Include path : F:/LLVM/9.0.0/dist/include
|
||||
-- - Binary path : F:/LLVM/9.0.0/dist/bin
|
||||
-- Found PythonInterp: C:/Program Files/Python37/python.exe (found suitable version "3.7.4", minimum required is "3.6")
|
||||
-- Found lit: C:/Program Files/Python36/Scripts/lit.exe
|
||||
-- Found FileCheck: F:/LLVM/6.0.1/dist/bin/FileCheck.exe
|
||||
-- Found CUDA: C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0 (found version "9.0")
|
||||
-- Found FileCheck: F:/LLVM/9.0.0/dist/bin/FileCheck.exe
|
||||
-- Found CUDA: C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v10.1 (found version "10.1")
|
||||
-- Configuring done
|
||||
-- Generating done
|
||||
-- Build files have been written to: f:/HIP/hipify-clang/build
|
||||
@@ -370,7 +369,7 @@ To process a file, `hipify-clang` needs access to the same headers that would be
|
||||
For example:
|
||||
|
||||
```shell
|
||||
./hipify-clang square.cu --cuda-path=/usr/local/cuda-10.0 -I /usr/local/cuda-10.0/samples/common/inc
|
||||
./hipify-clang square.cu --cuda-path=/usr/local/cuda-10.1 -I /usr/local/cuda-10.1/samples/common/inc
|
||||
```
|
||||
|
||||
`hipify-clang` arguments are given first, followed by a separator, and then the arguments you'd pass to `clang` if you
|
||||
|
||||
@@ -25,6 +25,181 @@ THE SOFTWARE.
|
||||
// Maps CUDA header names to HIP header names
|
||||
const std::map<llvm::StringRef, hipCounter> CUDA_DEVICE_FUNC_MAP{
|
||||
// math functions
|
||||
{"abs", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"labs", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"llabs", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"fabs", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"fabsf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"min", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"fminf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"fmin", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"max", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"fmaxf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"fmax", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"sin", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"cos", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"sincos", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"sincosf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"tan", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"sqrt", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"rsqrt", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"rsqrtf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"log2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"exp2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"exp2f", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"exp10", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"exp10f", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"expm1", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"expm1f", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"log2f", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"log10", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"log", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"log1p", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"log1pf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"floor", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"exp", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"cosh", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"sinh", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"tanh", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"acosh", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"acoshf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"asinh", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"asinhf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atanh", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atanhf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"ldexp", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"ldexpf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"logb", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"logbf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"ilogb", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"ilogbf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"scalbn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"scalbnf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"scalbln", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"scalblnf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"frexp", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"frexpf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"round", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"roundf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"lround", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"lroundf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"llround", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"llroundf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"rint", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"rintf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"lrint", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"lrintf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"llrint", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"llrintf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"nearbyint", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"nearbyintf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"ceil", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"trunc", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"truncf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"fdim", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"fdimf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atan2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atan", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"acos", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"asin", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"hypot", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"rhypot", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"hypotf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"rhypotf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"norm3d", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"rnorm3d", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"norm4d", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"rnorm4d", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"norm", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"rnorm", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"rnormf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"normf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"norm3df", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"rnorm3df", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"norm4df", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"rnorm4df", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"cbrt", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"cbrtf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"rcbrt", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"rcbrtf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"sinpi", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"sinpif", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"cospi", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"cospif", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"sincospi", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"sincospif", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"pow", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"modf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"fmod", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"remainder", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"remainderf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"remquo", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"remquof", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"j0", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"j0f", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"j1", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"j1f", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"jn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"jnf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"y0", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"y0f", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"y1", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"y1f", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"yn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"ynf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"cyl_bessel_i0", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"cyl_bessel_i0f", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"cyl_bessel_i1", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"cyl_bessel_i1f", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"erf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"erff", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"erfinv", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"erfinvf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"erfc", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"erfcf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"lgamma", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"erfcinv", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"erfcinvf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"normcdfinv", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"normcdfinvf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"normcdf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"normcdff", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"erfcx", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"erfcxf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"lgammaf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"tgamma", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"tgammaf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"copysign", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"copysignf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"nextafter", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"nextafterf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"nan", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"nanf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"fma", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"fmaf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"acosf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"asinf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atanf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atan2f", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"cosf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"sinf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"tanf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"coshf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"sinhf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"tanhf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"expf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"logf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"log10f", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"modff", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"powf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"sqrtf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"ceilf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"floorf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"fmodf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"signbit", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"isfinite", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"isnan", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"isinf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"umin", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"llmin", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"ullmin", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
@@ -60,13 +235,120 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DEVICE_FUNC_MAP{
|
||||
{"int2float", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"uint2float", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
// device functions
|
||||
{"__prof_trigger", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__trap", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__brkpt", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__pm0", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__pm1", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__pm2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__pm3", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__mulhi", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__umulhi", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__mul64hi", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__umul64hi", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__int_as_float", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float_as_int", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__uint_as_float", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float_as_uint", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__syncthreads", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__threadfence", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__threadfence_block", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__saturatef", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__sad", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__usad", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__mul24", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__umul24", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"fdividef", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fdividef", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"fdivide", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__sinf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__cosf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__tanf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__sincosf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__expf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__exp10f", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__log2f", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__log10f", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__logf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__powf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float2int_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float2int_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float2int_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float2int_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float2uint_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float2uint_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float2uint_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float2uint_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__int2float_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__int2float_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__int2float_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__int2float_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__uint2float_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__uint2float_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__uint2float_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__uint2float_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float2ll_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float2ll_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float2ll_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float2ll_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float2ull_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float2ull_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float2ull_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__float2ull_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__ll2float_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__ll2float_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__ll2float_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__ll2float_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__ull2float_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__ull2float_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__ull2float_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__ull2float_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fadd_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fadd_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fadd_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fadd_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fsub_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fsub_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fsub_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fsub_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fmul_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fmul_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fmul_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fmul_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fmaf_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fmaf_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fmaf_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fmaf_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__frcp_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__frcp_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__frcp_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__frcp_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fsqrt_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fsqrt_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fsqrt_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fsqrt_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__frsqrt_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fdiv_rn", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fdiv_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fdiv_ru", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__fdiv_rd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__clz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__ffs", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__popc", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__brev", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__clzll", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__ffsll", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__popcll", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__brevll", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__byte_perm", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__hadd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__rhadd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__uhadd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__urhadd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__double2int_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__double2uint_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__double2ll_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__double2ull_rz", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__prof_trigger", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__trap", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__brkpt", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__pm0", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__pm1", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__pm2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__pm3", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
// SIMD functions
|
||||
{"__vabs2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__vabsss2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
|
||||
@@ -36,51 +36,95 @@ using namespace llvm;
|
||||
|
||||
namespace perl {
|
||||
|
||||
const std::string space = " ";
|
||||
const std::string double_space = space + space;
|
||||
const std::string triple_space = double_space + space;
|
||||
const std::string tab = " ";
|
||||
const std::string double_tab = tab + tab;
|
||||
const std::string triple_tab = double_tab + tab;
|
||||
const std::string sSub = "sub";
|
||||
const std::string sReturn_0 = "return 0;\n";
|
||||
const std::string sReturn_m = "return $m;\n";
|
||||
const std::string sForeach = "foreach $func (\n";
|
||||
const std::string sMy = "my $m = 0;\n";
|
||||
|
||||
void generateUnsupportedDeviceFunctions(std::unique_ptr<std::ostream>& perlStreamPtr) {
|
||||
unsigned int num = 0;
|
||||
void generateSymbolFunctions(std::unique_ptr<std::ostream>& perlStreamPtr) {
|
||||
*perlStreamPtr.get() << "\n" << sSub << " transformSymbolFunctions\n" << "{\n" << tab << sMy;
|
||||
std::set<std::string> &funcSet = DeviceSymbolFunctions0;
|
||||
for (int i = 0; i < 2; ++i) {
|
||||
*perlStreamPtr.get() << tab + sForeach;
|
||||
if (i == 1) funcSet = DeviceSymbolFunctions1;
|
||||
unsigned int count = 0;
|
||||
for (auto& f : funcSet) {
|
||||
const auto found = CUDA_RUNTIME_FUNCTION_MAP.find(f);
|
||||
if (found != CUDA_RUNTIME_FUNCTION_MAP.end()) {
|
||||
*perlStreamPtr.get() << (count ? ",\n" : "") << double_tab << "\"" << found->second.hipName.str() << "\"";
|
||||
count++;
|
||||
}
|
||||
}
|
||||
*perlStreamPtr.get() << "\n" << tab << ")\n" << tab << "{\n" << double_tab;
|
||||
if (i ==0) *perlStreamPtr.get() << "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,/$func\\(HIP_SYMBOL\\($2\\),/g\n";
|
||||
else *perlStreamPtr.get() << "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,\\s*([^,\\)]+)\\s*(,\\s*|\\))\\s*/$func\\($2, HIP_SYMBOL\\($3\\)$4/g;\n";
|
||||
*perlStreamPtr.get() << tab << "}\n";
|
||||
}
|
||||
*perlStreamPtr.get() << tab << sReturn_m << "}\n";
|
||||
}
|
||||
|
||||
void generateDeviceFunctions(std::unique_ptr<std::ostream>& perlStreamPtr) {
|
||||
unsigned int countUnsupported = 0;
|
||||
unsigned int countSupported = 0;
|
||||
std::stringstream sSupported;
|
||||
std::stringstream sUnsupported;
|
||||
for (auto& ma : CUDA_DEVICE_FUNC_MAP) {
|
||||
if (Statistics::isUnsupported(ma.second)) {
|
||||
sUnsupported << (num ? ",\n" : "") << double_space << "\"" << ma.first.str() << "\"";
|
||||
num++;
|
||||
}
|
||||
bool isUnsupported = Statistics::isUnsupported(ma.second);
|
||||
(isUnsupported ? sUnsupported : sSupported) << ((isUnsupported && countUnsupported) || (!isUnsupported && countSupported) ? ",\n" : "") << double_tab << "\"" << ma.first.str() << "\"";
|
||||
if (isUnsupported) countUnsupported++;
|
||||
else countSupported++;
|
||||
}
|
||||
if (num) {
|
||||
*perlStreamPtr.get() << "\nsub warnUnsupportedDeviceFunctions\n" << "{\n" << space << "my $line_num = shift;\n" << space << "my $m = 0;\n" << space << "foreach $func (\n";
|
||||
*perlStreamPtr.get() << sUnsupported.str() << "\n" << space << ")\n";
|
||||
*perlStreamPtr.get() << space << "{\n";
|
||||
*perlStreamPtr.get() << double_space << "# match device function from the list of unsupported, except those, which have a namespace prefix (aka somenamespace::umin(...));\n";
|
||||
*perlStreamPtr.get() << double_space << "# function with only global namespace qualifier '::' (aka ::umin(...)) should be treated as a device function (and warned as well as without such qualifier);\n";
|
||||
*perlStreamPtr.get() << double_space << "my $mt_namespace = m/(\\w+)::($func)\\s*\\(\\s*.*\\s*\\)/g;\n";
|
||||
*perlStreamPtr.get() << double_space << "my $mt = m/($func)\\s*\\(\\s*.*\\s*\\)/g;\n";
|
||||
*perlStreamPtr.get() << double_space << "if ($mt && !$mt_namespace) {\n";
|
||||
*perlStreamPtr.get() << triple_space << "$m += $mt;\n";
|
||||
*perlStreamPtr.get() << triple_space << "print STDERR \" warning: $fileName:$line_num: unsupported device function \\\"$func\\\": $_\\n\";\n";
|
||||
*perlStreamPtr.get() << double_space << "}\n";
|
||||
*perlStreamPtr.get() << space << "}\n";
|
||||
*perlStreamPtr.get() << space << "return $m;\n";
|
||||
*perlStreamPtr.get() << "}\n";
|
||||
std::stringstream subCountSupported;
|
||||
std::stringstream subWarnUnsupported;
|
||||
std::stringstream subCommon;
|
||||
std::string sCommon = tab + sMy + tab + sForeach;
|
||||
subCountSupported << "\n" << sSub << " countSupportedDeviceFunctions\n" << "{\n" << (countSupported ? sCommon : tab + sReturn_0);
|
||||
subWarnUnsupported << "\n" << sSub << " warnUnsupportedDeviceFunctions\n" << "{\n" << (countUnsupported ? tab + "my $line_num = shift;\n" + sCommon : tab + sReturn_0);
|
||||
if (countSupported) {
|
||||
subCountSupported << sSupported.str() << "\n" << tab << ")\n";
|
||||
}
|
||||
if (countUnsupported) {
|
||||
subWarnUnsupported << sUnsupported.str() << "\n" << tab << ")\n";
|
||||
}
|
||||
if (countSupported || countUnsupported) {
|
||||
subCommon << tab << "{\n";
|
||||
subCommon << double_tab << "# match device function from the list, except those, which have a namespace prefix (aka somenamespace::umin(...));\n";
|
||||
subCommon << double_tab << "# function with only global namespace qualifier '::' (aka ::umin(...)) should be treated as a device function (and warned as well as without such qualifier);\n";
|
||||
subCommon << double_tab << "my $mt_namespace = m/(\\w+)::($func)\\s*\\(\\s*.*\\s*\\)/g;\n";
|
||||
subCommon << double_tab << "my $mt = m/($func)\\s*\\(\\s*.*\\s*\\)/g;\n";
|
||||
subCommon << double_tab << "if ($mt && !$mt_namespace) {\n";
|
||||
subCommon << triple_tab << "$m += $mt;\n";
|
||||
}
|
||||
if (countSupported) {
|
||||
subCountSupported << subCommon.str();
|
||||
}
|
||||
if (countUnsupported) {
|
||||
subWarnUnsupported << subCommon.str();
|
||||
subWarnUnsupported << triple_tab << "print STDERR \" warning: $fileName:$line_num: unsupported device function \\\"$func\\\": $_\\n\";\n";
|
||||
}
|
||||
if (countSupported || countUnsupported) {
|
||||
sCommon = double_tab + "}\n" + tab + "}\n" + tab + sReturn_m;
|
||||
}
|
||||
if (countSupported) subCountSupported << sCommon;
|
||||
if (countUnsupported) subWarnUnsupported << sCommon;
|
||||
subCountSupported << "}\n";
|
||||
subWarnUnsupported << "}\n";
|
||||
*perlStreamPtr.get() << subCountSupported.str();
|
||||
*perlStreamPtr.get() << subWarnUnsupported.str();
|
||||
}
|
||||
|
||||
bool generate(bool Generate) {
|
||||
if (!Generate) {
|
||||
return true;
|
||||
}
|
||||
if (!Generate) return true;
|
||||
std::string dstPerlMap = OutputPerlMapFilename, dstPerlMapDir = OutputPerlMapDir;
|
||||
if (dstPerlMap.empty()) {
|
||||
dstPerlMap = "hipify-perl-map";
|
||||
}
|
||||
if (dstPerlMap.empty()) dstPerlMap = "hipify-perl-map";
|
||||
std::error_code EC;
|
||||
if (!dstPerlMapDir.empty()) {
|
||||
std::string sOutputPerlMapDirAbsPath = getAbsoluteDirectoryPath(OutputPerlMapDir, EC, "output hipify-perl map");
|
||||
if (EC) {
|
||||
return false;
|
||||
}
|
||||
if (EC) return false;
|
||||
dstPerlMap = sOutputPerlMapDirAbsPath + "/" + dstPerlMap;
|
||||
}
|
||||
SmallString<128> tmpFile;
|
||||
@@ -93,18 +137,16 @@ namespace perl {
|
||||
std::unique_ptr<std::ostream> perlStreamPtr = std::unique_ptr<std::ostream>(new std::ofstream(tmpFile.c_str(), std::ios_base::trunc));
|
||||
std::string sConv = "my $conversions = ";
|
||||
*perlStreamPtr.get() << "@statNames = (";
|
||||
for (int i = 0; i < NUM_CONV_TYPES - 1; i++) {
|
||||
for (int i = 0; i < NUM_CONV_TYPES - 1; ++i) {
|
||||
*perlStreamPtr.get() << "\"" << counterNames[i] << "\", ";
|
||||
sConv += "$ft{'" + std::string(counterNames[i]) + "'} + ";
|
||||
}
|
||||
*perlStreamPtr.get() << "\"" << counterNames[NUM_CONV_TYPES - 1] << "\");\n\n";
|
||||
*perlStreamPtr.get() << sConv << "$ft{'" << counterNames[NUM_CONV_TYPES - 1] << "'};\n\n";
|
||||
for (int i = 0; i < NUM_CONV_TYPES; i++) {
|
||||
for (int i = 0; i < NUM_CONV_TYPES; ++i) {
|
||||
if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) {
|
||||
for (auto& ma : CUDA_INCLUDE_MAP) {
|
||||
if (Statistics::isUnsupported(ma.second)) {
|
||||
continue;
|
||||
}
|
||||
if (Statistics::isUnsupported(ma.second)) continue;
|
||||
if (i == ma.second.type) {
|
||||
std::string sCUDA = ma.first.str();
|
||||
std::string sHIP = ma.second.hipName.str();
|
||||
@@ -113,19 +155,17 @@ namespace perl {
|
||||
*perlStreamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << sCUDA << "\\b/" << sHIP << "/g;\n";
|
||||
}
|
||||
}
|
||||
}
|
||||
else {
|
||||
} else {
|
||||
for (auto& ma : CUDA_RENAMES_MAP()) {
|
||||
if (Statistics::isUnsupported(ma.second)) {
|
||||
continue;
|
||||
}
|
||||
if (Statistics::isUnsupported(ma.second)) continue;
|
||||
if (i == ma.second.type) {
|
||||
*perlStreamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << ma.first.str() << "\\b/" << ma.second.hipName.str() << "/g;\n";
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
generateUnsupportedDeviceFunctions(perlStreamPtr);
|
||||
generateSymbolFunctions(perlStreamPtr);
|
||||
generateDeviceFunctions(perlStreamPtr);
|
||||
perlStreamPtr.get()->flush();
|
||||
bool ret = true;
|
||||
EC = sys::fs::copy_file(tmpFile, dstPerlMap);
|
||||
@@ -133,9 +173,7 @@ namespace perl {
|
||||
llvm::errs() << "\n" << sHipify << sError << EC.message() << ": while copying " << tmpFile << " to " << dstPerlMap << "\n";
|
||||
ret = false;
|
||||
}
|
||||
if (!SaveTemps) {
|
||||
sys::fs::remove(tmpFile);
|
||||
}
|
||||
if (!SaveTemps) sys::fs::remove(tmpFile);
|
||||
return ret;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -179,7 +179,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_FUNCTION_MAP{
|
||||
|
||||
// 5.7. Execution Control
|
||||
// no analogue
|
||||
{"cudaFuncGetAttributes", {"hipFuncGetAttributes", "", CONV_EXECUTION, API_RUNTIME, HIP_UNSUPPORTED}},
|
||||
{"cudaFuncGetAttributes", {"hipFuncGetAttributes", "", CONV_EXECUTION, API_RUNTIME}},
|
||||
// no analogue
|
||||
{"cudaFuncSetAttribute", {"hipFuncSetAttribute", "", CONV_EXECUTION, API_RUNTIME, HIP_UNSUPPORTED}},
|
||||
// no analogue
|
||||
|
||||
@@ -22,6 +22,9 @@ THE SOFTWARE.
|
||||
|
||||
#pragma once
|
||||
|
||||
extern std::set<std::string> DeviceSymbolFunctions0;
|
||||
extern std::set<std::string> DeviceSymbolFunctions1;
|
||||
|
||||
namespace perl {
|
||||
|
||||
bool generate(bool Generate = true);
|
||||
|
||||
@@ -20,6 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <algorithm>
|
||||
#include <set>
|
||||
#include "HipifyAction.h"
|
||||
#include "clang/Basic/SourceLocation.h"
|
||||
#include "clang/Frontend/CompilerInstance.h"
|
||||
@@ -34,6 +36,25 @@ THE SOFTWARE.
|
||||
namespace ct = clang::tooling;
|
||||
namespace mat = clang::ast_matchers;
|
||||
|
||||
const std::string sCudaMemcpyToSymbol = "cudaMemcpyToSymbol";
|
||||
const std::string sCudaMemcpyToSymbolAsync = "cudaMemcpyToSymbolAsync";
|
||||
const std::string sCudaGetSymbolSize = "cudaGetSymbolSize";
|
||||
const std::string sCudaGetSymbolAddress = "cudaGetSymbolAddress";
|
||||
const std::string sCudaMemcpyFromSymbol = "cudaMemcpyFromSymbol";
|
||||
const std::string sCudaMemcpyFromSymbolAsync = "cudaMemcpyFromSymbolAsync";
|
||||
|
||||
std::set<std::string> DeviceSymbolFunctions0 {
|
||||
{sCudaMemcpyToSymbol},
|
||||
{sCudaMemcpyToSymbolAsync}
|
||||
};
|
||||
|
||||
std::set<std::string> DeviceSymbolFunctions1 {
|
||||
{sCudaGetSymbolSize},
|
||||
{sCudaGetSymbolAddress},
|
||||
{sCudaMemcpyFromSymbol},
|
||||
{sCudaMemcpyFromSymbolAsync}
|
||||
};
|
||||
|
||||
void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
|
||||
clang::SourceManager& SM = getCompilerInstance().getSourceManager();
|
||||
size_t begin = 0;
|
||||
@@ -83,7 +104,7 @@ void HipifyAction::RewriteToken(const clang::Token& t) {
|
||||
|
||||
void HipifyAction::FindAndReplace(llvm::StringRef name,
|
||||
clang::SourceLocation sl,
|
||||
const std::map<llvm::StringRef, hipCounter>& repMap) {
|
||||
const std::map<llvm::StringRef, hipCounter>& repMap, bool bReplace) {
|
||||
const auto found = repMap.find(name);
|
||||
if (found == repMap.end()) {
|
||||
// So it's an identifier, but not CUDA? Boring.
|
||||
@@ -100,6 +121,9 @@ void HipifyAction::FindAndReplace(llvm::StringRef name,
|
||||
DE.Report(sl, ID) << sWarn;
|
||||
return;
|
||||
}
|
||||
if (!bReplace) {
|
||||
return;
|
||||
}
|
||||
StringRef repName = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName;
|
||||
clang::SourceManager& SM = getCompilerInstance().getSourceManager();
|
||||
ct::Replacement Rep(SM, sl, name.size(), repName.str());
|
||||
@@ -313,8 +337,12 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc
|
||||
|
||||
// Next up are the four kernel configuration parameters, the last two of which are optional and default to zero.
|
||||
// Copy the two dimensional arguments verbatim.
|
||||
OS << "dim3(" << readSourceText(*SM, config->getArg(0)->getSourceRange()) << "), ";
|
||||
OS << "dim3(" << readSourceText(*SM, config->getArg(1)->getSourceRange()) << "), ";
|
||||
std::string sDim3 = "dim3(";
|
||||
for (unsigned int i = 0; i < 2; ++i) {
|
||||
const std::string sArg = readSourceText(*SM, config->getArg(i)->getSourceRange()).str();
|
||||
bool bDim3 = std::equal(sDim3.begin(), sDim3.end(), sArg.c_str());
|
||||
OS << (bDim3 ? "" : sDim3) << sArg << (bDim3 ? "" : ")") << ", ";
|
||||
}
|
||||
// The stream/memory arguments default to zero if omitted.
|
||||
OS << stringifyZeroDefaultedArg(*SM, config->getArg(2)) << ", ";
|
||||
OS << stringifyZeroDefaultedArg(*SM, config->getArg(3));
|
||||
@@ -392,7 +420,46 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match
|
||||
bool HipifyAction::cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
|
||||
if (const clang::CallExpr *call = Result.Nodes.getNodeAs<clang::CallExpr>("cudaDeviceFuncCall")) {
|
||||
const clang::FunctionDecl *funcDcl = call->getDirectCallee();
|
||||
FindAndReplace(funcDcl->getDeclName().getAsString(), llcompat::getBeginLoc(call), CUDA_DEVICE_FUNC_MAP);
|
||||
if (!funcDcl) {
|
||||
return true;
|
||||
}
|
||||
FindAndReplace(funcDcl->getDeclName().getAsString(), llcompat::getBeginLoc(call), CUDA_DEVICE_FUNC_MAP, false);
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
bool HipifyAction::cudaSymbolFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
|
||||
if (const clang::CallExpr * call = Result.Nodes.getNodeAs<clang::CallExpr>("cudaSymbolFuncCall")) {
|
||||
if (!call->getNumArgs()) {
|
||||
return true;
|
||||
}
|
||||
const clang::FunctionDecl* funcDcl = call->getDirectCallee();
|
||||
if (!funcDcl) {
|
||||
return true;
|
||||
}
|
||||
std::string sName = funcDcl->getDeclName().getAsString();
|
||||
unsigned int argNum = 0;
|
||||
if (DeviceSymbolFunctions0.find(sName) != DeviceSymbolFunctions0.end()) {
|
||||
argNum = 0;
|
||||
} else if (call->getNumArgs() > 1 && DeviceSymbolFunctions1.find(sName) != DeviceSymbolFunctions1.end()) {
|
||||
argNum = 1;
|
||||
} else {
|
||||
return true;
|
||||
}
|
||||
clang::SmallString<40> XStr;
|
||||
llvm::raw_svector_ostream OS(XStr);
|
||||
clang::SourceRange sr = call->getArg(argNum)->getSourceRange();
|
||||
clang::SourceManager* SM = Result.SourceManager;
|
||||
const std::string sSymbol = "HIP_SYMBOL";
|
||||
OS << sSymbol << "(" << readSourceText(*SM, sr) << ")";
|
||||
clang::SourceRange replacementRange = getWriteRange(*SM, { sr.getBegin(), sr.getEnd() });
|
||||
clang::SourceLocation s = replacementRange.getBegin();
|
||||
clang::SourceLocation e = replacementRange.getEnd();
|
||||
clang::LangOptions DefaultLangOptions;
|
||||
size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(e, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(s);
|
||||
ct::Replacement Rep(*SM, s, length, OS.str());
|
||||
clang::FullSourceLoc fullSL(s, *SM);
|
||||
insertReplacement(Rep, fullSL);
|
||||
}
|
||||
return true;
|
||||
}
|
||||
@@ -420,6 +487,24 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
|
||||
).bind("cudaSharedIncompleteArrayVar"),
|
||||
this
|
||||
);
|
||||
Finder->addMatcher(
|
||||
mat::callExpr(
|
||||
mat::isExpansionInMainFile(),
|
||||
mat::callee(
|
||||
mat::functionDecl(
|
||||
mat::hasAnyName(
|
||||
sCudaGetSymbolAddress,
|
||||
sCudaGetSymbolSize,
|
||||
sCudaMemcpyFromSymbol,
|
||||
sCudaMemcpyFromSymbolAsync,
|
||||
sCudaMemcpyToSymbol,
|
||||
sCudaMemcpyToSymbolAsync
|
||||
)
|
||||
)
|
||||
)
|
||||
).bind("cudaSymbolFuncCall"),
|
||||
this
|
||||
);
|
||||
Finder->addMatcher(
|
||||
mat::callExpr(
|
||||
mat::isExpansionInMainFile(),
|
||||
@@ -557,5 +642,6 @@ void HipifyAction::ExecuteAction() {
|
||||
void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
|
||||
if (cudaLaunchKernel(Result)) return;
|
||||
if (cudaSharedIncompleteArrayVar(Result)) return;
|
||||
if (cudaSymbolFuncCall(Result)) return;
|
||||
if (cudaDeviceFuncCall(Result)) return;
|
||||
}
|
||||
|
||||
@@ -71,6 +71,7 @@ public:
|
||||
bool cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
||||
bool cudaSharedIncompleteArrayVar(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
||||
bool cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
||||
bool cudaSymbolFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
||||
// Called by the preprocessor for each include directive during the non-raw lexing pass.
|
||||
void InclusionDirective(clang::SourceLocation hash_loc,
|
||||
const clang::Token &include_token,
|
||||
@@ -100,5 +101,5 @@ protected:
|
||||
void run(const clang::ast_matchers::MatchFinder::MatchResult& Result) override;
|
||||
std::unique_ptr<clang::ASTConsumer> CreateASTConsumer(clang::CompilerInstance &CI, llvm::StringRef InFile) override;
|
||||
bool Exclude(const hipCounter & hipToken);
|
||||
void FindAndReplace(llvm::StringRef name, clang::SourceLocation sl, const std::map<llvm::StringRef, hipCounter>& repMap);
|
||||
void FindAndReplace(llvm::StringRef name, clang::SourceLocation sl, const std::map<llvm::StringRef, hipCounter>& repMap, bool bReplace = true);
|
||||
};
|
||||
|
||||
@@ -34,6 +34,9 @@ THE SOFTWARE.
|
||||
#include "ArgParse.h"
|
||||
#include "StringUtils.h"
|
||||
#include "llvm/Support/Debug.h"
|
||||
#if LLVM_VERSION_MAJOR < 8
|
||||
#include "llvm/Support/Path.h"
|
||||
#endif
|
||||
|
||||
constexpr auto DEBUG_TYPE = "cuda2hip";
|
||||
|
||||
|
||||
@@ -0,0 +1,152 @@
|
||||
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
|
||||
|
||||
/*
|
||||
Copyright (c) 2015-present 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.
|
||||
*/
|
||||
|
||||
// CHECK: #include <hip/hip_runtime.h>
|
||||
#include <cuda_runtime.h>
|
||||
#include <iostream>
|
||||
#include <assert.h>
|
||||
|
||||
#define NUM 1024
|
||||
#define SIZE 1024 * 4
|
||||
|
||||
__device__ int globalIn[NUM];
|
||||
__device__ int globalOut[NUM];
|
||||
|
||||
__global__ void Assign(int* Out) {
|
||||
int tid = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
Out[tid] = globalIn[tid];
|
||||
globalOut[tid] = globalIn[tid];
|
||||
}
|
||||
|
||||
__device__ __constant__ int globalConst[NUM];
|
||||
|
||||
__global__ void checkAddress(int* addr, bool* out) {
|
||||
*out = (globalConst == addr);
|
||||
}
|
||||
|
||||
int main() {
|
||||
int *A, *Am, *B, *Ad, *C, *Cm;
|
||||
A = new int[NUM];
|
||||
B = new int[NUM];
|
||||
C = new int[NUM];
|
||||
for (int i = 0; i < NUM; ++i) {
|
||||
A[i] = -1 * i;
|
||||
B[i] = 0;
|
||||
C[i] = 0;
|
||||
}
|
||||
// CHECK: hipMalloc((void**)&Ad, SIZE);
|
||||
cudaMalloc((void**)&Ad, SIZE);
|
||||
// CHECK: hipHostMalloc((void**)&Am, SIZE);
|
||||
cudaMallocHost((void**)&Am, SIZE);
|
||||
// CHECK: hipHostMalloc((void**)&Cm, SIZE);
|
||||
cudaMallocHost((void**)&Cm, SIZE);
|
||||
for (int i = 0; i < NUM; ++i) {
|
||||
Am[i] = -1 * i;
|
||||
Cm[i] = 0;
|
||||
}
|
||||
// CHECK: hipStream_t stream = NULL;
|
||||
cudaStream_t stream = NULL;
|
||||
// CHECK: hipStreamCreate(&stream);
|
||||
cudaStreamCreate(&stream);
|
||||
// CHECK: hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), Am, SIZE, 0, hipMemcpyHostToDevice, stream);
|
||||
cudaMemcpyToSymbolAsync(globalIn, Am, SIZE, 0, cudaMemcpyHostToDevice, stream);
|
||||
// CHECK: hipStreamSynchronize(stream);
|
||||
cudaStreamSynchronize(stream);
|
||||
// CHECK: hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
|
||||
Assign<<<dim3(1, 1, 1), dim3(NUM, 1, 1)>>>(Ad);
|
||||
// CHECK: hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost);
|
||||
cudaMemcpy(B, Ad, SIZE, cudaMemcpyDeviceToHost);
|
||||
// CHECK: hipMemcpyFromSymbolAsync(Cm, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost, stream);
|
||||
cudaMemcpyFromSymbolAsync(Cm, globalOut, SIZE, 0, cudaMemcpyDeviceToHost, stream);
|
||||
// CHECK: hipStreamSynchronize(stream);
|
||||
cudaStreamSynchronize(stream);
|
||||
for (int i = 0; i < NUM; ++i) {
|
||||
assert(Am[i] == B[i]);
|
||||
assert(Am[i] == Cm[i]);
|
||||
}
|
||||
for (int i = 0; i < NUM; ++i) {
|
||||
A[i] = -2 * i;
|
||||
B[i] = 0;
|
||||
}
|
||||
// CHECK: hipMemcpyToSymbol(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice);
|
||||
cudaMemcpyToSymbol(globalIn, A, SIZE, 0, cudaMemcpyHostToDevice);
|
||||
// CHECK: hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
|
||||
Assign<<<dim3(1, 1, 1), dim3(NUM, 1, 1)>>>(Ad);
|
||||
// CHECK: hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost);
|
||||
cudaMemcpy(B, Ad, SIZE, cudaMemcpyDeviceToHost);
|
||||
// CHECK: hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost);
|
||||
cudaMemcpyFromSymbol(C, globalOut, SIZE, 0, cudaMemcpyDeviceToHost);
|
||||
for (int i = 0; i < NUM; ++i) {
|
||||
assert(A[i] == B[i]);
|
||||
assert(A[i] == C[i]);
|
||||
}
|
||||
for (int i = 0; i < NUM; ++i) {
|
||||
A[i] = -3 * i;
|
||||
B[i] = 0;
|
||||
}
|
||||
// CHECK: hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice, stream);
|
||||
cudaMemcpyToSymbolAsync(globalIn, A, SIZE, 0, cudaMemcpyHostToDevice, stream);
|
||||
// CHECK: hipStreamSynchronize(stream);
|
||||
cudaStreamSynchronize(stream);
|
||||
// CHECK: hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
|
||||
Assign<<<dim3(1, 1, 1), dim3(NUM, 1, 1)>>>(Ad);
|
||||
// CHECK: hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost);
|
||||
cudaMemcpy(B, Ad, SIZE, cudaMemcpyDeviceToHost);
|
||||
// CHECK: hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost, stream);
|
||||
cudaMemcpyFromSymbolAsync(C, globalOut, SIZE, 0, cudaMemcpyDeviceToHost, stream);
|
||||
// CHECK: hipStreamSynchronize(stream);
|
||||
cudaStreamSynchronize(stream);
|
||||
for (int i = 0; i < NUM; ++i) {
|
||||
assert(A[i] == B[i]);
|
||||
assert(A[i] == C[i]);
|
||||
}
|
||||
bool *checkOkD;
|
||||
bool checkOk = false;
|
||||
size_t symbolSize = 0;
|
||||
int *symbolAddress;
|
||||
// CHECK: hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConst));
|
||||
cudaGetSymbolSize(&symbolSize, globalConst);
|
||||
// CHECK: hipGetSymbolAddress((void**) &symbolAddress, HIP_SYMBOL(globalConst));
|
||||
cudaGetSymbolAddress((void**) &symbolAddress, globalConst);
|
||||
// CHECK: hipMalloc((void**)&checkOkD, sizeof(bool));
|
||||
cudaMalloc((void**)&checkOkD, sizeof(bool));
|
||||
// CHECK: hipLaunchKernelGGL(checkAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, symbolAddress, checkOkD);
|
||||
checkAddress<<<dim3(1, 1, 1), dim3(1, 1, 1)>>>(symbolAddress, checkOkD);
|
||||
// CHECK: hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost);
|
||||
cudaMemcpy(&checkOk, checkOkD, sizeof(bool), cudaMemcpyDeviceToHost);
|
||||
// CHECK: hipFree(checkOkD);
|
||||
cudaFree(checkOkD);
|
||||
assert(checkOk);
|
||||
assert(symbolSize == SIZE);
|
||||
// CHECK: hipHostFree(Am);
|
||||
cudaFreeHost(Am);
|
||||
// CHECK: hipHostFree(Cm);
|
||||
cudaFreeHost(Cm);
|
||||
// CHECK: hipFree(Ad);
|
||||
cudaFree(Ad);
|
||||
delete[] A;
|
||||
delete[] B;
|
||||
delete[] C;
|
||||
return 0;
|
||||
}
|
||||
+2
-2
@@ -304,7 +304,7 @@ struct runner<curandStateSobol32_t>
|
||||
CUDA_CALL(cudaMemcpy(directions, h_directions, size, cudaMemcpyHostToDevice));
|
||||
|
||||
const size_t blocks_x = next_power2((blocks + dimensions - 1) / dimensions);
|
||||
// CHECK: hipLaunchKernelGGL((init_kernel), dim3(dim3(blocks_x, dimensions)), dim3(threads), 0, 0, states, directions, offset);
|
||||
// CHECK: hipLaunchKernelGGL((init_kernel), dim3(blocks_x, dimensions), dim3(threads), 0, 0, states, directions, offset);
|
||||
init_kernel<<<dim3(blocks_x, dimensions), threads>>>(states, directions, offset);
|
||||
// CHECK: CUDA_CALL(hipPeekAtLastError());
|
||||
// CHECK: CUDA_CALL(hipDeviceSynchronize());
|
||||
@@ -329,7 +329,7 @@ struct runner<curandStateSobol32_t>
|
||||
const Extra extra)
|
||||
{
|
||||
const size_t blocks_x = next_power2((blocks + dimensions - 1) / dimensions);
|
||||
// CHECK: hipLaunchKernelGGL((generate_kernel), dim3(dim3(blocks_x, dimensions)), dim3(threads), 0, 0, states, data, size / dimensions, generate_func, extra);
|
||||
// CHECK: hipLaunchKernelGGL((generate_kernel), dim3(blocks_x, dimensions), dim3(threads), 0, 0, states, data, size / dimensions, generate_func, extra);
|
||||
generate_kernel<<<dim3(blocks_x, dimensions), threads>>>(states, data, size / dimensions, generate_func, extra);
|
||||
}
|
||||
};
|
||||
|
||||
@@ -247,7 +247,7 @@ API_TYPE set_API_type()
|
||||
void settings()
|
||||
{
|
||||
add_cachiers(cashiers_load);
|
||||
// CHECK: hipMemcpyToSymbol("cashiers_load", cashiers_load_h,
|
||||
// CHECK: hipMemcpyToSymbol(HIP_SYMBOL("cashiers_load"), cashiers_load_h,
|
||||
// CHECK: HOURS * sizeof(int), 0, hipMemcpyHostToDevice);
|
||||
cudaMemcpyToSymbol("cashiers_load", cashiers_load_h,
|
||||
HOURS * sizeof(int), 0, cudaMemcpyHostToDevice);
|
||||
|
||||
In neuem Issue referenzieren
Einen Benutzer sperren