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

Merge 'master' into 'amd-master'

Browse Source 
Change-Id: I0f3e131eb9415f0111526828bb1010a0c2c0be2a
This commit is contained in:
Jenkins
2019-08-09 05:53:35 -04:00
parent a05c22cad3 20628bd03f
commit 43e0d65de5
19 changed files with 181 additions and 34 deletions
Download Patch File Download Diff File Expand all files Collapse all files
CMakeLists.txt
+3 -5
View File
@@ -221,7 +221,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -DHIP_VERSION_MAJOR=${HIP_VERSION_MAJOR} -DHIP_VERSION_MINOR=${HIP_VERSION_MINOR} -DHIP_VERSION_PATCH=${HIP_VERSION_PATCH}")
# Add remaining flags
execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --cxxflags OUTPUT_VARIABLE HCC_CXX_FLAGS)
set(HCC_CXX_FLAGS "-hc -fno-gpu-rdc --amdgpu-target=gfx803 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 ")
set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -fPIC ${HCC_CXX_FLAGS} -I${HSA_PATH}/include")
# Set compiler and compiler flags
@@ -252,9 +252,6 @@ if(HIP_PLATFORM STREQUAL "hcc")
src/env.cpp
src/h2f.cpp)
execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS)
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS}")
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx803 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908")
if(COMPILE_HIP_ATP_MARKER)
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L/opt/rocm/profiler/CXLActivityLogger/bin/x86_64 -lCXLActivityLogger")
endif()
@@ -270,7 +267,8 @@ if(HIP_PLATFORM STREQUAL "hcc")
PRIVATE ${PROJECT_SOURCE_DIR}/include ${HSA_PATH}/include)
target_link_libraries(hiprtc PUBLIC stdc++fs)
endif()
set_target_properties(hip_hcc PROPERTIES CXX_VISIBILITY_PRESET hidden)
set_target_properties(hip_hcc PROPERTIES VISIBILITY_INLINES_HIDDEN 1)
if(HIP_PLATFORM STREQUAL "hcc")
find_package(amd_comgr REQUIRED CONFIG
README.md
+9
View File
@@ -37,6 +37,15 @@ HIP releases are typically of two types. The tag naming convention is different
- [HIP Debugging](docs/markdown/hip_debugging.md)
- [HIP Terminology](docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenCL)
- [hipify-clang](hipify-clang/README.md)
- Supported CUDA APIs:
* [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)
* [cuFFT](docs/markdown/CUFFT_API_supported_by_HIP.md)
* [cuSPARSE](docs/markdown/CUSPARSE_API_supported_by_HIP.md)
- [Developer/CONTRIBUTING Info](CONTRIBUTING.md)
- [Release Notes](RELEASE.md)
bin/hipcc
+5 -5
View File
@@ -118,10 +118,10 @@ if ($HIP_RUNTIME eq "VDI" and !defined $HIP_VDI_HOME) {
if (defined $HIP_VDI_HOME) {
if (!defined $HIP_CLANG_PATH and (-e "$HIP_VDI_HOME/bin/clang" or -e "$HIP_VDI_HOME/bin/clang.exe")) {
$HIP_CLANG_PATH = "$HIP_VDI_HOME/bin";
}
# With HIP_VDI_HOME defined, assume the installation of clang components, including headers.
if (!defined $HIP_CLANG_INCLUDE_PATH) {
$HIP_CLANG_INCLUDE_PATH = "$HIP_VDI_HOME/include/clang";
# With HIP_VDI_HOME defined, assume the installation of clang components, including headers.
if (!defined $HIP_CLANG_INCLUDE_PATH) {
$HIP_CLANG_INCLUDE_PATH = "$HIP_VDI_HOME/include/clang";
}
}
if (!defined $DEVICE_LIB_PATH and -e "$HIP_VDI_HOME/lib/bitcode") {
$DEVICE_LIB_PATH = "$HIP_VDI_HOME/lib/bitcode";
@@ -624,7 +624,7 @@ foreach $arg (@ARGV)
#print "O: <$arg>\n";
} else {
# input files and libraries
if (($arg =~ /\.cpp$/) or ($arg =~ /\.c$/) or ($arg =~ /\.cc$/) ) {
if (($arg =~ /\.cpp$/) or ($arg =~ /\.cxx$/) or ($arg =~ /\.c$/) or ($arg =~ /\.cc$/) ) {
$hasC = 1;
$needCXXFLAGS = 1;
if ($HIP_PLATFORM eq 'clang') {
bin/hipify-perl
+2
View File
@@ -252,6 +252,8 @@ while (@ARGV) {
$ft{'mem'} += s/\bcudaHostAllocPortable\b/hipHostMallocPortable/g;
$ft{'mem'} += s/\bcudaHostAllocMapped\b/hipHostMallocMapped/g;
$ft{'mem'} += s/\bcudaHostAllocWriteCombined\b/hipHostMallocWriteCombined/g;
$ft{'mem'} += s/\bcudaHostRegisterDefault\b/hipHostRegisterDefault/g;
$ft{'mem'} += s/\bcudaHostRegisterPortable\b/hipHostRegisterPortable/g;
$ft{'mem'} += s/\bcudaHostRegisterMapped\b/hipHostRegisterMapped/g;
$ft{'mem'} += s/\bcudaHostRegister\b/hipHostRegister/g;
$ft{'mem'} += s/\bcudaHostUnregister\b/hipHostUnregister/g;
include/hip/hcc_detail/channel_descriptor.h
+2
View File
@@ -23,11 +23,13 @@ THE SOFTWARE.
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H
#define HIP_INCLUDE_HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H
#include <hip/hip_common.h>
#include <hip/hcc_detail/driver_types.h>
#include <hip/hcc_detail/hip_vector_types.h>
#ifdef __cplusplus
HIP_PUBLIC_API
hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f);
static inline hipChannelFormatDesc hipCreateChannelDescHalf() {
include/hip/hcc_detail/driver_types.h
+14
View File
@@ -319,4 +319,18 @@ static inline struct hipExtent make_hipExtent(size_t w, size_t h, size_t d) {
return e;
}
typedef enum hipFunction_attribute {
HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK,
HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES,
HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES,
HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES,
HIP_FUNC_ATTRIBUTE_NUM_REGS,
HIP_FUNC_ATTRIBUTE_PTX_VERSION,
HIP_FUNC_ATTRIBUTE_BINARY_VERSION,
HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA,
HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES,
HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT,
HIP_FUNC_ATTRIBUTE_MAX
}hipFunction_attribute;
#endif
include/hip/hcc_detail/functional_grid_launch.hpp
+6 -13
View File
@@ -25,25 +25,16 @@ THE SOFTWARE.
#include "concepts.hpp"
#include "helpers.hpp"
#include "program_state.hpp"
#include "hip_runtime_api.h"
#include "hc.hpp"
#include "hip/hip_hcc.h"
#include "hip_runtime.h"
#include <algorithm>
#include <cstddef>
#include <cstdint>
#include <cstring>
#include <functional>
#include <iostream>
#include <mutex>
#include <stdexcept>
#include <string>
#include <tuple>
#include <type_traits>
#include <unordered_map>
#include <utility>
#include <vector>
#pragma GCC visibility push(hidden)
namespace hip_impl {
template <typename T, typename std::enable_if<std::is_integral<T>{}>::type* = nullptr>
@@ -114,7 +105,7 @@ inline hip_impl::kernarg make_kernarg(
}
hsa_agent_t target_agent(hipStream_t stream);
HIP_INTERNAL_EXPORTED_API hsa_agent_t target_agent(hipStream_t stream);
inline
__attribute__((visibility("hidden")))
@@ -194,3 +185,5 @@ inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBloc
hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{},
std::move(args)...);
}
#pragma GCC visibility pop
include/hip/hcc_detail/hip_runtime_api.h
+21 -8
View File
@@ -67,19 +67,18 @@ THE SOFTWARE.
#define HIP_LAUNCH_PARAM_END ((void*)0x03)
#ifdef __cplusplus
#include <algorithm>
#include <mutex>
#include <string>
#include <unordered_map>
#include <vector>
#define __dparm(x) \
= x
#else
#define __dparm(x)
#endif
#ifdef __GNUC__
#pragma GCC visibility push (default)
#endif
#ifdef __cplusplus
namespace hip_impl {
hipError_t hip_init();
} // namespace hip_impl
@@ -2603,16 +2602,27 @@ hipError_t hipModuleUnload(hipModule_t module);
hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, const char* kname);
/**
* @bried Find out attributes for a given function.
* @brief Find out attributes for a given function.
*
* @param [out] attr
* @param [in] func
*
* @returns hipSuccess, hipErrorInvalidDeviceFunction
* @returns hipSuccess, hipErrorInvalidValue, hipErrorInvalidDeviceFunction
*/
hipError_t hipFuncGetAttributes(struct hipFuncAttributes* attr, const void* func);
/**
* @brief Find out a specific attribute for a given function.
*
* @param [out] value
* @param [in] attrib
* @param [in] hfunc
*
* @returns hipSuccess, hipErrorInvalidValue, hipErrorInvalidDeviceFunction
*/
hipError_t hipFuncGetAttribute(int* value, hipFunction_attribute attrib, hipFunction_t hfunc);
#if !__HIP_VDI__
#if defined(__cplusplus)
} // extern "C"
@@ -3222,6 +3232,9 @@ hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject);
#endif
#ifdef __GNUC__
#pragma GCC visibility pop
#endif
/**
*-------------------------------------------------------------------------------------------------
include/hip/hcc_detail/program_state.hpp
+12
View File
@@ -31,11 +31,19 @@ THE SOFTWARE.
#include <cstdint>
#include <cstdlib>
#include <hip/hip_common.h>
struct ihipModuleSymbol_t;
using hipFunction_t = ihipModuleSymbol_t*;
namespace hip_impl {
// This section contains internal APIs that
// needs to be exported
#ifdef __GNUC__
#pragma GCC visibility push (default)
#endif
struct kernarg_impl;
class kernarg {
public:
@@ -83,6 +91,10 @@ private:
friend kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t);
};
#ifdef __GNUC__
#pragma GCC visibility pop
#endif
inline
__attribute__((visibility("hidden")))
program_state& get_program_state() {
include/hip/hip_common.h
+8
View File
@@ -45,6 +45,14 @@ THE SOFTWARE.
#define __HIP_DEVICE_COMPILE__ 1
#endif
#ifdef __GNUC__
#define HIP_PUBLIC_API __attribute__ ((visibility ("default")))
#define HIP_INTERNAL_EXPORTED_API __attribute__ ((visibility ("default")))
#else
#define HIP_PUBLIC_API
#define HIP_INTERNAL_EXPORTED_API
#endif
#if __HIP_DEVICE_COMPILE__ == 0
// 32-bit Atomics
#define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (0)
include/hip/hip_hcc.h
+4
View File
@@ -48,6 +48,7 @@ class accelerator_view;
* @brief Return hc::accelerator associated with the specified deviceId
* @return #hipSuccess, #hipErrorInvalidDevice
*/
HIP_PUBLIC_API
hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc);
/**
@@ -56,6 +57,7 @@ hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc);
* If stream is 0, the accelerator_view for the default stream is returned.
* @return #hipSuccess
*/
HIP_PUBLIC_API
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** av);
@@ -89,6 +91,7 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** a
* HIP/ROCm actually updates the start event when the associated kernel completes.
*/
HIP_PUBLIC_API
hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t localWorkSizeX, uint32_t localWorkSizeY,
@@ -98,6 +101,7 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
hipEvent_t stopEvent = nullptr,
uint32_t flags = 0);
HIP_PUBLIC_API
hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t localWorkSizeX, uint32_t localWorkSizeY,
include/hip/hip_runtime_api.h
-2
View File
@@ -117,13 +117,11 @@ typedef struct hipDeviceProp_t {
int integrated; ///< APU vs dGPU
int cooperativeLaunch; ///< HIP device supports cooperative launch
int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple devices
#if !__HIP_VDI__ // Temporarily disable the following five new fields for HIP/VDI runtime
int maxTexture1D; ///< Maximum number of elements in 1D images
int maxTexture2D[2]; ///< Maximum dimensions (width, height) of 2D images, in image elements
int maxTexture3D[3]; ///< Maximum dimensions (width, height, depth) of 3D images, in image elements
unsigned int* hdpMemFlushCntl; ///< Addres of HDP_MEM_COHERENCY_FLUSH_CNTL register
unsigned int* hdpRegFlushCntl; ///< Addres of HDP_REG_COHERENCY_FLUSH_CNTL register
#endif
} hipDeviceProp_t;
include/hip/nvcc_detail/hip_runtime_api.h
+18
View File
@@ -162,6 +162,7 @@ typedef CUdeviceptr hipDeviceptr_t;
typedef struct cudaArray hipArray;
typedef struct cudaArray* hipArray_const_t;
typedef cudaFuncAttributes hipFuncAttributes;
typedef CUfunction_attribute hipFunction_attribute;
#define hip_Memcpy2D CUDA_MEMCPY2D
#define hipMemcpy3DParms cudaMemcpy3DParms
#define hipArrayDefault cudaArrayDefault
@@ -197,6 +198,19 @@ typedef cudaSurfaceObject_t hipSurfaceObject_t;
#define hipSharedMemBankSizeFourByte cudaSharedMemBankSizeFourByte
#define hipSharedMemBankSizeEightByte cudaSharedMemBankSizeEightByte
//Function Attributes
#define HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK
#define HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES
#define HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES CU_FUNC_ATTRIBUTE_CONST_SIZE_BYTES
#define HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES
#define HIP_FUNC_ATTRIBUTE_NUM_REGS CU_FUNC_ATTRIBUTE_NUM_REGS
#define HIP_FUNC_ATTRIBUTE_PTX_VERSION CU_FUNC_ATTRIBUTE_PTX_VERSION
#define HIP_FUNC_ATTRIBUTE_BINARY_VERSION CU_FUNC_ATTRIBUTE_BINARY_VERSION
#define HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA CU_FUNC_ATTRIBUTE_CACHE_MODE_CA
#define HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES
#define HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT
#define HIP_FUNC_ATTRIBUTE_MAX CU_FUNC_ATTRIBUTE_MAX
inline static hipError_t hipCUDAErrorTohipError(cudaError_t cuError) {
switch (cuError) {
case cudaSuccess:
@@ -1197,6 +1211,10 @@ inline static hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const voi
return hipCUDAErrorTohipError(cudaFuncGetAttributes(attr, func));
}
inline static hipError_t hipFuncGetAttribute (int* value, hipFunction_attribute attrib, hipFunction_t hfunc) {
return hipCUResultTohipError(cuFuncGetAttribute(value, attrib, hfunc));
}
inline static hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod,
const char* name) {
return hipCUResultTohipError(cuModuleGetGlobal(dptr, bytes, hmod, name));
samples/0_Intro/module_api_global/runKernel.cpp
+5
View File
@@ -124,6 +124,11 @@ int main() {
{
hipFunction_t Function;
HIP_CHECK(hipModuleGetFunction(&Function, Module, "test_globals"));
int val =-1;
HIP_CHECK(hipFuncGetAttribute(&val, HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES,Function));
printf("Shared Size Bytes = %d\n",val);
HIP_CHECK(hipFuncGetAttribute(&val, HIP_FUNC_ATTRIBUTE_NUM_REGS, Function));
printf("Num Regs = %d\n",val);
HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config));
hipMemcpyDtoH(B, Bd, SIZE);
src/functional_grid_launch.inl
+1 -1
View File
@@ -41,7 +41,7 @@ using namespace std;
namespace hip_impl
{
hsa_agent_t target_agent(hipStream_t stream)
HIP_INTERNAL_EXPORTED_API hsa_agent_t target_agent(hipStream_t stream)
{
if (stream) {
return *static_cast<hsa_agent_t*>(
src/hip_hcc.cpp
+1
View File
@@ -46,6 +46,7 @@ THE SOFTWARE.
#include "hsa/hsa_ext_image.h"
#include "hip/hip_runtime.h"
#include "hip_hcc_internal.h"
#include "hip/hip_hcc.h"
#include "trace_helper.h"
#include "env.h"
src/hip_module.cpp
+46
View File
@@ -803,6 +803,52 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func)
return ihipLogStatus(hipSuccess);
}
hipError_t hipFuncGetAttribute(int* value, hipFunction_attribute attrib, hipFunction_t hfunc)
{
HIP_INIT_API(hipFuncGetAttribute, value, attrib, hfunc);
using namespace hip_impl;
hipError_t retVal = hipSuccess;
if (!value) return ihipLogStatus(hipErrorInvalidValue);
hipFuncAttributes attr{};
attr = make_function_attributes(tls, *hfunc);
switch(attrib) {
case HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES:
*value = (int) attr.sharedSizeBytes;
break;
case HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK:
*value = attr.maxThreadsPerBlock;
break;
case HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES:
*value = (int) attr.constSizeBytes;
break;
case HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES:
*value = (int) attr.localSizeBytes;
break;
case HIP_FUNC_ATTRIBUTE_NUM_REGS:
*value = attr.numRegs;
break;
case HIP_FUNC_ATTRIBUTE_PTX_VERSION:
*value = attr.ptxVersion;
break;
case HIP_FUNC_ATTRIBUTE_BINARY_VERSION:
*value = attr.binaryVersion;
break;
case HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA:
*value = attr.cacheModeCA;
break;
case HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES:
*value = attr.maxDynamicSharedSizeBytes;
break;
case HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT:
*value = attr.preferredShmemCarveout;
break;
default:
retVal = hipErrorInvalidValue;
}
return ihipLogStatus(retVal);
}
hipError_t ihipModuleLoadData(TlsData *tls, hipModule_t* module, const void* image) {
using namespace hip_impl;
tests/src/deviceLib/hipDoublePrecisionMathHost.cpp
+16
View File
@@ -59,7 +59,9 @@ __host__ void double_precision_math_functions() {
// erfcx(0.0);
// erfinv(1.0);
exp(0.0);
#ifdef __unix__
exp10(0.0);
#endif
exp2(0.0);
expm1(0.0);
fabs(1.0);
@@ -75,9 +77,15 @@ __host__ void double_precision_math_functions() {
std::isfinite(0.0);
std::isinf(0.0);
std::isnan(0.0);
#ifdef __unix__
j0(0.0);
j1(0.0);
jn(-1.0, 1.0);
#elif _WIN64
_j0(0.0);
_j1(0.0);
_jn(-1.0, 1.0);
#endif
ldexp(0.0, 0);
// lgamma(1.0);
llrint(0.0);
@@ -120,7 +128,9 @@ __host__ void double_precision_math_functions() {
scalbn(0.0, 1);
std::signbit(1.0);
sin(0.0);
#ifdef _unix__
sincos(0.0, &fX, &fY);
#endif
// sincospi(0.0, &fX, &fY);
sinh(0.0);
// sinpi(0.0);
@@ -129,9 +139,15 @@ __host__ void double_precision_math_functions() {
tanh(0.0);
tgamma(2.0);
trunc(0.0);
#ifdef __unix__
y0(1.0);
y1(1.0);
yn(1, 1.0);
#elif _WIN64
_y0(1.0);
_y1(1.0);
_yn(1, 1.0);
#endif
}
static void compileOnHost() { double_precision_math_functions(); }
tests/src/deviceLib/hipSinglePrecisionMathHost.cpp
+8
View File
@@ -58,7 +58,9 @@ __host__ void single_precision_math_functions() {
// erfcxf(0.0f);
erff(0.0f);
// erfinvf(1.0f);
#ifdef __unix__
exp10f(0.0f);
#endif
exp2f(0.0f);
expf(0.0f);
expm1f(0.0f);
@@ -78,9 +80,11 @@ __host__ void single_precision_math_functions() {
std::isfinite(0.0f);
std::isinf(0.0f);
std::isnan(0.0f);
#ifdef __unix__
j0f(0.0f);
j1f(0.0f);
jnf(-1.0f, 1.0f);
#endif
ldexpf(0.0f, 0);
lgammaf(1.0f);
llrintf(0.0f);
@@ -121,7 +125,9 @@ __host__ void single_precision_math_functions() {
scalblnf(0.0f, 1);
scalbnf(0.0f, 1);
std::signbit(1.0f);
#ifdef __unix__
sincosf(0.0f, &fX, &fY);
#endif
// sincospif(0.0f, &fX, &fY);
sinf(0.0f);
sinhf(0.0f);
@@ -131,9 +137,11 @@ __host__ void single_precision_math_functions() {
tanhf(0.0f);
tgammaf(2.0f);
truncf(0.0f);
#ifdef __unix__
y0f(1.0f);
y1f(1.0f);
ynf(1, 1.0f);
#endif
}
static void compileOnHost() { single_precision_math_functions(); }