diff --git a/bin/hipcc b/bin/hipcc index 4be1c084fb..c759802e2e 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -383,7 +383,6 @@ if($HIP_PLATFORM eq "nvcc"){ my $toolArgs = ""; # arguments to pass to the hcc or nvcc tool my $optArg = ""; # -O args -my $gArg = ""; # -g args foreach $arg (@ARGV) { @@ -494,10 +493,6 @@ foreach $arg (@ARGV) { $optArg = $arg; } - if($arg =~ m/^-g/) - { - $gArg = $arg; - } ## process linker response file for hip-clang ## extract object files from static library and pass them directly to @@ -876,7 +871,7 @@ if ($needHipHcc) { if ($HIP_PLATFORM eq "clang") { # Set default optimization level to -O3 for hip-clang. - if ($optArg eq "" and $gArg ne "-g") { + if ($optArg eq "") { $HIPCXXFLAGS .= " -O3"; $HIPLDFLAGS .= " -O3"; } diff --git a/docs/markdown/CUDNN_API_supported_by_HIP.md b/docs/markdown/CUDNN_API_supported_by_HIP.md index dfecc96f1b..7efa922a21 100644 --- a/docs/markdown/CUDNN_API_supported_by_HIP.md +++ b/docs/markdown/CUDNN_API_supported_by_HIP.md @@ -233,9 +233,11 @@ | define |`CUDNN_SEQDATA_DIM_COUNT` | 9.0 | | | struct |`cudnnSeqDataStruct` | 9.0 | | | struct* |`cudnnSeqDataDescriptor_t` | 9.0 | | -| enum |***`cudnnAttnQueryMap_t`*** | 9.0 | | +| unsigned |***`cudnnAttnQueryMap_t`*** | 9.0 | | | 0 |*`CUDNN_ATTN_QUERYMAP_ALL_TO_ONE`* | 9.0 | | -| 1 |*`CUDNN_ATTN_QUERYMAP_ONE_TO_ONE`* | 9.0 | | +| 1U << 0 |*`CUDNN_ATTN_QUERYMAP_ONE_TO_ONE`* | 9.0 | | +| 1 |*`CUDNN_ATTN_DISABLE_PROJ_BIASES`* | 10.1 Update 2 | | +| 1U << 1 |*`CUDNN_ATTN_ENABLE_PROJ_BIASES`* | 10.1 Update 2 | | | struct |`cudnnAttnStruct` | 9.0 | | | struct* |`cudnnAttnDescriptor_t` | 9.0 | | | enum |***`cudnnMultiHeadAttnWeightKind_t`*** | 9.0 | | @@ -243,6 +245,11 @@ | 1 |*`CUDNN_MH_ATTN_K_WEIGHTS`* | 9.0 | | | 2 |*`CUDNN_MH_ATTN_V_WEIGHTS`* | 9.0 | | | 3 |*`CUDNN_MH_ATTN_O_WEIGHTS`* | 9.0 | | +| 4 |*`CUDNN_MH_ATTN_Q_BIASES`* | 10.1 Update 2 | | +| 5 |*`CUDNN_MH_ATTN_K_BIASES`* | 10.1 Update 2 | | +| 6 |*`CUDNN_MH_ATTN_V_BIASES`* | 10.1 Update 2 | | +| 7 |*`CUDNN_MH_ATTN_O_BIASES`* | 10.1 Update 2 | | +| define 8 |`CUDNN_ATTN_WKIND_COUNT` | 10.1 Update 2 | | | enum |***`cudnnWgradMode_t`*** | 9.0 | | | 0 |*`CUDNN_WGRAD_MODE_ADD`* | 9.0 | | | 1 |*`CUDNN_WGRAD_MODE_SET`* | 9.0 | | diff --git a/hipify-clang/CMakeLists.txt b/hipify-clang/CMakeLists.txt index 7ba92cd93d..e06b33d103 100644 --- a/hipify-clang/CMakeLists.txt +++ b/hipify-clang/CMakeLists.txt @@ -65,24 +65,31 @@ target_link_libraries(hipify-clang PRIVATE LLVMOption LLVMCore) -if(WIN32) - target_link_libraries(hipify-clang PRIVATE version) -endif() - -if ((LLVM_PACKAGE_VERSION VERSION_EQUAL "7") OR (LLVM_PACKAGE_VERSION VERSION_GREATER "7")) +if (LLVM_PACKAGE_VERSION VERSION_GREATER_EQUAL "7") target_link_libraries(hipify-clang PRIVATE clangToolingInclusions) endif() -set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}") -set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS} -std=c++14") -if(MSVC) +if (MSVC) + target_link_libraries(hipify-clang PRIVATE version) target_compile_options(hipify-clang PRIVATE "/Od /GR- /EHs- /EHc-") set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} /SUBSYSTEM:WINDOWS") + set(StdCpp "/std:c++") else() set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -pthread -fno-rtti -fvisibility-inlines-hidden") + set(StdCpp "-std=c++") endif() -set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DIRS}/clang/${LLVM_VERSION_MAJOR}.${LLVM_VERSION_MINOR}.${LLVM_VERSION_PATCH}\\\"") +if (LLVM_PACKAGE_VERSION VERSION_GREATER_EQUAL "10.0") + string(APPEND StdCpp "14") +# MSVC starting from 1900 (VS 2015) supports only the following c++ std values: c++14|c++17|c++latest +elseif (MSVC) + set(StdCpp "") +else() + string(APPEND StdCpp "11") +endif() + +set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS} ${StdCpp} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DIRS}/clang/${LLVM_VERSION_MAJOR}.${LLVM_VERSION_MINOR}.${LLVM_VERSION_PATCH}\\\"") install(TARGETS hipify-clang DESTINATION bin) @@ -111,15 +118,15 @@ if (HIPIFY_CLANG_TESTS) (CUDA_VERSION VERSION_GREATER "9.2" AND LLVM_PACKAGE_VERSION VERSION_LESS "8.0") OR (CUDA_VERSION VERSION_GREATER "10.0" AND LLVM_PACKAGE_VERSION VERSION_LESS "9.0")) message(SEND_ERROR "CUDA ${CUDA_VERSION} is not supported by clang ${LLVM_PACKAGE_VERSION}.") - if (CUDA_VERSION VERSION_LESS "7.0") + if (CUDA_VERSION_MAJOR VERSION_LESS "7") message(STATUS "Please install CUDA 7.0 or higher.") - elseif ((CUDA_VERSION VERSION_EQUAL "7.0") OR (CUDA_VERSION VERSION_EQUAL "7.5")) + elseif (CUDA_VERSION_MAJOR VERSION_LESS "8") message(STATUS "Please install clang 3.8 or higher.") - elseif (CUDA_VERSION VERSION_EQUAL "8.0") + elseif (CUDA_VERSION_MAJOR VERSION_LESS "9") message(STATUS "Please install clang 4.0 or higher.") elseif (CUDA_VERSION VERSION_EQUAL "9.0") message(STATUS "Please install clang 6.0 or higher.") - elseif ((CUDA_VERSION VERSION_EQUAL "9.1") OR (CUDA_VERSION VERSION_EQUAL "9.2")) + elseif (CUDA_VERSION_MAJOR VERSION_LESS "10") message(STATUS "Please install clang 7.0 or higher.") elseif (CUDA_VERSION VERSION_EQUAL "10.0") message(STATUS "Please install clang 8.0 or higher.") diff --git a/hipify-clang/src/ArgParse.cpp b/hipify-clang/src/ArgParse.cpp index 8a5a5b6e6c..751de0d028 100644 --- a/hipify-clang/src/ArgParse.cpp +++ b/hipify-clang/src/ArgParse.cpp @@ -138,4 +138,9 @@ cl::list MacroNames("D", cl::Prefix, cl::cat(ToolTemplateCategory)); +cl::opt SkipExcludedPPConditionalBlocks("skip-excluded-preprocessor-conditional-blocks", + cl::desc("Enable default preprocessor behaviour by skipping undefined conditional blocks"), + cl::value_desc("skip-excluded-preprocessor-conditional-blocks"), + cl::cat(ToolTemplateCategory)); + cl::extrahelp CommonHelp(ct::CommonOptionsParser::HelpMessage); diff --git a/hipify-clang/src/ArgParse.h b/hipify-clang/src/ArgParse.h index 5b8c763647..64fb2e9678 100644 --- a/hipify-clang/src/ArgParse.h +++ b/hipify-clang/src/ArgParse.h @@ -52,3 +52,4 @@ extern cl::opt Examine; extern cl::extrahelp CommonHelp; extern cl::opt TranslateToRoc; extern cl::opt DashDash; +extern cl::opt SkipExcludedPPConditionalBlocks; diff --git a/hipify-clang/src/CUDA2HIP_DNN_API_types.cpp b/hipify-clang/src/CUDA2HIP_DNN_API_types.cpp index a7d277e4c2..cc372067b1 100644 --- a/hipify-clang/src/CUDA2HIP_DNN_API_types.cpp +++ b/hipify-clang/src/CUDA2HIP_DNN_API_types.cpp @@ -229,14 +229,21 @@ const std::map CUDA_DNN_TYPE_NAME_MAP{ {"CUDNN_SEQDATA_BATCH_DIM", {"HIPDNN_SEQDATA_BATCH_DIM", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1 {"CUDNN_SEQDATA_BEAM_DIM", {"HIPDNN_SEQDATA_BEAM_DIM", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 2 {"CUDNN_SEQDATA_VECT_DIM", {"HIPDNN_SEQDATA_VECT_DIM", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 3 - {"cudnnAttnQueryMap_t", {"hipdnnAttnQueryMap_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnAttnQueryMap_t", {"hipdnnAttnQueryMap_t", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, {"CUDNN_ATTN_QUERYMAP_ALL_TO_ONE", {"HIPDNN_ATTN_QUERYMAP_ALL_TO_ONE", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 - {"CUDNN_ATTN_QUERYMAP_ONE_TO_ONE", {"HIPDNN_ATTN_QUERYMAP_ONE_TO_ONE", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1 + {"CUDNN_ATTN_QUERYMAP_ONE_TO_ONE", {"HIPDNN_ATTN_QUERYMAP_ONE_TO_ONE", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1U << 0 + {"CUDNN_ATTN_DISABLE_PROJ_BIASES", {"HIPDNN_ATTN_DISABLE_PROJ_BIASES", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 + {"CUDNN_ATTN_ENABLE_PROJ_BIASES", {"HIPDNN_ATTN_ENABLE_PROJ_BIASES", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1U << 1 {"cudnnMultiHeadAttnWeightKind_t", {"hipdnnMultiHeadAttnWeightKind_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, {"CUDNN_MH_ATTN_Q_WEIGHTS", {"HIPDNN_MH_ATTN_Q_WEIGHTS", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 {"CUDNN_MH_ATTN_K_WEIGHTS", {"HIPDNN_MH_ATTN_K_WEIGHTS", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1 {"CUDNN_MH_ATTN_V_WEIGHTS", {"HIPDNN_MH_ATTN_V_WEIGHTS", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 2 {"CUDNN_MH_ATTN_O_WEIGHTS", {"HIPDNN_MH_ATTN_O_WEIGHTS", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 3 + {"CUDNN_MH_ATTN_Q_BIASES", {"HIPDNN_MH_ATTN_Q_BIASES", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 4 + {"CUDNN_MH_ATTN_K_BIASES", {"HIPDNN_MH_ATTN_K_BIASES", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 5 + {"CUDNN_MH_ATTN_V_BIASES", {"HIPDNN_MH_ATTN_V_BIASES", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 6 + {"CUDNN_MH_ATTN_O_BIASES", {"HIPDNN_MH_ATTN_O_BIASES", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 7 + {"CUDNN_ATTN_WKIND_COUNT", {"HIPDNN_ATTN_WKIND_COUNT", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 8 {"cudnnWgradMode_t", {"hipdnnWgradMode_t", "", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, {"CUDNN_WGRAD_MODE_ADD", {"HIPDNN_WGRAD_MODE_ADD", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 {"CUDNN_WGRAD_MODE_SET", {"HIPDNN_WGRAD_MODE_SET", "", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1 diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index 52ba2eeaaf..e223ab7f01 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/hipify-clang/src/HipifyAction.cpp @@ -484,6 +484,11 @@ public: } +bool HipifyAction::BeginInvocation(clang::CompilerInstance &CI) { + llcompat::RetainExcludedConditionalBlocks(CI); + return true; +} + void HipifyAction::ExecuteAction() { clang::Preprocessor& PP = getCompilerInstance().getPreprocessor(); clang::SourceManager& SM = getCompilerInstance().getSourceManager(); diff --git a/hipify-clang/src/HipifyAction.h b/hipify-clang/src/HipifyAction.h index adafba2df9..31ccc0b648 100644 --- a/hipify-clang/src/HipifyAction.h +++ b/hipify-clang/src/HipifyAction.h @@ -91,6 +91,8 @@ protected: void insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL); // FrontendAction entry point. void ExecuteAction() override; + // Callback before starting processing a single input; used by hipify-clang for setting Preprocessor options. + bool BeginInvocation(clang::CompilerInstance &CI) override; // Called at the start of each new file to process. void EndSourceFileAction() override; // MatchCallback API entry point. Called by the AST visitor while searching the AST for things we registered an interest for. diff --git a/hipify-clang/src/LLVMCompat.cpp b/hipify-clang/src/LLVMCompat.cpp index d2573ecf22..8bb3eeda25 100644 --- a/hipify-clang/src/LLVMCompat.cpp +++ b/hipify-clang/src/LLVMCompat.cpp @@ -20,8 +20,13 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +#include "ArgParse.h" #include "LLVMCompat.h" #include "llvm/Support/Path.h" +#include "clang/Lex/PreprocessorOptions.h" +#include "clang/Frontend/CompilerInstance.h" + +const std::string sHipify = "[HIPIFY] ", sConflict = "conflict: ", sError = "error: ", sWarning = "warning: "; namespace llcompat { @@ -122,4 +127,20 @@ bool pragma_once_outside_header() { #endif } +void RetainExcludedConditionalBlocks(clang::CompilerInstance &CI) { +#if LLVM_VERSION_MAJOR > 9 + clang::PreprocessorOptions &PPOpts = CI.getPreprocessorOpts(); + PPOpts.RetainExcludedConditionalBlocks = !SkipExcludedPPConditionalBlocks; +#endif +} + +bool CheckCompatibility() { +#if LLVM_VERSION_MAJOR < 10 + if (SkipExcludedPPConditionalBlocks) { + llvm::errs() << "\n" << sHipify << sWarning << "Option '" << SkipExcludedPPConditionalBlocks.ArgStr.str() << "' is supported starting from LLVM version 10.0\n"; + } +#endif + return true; +} + } // namespace llcompat diff --git a/hipify-clang/src/LLVMCompat.h b/hipify-clang/src/LLVMCompat.h index 069aeb91a3..c91e2815e2 100644 --- a/hipify-clang/src/LLVMCompat.h +++ b/hipify-clang/src/LLVMCompat.h @@ -30,6 +30,8 @@ THE SOFTWARE. namespace ct = clang::tooling; +extern const std::string sHipify, sConflict, sError, sWarning; + // Things for papering over the differences between different LLVM versions. namespace llcompat { @@ -83,4 +85,8 @@ std::error_code real_path(const Twine &path, SmallVectorImpl &output, bool pragma_once_outside_header(); +void RetainExcludedConditionalBlocks(clang::CompilerInstance &CI); + +bool CheckCompatibility(); + } // namespace llcompat diff --git a/hipify-clang/src/main.cpp b/hipify-clang/src/main.cpp index 3b8f454fd7..d6a53b170d 100644 --- a/hipify-clang/src/main.cpp +++ b/hipify-clang/src/main.cpp @@ -37,7 +37,6 @@ THE SOFTWARE. #define DEBUG_TYPE "cuda2hip" -std::string sHipify = "[HIPIFY] ", sConflict = "conflict: ", sError = "error: "; namespace ct = clang::tooling; std::string getAbsoluteFilePath(const std::string& sFile, std::error_code& EC) { @@ -240,13 +239,16 @@ int main(int argc, const char **argv) { } llcompat::PrintStackTraceOnErrorSignal(); ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::ZeroOrMore); + if (!llcompat::CheckCompatibility()) { + return 1; + } std::vector fileSources = OptionsParser.getSourcePathList(); if (fileSources.empty() && !GeneratePerl && !GeneratePython) { - llvm::errs() << "\n" << sHipify << sError << "Must specify at least 1 positional argument for source file." << "\n"; + llvm::errs() << "\n" << sHipify << sError << "Must specify at least 1 positional argument for source file" << "\n"; return 1; } if (!generatePerl(GeneratePerl)) { - llvm::errs() << "\n" << sHipify << sError << "hipify-perl generating failed." << "\n"; + llvm::errs() << "\n" << sHipify << sError << "hipify-perl generating failed" << "\n"; return 1; } bool bToRoc = TranslateToRoc; @@ -254,7 +256,7 @@ int main(int argc, const char **argv) { bool bToPython = generatePython(GeneratePython); TranslateToRoc = bToRoc; if (!bToPython) { - llvm::errs() << "\n" << sHipify << sError << "hipify-python generating failed." << "\n"; + llvm::errs() << "\n" << sHipify << sError << "hipify-python generating failed" << "\n"; return 1; } if (fileSources.empty()) { @@ -268,15 +270,15 @@ int main(int argc, const char **argv) { } if (!dst.empty()) { if (fileSources.size() > 1) { - llvm::errs() << sHipify << sConflict << "-o and multiple source files are specified.\n"; + llvm::errs() << sHipify << sConflict << "-o and multiple source files are specified\n"; return 1; } if (Inplace) { - llvm::errs() << sHipify << sConflict << "both -o and -inplace options are specified.\n"; + llvm::errs() << sHipify << sConflict << "both -o and -inplace options are specified\n"; return 1; } if (NoOutput) { - llvm::errs() << sHipify << sConflict << "both -no-output and -o options are specified.\n"; + llvm::errs() << sHipify << sConflict << "both -no-output and -o options are specified\n"; return 1; } if (!dstDir.empty()) { @@ -284,11 +286,11 @@ int main(int argc, const char **argv) { } } if (NoOutput && Inplace) { - llvm::errs() << sHipify << sConflict << "both -no-output and -inplace options are specified.\n"; + llvm::errs() << sHipify << sConflict << "both -no-output and -inplace options are specified\n"; return 1; } if (!dstDir.empty() && Inplace) { - llvm::errs() << sHipify << sConflict << "both -o-dir and -inplace options are specified.\n"; + llvm::errs() << sHipify << sConflict << "both -o-dir and -inplace options are specified\n"; return 1; } if (Examine) { @@ -425,8 +427,8 @@ int main(int argc, const char **argv) { // Hipify _all_ the things! if (Tool.runAndSave(&actionFactory)) { currentStat.hasErrors = true; - LLVM_DEBUG(llvm::dbgs() << "Skipped some replacements.\n"); Result = 1; + LLVM_DEBUG(llvm::dbgs() << "Skipped some replacements.\n"); } // Copy the tmpfile to the output if (!NoOutput && !currentStat.hasErrors) { diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 684449aa6b..5ad1ced41c 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -748,28 +748,36 @@ unsigned long long int __ballot64(int predicate) { // hip.amdgcn.bc - lanemask __device__ inline -int64_t __lanemask_gt() +uint64_t __lanemask_gt() { - int32_t activelane = __ockl_activelane_u32(); - int64_t ballot = __ballot64(1); - if (activelane != 63) { - int64_t tmp = (~0ULL) << (activelane + 1); - return tmp & ballot; - } - return 0; + uint32_t lane = __ockl_lane_u32(); + if (lane == 63) + return 0; + uint64_t ballot = __ballot64(1); + uint64_t mask = (~((uint64_t)0)) << (lane + 1); + return mask & ballot; } __device__ inline -int64_t __lanemask_lt() +uint64_t __lanemask_lt() { - int32_t activelane = __ockl_activelane_u32(); + uint32_t lane = __ockl_lane_u32(); int64_t ballot = __ballot64(1); - if (activelane == 0) - return 0; - return ballot; + uint64_t mask = ((uint64_t)1 << lane) - (uint64_t)1; + return mask & ballot; } +__device__ +inline +uint64_t __lanemask_eq() +{ + uint32_t lane = __ockl_lane_u32(); + int64_t mask = ((uint64_t)1 << lane); + return mask; +} + + __device__ inline void* __local_to_generic(void* p) { return p; } #ifdef __HIP_DEVICE_COMPILE__ diff --git a/include/hip/hcc_detail/device_library_decls.h b/include/hip/hcc_detail/device_library_decls.h index 9a7636fc4a..8cfb020caa 100644 --- a/include/hip/hcc_detail/device_library_decls.h +++ b/include/hip/hcc_detail/device_library_decls.h @@ -62,6 +62,16 @@ extern "C" __device__ __attribute__((const)) float __ocml_fmax_f32(float, float) extern "C" __device__ __attribute__((convergent)) void __ockl_gws_init(uint nwm1, uint rid); extern "C" __device__ __attribute__((convergent)) void __ockl_gws_barrier(uint nwm1, uint rid); +extern "C" __device__ __attribute__((const)) uint32_t __ockl_lane_u32(); +extern "C" __device__ __attribute__((const)) int __ockl_grid_is_valid(void); +extern "C" __device__ __attribute__((convergent)) void __ockl_grid_sync(void); +extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_num_grids(void); +extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_grid_rank(void); +extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_size(void); +extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_thread_rank(void); +extern "C" __device__ __attribute__((const)) int __ockl_multi_grid_is_valid(void); +extern "C" __device__ __attribute__((convergent)) void __ockl_multi_grid_sync(void); + // Introduce local address space #define __local __attribute__((address_space(3))) diff --git a/include/hip/hcc_detail/functional_grid_launch.hpp b/include/hip/hcc_detail/functional_grid_launch.hpp index 883470b1fb..c493eec933 100644 --- a/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/include/hip/hcc_detail/functional_grid_launch.hpp @@ -129,7 +129,7 @@ void hipLaunchKernelGGLImpl( template inline -void hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* blockSize, +hipError_t hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* blockSize, F kernel, size_t dynSharedMemPerBlk, uint32_t blockSizeLimit) { using namespace hip_impl; @@ -138,13 +138,13 @@ void hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* blockSize, auto f = get_program_state().kernel_descriptor(reinterpret_cast(kernel), target_agent(0)); - hipOccupancyMaxPotentialBlockSize(gridSize, blockSize, f, + return hipOccupancyMaxPotentialBlockSize(gridSize, blockSize, f, dynSharedMemPerBlk, blockSizeLimit); } template inline -void hipOccupancyMaxActiveBlocksPerMultiprocessor(uint32_t* numBlocks, F kernel, +hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(uint32_t* numBlocks, F kernel, uint32_t blockSize, size_t dynSharedMemPerBlk) { using namespace hip_impl; @@ -153,7 +153,7 @@ void hipOccupancyMaxActiveBlocksPerMultiprocessor(uint32_t* numBlocks, F kernel, auto f = get_program_state().kernel_descriptor(reinterpret_cast(kernel), target_agent(0)); - hipOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, f, blockSize, dynSharedMemPerBlk); + return hipOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, f, blockSize, dynSharedMemPerBlk); } template @@ -177,13 +177,4 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, stream, &config[0]); } -template -[[deprecated("hipLaunchKernel is deprecated and will be removed in the next " - "version of HIP; please upgrade to hipLaunchKernelGGL.")]] -inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBlocks, - std::uint32_t groupMemBytes, hipStream_t stream, Args... args) { - hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{}, - std::move(args)...); -} - #pragma GCC visibility pop diff --git a/include/hip/hcc_detail/hip_cooperative_groups.h b/include/hip/hcc_detail/hip_cooperative_groups.h new file mode 100644 index 0000000000..27ce887723 --- /dev/null +++ b/include/hip/hcc_detail/hip_cooperative_groups.h @@ -0,0 +1,216 @@ +/* +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. +*/ + +/** + * @file hcc_detail/hip_cooperative_groups.h + * + * @brief Device side implementation of `Cooperative Group` feature. + * + * Defines new types and device API wrappers related to `Cooperative Group` + * feature, which the programmer can directly use in his kernel(s) in order to + * make use of this feature. + */ +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COOPERATIVE_GROUPS_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COOPERATIVE_GROUPS_H + +#if __cplusplus +#include + +namespace cooperative_groups { + +/** \brief The base type of all cooperative group types + * + * \details Holds the key properties of a constructed cooperative group type + * object, like the group type, its size, etc + */ +class thread_group { + protected: + uint32_t _type; // thread_group type + uint32_t _size; // total number of threads in the tread_group + uint64_t _mask; // Lanemask for coalesced and tiled partitioned group types, + // LSB represents lane 0, and MSB represents lane 63 + + // Construct a thread group, and set thread group type and other essential + // thread group properties. This generic thread group is directly constructed + // only when the group is supposed to contain only the calling the thread + // (throurh the API - `this_thread()`), and in all other cases, this thread + // group object is a sub-object of some other derived thread group object + __CG_QUALIFIER__ thread_group(internal::group_type type, uint32_t size, + uint64_t mask = (uint64_t)0) { + _type = type; + _size = size; + _mask = mask; + } + + public: + // Total number of threads in the thread group, and this serves the purpose + // for all derived cooperative group types since their `size` is directly + // saved during the construction + __CG_QUALIFIER__ uint32_t size() const { + return _size; + } + // Rank of the calling thread within [0, size()) + __CG_QUALIFIER__ uint32_t thread_rank() const; + // Is this cooperative group type valid? + __CG_QUALIFIER__ bool is_valid() const; + // synchronize the threads in the thread group + __CG_QUALIFIER__ void sync() const; +}; + +/** \brief The multi-grid cooperative group type + * + * \details Represents an inter-device cooperative group type where the + * participating threads within the group spans across multple + * devices, running the (same) kernel on these devices + */ +class multi_grid_group : public thread_group { + // Only these friend functions are allowed to construct an object of this class + // and access its resources + friend __CG_QUALIFIER__ multi_grid_group this_multi_grid(); + + protected: + // Construct mutli-grid thread group (through the API this_multi_grid()) + explicit __CG_QUALIFIER__ multi_grid_group(uint32_t size) + : thread_group(internal::cg_multi_grid, size) { } + + public: + // Number of invocations participating in this multi-grid group. In other + // words, the number of GPUs + __CG_QUALIFIER__ uint32_t num_grids() { + return internal::multi_grid::num_grids(); + } + // Rank of this invocation. In other words, an ID number within the range + // [0, num_grids()) of the GPU, this kernel is running on + __CG_QUALIFIER__ uint32_t grid_rank() { + return internal::multi_grid::grid_rank(); + } + __CG_QUALIFIER__ uint32_t thread_rank() const { + return internal::multi_grid::thread_rank(); + } + __CG_QUALIFIER__ bool is_valid() const { + return internal::multi_grid::is_valid(); + } + __CG_QUALIFIER__ void sync() const { + internal::multi_grid::sync(); + } +}; + +/** \brief User exposed API interface to construct multi-grid cooperative + * group type object - `multi_grid_group` + * + * \details User is not allowed to directly construct an object of type + * `multi_grid_group`. Instead, he should construct it through this + * API function + */ +__CG_QUALIFIER__ multi_grid_group +this_multi_grid() { + return multi_grid_group(internal::multi_grid::size()); +} + +/** \brief The grid cooperative group type + * + * \details Represents an inter-workgroup cooperative group type where the + * participating threads within the group spans across multiple + * workgroups running the (same) kernel on the same device + */ +class grid_group : public thread_group { + // Only these friend functions are allowed to construct an object of this class + // and access its resources + friend __CG_QUALIFIER__ grid_group this_grid(); + + protected: + // Construct grid thread group (through the API this_grid()) + explicit __CG_QUALIFIER__ grid_group(uint32_t size) + : thread_group(internal::cg_grid, size) { } + + public: + __CG_QUALIFIER__ uint32_t thread_rank() const { + return internal::grid::thread_rank(); + } + __CG_QUALIFIER__ bool is_valid() const { + return internal::grid::is_valid(); + } + __CG_QUALIFIER__ void sync() const { + internal::grid::sync(); + } +}; + +/** \brief User exposed API interface to construct grid cooperative group type + * object - `grid_group` + * + * \details User is not allowed to directly construct an object of type + * `multi_grid_group`. Instead, he should construct it through this + * API function + */ +__CG_QUALIFIER__ grid_group +this_grid() { + return grid_group(internal::grid::size()); +} + +/** + * Implemenation of all publicly exposed base class APIs + */ +__CG_QUALIFIER__ uint32_t thread_group::thread_rank() const { + switch (this->_type) { + case internal::cg_multi_grid: { + return (static_cast(this)->thread_rank()); + } + case internal::cg_grid: { + return (static_cast(this)->thread_rank()); + } + default: { + return 0; //TODO(mahesha) + } + } +} + +__CG_QUALIFIER__ bool thread_group::is_valid() const { + switch (this->_type) { + case internal::cg_multi_grid: { + return (static_cast(this)->is_valid()); + } + case internal::cg_grid: { + return (static_cast(this)->is_valid()); + } + default: { + return false; + } + } +} + +__CG_QUALIFIER__ void thread_group::sync() const { + switch (this->_type) { + case internal::cg_multi_grid: { + static_cast(this)->sync(); + break; + } + case internal::cg_grid: { + static_cast(this)->sync(); + break; + } + } +} + +} // namespace cooperative_groups + +#endif // __cplusplus +#endif // HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COOPERATIVE_GROUPS_H diff --git a/include/hip/hcc_detail/hip_cooperative_groups_helper.h b/include/hip/hcc_detail/hip_cooperative_groups_helper.h new file mode 100644 index 0000000000..b74d16d23b --- /dev/null +++ b/include/hip/hcc_detail/hip_cooperative_groups_helper.h @@ -0,0 +1,144 @@ +/* +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. +*/ + +/** + * @file hcc_detail/hip_cooperative_groups_helper.h + * + * @brief Device side implementation of cooperative group feature. + * + * Defines helper constructs and APIs which aid the types and device API + * wrappers defined within `hcc_detail/hip_cooperative_groups.h`. + */ +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COOPERATIVE_GROUPS_HELPER_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COOPERATIVE_GROUPS_HELPER_H + +#if __cplusplus +#include +#include + +#if !defined(__align__) +#define __align__(x) __attribute__((aligned(x))) +#endif + +#if !defined(__CG_QUALIFIER__) +#define __CG_QUALIFIER__ __device__ __forceinline__ +#endif + +#if !defined(__CG_STATIC_QUALIFIER__) +#define __CG_STATIC_QUALIFIER__ __device__ static __forceinline__ +#endif + +#if !defined(WAVEFRONT_SIZE) +#define WAVEFRONT_SIZE 64 +#endif + +namespace cooperative_groups { + +namespace internal { + +/** \brief Enums representing different cooperative group types + */ +typedef enum { + cg_invalid, + cg_multi_grid, + cg_grid +} group_type; + +/** + * Functionalities related to multi-grid cooperative group type + */ +namespace multi_grid { + +__CG_STATIC_QUALIFIER__ uint32_t num_grids() { + return (uint32_t)__ockl_multi_grid_num_grids(); +} + +__CG_STATIC_QUALIFIER__ uint32_t grid_rank() { + return (uint32_t)__ockl_multi_grid_grid_rank(); +} + +__CG_STATIC_QUALIFIER__ uint32_t size() { + return (uint32_t)__ockl_multi_grid_size(); +} + +__CG_STATIC_QUALIFIER__ uint32_t thread_rank() { + return (uint32_t)__ockl_multi_grid_thread_rank(); +} + +__CG_STATIC_QUALIFIER__ bool is_valid() { + return (bool)__ockl_multi_grid_is_valid(); +} + +__CG_STATIC_QUALIFIER__ void sync() { + __ockl_multi_grid_sync(); +} + +} // namespace multi_grid + +/** + * Functionalities related to grid cooperative group type + */ +namespace grid { + +__CG_STATIC_QUALIFIER__ uint32_t size() { + return (uint32_t)((hipBlockDim_z * hipGridDim_z) * + (hipBlockDim_y * hipGridDim_y) * + (hipBlockDim_x * hipGridDim_x)); +} + +__CG_STATIC_QUALIFIER__ uint32_t thread_rank() { + // Compute global id of the workgroup to which the current threads belongs to + uint32_t blkIdx = + (uint32_t)((hipBlockIdx_z * hipGridDim_y * hipGridDim_x) + + (hipBlockIdx_y * hipGridDim_x) + + (hipBlockIdx_x)); + + // Compute total number of threads being passed to reach current workgroup + // within grid + uint32_t num_threads_till_current_workgroup = + (uint32_t)(blkIdx * (hipBlockIdx_x * hipBlockIdx_y * hipBlockIdx_z)); + + // Compute thread local rank within current workgroup + uint32_t local_thread_rank = + (uint32_t)((hipThreadIdx_z * hipBlockDim_y * hipBlockDim_x) + + (hipThreadIdx_y * hipBlockDim_x) + + (hipThreadIdx_x)); + + return (num_threads_till_current_workgroup + local_thread_rank); +} + +__CG_STATIC_QUALIFIER__ bool is_valid() { + return (bool)__ockl_grid_is_valid(); +} + +__CG_STATIC_QUALIFIER__ void sync() { + __ockl_grid_sync(); +} + +} // namespace grid + +} // namespace internal + +} // namespace cooperative_groups + +#endif // __cplusplus +#endif // HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COOPERATIVE_GROUPS_HELPER_H diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 97ff4bd1d7..c93aa67994 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -1520,6 +1520,34 @@ hipError_t hipMemcpyToSymbol(void*, const void*, size_t, size_t, hipMemcpyKind, } // Namespace hip_impl. #endif +#ifdef __cplusplus +extern "C" { +#endif + +/** + * @brief C compliant kernel launch API + * + * @param [in] function_address - kernel function pointer. + * @param [in] numBlocks - number of blocks + * @param [in] dimBlocks - dimension of a block + * @param [in] args - kernel arguments + * @param [in] sharedMemBytes - Amount of dynamic shared memory to allocate for this kernel. The + * Kernel can access this with HIP_DYNAMIC_SHARED. + * @param [in] stream - Stream where the kernel should be dispatched. May be 0, in which case th + * default stream is used with associated synchronization rules. + * + * @returns #hipSuccess, #hipErrorInvalidValue, hipInvalidDevice + * + */ + +hipError_t hipLaunchKernel(const void* function_address, + dim3 numBlocks, dim3 dimBlocks, void** args, + size_t sharedMemBytes, hipStream_t stream); + +#ifdef __cplusplus +} +#endif + #if defined(__cplusplus) extern "C" { #endif diff --git a/include/hip/hip_cooperative_groups.h b/include/hip/hip_cooperative_groups.h new file mode 100644 index 0000000000..d919e83c7f --- /dev/null +++ b/include/hip/hip_cooperative_groups.h @@ -0,0 +1,43 @@ +/* +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. +*/ + +/** + * @file hip_cooperative_groups.h + * + * @brief Defines new types and device API wrappers for `Cooperative Group` + * feature. + */ + +#ifndef HIP_INCLUDE_HIP_HIP_COOPERATIVE_GROUP_H +#define HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H + +#if defined(__HIP_PLATFORM_HCC__) && !defined(__HIP_PLATFORM_NVCC__) +#if __cplusplus +#include +#endif +#elif defined(__HIP_PLATFORM_NVCC__) && !defined(__HIP_PLATFORM_HCC__) +#include +#else +#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); +#endif + +#endif // HIP_INCLUDE_HIP_HIP_COOPERATIVE_GROUP_H diff --git a/include/hip/nvcc_detail/hip_runtime.h b/include/hip/nvcc_detail/hip_runtime.h index 3a18f11614..d159e116c0 100644 --- a/include/hip/nvcc_detail/hip_runtime.h +++ b/include/hip/nvcc_detail/hip_runtime.h @@ -31,11 +31,6 @@ THE SOFTWARE. typedef int hipLaunchParm; -#define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ - do { \ - kernelName<<>>(0, ##__VA_ARGS__); \ - } while (0) - #define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ do { \ kernelName<<>>(__VA_ARGS__); \ diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index 532e1d827e..095943313c 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -216,38 +216,80 @@ inline static hipError_t hipCUDAErrorTohipError(cudaError_t cuError) { switch (cuError) { case cudaSuccess: return hipSuccess; + case cudaErrorProfilerDisabled: + return hipErrorProfilerDisabled; + case cudaErrorProfilerNotInitialized: + return hipErrorProfilerNotInitialized; + case cudaErrorProfilerAlreadyStarted: + return hipErrorProfilerAlreadyStarted; + case cudaErrorProfilerAlreadyStopped: + return hipErrorProfilerAlreadyStopped; + case cudaErrorInsufficientDriver: + return hipErrorInsufficientDriver; + case cudaErrorUnsupportedLimit: + return hipErrorUnsupportedLimit; + case cudaErrorPeerAccessUnsupported: + return hipErrorPeerAccessUnsupported; + case cudaErrorInvalidGraphicsContext: + return hipErrorInvalidGraphicsContext; + case cudaErrorSharedObjectSymbolNotFound: + return hipErrorSharedObjectSymbolNotFound; + case cudaErrorSharedObjectInitFailed: + return hipErrorSharedObjectInitFailed; + case cudaErrorOperatingSystem: + return hipErrorOperatingSystem; + case cudaErrorSetOnActiveProcess: + return hipErrorSetOnActiveProcess; + case cudaErrorIllegalAddress: + return hipErrorIllegalAddress; + case cudaErrorInvalidSymbol: + return hipErrorInvalidSymbol; + case cudaErrorMissingConfiguration: + return hipErrorMissingConfiguration; case cudaErrorMemoryAllocation: return hipErrorMemoryAllocation; - case cudaErrorLaunchOutOfResources: - return hipErrorLaunchOutOfResources; - case cudaErrorInvalidValue: - return hipErrorInvalidValue; - case cudaErrorInvalidResourceHandle: - return hipErrorInvalidResourceHandle; - case cudaErrorInvalidDevice: - return hipErrorInvalidDevice; - case cudaErrorInvalidMemcpyDirection: - return hipErrorInvalidMemcpyDirection; - case cudaErrorInvalidDevicePointer: - return hipErrorInvalidDevicePointer; case cudaErrorInitializationError: return hipErrorInitializationError; - case cudaErrorNoDevice: - return hipErrorNoDevice; - case cudaErrorNotReady: - return hipErrorNotReady; + case cudaErrorLaunchFailure: + return hipErrorLaunchFailure; + case cudaErrorPriorLaunchFailure: + return hipErrorPriorLaunchFailure; + case cudaErrorLaunchOutOfResources: + return hipErrorLaunchOutOfResources; + case cudaErrorInvalidDeviceFunction: + return hipErrorInvalidDeviceFunction; + case cudaErrorInvalidConfiguration: + return hipErrorInvalidConfiguration; + case cudaErrorInvalidDevice: + return hipErrorInvalidDevice; + case cudaErrorInvalidValue: + return hipErrorInvalidValue; + case cudaErrorInvalidDevicePointer: + return hipErrorInvalidDevicePointer; + case cudaErrorInvalidMemcpyDirection: + return hipErrorInvalidMemcpyDirection; case cudaErrorUnknown: return hipErrorUnknown; - case cudaErrorPeerAccessNotEnabled: - return hipErrorPeerAccessNotEnabled; + case cudaErrorInvalidResourceHandle: + return hipErrorInvalidResourceHandle; + case cudaErrorNotReady: + return hipErrorNotReady; + case cudaErrorNoDevice: + return hipErrorNoDevice; case cudaErrorPeerAccessAlreadyEnabled: return hipErrorPeerAccessAlreadyEnabled; + case cudaErrorPeerAccessNotEnabled: + return hipErrorPeerAccessNotEnabled; case cudaErrorHostMemoryAlreadyRegistered: return hipErrorHostMemoryAlreadyRegistered; case cudaErrorHostMemoryNotRegistered: return hipErrorHostMemoryNotRegistered; - case cudaErrorUnsupportedLimit: - return hipErrorUnsupportedLimit; + case cudaErrorMapBufferObjectFailed: + return hipErrorMapBufferObjectFailed; + case cudaErrorAssert: + return hipErrorAssert; + case cudaErrorNotSupported: + return hipErrorNotSupported; default: return hipErrorUnknown; // Note - translated error. } @@ -1241,6 +1283,13 @@ inline static hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* im cuModuleLoadDataEx(module, image, numOptions, options, optionValues)); } +inline static hipError_t hipLaunchKernel(const void* function_address, dim3 numBlocks, + dim3 dimBlocks, void** args, size_t sharedMemBytes, + hipStream_t stream) +{ + return hipCUDAErrorTohipError(cudaLaunchKernel(function_address,numBlocks,dimBlocks,args,sharedMemBytes,stream)); +} + inline static hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, diff --git a/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp index 90645f6b72..6fd49fdb0f 100644 --- a/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp +++ b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp @@ -21,7 +21,12 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" -extern texture tex; +#if __HIP__ +__hip_pinned_shadow__ +#else +extern +#endif +texture tex; extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) { int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; diff --git a/samples/2_Cookbook/13_occupancy/Makefile b/samples/2_Cookbook/13_occupancy/Makefile new file mode 100644 index 0000000000..3ec069d9e6 --- /dev/null +++ b/samples/2_Cookbook/13_occupancy/Makefile @@ -0,0 +1,20 @@ +HIP_PATH?= $(wildcard /opt/rocm/hip) +ifeq (,$(HIP_PATH)) + HIP_PATH=../../.. +endif +HIPCC=$(HIP_PATH)/bin/hipcc + +EXE=./occupancy + +.PHONY: test + +all: test + +$(EXE): occupancy.cpp + $(HIPCC) $^ -o $@ + +test: $(EXE) + $(EXE) + +clean: + rm -f *.o $(EXE) diff --git a/samples/2_Cookbook/13_occupancy/occupancy.cpp b/samples/2_Cookbook/13_occupancy/occupancy.cpp new file mode 100644 index 0000000000..a9f4e198b0 --- /dev/null +++ b/samples/2_Cookbook/13_occupancy/occupancy.cpp @@ -0,0 +1,176 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hip/hip_runtime.h" +#include +#define NUM 1000000 + +#define HIP_CHECK(status) \ + if (status != hipSuccess) { \ + std::cout << "Got Status: " << status << " at Line: " << __LINE__ << std::endl; \ + exit(0); \ + } + +// Device (Kernel) function +__global__ void multiply(float* C, float* A, float* B, int N){ + + int tx = blockDim.x*blockIdx.x+threadIdx.x; + + if (tx < N){ + C[tx] = A[tx] * B[tx]; + } +} +// CPU implementation +void multiplyCPU(float* C, float* A, float* B, int N){ + + for(unsigned int i=0; i eps) { + errors++; + } + } + + if (errors != 0){ + printf("\nManual Test FAILED: %d errors\n", errors); + errors=0; + } else { + printf("\nManual Test PASSED!\n"); + } + + for (i = 0; i < NUM; i++) { + if (std::abs(C1[i] - cpuC[i]) > eps) { + errors++; + } + } + + if (errors != 0){ + printf("\n Automatic Test FAILED: %d errors\n", errors); + } else { + printf("\nAutomatic Test PASSED!\n"); + } + + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Bd)); + HIP_CHECK(hipFree(C0d)); + HIP_CHECK(hipFree(C1d)); + + free(A); + free(B); + free(C0); + free(C1); + free(cpuC); +} diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index da0b38ed4f..96b8e84298 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -1671,6 +1671,8 @@ const char* ihipErrorString(hipError_t hip_error) { return "hipErrorProfilerAlreadyStarted"; case hipErrorProfilerAlreadyStopped: return "hipErrorProfilerAlreadyStopped"; + case hipErrorInsufficientDriver: + return "hipErrorInsufficientDriver"; case hipErrorInvalidImage: return "hipErrorInvalidImage"; case hipErrorInvalidContext: @@ -1725,6 +1727,8 @@ const char* ihipErrorString(hipError_t hip_error) { return "hipErrorNotFound"; case hipErrorIllegalAddress: return "hipErrorIllegalAddress"; + case hipErrorInvalidSymbol: + return "hipErrorInvalidSymbol"; case hipErrorMissingConfiguration: return "hipErrorMissingConfiguration"; @@ -1773,6 +1777,12 @@ const char* ihipErrorString(hipError_t hip_error) { return "hipErrorHostMemoryAlreadyRegistered"; case hipErrorHostMemoryNotRegistered: return "hipErrorHostMemoryNotRegistered"; + case hipErrorMapBufferObjectFailed: + return "hipErrorMapBufferObjectFailed"; + case hipErrorAssert: + return "hipErrorAssert"; + case hipErrorNotSupported: + return "hipErrorNotSupported"; case hipErrorTbd: return "hipErrorTbd"; default: diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 50b3cc9b04..2548660a72 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -1238,3 +1238,31 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor( tls, numBlocks, f, blockSize, dynSharedMemPerBlk)); } + +hipError_t hipLaunchKernel( + const void* func_addr, dim3 numBlocks, dim3 dimBlocks, void** args, + size_t sharedMemBytes, hipStream_t stream) +{ + HIP_INIT_API(hipLaunchKernel,func_addr,numBlocks,dimBlocks,args,sharedMemBytes,stream); + + hipFunction_t kd = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)func_addr, + hip_impl::target_agent(stream)); + + if(kd == nullptr || kd->_header == nullptr) + return ihipLogStatus(hipErrorInvalidValue); + + size_t szKernArg = kd->_header->kernarg_segment_byte_size; + + if(args == NULL && szKernArg != 0) + return ihipLogStatus(hipErrorInvalidValue); + + void* config[]{ + HIP_LAUNCH_PARAM_BUFFER_POINTER, + args, + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &szKernArg, + HIP_LAUNCH_PARAM_END}; + + return ihipLogStatus(ihipModuleLaunchKernel(tls, kd, numBlocks.x * dimBlocks.x, numBlocks.y * dimBlocks.y, numBlocks.z * dimBlocks.z, + dimBlocks.x, dimBlocks.y, dimBlocks.z, sharedMemBytes, stream, nullptr, (void**)&config, nullptr, nullptr, 0)); +} diff --git a/tests/hipify-clang/lit.cfg b/tests/hipify-clang/lit.cfg index 594d5fcf69..64f82e57fa 100644 --- a/tests/hipify-clang/lit.cfg +++ b/tests/hipify-clang/lit.cfg @@ -42,6 +42,10 @@ if config.cuda_version_major < 10: config.excludes.append('cuSPARSE_10.cu') config.excludes.append('cuSPARSE_11.cu') +if config.llvm_version_major < 10: + config.excludes.append('pp_if_else_conditionals_LLVM_10.cu') + config.excludes.append('pp_if_else_conditionals_01_LLVM_10.cu') + # name: The name of this test suite. config.name = 'hipify' diff --git a/tests/hipify-clang/lit.site.cfg.in b/tests/hipify-clang/lit.site.cfg.in index 71cc3e08d3..6ef4dc007a 100644 --- a/tests/hipify-clang/lit.site.cfg.in +++ b/tests/hipify-clang/lit.site.cfg.in @@ -3,6 +3,7 @@ import os config.pointer_size = @CMAKE_SIZEOF_VOID_P@ config.llvm_version = "@LLVM_PACKAGE_VERSION@" +config.llvm_version_major = int("@LLVM_VERSION_MAJOR@") config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@" config.obj_root = "@CMAKE_CURRENT_BINARY_DIR@" config.cuda_root = "@CUDA_TOOLKIT_ROOT_DIR@" diff --git a/tests/hipify-clang/unit_tests/libraries/cuComplex/cuComplex_Julia.cu b/tests/hipify-clang/unit_tests/libraries/cuComplex/cuComplex_Julia.cu index 196a94197b..5c2cd3b1e0 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuComplex/cuComplex_Julia.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuComplex/cuComplex_Julia.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args "--skip-excluded-preprocessor-conditional-blocks" %clang_args // CHECK: #include // CHECK: #include "hip/hip_complex.h" diff --git a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu index 349d6471db..8618be4748 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args "--skip-excluded-preprocessor-conditional-blocks" %clang_args #include #include #include diff --git a/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals.cu b/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals.cu new file mode 100644 index 0000000000..51bfeb6017 --- /dev/null +++ b/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals.cu @@ -0,0 +1,30 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args "--skip-excluded-preprocessor-conditional-blocks" %clang_args +// CHECK: #include + +#include + +__global__ void axpy_kernel(float a, float* x, float* y) { + y[threadIdx.x] = a * x[threadIdx.x]; +} + +void axpy(float a, float* x, float* y) { + +#ifdef SOME_MACRO + // CHECK: axpy_kernel <<<1, 1>>> (a, y, x); + axpy_kernel <<<1, 1>>> (a, y, x); +#endif + +#ifndef SOME_MACRO + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(2), 0, 0, a, y, x); + axpy_kernel <<<1, 2>>> (a, y, x); +#endif + +#ifdef SOME_MACRO + // CHECK: axpy_kernel <<<1, 3>>> (a, y, x); + axpy_kernel <<<1, 3>>> (a, y, x); +#else + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(4), 0, 0, a, x, y); + axpy_kernel <<<1, 4>>> (a, x, y); +#endif + +} \ No newline at end of file diff --git a/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_01.cu b/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_01.cu new file mode 100644 index 0000000000..310d896054 --- /dev/null +++ b/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_01.cu @@ -0,0 +1,52 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args "--skip-excluded-preprocessor-conditional-blocks" %clang_args +// CHECK: #include + +__global__ void axpy_kernel(float a, float* x, float* y) { + y[threadIdx.x] = a * x[threadIdx.x]; +} + +void axpy(float a, float* x, float* y) { +float* y_new = nullptr; +#ifdef SOME_MACRO + y_new = x; + // CHECK: axpy_kernel <<<1, 1>>> (a, y_new, x); + axpy_kernel <<<1, 1>>> (a, y_new, x); +#endif + +#ifndef SOME_MACRO + y_new = y; + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(2), 0, 0, a, y_new, x); + axpy_kernel <<<1, 2>>> (a, y_new, x); +#endif + +#ifdef SOME_MACRO + // CHECK: axpy_kernel <<<1, 3>>> (a, y, x); + axpy_kernel <<<1, 3>>> (a, y, x); +#else + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(4), 0, 0, a, x, y); + axpy_kernel <<<1, 4>>> (a, x, y); +#endif + +#ifdef SOME_MACRO + // CHECK: axpy_kernel <<<1, 5>>> (a, y, x); + axpy_kernel <<<1, 5>>> (a, y, x); +#elif defined SOME_MACRO_1 + // CHECK: axpy_kernel <<<1, 6>>> (a, x, y); + axpy_kernel <<<1, 6>>> (a, x, y); +#else + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(7), 0, 0, a, x, y); + axpy_kernel <<<1, 7>>> (a, x, y); +#endif + +#ifndef SOME_MACRO + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(8), 0, 0, a, y, x); + axpy_kernel <<<1, 8>>> (a, y, x); +#elif !defined(SOME_MACRO_1) + // CHECK: axpy_kernel <<<1, 9>>> (a, x, y); + axpy_kernel <<<1, 9>>> (a, x, y); +#else + // CHECK: axpy_kernel <<<1, 10>>> (a, x, y); + axpy_kernel <<<1, 10>>> (a, x, y); +#endif + +} \ No newline at end of file diff --git a/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_01_LLVM_10.cu b/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_01_LLVM_10.cu new file mode 100644 index 0000000000..a5c7c41745 --- /dev/null +++ b/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_01_LLVM_10.cu @@ -0,0 +1,52 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// CHECK: #include + +__global__ void axpy_kernel(float a, float* x, float* y) { + y[threadIdx.x] = a * x[threadIdx.x]; +} + +void axpy(float a, float* x, float* y) { +float* y_new = nullptr; +#ifdef SOME_MACRO + y_new = x; + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(1), 0, 0, a, y_new, x); + axpy_kernel <<<1, 1>>> (a, y_new, x); +#endif + +#ifndef SOME_MACRO + y_new = y; + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(2), 0, 0, a, y_new, x); + axpy_kernel <<<1, 2>>> (a, y_new, x); +#endif + +#ifdef SOME_MACRO + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(3), 0, 0, a, y, x); + axpy_kernel <<<1, 3>>> (a, y, x); +#else + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(4), 0, 0, a, x, y); + axpy_kernel <<<1, 4>>> (a, x, y); +#endif + +#ifdef SOME_MACRO + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(5), 0, 0, a, y, x); + axpy_kernel <<<1, 5>>> (a, y, x); +#elif defined SOME_MACRO_1 + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(6), 0, 0, a, x, y); + axpy_kernel <<<1, 6>>> (a, x, y); +#else + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(7), 0, 0, a, x, y); + axpy_kernel <<<1, 7>>> (a, x, y); +#endif + +#ifndef SOME_MACRO + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(8), 0, 0, a, y, x); + axpy_kernel <<<1, 8>>> (a, y, x); +#elif !defined(SOME_MACRO_1) + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(9), 0, 0, a, x, y); + axpy_kernel <<<1, 9>>> (a, x, y); +#else + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(10), 0, 0, a, x, y); + axpy_kernel <<<1, 10>>> (a, x, y); +#endif + +} \ No newline at end of file diff --git a/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_LLVM_10.cu b/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_LLVM_10.cu new file mode 100644 index 0000000000..06ce48ebef --- /dev/null +++ b/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_LLVM_10.cu @@ -0,0 +1,30 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// CHECK: #include + +#include + +__global__ void axpy_kernel(float a, float* x, float* y) { + y[threadIdx.x] = a * x[threadIdx.x]; +} + +void axpy(float a, float* x, float* y) { + +#ifdef SOME_MACRO + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(1), 0, 0, a, y, x); + axpy_kernel <<<1, 1>>> (a, y, x); +#endif + +#ifndef SOME_MACRO + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(2), 0, 0, a, y, x); + axpy_kernel <<<1, 2>>> (a, y, x); +#endif + +#ifdef SOME_MACRO + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(3), 0, 0, a, y, x); + axpy_kernel <<<1, 3>>> (a, y, x); +#else + // CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(4), 0, 0, a, x, y); + axpy_kernel <<<1, 4>>> (a, x, y); +#endif + +} \ No newline at end of file diff --git a/tests/src/gcc/LaunchKernel.c b/tests/src/gcc/LaunchKernel.c new file mode 100644 index 0000000000..4ebe59d7c5 --- /dev/null +++ b/tests/src/gcc/LaunchKernel.c @@ -0,0 +1,165 @@ + +/* Copyright (c) 2019-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 WARRANNTY OF ANY KIND, EXPRESS OR + * IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + + +/* HIT_START + * BUILD_CMD: gpu.o %hc -I%hip-path/include -g -c %S/gpu.cpp -o %T/gpu.o EXCLUDE_HIP_PLATFORM nvcc + * BUILD_CMD: launchkernel.o %cc -D__HIP_PLATFORM_HCC__ -g -I%hip-path/include -c %S/LaunchKernel.c -o %T/launchkernel.o EXCLUDE_HIP_PLATFORM nvcc + * BUILD_CMD: LaunchKernel %hc %T/launchkernel.o %T/gpu.o -g -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/%t DEPENDS gpu.o launchkernel.o EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t EXCLUDE_HIP_PLATFORM nvcc + * HIT_END + */ + + +#include "hip/hip_runtime.h" +#include +#include "LaunchKernel.h" + +#define HIPCHECK(error) \ + { \ + hipError_t localError = error; \ + if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \ + printf("%serror: '%s'(%d) from %s at %s:%d%s\n", "\x1B[31m", hipGetErrorString(localError), \ + localError, #error, __FILE__, __LINE__, "\x1B[0m"); \ + return false; \ + } \ + } + + +bool LaunchKernelArg() +{ + dim3 blocks = {1,1,1}; + dim3 threads = {1,1,1}; + + HIPCHECK(hipLaunchKernel(kernel, blocks, threads,NULL, 0, 0)); + + return true; +} + +bool LaunchKernelArg1() +{ + int A = 0; + int *A_d = NULL; + dim3 blocks = {1,1,1}; + dim3 threads = {1,1,1}; + + // Allocate Device memory + HIPCHECK(hipMalloc((void**)&A_d, sizeof(int))); + + void* Args[]={A_d}; + HIPCHECK(hipLaunchKernel(kernel1, blocks, threads, Args,0,0)); + + // Get the result back to host memory + HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost)); + + HIPCHECK(hipFree(A_d)); + + if(A != 333) + return false; + + return true; +} + +bool LaunchKernelArg2() +{ + int A = 0; + int B = 123; + int *A_d = NULL; + int *B_d = NULL; + + dim3 blocks = {1,1,1}; + dim3 threads = {1,1,1}; + + // Allocate Device memory + HIPCHECK(hipMalloc((void**)&A_d, sizeof(int))); + + HIPCHECK(hipMalloc((void**)&B_d, sizeof(int))); + + // Copy data from host memory to device memory + HIPCHECK(hipMemcpy(B_d,&B, sizeof(int), hipMemcpyHostToDevice)); + + void* Args[]={A_d,B_d}; + HIPCHECK(hipLaunchKernel(kernel2, blocks, threads, Args,0,0)); + + // Get the result back to host memory + HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost)); + + HIPCHECK(hipFree(A_d)); + HIPCHECK(hipFree(B_d)); + + if(A != 123) + return false; + + return true; +} + +bool LaunchKernelArg3() +{ + int A = 321; + int B = 123; + int C = 0; + int *A_d = NULL; + int *B_d = NULL; + int *C_d = NULL; + + dim3 blocks = {1,1,1}; + dim3 threads = {1,1,1}; + + // Allocate Device memory + HIPCHECK(hipMalloc((void**)&A_d, sizeof(int))); + + HIPCHECK(hipMalloc((void**)&B_d, sizeof(int))); + + HIPCHECK(hipMalloc((void**)&C_d, sizeof(int))); + + // Copy data from host memory to device memory + HIPCHECK(hipMemcpy(A_d,&A, sizeof(int), hipMemcpyHostToDevice)); + + HIPCHECK(hipMemcpy(B_d,&B, sizeof(int), hipMemcpyHostToDevice)); + + void* Args[]={A_d,B_d,C_d}; + HIPCHECK(hipLaunchKernel(kernel3, blocks, threads, Args,0,0)); + + // Get the result back to host memory + HIPCHECK(hipMemcpy(&C, C_d, sizeof(int), hipMemcpyDeviceToHost)); + + HIPCHECK(hipFree(A_d)); + HIPCHECK(hipFree(B_d)); + HIPCHECK(hipFree(C_d)); + + if(C != 444) + return false; + + return true; +} + + +int main() +{ + + if( LaunchKernelArg() && + LaunchKernelArg1() && + LaunchKernelArg2() && + LaunchKernelArg3()) + { + printf("PASSED!\n"); + } + else + printf("FAILED\n"); +} diff --git a/tests/src/gcc/LaunchKernel.h b/tests/src/gcc/LaunchKernel.h new file mode 100644 index 0000000000..b7424d8b46 --- /dev/null +++ b/tests/src/gcc/LaunchKernel.h @@ -0,0 +1,23 @@ + +/* Copyright (c) 2019-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 WARRANNTY OF ANY KIND, EXPRESS OR + * IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +extern __global__ void kernel(); +extern __global__ void kernel1(int*); +extern __global__ void kernel2(int*,int*); +extern __global__ void kernel3(int*,int*,int*); diff --git a/tests/src/gcc/gpu.cpp b/tests/src/gcc/gpu.cpp new file mode 100644 index 0000000000..00223478fa --- /dev/null +++ b/tests/src/gcc/gpu.cpp @@ -0,0 +1,46 @@ +/* + * Copyright (c) 2019-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 WARRANNTY OF ANY KIND, EXPRESS OR + * IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * */ + + +#include + +extern "C" +{ + +__global__ void kernel() +{ + int a = 20; +} + +__global__ void kernel1(int *a) +{ + *a = 333; +} + +__global__ void kernel2(int *a, int*b) +{ + *a = *b; +} + +__global__ void kernel3(int *a, int*b, int* c) +{ + *c = *a+*b; +} + +}//extern "C"