Merge 'master' into 'amd-master'
Change-Id: Ie207caaa50ef3b2a01ad871a8cf9c7336b81b74a
This commit is contained in:
Jenkins
+20
-1
@@ -58,6 +58,21 @@ if(NOT DEFINED HIP_PLATFORM)
|
||||
endif()
|
||||
message(STATUS "HIP Platform: " ${HIP_PLATFORM})
|
||||
|
||||
# Determine HIP_COMPILER
|
||||
# Either hcc or clang; default is hcc
|
||||
if(NOT DEFINED HIP_COMPILER)
|
||||
if(NOT DEFINED ENV{HIP_COMPILER})
|
||||
set(HIP_COMPILER "hcc" CACHE STRING "HIP Compiler")
|
||||
else()
|
||||
set(HIP_COMPILER $ENV{HIP_COMPILER} CACHE STRING "HIP Compiler")
|
||||
endif()
|
||||
endif()
|
||||
if(NOT (HIP_COMPILER STREQUAL "hcc" OR HIP_COMPILER STREQUAL "clang"))
|
||||
message(FATAL_ERROR "Must use HIP_COMPILER as hcc or clang")
|
||||
endif()
|
||||
message(STATUS "HIP Compiler: " ${HIP_COMPILER})
|
||||
|
||||
|
||||
# If HIP_PLATFORM is hcc, we need HCC_HOME and HSA_PATH to be defined
|
||||
if(HIP_PLATFORM STREQUAL "hcc")
|
||||
# Determine HCC_HOME
|
||||
@@ -214,7 +229,11 @@ if(HIP_PLATFORM STREQUAL "hcc")
|
||||
add_library(host INTERFACE)
|
||||
target_link_libraries(host INTERFACE hip_hcc)
|
||||
add_library(device INTERFACE)
|
||||
target_link_libraries(device INTERFACE host hip_device hcc::hccrt hcc::hc_am)
|
||||
if(HIP_COMPILER STREQUAL "hcc")
|
||||
target_link_libraries(device INTERFACE host hip_device hcc::hccrt hcc::hc_am)
|
||||
elseif(HIP_COMPILER STREQUAL "clang")
|
||||
target_link_libraries(device INTERFACE host hip_device)
|
||||
endif()
|
||||
|
||||
# Generate .hipInfo
|
||||
file(WRITE "${PROJECT_BINARY_DIR}/.hipInfo" ${_buildInfo})
|
||||
|
||||
@@ -57,6 +57,7 @@ $HIP_PATH=$ENV{'HIP_PATH'} // dirname (dirname $0); # use parent directory of hi
|
||||
$HIP_VDI_HOME=$ENV{'HIP_VDI_HOME'};
|
||||
$HIP_CLANG_PATH=$ENV{'HIP_CLANG_PATH'};
|
||||
$DEVICE_LIB_PATH=$ENV{'DEVICE_LIB_PATH'};
|
||||
$HIP_CLANG_HCC_COMPAT_MODE=$ENV{'HIP_CLANG_HCC_COMPAT_MODE'}; # HCC compatibility mode
|
||||
|
||||
#---
|
||||
# Read .hipInfo
|
||||
@@ -158,6 +159,10 @@ if ($HIP_PLATFORM eq "clang") {
|
||||
|
||||
$HIPCXXFLAGS .= " -std=c++11 -isystem $HIP_CLANG_INCLUDE_PATH";
|
||||
$HIPLDFLAGS .= " --hip-device-lib-path=$DEVICE_LIB_PATH -L$HIP_LIB_PATH -Wl,--rpath=$HIP_LIB_PATH -lhip_hcc";
|
||||
if ($HIP_CLANG_HCC_COMPAT_MODE) {
|
||||
## Allow __fp16 as function parameter and return type.
|
||||
$HIPCXXFLAGS .= " -Xclang -fallow-half-arguments-and-returns -D__HIP_HCC_COMPAT_MODE__=1";
|
||||
}
|
||||
} elsif ($HIP_PLATFORM eq "hcc") {
|
||||
$HIP_INCLUDE_PATH = "$HIP_PATH/include";
|
||||
$HSA_PATH=$ENV{'HSA_PATH'} // "/opt/rocm/hsa";
|
||||
@@ -455,6 +460,53 @@ foreach $arg (@ARGV)
|
||||
close $in;
|
||||
close $out;
|
||||
$arg = "$new_arg -Wl,\@$new_file";
|
||||
} elsif (($arg =~ m/\.a$/ || $arg =~ m/\.lo$/) &&
|
||||
$HIP_PLATFORM eq 'clang') {
|
||||
## process static library for hip-clang
|
||||
## extract object files from static library and pass them directly to
|
||||
## hip-clang.
|
||||
## ToDo: Remove this after hip-clang switch to lto and lld is able to
|
||||
## handle clang-offload-bundler bundles.
|
||||
|
||||
my $new_arg = "";
|
||||
my $tmpdir = get_temp_dir ();
|
||||
my $libFile = $arg;
|
||||
my $path = abs_path($arg);
|
||||
my @objs = split ('\n', `cd $tmpdir; ar xv $path`);
|
||||
## Check if all files in .a are object files.
|
||||
my $allIsObj = 1;
|
||||
my $realObjs = "";
|
||||
foreach my $obj (@objs) {
|
||||
chomp $obj;
|
||||
$obj =~ s/^x - //;
|
||||
$obj = "$tmpdir/$obj";
|
||||
my $fileType = `file $obj`;
|
||||
my $isObj = ($fileType =~ m/ELF/ or $fileType =~ m/COFF/);
|
||||
$allIsObj = ($allIsObj and $isObj);
|
||||
if ($isObj) {
|
||||
$realObjs = ($realObjs . " " . $obj);
|
||||
} else {
|
||||
push (@inputs, $obj);
|
||||
if ($new_arg ne "") {
|
||||
$new_arg .= " ";
|
||||
}
|
||||
$new_arg .= "$obj";
|
||||
}
|
||||
}
|
||||
chomp $realObjs;
|
||||
if ($allIsObj) {
|
||||
$new_arg = $arg;
|
||||
} elsif ($realObjs) {
|
||||
my($libBaseName, $libDir, $libExt) = fileparse($libFile);
|
||||
$libBaseName = mktemp($libBaseName . "XXXX") . $libExt;
|
||||
system("cd $tmpdir; ar c $libBaseName $realObjs");
|
||||
$new_arg .= " $tmpdir/$libBaseName";
|
||||
}
|
||||
$arg = "$new_arg";
|
||||
if ($toolArgs =~ m/-Xlinker$/) {
|
||||
$toolArgs = substr $toolArgs, 0, -8;
|
||||
chomp $toolArgs;
|
||||
}
|
||||
} elsif ($arg =~ m/^-/) {
|
||||
# options start with -
|
||||
|
||||
|
||||
@@ -0,0 +1,37 @@
|
||||
# cuComplex API supported by HIP
|
||||
|
||||
## **1. cuComplex Data types**
|
||||
|
||||
| **type** | **CUDA** | **HIP** |**HIP value** (if differs) |
|
||||
|-------------:|---------------------------------------------------------------|------------------------------------------------------------|---------------------------|
|
||||
| float2 |***`cuFloatComplex`*** |***`hipFloatComplex`*** | struct |
|
||||
| double2 |***`cuDoubleComplex`*** |***`hipDoubleComplex`*** | struct |
|
||||
| float2 |***`cuComplex`*** |***`hipComplex`*** | struct |
|
||||
|
||||
## **2. cuComplex API functions**
|
||||
|
||||
| **CUDA** | **HIP** |
|
||||
|-----------------------------------------------------------|-------------------------------------------------|
|
||||
|`cuCrealf` |`hipCrealf` |
|
||||
|`cuCimagf` |`hipCimagf` |
|
||||
|`make_cuFloatComplex` |`make_hipFloatComplex` |
|
||||
|`cuConjf` |`hipConjf` |
|
||||
|`cuCaddf` |`hipCaddf` |
|
||||
|`cuCsubf` |`hipCsubf` |
|
||||
|`cuCmulf` |`hipCmulf` |
|
||||
|`cuCdivf` |`hipCdivf` |
|
||||
|`cuCabsf` |`hipCabsf` |
|
||||
|`cuCreal` |`hipCreal` |
|
||||
|`cuCimag` |`hipCimag` |
|
||||
|`make_cuDoubleComplex` |`make_hipDoubleComplex` |
|
||||
|`cuConj` |`hipConj` |
|
||||
|`cuCadd` |`hipCadd` |
|
||||
|`cuCsub` |`hipCsub` |
|
||||
|`cuCmul` |`hipCmul` |
|
||||
|`cuCdiv` |`hipCdiv` |
|
||||
|`cuCabs` |`hipCabs` |
|
||||
|`make_cuComplex` |`make_hipComplex` |
|
||||
|`cuComplexFloatToDouble` |`hipComplexFloatToDouble` |
|
||||
|`cuComplexDoubleToFloat` |`hipComplexDoubleToFloat` |
|
||||
|`cuCfmaf` |`hipCfmaf` |
|
||||
|`cuCfma` |`hipCfma` |
|
||||
@@ -21,6 +21,7 @@
|
||||
|
||||
- [Runtime API](../docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md)
|
||||
- [Driver API](../docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md)
|
||||
- [cuComplex API](../docs/markdown/cuComplex_API_supported_by_HIP.md)
|
||||
- [cuBLAS](../docs/markdown/CUBLAS_API_supported_by_HIP.md)
|
||||
- [cuRAND](../docs/markdown/CURAND_API_supported_by_HIP.md)
|
||||
- [cuDNN](../docs/markdown/CUDNN_API_supported_by_HIP.md)
|
||||
|
||||
@@ -366,7 +366,6 @@ const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP{
|
||||
{"channel_descriptor.h", {"hip/channel_descriptor.h", CONV_INCLUDE, API_RUNTIME}},
|
||||
{"device_functions.h", {"hip/device_functions.h", CONV_INCLUDE, API_RUNTIME}},
|
||||
{"driver_types.h", {"hip/driver_types.h", CONV_INCLUDE, API_RUNTIME}},
|
||||
{"cuComplex.h", {"hip/hip_complex.h", CONV_INCLUDE, API_RUNTIME}},
|
||||
{"cuda_fp16.h", {"hip/hip_fp16.h", CONV_INCLUDE, API_RUNTIME}},
|
||||
{"cuda_texture_types.h", {"hip/hip_texture_types.h", CONV_INCLUDE, API_RUNTIME}},
|
||||
{"vector_types.h", {"hip/hip_vector_types.h", CONV_INCLUDE, API_RUNTIME}},
|
||||
@@ -397,12 +396,11 @@ const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP{
|
||||
// CUDNN includes
|
||||
{"cudnn.h", {"hipDNN.h", CONV_INCLUDE_CUDA_MAIN_H, API_DNN}},
|
||||
|
||||
// CUDNN includes
|
||||
// CUFFT includes
|
||||
{"cufft.h", {"hipfft.h", CONV_INCLUDE_CUDA_MAIN_H, API_FFT}},
|
||||
|
||||
// HIP includes
|
||||
// TODO: uncomment this when hip/cudacommon.h will be renamed to hip/hipcommon.h
|
||||
// {"cudacommon.h", {"hipcommon.h", CONV_INCLUDE, API_RUNTIME}},
|
||||
// cuComplex includes
|
||||
{"cuComplex.h", {"hip/hip_complex.h", CONV_INCLUDE_CUDA_MAIN_H, API_COMPLEX}},
|
||||
};
|
||||
|
||||
/// All other identifiers: function and macro names.
|
||||
@@ -1387,6 +1385,36 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
|
||||
{"cuGraphicsEGLRegisterImage", {"hipGraphicsEGLRegisterImage", CONV_EGL, API_DRIVER, HIP_UNSUPPORTED}}, // API_Runtime ANALOGUE (cudaGraphicsEGLRegisterImage)
|
||||
{"cuGraphicsResourceGetMappedEglFrame", {"hipGraphicsResourceGetMappedEglFrame", CONV_EGL, API_DRIVER, HIP_UNSUPPORTED}}, // API_Runtime ANALOGUE (cudaGraphicsResourceGetMappedEglFrame)
|
||||
|
||||
|
||||
////////////////////////////// cuComplex API //////////////////////////////
|
||||
{"cuFloatComplex", {"hipFloatComplex", CONV_TYPE, API_COMPLEX}},
|
||||
{"cuDoubleComplex", {"hipDoubleComplex", CONV_TYPE, API_COMPLEX}},
|
||||
{"cuComplex", {"hipComplex", CONV_TYPE, API_COMPLEX}},
|
||||
|
||||
{"cuCrealf", {"hipCrealf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCimagf", {"hipCimagf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"make_cuFloatComplex", {"make_hipFloatComplex", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuConjf", {"hipConjf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCaddf", {"hipCaddf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCsubf", {"hipCsubf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCmulf", {"hipCmulf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCdivf", {"hipCdivf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCabsf", {"hipCabsf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCreal", {"hipCreal", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCimag", {"hipCimag", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"make_cuDoubleComplex", {"make_hipDoubleComplex", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuConj", {"hipConj", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCadd", {"hipCadd", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCsub", {"hipCsub", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCmul", {"hipCmul", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCdiv", {"hipCdiv", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCabs", {"hipCabs", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"make_cuComplex", {"make_hipComplex", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuComplexFloatToDouble", {"hipComplexFloatToDouble", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuComplexDoubleToFloat", {"hipComplexDoubleToFloat", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCfmaf", {"hipCfmaf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCfma", {"hipCfma", CONV_COMPLEX, API_COMPLEX}},
|
||||
|
||||
/////////////////////////////// CUDA RT API ///////////////////////////////
|
||||
// Data types
|
||||
{"cudaDataType_t", {"hipDataType_t", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}},
|
||||
|
||||
@@ -168,6 +168,10 @@ bool HipifyAction::Exclude(const hipCounter & hipToken) {
|
||||
if (insertedFFTHeader) { return true; }
|
||||
insertedFFTHeader = true;
|
||||
return false;
|
||||
case API_COMPLEX:
|
||||
if (insertedComplexHeader) { return true; }
|
||||
insertedComplexHeader = true;
|
||||
return false;
|
||||
default:
|
||||
return false;
|
||||
}
|
||||
|
||||
@@ -29,6 +29,7 @@ private:
|
||||
bool insertedRAND_kernelHeader = false;
|
||||
bool insertedDNNHeader = false;
|
||||
bool insertedFFTHeader = false;
|
||||
bool insertedComplexHeader = false;
|
||||
bool firstHeader = false;
|
||||
bool pragmaOnce = false;
|
||||
clang::SourceLocation firstHeaderLoc;
|
||||
|
||||
@@ -8,13 +8,13 @@ const char *counterNames[NUM_CONV_TYPES] = {
|
||||
"version", "init", "device", "mem", "kern", "coord_func", "math_func", "device_func",
|
||||
"special_func", "stream", "event", "occupancy", "ctx", "peer", "module",
|
||||
"cache", "exec", "err", "def", "tex", "gl", "graphics",
|
||||
"surface", "jit", "d3d9", "d3d10", "d3d11", "vdpau", "egl",
|
||||
"surface", "jit", "d3d9", "d3d10", "d3d11", "vdpau", "egl", "complex",
|
||||
"thread", "other", "include", "include_cuda_main_header", "type", "literal",
|
||||
"numeric_literal"
|
||||
};
|
||||
|
||||
const char *apiNames[NUM_API_TYPES] = {
|
||||
"CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API", "CUDNN API", "CUFFT API"
|
||||
"CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API", "CUDNN API", "CUFFT API", "cuComplex API"
|
||||
};
|
||||
|
||||
namespace {
|
||||
|
||||
@@ -40,6 +40,7 @@ enum ConvTypes {
|
||||
CONV_D3D11,
|
||||
CONV_VDPAU,
|
||||
CONV_EGL,
|
||||
CONV_COMPLEX,
|
||||
CONV_THREAD,
|
||||
CONV_OTHER,
|
||||
CONV_INCLUDE,
|
||||
@@ -58,6 +59,7 @@ enum ApiTypes {
|
||||
API_RAND,
|
||||
API_DNN,
|
||||
API_FFT,
|
||||
API_COMPLEX,
|
||||
API_LAST
|
||||
};
|
||||
constexpr int NUM_API_TYPES = (int) ApiTypes::API_LAST;
|
||||
|
||||
@@ -436,6 +436,34 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim;
|
||||
#define hipGridDim_y gridDim.y
|
||||
#define hipGridDim_z gridDim.z
|
||||
|
||||
#if __HIP_HCC_COMPAT_MODE__
|
||||
// Define HCC work item functions in terms of HIP builtin variables.
|
||||
#pragma push_macro("__DEFINE_HCC_FUNC")
|
||||
#define __DEFINE_HCC_FUNC(hc_fun,hip_var) \
|
||||
inline __device__ __attribute__((always_inline)) uint hc_get_##hc_fun(uint i) { \
|
||||
if (i==0) \
|
||||
return hip_var.x; \
|
||||
else if(i==1) \
|
||||
return hip_var.y; \
|
||||
else \
|
||||
return hip_var.z; \
|
||||
}
|
||||
|
||||
__DEFINE_HCC_FUNC(workitem_id, threadIdx)
|
||||
__DEFINE_HCC_FUNC(group_id, blockIdx)
|
||||
__DEFINE_HCC_FUNC(group_size, blockDim)
|
||||
__DEFINE_HCC_FUNC(num_groups, gridDim)
|
||||
#pragma pop_macro("__DEFINE_HCC_FUNC")
|
||||
|
||||
extern "C" __device__ __attribute__((const)) size_t __ockl_get_global_id(uint);
|
||||
inline __device__ __attribute__((always_inline)) uint
|
||||
hc_get_workitem_absolute_id(int dim)
|
||||
{
|
||||
return (uint)__ockl_get_global_id(dim);
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
// Support std::complex.
|
||||
#pragma push_macro("__CUDA__")
|
||||
#define __CUDA__
|
||||
|
||||
@@ -31,6 +31,12 @@ THE SOFTWARE.
|
||||
#include <limits>
|
||||
#include <stdint.h>
|
||||
|
||||
// HCC's own math functions should be included first, otherwise there will
|
||||
// be conflicts when hip/math_functions.h is included before hip/hip_runtime.h.
|
||||
#ifdef __HCC__
|
||||
#include "kalmar_math.h"
|
||||
#endif
|
||||
|
||||
#pragma push_macro("__DEVICE__")
|
||||
#pragma push_macro("__RETURN_TYPE")
|
||||
|
||||
@@ -1332,3 +1338,8 @@ __HIP_OVERLOAD2(double, min)
|
||||
#pragma pop_macro("__HIP_OVERLOAD2")
|
||||
#pragma pop_macro("__DEVICE__")
|
||||
#pragma pop_macro("__RETURN_TYPE")
|
||||
|
||||
// For backward compatibility.
|
||||
// There are HIP applications e.g. TensorFlow, expecting __HIP_ARCH_* macros
|
||||
// defined after including math_functions.h.
|
||||
#include <hip/hcc_detail/hip_runtime.h>
|
||||
|
||||
@@ -93,12 +93,11 @@ public:
|
||||
}
|
||||
};
|
||||
|
||||
const std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>& executables(
|
||||
bool rebuild = false);
|
||||
const std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>& executables();
|
||||
const std::unordered_map<std::uintptr_t, std::vector<std::pair<hsa_agent_t, Kernel_descriptor>>>&
|
||||
functions(bool rebuild = false);
|
||||
const std::unordered_map<std::uintptr_t, std::string>& function_names(bool rebuild = false);
|
||||
std::unordered_map<std::string, void*>& globals(bool rebuild = false);
|
||||
functions();
|
||||
const std::unordered_map<std::uintptr_t, std::string>& function_names();
|
||||
std::unordered_map<std::string, void*>& globals();
|
||||
|
||||
hsa_executable_t load_executable(const std::string& file, hsa_executable_t executable,
|
||||
hsa_agent_t agent);
|
||||
|
||||
@@ -27,6 +27,7 @@ THE SOFTWARE.
|
||||
// paths to provide a consistent include env and avoid "missing symbol" errors that only appears
|
||||
// on NVCC path:
|
||||
|
||||
#include <hip/hip_common.h>
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) && !defined(__HIP_PLATFORM_NVCC__)
|
||||
#include <hip/hcc_detail/math_functions.h>
|
||||
|
||||
+3
-3
@@ -80,7 +80,7 @@ inline void copy_kernel_section_to_fat_binary(const std::string& tmp, const std:
|
||||
std::find_if(reader.sections.begin(), reader.sections.end(),
|
||||
[](const ELFIO::section* x) { return x->get_name() == kernel_section(); });
|
||||
|
||||
std::ofstream out{output + fat_binary_extension()};
|
||||
std::ofstream out{output};
|
||||
|
||||
if (it == reader.sections.end()) {
|
||||
std::cerr << "Warning: no kernels were generated; fat binary shall "
|
||||
@@ -95,8 +95,8 @@ inline void generate_fat_binary(const std::vector<std::string>& sources,
|
||||
const std::vector<std::string>& targets, const std::string& flags,
|
||||
const std::string& output) {
|
||||
static const auto d = [](const std::string* f) { remove(f->c_str()); };
|
||||
|
||||
std::unique_ptr<const std::string, decltype(d)> tmp{&output, d};
|
||||
std::string temp_str = output + ".tmp";
|
||||
std::unique_ptr<const std::string, decltype(d)> tmp{&temp_str, d};
|
||||
|
||||
redi::ipstream hipcc{make_hipcc_call(sources, targets, flags, *tmp), redi::pstream::pstderr};
|
||||
|
||||
|
||||
@@ -30,7 +30,7 @@ THE SOFTWARE.
|
||||
#define LEN 64
|
||||
#define SIZE LEN << 2
|
||||
|
||||
#define fileName "vcpy_kernel.code.adipose"
|
||||
#define fileName "vcpy_kernel.code"
|
||||
#define kernel_name "hello_world"
|
||||
|
||||
#define HIP_CHECK(status) \
|
||||
@@ -66,32 +66,13 @@ int main() {
|
||||
HIP_CHECK(hipModuleLoad(&Module, fileName));
|
||||
HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name));
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
uint32_t len = LEN;
|
||||
uint32_t one = 1;
|
||||
|
||||
struct {
|
||||
void* _Ad;
|
||||
void* _Bd;
|
||||
} args;
|
||||
|
||||
args._Ad = Ad;
|
||||
args._Bd = Bd;
|
||||
|
||||
#endif
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
struct {
|
||||
uint32_t _hidden[1];
|
||||
void* _Ad;
|
||||
void* _Bd;
|
||||
} args;
|
||||
|
||||
args._hidden[0] = 0;
|
||||
args._Ad = Ad;
|
||||
args._Bd = Bd;
|
||||
#endif
|
||||
|
||||
args._Ad = (void*) Ad;
|
||||
args._Bd = (void*) Bd;
|
||||
|
||||
size_t size = sizeof(args);
|
||||
|
||||
|
||||
@@ -30,7 +30,7 @@ THE SOFTWARE.
|
||||
#define LEN 64
|
||||
#define SIZE LEN * sizeof(float)
|
||||
|
||||
#define fileName "vcpy_kernel.code.adipose"
|
||||
#define fileName "vcpy_kernel.code"
|
||||
float myDeviceGlobal;
|
||||
float myDeviceGlobalArray[16];
|
||||
#define HIP_CHECK(cmd) \
|
||||
@@ -79,32 +79,13 @@ int main() {
|
||||
myDeviceGlobalArray[i] = i * 1000.0f;
|
||||
}
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
uint32_t len = LEN;
|
||||
uint32_t one = 1;
|
||||
|
||||
struct {
|
||||
void* _Ad;
|
||||
void* _Bd;
|
||||
} args;
|
||||
|
||||
args._Ad = Ad;
|
||||
args._Bd = Bd;
|
||||
|
||||
#endif
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
struct {
|
||||
uint32_t _hidden[1];
|
||||
void* _Ad;
|
||||
void* _Bd;
|
||||
} args;
|
||||
|
||||
args._hidden[0] = 0;
|
||||
args._Ad = Ad;
|
||||
args._Bd = Bd;
|
||||
#endif
|
||||
|
||||
args._Ad = (void*) Ad;
|
||||
args._Bd = (void*) Bd;
|
||||
|
||||
size_t size = sizeof(args);
|
||||
|
||||
|
||||
@@ -21,13 +21,13 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip/hip_runtime_api.h"
|
||||
//#include "hip/hip_runtime_api.h"
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <vector>
|
||||
#include <hip/hip_hcc.h>
|
||||
//#include <hip/hip_hcc.h>
|
||||
|
||||
#define fileName "tex2dKernel.code.adipose"
|
||||
#define fileName "tex2dKernel.code"
|
||||
|
||||
texture<float, 2, hipReadModeElementType> tex;
|
||||
bool testResult = false;
|
||||
@@ -87,34 +87,15 @@ bool runTest(int argc, char** argv) {
|
||||
float* dData = NULL;
|
||||
hipMalloc((void**)&dData, size);
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
|
||||
struct {
|
||||
void* _Ad;
|
||||
unsigned int _Bd;
|
||||
unsigned int _Cd;
|
||||
} args;
|
||||
args._Ad = dData;
|
||||
args._Ad = (void*) dData;
|
||||
args._Bd = width;
|
||||
args._Cd = height;
|
||||
|
||||
#endif
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
struct {
|
||||
uint32_t _hidden[1];
|
||||
void* _Ad;
|
||||
unsigned int _Bd;
|
||||
unsigned int _Cd;
|
||||
} args;
|
||||
|
||||
args._hidden[0] = 0;
|
||||
args._Ad = dData;
|
||||
args._Bd = width;
|
||||
args._Cd = height;
|
||||
#endif
|
||||
|
||||
|
||||
size_t sizeTemp = sizeof(args);
|
||||
|
||||
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE,
|
||||
|
||||
@@ -92,19 +92,13 @@ namespace hip_impl
|
||||
hipStream_t stream,
|
||||
void** kernarg)
|
||||
{
|
||||
auto it0 = functions().find(function_address);
|
||||
const auto it0 = functions().find(function_address);
|
||||
|
||||
if (it0 == functions().cend()) {
|
||||
// Re-init device code maps once again to help locate kernels
|
||||
// loaded after HIP runtime initialization via means such as
|
||||
// dlopen().
|
||||
it0 = functions(true).find(function_address);
|
||||
if (it0 == functions().cend()) {
|
||||
throw runtime_error{
|
||||
"No device code available for function: " +
|
||||
name(function_address)
|
||||
};
|
||||
}
|
||||
throw runtime_error{
|
||||
"No device code available for function: " +
|
||||
name(function_address)
|
||||
};
|
||||
}
|
||||
|
||||
auto agent = target_agent(stream);
|
||||
|
||||
+37
-131
@@ -74,15 +74,11 @@ vector<string> copy_names_of_undefined_symbols(const symbol_section_accessor& se
|
||||
}
|
||||
|
||||
const std::unordered_map<std::string, std::pair<ELFIO::Elf64_Addr, ELFIO::Elf_Xword>>&
|
||||
symbol_addresses(bool rebuild = false) {
|
||||
symbol_addresses() {
|
||||
static unordered_map<string, pair<Elf64_Addr, Elf_Xword>> r;
|
||||
static once_flag f;
|
||||
|
||||
auto cons = [rebuild]() {
|
||||
if (rebuild) {
|
||||
r.clear();
|
||||
}
|
||||
|
||||
call_once(f, []() {
|
||||
dl_iterate_phdr(
|
||||
[](dl_phdr_info* info, size_t, void*) {
|
||||
static constexpr const char self[] = "/proc/self/exe";
|
||||
@@ -112,12 +108,7 @@ symbol_addresses(bool rebuild = false) {
|
||||
return 0;
|
||||
},
|
||||
nullptr);
|
||||
};
|
||||
|
||||
call_once(f, cons);
|
||||
if (rebuild) {
|
||||
cons();
|
||||
}
|
||||
});
|
||||
|
||||
return r;
|
||||
}
|
||||
@@ -175,34 +166,21 @@ vector<char> code_object_blob_for_process() {
|
||||
return r;
|
||||
}
|
||||
|
||||
const unordered_map<hsa_isa_t, vector<vector<char>>>& code_object_blobs(bool rebuild = false) {
|
||||
const unordered_map<hsa_isa_t, vector<vector<char>>>& code_object_blobs() {
|
||||
static unordered_map<hsa_isa_t, vector<vector<char>>> r;
|
||||
static once_flag f;
|
||||
|
||||
auto cons = [rebuild]() {
|
||||
// names of shared libraries who .kernel sections already loaded
|
||||
static unordered_set<string> lib_names;
|
||||
call_once(f, []() {
|
||||
static vector<vector<char>> blobs{code_object_blob_for_process()};
|
||||
|
||||
if (rebuild) {
|
||||
r.clear();
|
||||
blobs.clear();
|
||||
}
|
||||
|
||||
dl_iterate_phdr(
|
||||
[](dl_phdr_info* info, std::size_t, void*) {
|
||||
elfio tmp;
|
||||
if ((lib_names.find(info->dlpi_name) == lib_names.end()) &&
|
||||
(tmp.load(info->dlpi_name))) {
|
||||
if (tmp.load(info->dlpi_name)) {
|
||||
const auto it = find_section_if(
|
||||
tmp, [](const section* x) { return x->get_name() == ".kernel"; });
|
||||
|
||||
if (it) {
|
||||
blobs.emplace_back(
|
||||
it->get_data(), it->get_data() + it->get_size());
|
||||
// register the shared library as already loaded
|
||||
lib_names.emplace(info->dlpi_name);
|
||||
}
|
||||
if (it) blobs.emplace_back(it->get_data(), it->get_data() + it->get_size());
|
||||
}
|
||||
return 0;
|
||||
},
|
||||
@@ -216,13 +194,7 @@ const unordered_map<hsa_isa_t, vector<vector<char>>>& code_object_blobs(bool reb
|
||||
}
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
|
||||
call_once(f, cons);
|
||||
if (rebuild) {
|
||||
cons();
|
||||
}
|
||||
});
|
||||
|
||||
return r;
|
||||
}
|
||||
@@ -244,13 +216,13 @@ vector<pair<uintptr_t, string>> function_names_for(const elfio& reader, section*
|
||||
return r;
|
||||
}
|
||||
|
||||
const vector<pair<uintptr_t, string>>& function_names_for_process(bool rebuild = false) {
|
||||
const vector<pair<uintptr_t, string>>& function_names_for_process() {
|
||||
static constexpr const char self[] = "/proc/self/exe";
|
||||
|
||||
static vector<pair<uintptr_t, string>> r;
|
||||
static once_flag f;
|
||||
|
||||
auto cons = [rebuild]() {
|
||||
call_once(f, []() {
|
||||
elfio reader;
|
||||
|
||||
if (!reader.load(self)) {
|
||||
@@ -261,26 +233,16 @@ const vector<pair<uintptr_t, string>>& function_names_for_process(bool rebuild =
|
||||
find_section_if(reader, [](const section* x) { return x->get_type() == SHT_SYMTAB; });
|
||||
|
||||
if (symtab) r = function_names_for(reader, symtab);
|
||||
};
|
||||
|
||||
call_once(f, cons);
|
||||
if (rebuild) {
|
||||
cons();
|
||||
}
|
||||
});
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
const unordered_map<string, vector<hsa_executable_symbol_t>>& kernels(bool rebuild = false) {
|
||||
const unordered_map<string, vector<hsa_executable_symbol_t>>& kernels() {
|
||||
static unordered_map<string, vector<hsa_executable_symbol_t>> r;
|
||||
static once_flag f;
|
||||
|
||||
auto cons = [rebuild]() {
|
||||
if (rebuild) {
|
||||
r.clear();
|
||||
executables(rebuild);
|
||||
}
|
||||
|
||||
call_once(f, []() {
|
||||
static const auto copy_kernels = [](hsa_executable_t, hsa_agent_t,
|
||||
hsa_executable_symbol_t s, void*) {
|
||||
if (type(s) == HSA_SYMBOL_KIND_KERNEL) r[name(s)].push_back(s);
|
||||
@@ -294,12 +256,7 @@ const unordered_map<string, vector<hsa_executable_symbol_t>>& kernels(bool rebui
|
||||
copy_kernels, nullptr);
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
call_once(f, cons);
|
||||
if (rebuild) {
|
||||
cons();
|
||||
}
|
||||
});
|
||||
|
||||
return r;
|
||||
}
|
||||
@@ -338,19 +295,13 @@ void load_code_object_and_freeze_executable(
|
||||
|
||||
namespace hip_impl {
|
||||
const unordered_map<hsa_agent_t, vector<hsa_executable_t>>&
|
||||
executables(bool rebuild) { // TODO: This leaks the hsa_executable_ts, it should use RAII.
|
||||
executables() { // TODO: This leaks the hsa_executable_ts, it should use RAII.
|
||||
static unordered_map<hsa_agent_t, vector<hsa_executable_t>> r;
|
||||
static once_flag f;
|
||||
|
||||
auto cons = [rebuild]() {
|
||||
call_once(f, []() {
|
||||
static const auto accelerators = hc::accelerator::get_all();
|
||||
|
||||
if (rebuild) {
|
||||
// do NOT clear r so we reuse instances of hsa_executable_t
|
||||
// created previously
|
||||
code_object_blobs(rebuild);
|
||||
}
|
||||
|
||||
for (auto&& acc : accelerators) {
|
||||
auto agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
|
||||
|
||||
@@ -384,29 +335,17 @@ executables(bool rebuild) { // TODO: This leaks the hsa_executable_ts, it shoul
|
||||
},
|
||||
agent);
|
||||
}
|
||||
};
|
||||
|
||||
call_once(f, cons);
|
||||
if (rebuild) {
|
||||
cons();
|
||||
}
|
||||
});
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
const unordered_map<uintptr_t, string>& function_names(bool rebuild) {
|
||||
const unordered_map<uintptr_t, string>& function_names() {
|
||||
static unordered_map<uintptr_t, string> r{function_names_for_process().cbegin(),
|
||||
function_names_for_process().cend()};
|
||||
static once_flag f;
|
||||
|
||||
auto cons = [rebuild]() {
|
||||
if (rebuild) {
|
||||
r.clear();
|
||||
function_names_for_process(rebuild);
|
||||
r.insert(function_names_for_process().cbegin(),
|
||||
function_names_for_process().cend());
|
||||
}
|
||||
|
||||
call_once(f, []() {
|
||||
dl_iterate_phdr(
|
||||
[](dl_phdr_info* info, size_t, void*) {
|
||||
elfio tmp;
|
||||
@@ -426,32 +365,16 @@ const unordered_map<uintptr_t, string>& function_names(bool rebuild) {
|
||||
return 0;
|
||||
},
|
||||
nullptr);
|
||||
};
|
||||
|
||||
call_once(f, cons);
|
||||
if (rebuild) {
|
||||
static mutex mtx;
|
||||
lock_guard<mutex> lck{mtx};
|
||||
cons();
|
||||
}
|
||||
});
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
const unordered_map<uintptr_t, vector<pair<hsa_agent_t, Kernel_descriptor>>>& functions(bool rebuild) {
|
||||
const unordered_map<uintptr_t, vector<pair<hsa_agent_t, Kernel_descriptor>>>& functions() {
|
||||
static unordered_map<uintptr_t, vector<pair<hsa_agent_t, Kernel_descriptor>>> r;
|
||||
static once_flag f;
|
||||
|
||||
auto cons = [rebuild]() {
|
||||
if (rebuild) {
|
||||
// do NOT clear r so we reuse instances of pair<hsa_agent_t, Kernel_descriptor>
|
||||
// created previously
|
||||
|
||||
function_names(rebuild);
|
||||
kernels(rebuild);
|
||||
globals(rebuild);
|
||||
}
|
||||
|
||||
call_once(f, []() {
|
||||
for (auto&& function : function_names()) {
|
||||
const auto it = kernels().find(function.second);
|
||||
|
||||
@@ -463,34 +386,15 @@ const unordered_map<uintptr_t, vector<pair<hsa_agent_t, Kernel_descriptor>>>& fu
|
||||
}
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
call_once(f, cons);
|
||||
if (rebuild) {
|
||||
static mutex mtx;
|
||||
lock_guard<mutex> lck{mtx};
|
||||
cons();
|
||||
}
|
||||
});
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
unordered_map<string, void*>& globals(bool rebuild) {
|
||||
unordered_map<string, void*>& globals() {
|
||||
static unordered_map<string, void*> r;
|
||||
static once_flag f;
|
||||
auto cons =[rebuild]() {
|
||||
if (rebuild) {
|
||||
r.clear();
|
||||
symbol_addresses(rebuild);
|
||||
}
|
||||
|
||||
r.reserve(symbol_addresses().size());
|
||||
};
|
||||
|
||||
call_once(f, cons);
|
||||
if (rebuild) {
|
||||
cons();
|
||||
}
|
||||
call_once(f, []() { r.reserve(symbol_addresses().size()); });
|
||||
|
||||
return r;
|
||||
}
|
||||
@@ -513,15 +417,17 @@ hsa_executable_t load_executable(const string& file, hsa_executable_t executable
|
||||
return executable;
|
||||
}
|
||||
|
||||
// To force HIP to load the kernels and to setup the function
|
||||
// symbol map on program startup
|
||||
class startup_kernel_loader {
|
||||
private:
|
||||
startup_kernel_loader() { functions(); }
|
||||
startup_kernel_loader(const startup_kernel_loader&) = delete;
|
||||
startup_kernel_loader& operator=(const startup_kernel_loader&) = delete;
|
||||
static startup_kernel_loader skl;
|
||||
};
|
||||
startup_kernel_loader startup_kernel_loader::skl;
|
||||
// HIP startup kernel loader logic
|
||||
// When enabled HIP_STARTUP_LOADER, HIP will load the kernels and setup
|
||||
// the function symbol map on program startup
|
||||
extern "C" void __attribute__((constructor)) __startup_kernel_loader_init() {
|
||||
int hip_startup_loader=0;
|
||||
if (std::getenv("HIP_STARTUP_LOADER"))
|
||||
hip_startup_loader = atoi(std::getenv("HIP_STARTUP_LOADER"));
|
||||
if (hip_startup_loader) functions();
|
||||
}
|
||||
|
||||
extern "C" void __attribute__((destructor)) __startup_kernel_loader_fini() {
|
||||
}
|
||||
|
||||
} // Namespace hip_impl.
|
||||
|
||||
@@ -0,0 +1,58 @@
|
||||
// RUN: %run_test hipify "%s" "%t" %cuda_args
|
||||
|
||||
// CHECK: #include <hip/hip_runtime.h>
|
||||
// CHECK: #include "hip/hip_complex.h"
|
||||
#include "cuComplex.h"
|
||||
|
||||
#define TYPEFLOAT
|
||||
#define DIMX 100
|
||||
#define DIMY 40
|
||||
#define moveX 2
|
||||
#define moveY 1
|
||||
|
||||
#define MAXITERATIONS 10
|
||||
|
||||
#ifdef TYPEFLOAT
|
||||
#define TYPE float
|
||||
// CHECK: #define cTYPE hipFloatComplex
|
||||
#define cTYPE cuFloatComplex
|
||||
// CHECK: #define cMakecuComplex(re,i) make_hipFloatComplex(re,i)
|
||||
#define cMakecuComplex(re,i) make_cuFloatComplex(re,i)
|
||||
#endif
|
||||
#ifdef TYPEDOUBLE
|
||||
// CHECK: #define TYPE hipDoubleComplex
|
||||
#define TYPE cuDoubleComplex
|
||||
// CHECK: #define cMakecuComplex(re,i) make_hipDoubleComplex(re,i)
|
||||
#define cMakecuComplex(re,i) make_cuDoubleComplex(re,i)
|
||||
#endif
|
||||
|
||||
__device__ cTYPE juliaFunctor(cTYPE p, cTYPE c) {
|
||||
// CHECK: return hipCaddf(hipCmulf(p, p), c);
|
||||
return cuCaddf(cuCmulf(p, p), c);
|
||||
}
|
||||
|
||||
__device__ cTYPE convertToComplex(int x, int y, float zoom) {
|
||||
TYPE jx = 1.5 * (x - DIMX / 2) / (0.5 * zoom * DIMX) + moveX;
|
||||
TYPE jy = (y - DIMY / 2) / (0.5 * zoom * DIMY) + moveY;
|
||||
return cMakecuComplex(jx, jy);
|
||||
}
|
||||
|
||||
__device__ int evolveComplexPoint(cTYPE p, cTYPE c) {
|
||||
int it = 1;
|
||||
// CHECK: while (it <= MAXITERATIONS && hipCabsf(p) <= 4) {
|
||||
while (it <= MAXITERATIONS && cuCabsf(p) <= 4) {
|
||||
p = juliaFunctor(p, c);
|
||||
it++;
|
||||
}
|
||||
return it;
|
||||
}
|
||||
|
||||
__global__ void computeJulia(int* data, cTYPE c, float zoom) {
|
||||
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int j = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
|
||||
if (i<DIMX && j<DIMY) {
|
||||
cTYPE p = convertToComplex(i, j, zoom);
|
||||
data[i*DIMY + j] = evolveComplexPoint(p, c);
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,158 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp
|
||||
* RUN: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
// Test include math_functions.h then hip_runtime.h.
|
||||
// Incorrect implementation causes compilation failure due to conflict
|
||||
// declartions.
|
||||
|
||||
#include <hip/math_functions.h>
|
||||
|
||||
// Test __HIP_DEVICE_COMPILE__ is defined after math_functions.h
|
||||
// is included.
|
||||
//
|
||||
__device__ __host__ inline void throw_std_bad_alloc()
|
||||
{
|
||||
#ifndef __HIP_DEVICE_COMPILE__
|
||||
throw std::bad_alloc();
|
||||
#else
|
||||
std::size_t huge = static_cast<std::size_t>(-1);
|
||||
new int[huge];
|
||||
#endif
|
||||
}
|
||||
|
||||
// Test __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ and __HIP_ARCH_HAS_DYNAMIC_PARALLEL__
|
||||
// is defined. Eigen HIP/hcc/Half.h __ldg depends on this.
|
||||
#if !defined(__HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__) || \
|
||||
!defined(__HIP_ARCH_HAS_DYNAMIC_PARALLEL__)
|
||||
#error \
|
||||
"__HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ or __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ not defined"
|
||||
#endif
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include "test_common.h"
|
||||
|
||||
__global__ void FloatMathPrecise() {
|
||||
int iX;
|
||||
float fX, fY;
|
||||
|
||||
acosf(1.0f);
|
||||
acoshf(1.0f);
|
||||
asinf(0.0f);
|
||||
asinhf(0.0f);
|
||||
atan2f(0.0f, 1.0f);
|
||||
atanf(0.0f);
|
||||
atanhf(0.0f);
|
||||
cbrtf(0.0f);
|
||||
fX = ceilf(0.0f);
|
||||
fX = copysignf(1.0f, -2.0f);
|
||||
cosf(0.0f);
|
||||
coshf(0.0f);
|
||||
cospif(0.0f);
|
||||
cyl_bessel_i0f(0.0f);
|
||||
cyl_bessel_i1f(0.0f);
|
||||
erfcf(0.0f);
|
||||
erfcinvf(2.0f);
|
||||
erfcxf(0.0f);
|
||||
erff(0.0f);
|
||||
erfinvf(1.0f);
|
||||
exp10f(0.0f);
|
||||
exp2f(0.0f);
|
||||
expf(0.0f);
|
||||
expm1f(0.0f);
|
||||
fX = fabsf(1.0f);
|
||||
fdimf(1.0f, 0.0f);
|
||||
fdividef(0.0f, 1.0f);
|
||||
fX = floorf(0.0f);
|
||||
fmaf(1.0f, 2.0f, 3.0f);
|
||||
fX = fmaxf(0.0f, 0.0f);
|
||||
fX = fminf(0.0f, 0.0f);
|
||||
fmodf(0.0f, 1.0f);
|
||||
frexpf(0.0f, &iX);
|
||||
hypotf(1.0f, 0.0f);
|
||||
ilogbf(1.0f);
|
||||
isfinite(0.0f);
|
||||
fX = isinf(0.0f);
|
||||
fX = isnan(0.0f);
|
||||
j0f(0.0f);
|
||||
j1f(0.0f);
|
||||
jnf(-1.0f, 1.0f);
|
||||
ldexpf(0.0f, 0);
|
||||
lgammaf(1.0f);
|
||||
llrintf(0.0f);
|
||||
llroundf(0.0f);
|
||||
log10f(1.0f);
|
||||
log1pf(-1.0f);
|
||||
log2f(1.0f);
|
||||
logbf(1.0f);
|
||||
logf(1.0f);
|
||||
lrintf(0.0f);
|
||||
lroundf(0.0f);
|
||||
modff(0.0f, &fX);
|
||||
fX = nanf("1");
|
||||
fX = nearbyintf(0.0f);
|
||||
nextafterf(0.0f, 0.0f);
|
||||
norm3df(1.0f, 0.0f, 0.0f);
|
||||
norm4df(1.0f, 0.0f, 0.0f, 0.0f);
|
||||
normcdff(0.0f);
|
||||
normcdfinvf(1.0f);
|
||||
fX = 1.0f;
|
||||
normf(1, &fX);
|
||||
powf(1.0f, 0.0f);
|
||||
rcbrtf(1.0f);
|
||||
remainderf(2.0f, 1.0f);
|
||||
remquof(1.0f, 2.0f, &iX);
|
||||
rhypotf(0.0f, 1.0f);
|
||||
fY = rintf(1.0f);
|
||||
rnorm3df(0.0f, 0.0f, 1.0f);
|
||||
rnorm4df(0.0f, 0.0f, 0.0f, 1.0f);
|
||||
fX = 1.0f;
|
||||
rnormf(1, &fX);
|
||||
fY = roundf(0.0f);
|
||||
rsqrtf(1.0f);
|
||||
scalblnf(0.0f, 1);
|
||||
scalbnf(0.0f, 1);
|
||||
signbit(1.0f);
|
||||
sincosf(0.0f, &fX, &fY);
|
||||
sincospif(0.0f, &fX, &fY);
|
||||
sinf(0.0f);
|
||||
sinhf(0.0f);
|
||||
sinpif(0.0f);
|
||||
sqrtf(0.0f);
|
||||
tanf(0.0f);
|
||||
tanhf(0.0f);
|
||||
tgammaf(2.0f);
|
||||
fY = truncf(0.0f);
|
||||
y0f(1.0f);
|
||||
y1f(1.0f);
|
||||
ynf(1, 1.0f);
|
||||
}
|
||||
|
||||
int main() {
|
||||
hipLaunchKernelGGL(FloatMathPrecise, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0);
|
||||
passed();
|
||||
}
|
||||
Reference in New Issue
Block a user