From 65ba03dd4cf23f268104b4cdcaa535ffed1337c7 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 28 Aug 2019 21:17:35 +0300 Subject: [PATCH] [HIPIFY][#207][fix] Translate all preprocessor's conditional blocks + Start to translate preprocessor's false conditional blocks too: based on clang's https://reviews.llvm.org/D66597; available only starting from LLVM 10.0 or trunk. + Option -skip-excluded-preprocessor-conditional-blocks for skipping excluded conditional blocks: the default behavior for hipify-clang built with LLVM < 10.0; false by default for hipify-clang built with LLVM 10 or trunk. + Add 4 preprocessor unit tests, 2 of which are LLVM 10.0 only + Update couple of existing tests by setting -skip-excluded-preprocessor-conditional-blocks option: update lit testing accordingly [ROCm/clr commit: 32f22a0117348c3c60b626e0e1aa640276b71b4a] --- .../clr/hipamd/hipify-clang/src/ArgParse.cpp | 5 ++ .../clr/hipamd/hipify-clang/src/ArgParse.h | 1 + .../hipamd/hipify-clang/src/HipifyAction.cpp | 5 ++ .../hipamd/hipify-clang/src/HipifyAction.h | 2 + .../hipamd/hipify-clang/src/LLVMCompat.cpp | 21 ++++++++ .../clr/hipamd/hipify-clang/src/LLVMCompat.h | 6 +++ projects/clr/hipamd/hipify-clang/src/main.cpp | 22 ++++---- .../clr/hipamd/tests/hipify-clang/lit.cfg | 4 ++ .../hipamd/tests/hipify-clang/lit.site.cfg.in | 1 + .../libraries/cuComplex/cuComplex_Julia.cu | 2 +- .../libraries/cuSPARSE/cuSPARSE_03.cu | 2 +- .../unit_tests/pp/pp_if_else_conditionals.cu | 30 +++++++++++ .../pp/pp_if_else_conditionals_01.cu | 52 +++++++++++++++++++ .../pp/pp_if_else_conditionals_01_LLVM_10.cu | 52 +++++++++++++++++++ .../pp/pp_if_else_conditionals_LLVM_10.cu | 30 +++++++++++ 15 files changed, 223 insertions(+), 12 deletions(-) create mode 100644 projects/clr/hipamd/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals.cu create mode 100644 projects/clr/hipamd/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_01.cu create mode 100644 projects/clr/hipamd/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_01_LLVM_10.cu create mode 100644 projects/clr/hipamd/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_LLVM_10.cu diff --git a/projects/clr/hipamd/hipify-clang/src/ArgParse.cpp b/projects/clr/hipamd/hipify-clang/src/ArgParse.cpp index 8a5a5b6e6c..751de0d028 100644 --- a/projects/clr/hipamd/hipify-clang/src/ArgParse.cpp +++ b/projects/clr/hipamd/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/projects/clr/hipamd/hipify-clang/src/ArgParse.h b/projects/clr/hipamd/hipify-clang/src/ArgParse.h index 5b8c763647..64fb2e9678 100644 --- a/projects/clr/hipamd/hipify-clang/src/ArgParse.h +++ b/projects/clr/hipamd/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/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp index 52ba2eeaaf..e223ab7f01 100644 --- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp +++ b/projects/clr/hipamd/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/projects/clr/hipamd/hipify-clang/src/HipifyAction.h b/projects/clr/hipamd/hipify-clang/src/HipifyAction.h index adafba2df9..31ccc0b648 100644 --- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.h +++ b/projects/clr/hipamd/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/projects/clr/hipamd/hipify-clang/src/LLVMCompat.cpp b/projects/clr/hipamd/hipify-clang/src/LLVMCompat.cpp index d2573ecf22..8bb3eeda25 100644 --- a/projects/clr/hipamd/hipify-clang/src/LLVMCompat.cpp +++ b/projects/clr/hipamd/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/projects/clr/hipamd/hipify-clang/src/LLVMCompat.h b/projects/clr/hipamd/hipify-clang/src/LLVMCompat.h index 069aeb91a3..c91e2815e2 100644 --- a/projects/clr/hipamd/hipify-clang/src/LLVMCompat.h +++ b/projects/clr/hipamd/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/projects/clr/hipamd/hipify-clang/src/main.cpp b/projects/clr/hipamd/hipify-clang/src/main.cpp index 3b8f454fd7..d6a53b170d 100644 --- a/projects/clr/hipamd/hipify-clang/src/main.cpp +++ b/projects/clr/hipamd/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/projects/clr/hipamd/tests/hipify-clang/lit.cfg b/projects/clr/hipamd/tests/hipify-clang/lit.cfg index 594d5fcf69..64f82e57fa 100644 --- a/projects/clr/hipamd/tests/hipify-clang/lit.cfg +++ b/projects/clr/hipamd/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/projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in b/projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in index 71cc3e08d3..6ef4dc007a 100644 --- a/projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in +++ b/projects/clr/hipamd/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/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuComplex/cuComplex_Julia.cu b/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuComplex/cuComplex_Julia.cu index 196a94197b..5c2cd3b1e0 100644 --- a/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuComplex/cuComplex_Julia.cu +++ b/projects/clr/hipamd/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/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu b/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu index 349d6471db..8618be4748 100644 --- a/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu +++ b/projects/clr/hipamd/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/projects/clr/hipamd/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals.cu b/projects/clr/hipamd/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals.cu new file mode 100644 index 0000000000..51bfeb6017 --- /dev/null +++ b/projects/clr/hipamd/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/projects/clr/hipamd/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_01.cu b/projects/clr/hipamd/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_01.cu new file mode 100644 index 0000000000..310d896054 --- /dev/null +++ b/projects/clr/hipamd/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/projects/clr/hipamd/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_01_LLVM_10.cu b/projects/clr/hipamd/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_01_LLVM_10.cu new file mode 100644 index 0000000000..a5c7c41745 --- /dev/null +++ b/projects/clr/hipamd/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/projects/clr/hipamd/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_LLVM_10.cu b/projects/clr/hipamd/tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_LLVM_10.cu new file mode 100644 index 0000000000..06ce48ebef --- /dev/null +++ b/projects/clr/hipamd/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