diff --git a/CMakeLists.txt b/CMakeLists.txt index a4da3b1920..818c75783e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -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}) diff --git a/bin/hipcc b/bin/hipcc index 9df3459fd1..53ef35a558 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -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 - diff --git a/docs/markdown/cuComplex_API_supported_by_HIP.md b/docs/markdown/cuComplex_API_supported_by_HIP.md new file mode 100644 index 0000000000..621973ba61 --- /dev/null +++ b/docs/markdown/cuComplex_API_supported_by_HIP.md @@ -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` | diff --git a/hipify-clang/README.md b/hipify-clang/README.md index d5a825033d..7f55dea800 100644 --- a/hipify-clang/README.md +++ b/hipify-clang/README.md @@ -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) diff --git a/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp index ca53738abb..4bd7263bad 100644 --- a/hipify-clang/src/CUDA2HipMap.cpp +++ b/hipify-clang/src/CUDA2HipMap.cpp @@ -366,7 +366,6 @@ const std::map 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 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 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}}, diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index 20bc9dcc2c..a9bd1aa085 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/hipify-clang/src/HipifyAction.cpp @@ -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; } diff --git a/hipify-clang/src/HipifyAction.h b/hipify-clang/src/HipifyAction.h index ad987c921e..1262142cfc 100644 --- a/hipify-clang/src/HipifyAction.h +++ b/hipify-clang/src/HipifyAction.h @@ -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; diff --git a/hipify-clang/src/Statistics.cpp b/hipify-clang/src/Statistics.cpp index 595f243857..2115c567a5 100644 --- a/hipify-clang/src/Statistics.cpp +++ b/hipify-clang/src/Statistics.cpp @@ -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 { diff --git a/hipify-clang/src/Statistics.h b/hipify-clang/src/Statistics.h index 00ecce3eda..1ded45f0e4 100644 --- a/hipify-clang/src/Statistics.h +++ b/hipify-clang/src/Statistics.h @@ -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; diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 07feddc1cd..41710e4165 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -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__ diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index 629ecf7c50..10967605d4 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/include/hip/hcc_detail/math_functions.h @@ -31,6 +31,12 @@ THE SOFTWARE. #include #include +// 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 diff --git a/include/hip/hcc_detail/program_state.hpp b/include/hip/hcc_detail/program_state.hpp index bdb87b3509..ac689fdb89 100644 --- a/include/hip/hcc_detail/program_state.hpp +++ b/include/hip/hcc_detail/program_state.hpp @@ -93,12 +93,11 @@ public: } }; -const std::unordered_map>& executables( - bool rebuild = false); +const std::unordered_map>& executables(); const std::unordered_map>>& -functions(bool rebuild = false); -const std::unordered_map& function_names(bool rebuild = false); -std::unordered_map& globals(bool rebuild = false); +functions(); +const std::unordered_map& function_names(); +std::unordered_map& globals(); hsa_executable_t load_executable(const std::string& file, hsa_executable_t executable, hsa_agent_t agent); diff --git a/include/hip/math_functions.h b/include/hip/math_functions.h index 23dddf9f9d..2dfec4551b 100644 --- a/include/hip/math_functions.h +++ b/include/hip/math_functions.h @@ -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 #if defined(__HIP_PLATFORM_HCC__) && !defined(__HIP_PLATFORM_NVCC__) #include diff --git a/lpl_ca/lpl.hpp b/lpl_ca/lpl.hpp index cbd7fe8386..c3992e43c0 100644 --- a/lpl_ca/lpl.hpp +++ b/lpl_ca/lpl.hpp @@ -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& sources, const std::vector& targets, const std::string& flags, const std::string& output) { static const auto d = [](const std::string* f) { remove(f->c_str()); }; - - std::unique_ptr tmp{&output, d}; + std::string temp_str = output + ".tmp"; + std::unique_ptr tmp{&temp_str, d}; redi::ipstream hipcc{make_hipcc_call(sources, targets, flags, *tmp), redi::pstream::pstderr}; diff --git a/samples/0_Intro/module_api/runKernel.cpp b/samples/0_Intro/module_api/runKernel.cpp index 129c458dd0..ed74a7e4eb 100644 --- a/samples/0_Intro/module_api/runKernel.cpp +++ b/samples/0_Intro/module_api/runKernel.cpp @@ -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); diff --git a/samples/0_Intro/module_api_global/runKernel.cpp b/samples/0_Intro/module_api_global/runKernel.cpp index ebae029a16..0047f79b7f 100644 --- a/samples/0_Intro/module_api_global/runKernel.cpp +++ b/samples/0_Intro/module_api_global/runKernel.cpp @@ -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); diff --git a/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp index 2485d455ef..385ed4775a 100644 --- a/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp +++ b/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp @@ -21,13 +21,13 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" -#include "hip/hip_runtime_api.h" +//#include "hip/hip_runtime_api.h" #include #include #include -#include +//#include -#define fileName "tex2dKernel.code.adipose" +#define fileName "tex2dKernel.code" texture 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, diff --git a/src/functional_grid_launch.inl b/src/functional_grid_launch.inl index 6283d1aaba..704cf9ffeb 100644 --- a/src/functional_grid_launch.inl +++ b/src/functional_grid_launch.inl @@ -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); diff --git a/src/program_state.cpp b/src/program_state.cpp index fab71b1def..e8aff28faf 100644 --- a/src/program_state.cpp +++ b/src/program_state.cpp @@ -74,15 +74,11 @@ vector copy_names_of_undefined_symbols(const symbol_section_accessor& se } const std::unordered_map>& -symbol_addresses(bool rebuild = false) { +symbol_addresses() { static unordered_map> 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 code_object_blob_for_process() { return r; } -const unordered_map>>& code_object_blobs(bool rebuild = false) { +const unordered_map>>& code_object_blobs() { static unordered_map>> r; static once_flag f; - auto cons = [rebuild]() { - // names of shared libraries who .kernel sections already loaded - static unordered_set lib_names; + call_once(f, []() { static vector> 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>>& code_object_blobs(bool reb } } } - }; - - - call_once(f, cons); - if (rebuild) { - cons(); - } + }); return r; } @@ -244,13 +216,13 @@ vector> function_names_for(const elfio& reader, section* return r; } -const vector>& function_names_for_process(bool rebuild = false) { +const vector>& function_names_for_process() { static constexpr const char self[] = "/proc/self/exe"; static vector> r; static once_flag f; - auto cons = [rebuild]() { + call_once(f, []() { elfio reader; if (!reader.load(self)) { @@ -261,26 +233,16 @@ const vector>& 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>& kernels(bool rebuild = false) { +const unordered_map>& kernels() { static unordered_map> 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>& 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>& -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> 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(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& function_names(bool rebuild) { +const unordered_map& function_names() { static unordered_map 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& function_names(bool rebuild) { return 0; }, nullptr); - }; - - call_once(f, cons); - if (rebuild) { - static mutex mtx; - lock_guard lck{mtx}; - cons(); - } + }); return r; } -const unordered_map>>& functions(bool rebuild) { +const unordered_map>>& functions() { static unordered_map>> r; static once_flag f; - auto cons = [rebuild]() { - if (rebuild) { - // do NOT clear r so we reuse instances of pair - // 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>>& fu } } } - }; - - call_once(f, cons); - if (rebuild) { - static mutex mtx; - lock_guard lck{mtx}; - cons(); - } + }); return r; } -unordered_map& globals(bool rebuild) { +unordered_map& globals() { static unordered_map 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. diff --git a/tests/hipify-clang/cuComplex/cuComplex_Julia.cu b/tests/hipify-clang/cuComplex/cuComplex_Julia.cu new file mode 100644 index 0000000000..8bf9587b94 --- /dev/null +++ b/tests/hipify-clang/cuComplex/cuComplex_Julia.cu @@ -0,0 +1,58 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// CHECK: #include +// 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 + +// 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(-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 +#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(); +}