diff --git a/projects/hip/hipify-clang/CMakeLists.txt b/projects/hip/hipify-clang/CMakeLists.txt index 1aaa68097f..1a0648f9a0 100644 --- a/projects/hip/hipify-clang/CMakeLists.txt +++ b/projects/hip/hipify-clang/CMakeLists.txt @@ -1,63 +1,70 @@ cmake_minimum_required(VERSION 2.8.8) project(hipify-clang) -if (PARENT_SCOPE) - set(BUILD_HIPIFY_CLANG 0 PARENT_SCOPE) +set(BUILD_HIPIFY_CLANG 0 CACHE INTERNAL "") + +if (HIPIFY_CLANG_LLVM_DIR) + find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED NO_DEFAULT_PATH) +else() + message(STATUS "hipify-clang will not be built. To build it please specify absolute path to LLVM 3.8 or higher using HIPIFY_CLANG_LLVM_DIR") + return() endif() -# Find LLVM package -find_package(LLVM 3.8 QUIET PATHS ${HIPIFY_CLANG_LLVM_DIR} NO_DEFAULT_PATH) -if (NOT ${LLVM_FOUND}) - find_package(LLVM 3.9 QUIET PATHS ${HIPIFY_CLANG_LLVM_DIR} NO_DEFAULT_PATH) - if (NOT ${LLVM_FOUND}) - message(STATUS "hipify-clang will not be built. To build it please specify absolute path to LLVM 3.8 or LLVM 3.9 package/dist using -DHIPIFY_CLANG_LLVM_DIR") - endif() +option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" ON) + +# Disable the tests if `lit` or `FileCheck` is not installed. +find_program(LIT_COMMAND lit) +find_program(FILECHECK_COMMAND FileCheck) +find_program(SOCAT_COMMAND socat) +if (NOT LIT_COMMAND OR NOT FILECHECK_COMMAND OR NOT SOCAT_COMMAND) + set(HIPIFY_CLANG_TESTS OFF CACHE INTERNAL "") + message(STATUS "hipify-clang's tests are not being built because `lit`,`FileCheck` or `socat` could not be found.") endif() -if (${LLVM_FOUND}) - list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR}) - include(AddLLVM) - message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}") +list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR}) +include(AddLLVM) - include_directories(${LLVM_INCLUDE_DIRS}) - link_directories(${LLVM_LIBRARY_DIRS}) - add_definitions(${LLVM_DEFINITIONS}) - add_llvm_executable(hipify-clang src/Cuda2Hip.cpp) +message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}") - set(CMAKE_CXX_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang++) - set(CMAKE_C_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang) +include_directories(${LLVM_INCLUDE_DIRS}) +link_directories(${LLVM_LIBRARY_DIRS}) +add_definitions(${LLVM_DEFINITIONS}) +add_llvm_executable(hipify-clang src/Cuda2Hip.cpp) - # Link against LLVM and CLANG libraries - target_link_libraries(hipify-clang - clangASTMatchers - clangFrontend - clangTooling - clangParse - clangSerialization - clangSema - clangEdit - clangFormat - clangLex - clangAnalysis - clangDriver - clangAST - clangToolingCore - clangRewrite - clangBasic - LLVMProfileData - LLVMSupport - LLVMMCParser - LLVMMC - LLVMBitReader - LLVMOption - LLVMCore) +set(CMAKE_CXX_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang++) +set(CMAKE_C_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang) + +# Link against LLVM and CLANG libraries +target_link_libraries(hipify-clang + clangASTMatchers + clangFrontend + clangTooling + clangParse + clangSerialization + clangSema + clangEdit + clangFormat + clangLex + clangAnalysis + clangDriver + clangAST + clangToolingCore + clangRewrite + clangBasic + LLVMProfileData + LLVMSupport + LLVMMCParser + LLVMMC + LLVMBitReader + LLVMOption + LLVMCore) if(WIN32) target_link_libraries(hipify-clang version) endif() - set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}") - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}") +set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}") if(MSVC) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /GR- /EHs- /EHc-") set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} /SUBSYSTEM:WINDOWS") @@ -65,42 +72,33 @@ else() set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden") endif() - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DIRS}/clang/${LLVM_VERSION_MAJOR}.${LLVM_VERSION_MINOR}.${LLVM_VERSION_PATCH}\\\"") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DIRS}/clang/${LLVM_VERSION_MAJOR}.${LLVM_VERSION_MINOR}.${LLVM_VERSION_PATCH}\\\"") - install(TARGETS hipify-clang DESTINATION bin) +install(TARGETS hipify-clang DESTINATION bin) if (HIPIFY_CLANG_TESTS) - # tests - set(Python_ADDITIONAL_VERSIONS 2.7) - include(FindPythonInterp) - if(NOT PYTHONINTERP_FOUND) - message(FATAL_ERROR - "Unable to find Python interpreter, required for builds and testing.\n\n" - "Please install Python or specify the PYTHON_EXECUTABLE CMake variable.") - endif() + find_package(PythonInterp 2.7 REQUIRED EXACT) - find_program(LIT_COMMAND lit) - - set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) + # Populates CUDA_TOOLKIT_ROOT_DIR, which is then applied to the lit config to give the + # value of --cuda-path for the test runs. + find_package(CUDA REQUIRED) configure_file( - ${CMAKE_SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in + ${CMAKE_CURRENT_LIST_DIR}/../tests/hipify-clang/lit.site.cfg.in ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg - @ONLY) + @ONLY + ) add_lit_testsuite(test-hipify "Running HIPify regression tests" - ${CMAKE_SOURCE_DIR}/tests/hipify-clang + ${CMAKE_CURRENT_LIST_DIR}/../tests/hipify-clang PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg - DEPENDS hipify-clang lit - ) + ARGS -v + DEPENDS hipify-clang + ) add_custom_target(test-hipify-clang) add_dependencies(test-hipify-clang test-hipify) set_target_properties(test-hipify-clang PROPERTIES FOLDER "Tests") endif() -if (PARENT_SCOPE) - set(BUILD_HIPIFY_CLANG 1 PARENT_SCOPE) -endif() - -endif() +set(BUILD_HIPIFY_CLANG 1 CACHE INTERNAL "") diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp index 907e78e54f..1f0bcdba65 100644 --- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp @@ -2928,7 +2928,13 @@ protected: std::string mainFileName; virtual void insertReplacement(const Replacement &rep, const FullSourceLoc &fullSL) { +#if LLVM_VERSION_MAJOR > 3 + // New clang added error checking to Replacements, and *insists* that you explicitly check it. + llvm::Error e = Replace->add(rep); +#else + // In older versions, it's literally an std::set Replace->insert(rep); +#endif if (PrintStats) { LOCs.insert(fullSL.getExpansionLineNumber()); } @@ -3022,14 +3028,17 @@ class Cuda2HipCallback; class HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks, public Cuda2Hip { public: HipifyPPCallbacks(Replacements *R, const std::string &mainFileName) - : Cuda2Hip(R, mainFileName), SeenEnd(false), _sm(nullptr), _pp(nullptr) {} + : Cuda2Hip(R, mainFileName) {} - virtual bool handleBeginSource(CompilerInstance &CI, StringRef Filename) override { + virtual bool handleBeginSource(CompilerInstance &CI +#if LLVM_VERSION_MAJOR <= 4 + , StringRef Filename +#endif + ) override { Preprocessor &PP = CI.getPreprocessor(); SourceManager &SM = CI.getSourceManager(); setSourceManager(&SM); PP.addPPCallbacks(std::unique_ptr(this)); - PP.Retain(); setPreprocessor(&PP); return true; } @@ -3103,16 +3112,23 @@ public: const MacroDefinition &MD, SourceRange Range, const MacroArgs *Args) override { if (_sm->isWrittenInMainFile(MacroNameTok.getLocation())) { - for (unsigned int i = 0; Args && i < MD.getMacroInfo()->getNumArgs(); i++) { + // The getNumArgs function was rather unhelpfully renamed in clang 4.0. Its semantics + // remain unchanged. +#if LLVM_VERSION_MAJOR > 4 + #define GET_NUM_ARGS() getNumParams() +#else + #define GET_NUM_ARGS() getNumArgs() +#endif + for (unsigned int i = 0; Args && i < MD.getMacroInfo()->GET_NUM_ARGS(); i++) { std::vector toks; // Code below is a kind of stolen from 'MacroArgs::getPreExpArgument' // to workaround the 'const' MacroArgs passed into this hook. const Token *start = Args->getUnexpArgument(i); size_t len = Args->getArgLength(start) + 1; -#if (LLVM_VERSION_MAJOR >= 3) && (LLVM_VERSION_MINOR >= 9) - _pp->EnterTokenStream(ArrayRef(start, len), false); -#else +#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8) _pp->EnterTokenStream(start, len, false, false); +#else + _pp->EnterTokenStream(ArrayRef(start, len), false); #endif do { toks.push_back(Token()); @@ -3185,15 +3201,15 @@ public: void EndOfMainFile() override {} - bool SeenEnd; + bool SeenEnd = false; void setSourceManager(SourceManager *sm) { _sm = sm; } void setPreprocessor(Preprocessor *pp) { _pp = pp; } void setMatch(Cuda2HipCallback *match) { Match = match; } private: - SourceManager *_sm; - Preprocessor *_pp; - Cuda2HipCallback *Match; + SourceManager *_sm = nullptr; + Preprocessor *_pp = nullptr; + Cuda2HipCallback *Match = nullptr; }; class Cuda2HipCallback : public MatchFinder::MatchCallback, public Cuda2Hip { @@ -3799,7 +3815,7 @@ private: } bool stringLiteral(const MatchFinder::MatchResult &Result) { - if (const StringLiteral *sLiteral = Result.Nodes.getNodeAs("stringLiteral")) { + if (const clang::StringLiteral *sLiteral = Result.Nodes.getNodeAs("stringLiteral")) { if (sLiteral->getCharByteWidth() == 1) { StringRef s = sLiteral->getString(); SourceManager *SM = Result.SourceManager; @@ -4207,11 +4223,15 @@ void copyFile(const std::string& src, const std::string& dst) { int main(int argc, const char **argv) { auto start = std::chrono::steady_clock::now(); auto begin = start; -#if (LLVM_VERSION_MAJOR >= 3) && (LLVM_VERSION_MINOR >= 9) - llvm::sys::PrintStackTraceOnErrorSignal(StringRef()); -#else + + // The signature of PrintStackTraceOnErrorSignal changed in llvm 3.9. We don't support + // anything older than 3.8, so let's specifically detect the one old version we support. +#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8) llvm::sys::PrintStackTraceOnErrorSignal(); +#else + llvm::sys::PrintStackTraceOnErrorSignal(StringRef()); #endif + CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::OneOrMore); std::vector fileSources = OptionsParser.getSourcePathList(); std::string dst = OutputFilename; @@ -4271,12 +4291,23 @@ int main(int argc, const char **argv) { // So what we do is operate on a copy, which we then move to the output. RefactoringTool Tool(OptionsParser.getCompilations(), tmpFile); ast_matchers::MatchFinder Finder; - HipifyPPCallbacks PPCallbacks(&Tool.getReplacements(), tmpFile); - Cuda2HipCallback Callback(&Tool.getReplacements(), &Finder, &PPCallbacks, tmpFile); + + // The Replacements to apply to the file `src`. + Replacements* replacementsToUse; +#if LLVM_VERSION_MAJOR > 3 + // getReplacements() now returns a map from filename to Replacements - so create an entry + // for this source file and return a pointer to it. + replacementsToUse = &(Tool.getReplacements()[tmpFile]); +#else + replacementsToUse = &Tool.getReplacements(); +#endif + + HipifyPPCallbacks* PPCallbacks = new HipifyPPCallbacks(replacementsToUse, tmpFile); + Cuda2HipCallback Callback(replacementsToUse, &Finder, PPCallbacks, tmpFile); addAllMatchers(Finder, &Callback); - auto action = newFrontendActionFactory(&Finder, &PPCallbacks); + auto action = newFrontendActionFactory(&Finder, PPCallbacks); Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("--cuda-host-only", ArgumentInsertPosition::BEGIN)); @@ -4300,9 +4331,14 @@ int main(int argc, const char **argv) { SourceManager SM(Diagnostics, Tool.getFiles()); if (PrintStats) { DEBUG(dbgs() << "Replacements collected by the tool:\n"); - for (const auto &r : Tool.getReplacements()) { - DEBUG(dbgs() << r.toString() << "\n"); - repBytes += r.getLength(); +#if LLVM_VERSION_MAJOR > 3 + Replacements& replacements = Tool.getReplacements().begin()->second; +#else + Replacements& replacements = Tool.getReplacements(); +#endif + for (const auto &replacement : replacements) { + DEBUG(dbgs() << replacement.toString() << "\n"); + repBytes += replacement.getLength(); } std::ifstream src_file(dst, std::ios::binary | std::ios::ate); src_file.clear(); @@ -4329,13 +4365,13 @@ int main(int argc, const char **argv) { } std::remove(csv.c_str()); } - if (0 == printStats(csv, src, PPCallbacks, Callback, repBytes, bytes, lines, start)) { + if (0 == printStats(csv, src, *PPCallbacks, Callback, repBytes, bytes, lines, start)) { filesTranslated--; } start = std::chrono::steady_clock::now(); repBytesTotal += repBytes; bytesTotal += bytes; - changedLinesTotal += PPCallbacks.LOCs.size() + Callback.LOCs.size(); + changedLinesTotal += PPCallbacks->LOCs.size() + Callback.LOCs.size(); linesTotal += lines; } dst.clear(); diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 2c98ac804f..2d67c31fe7 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -614,6 +614,11 @@ void ihipDevice_t::locked_reset() _state = 0; am_memtracker_reset(_acc); + //FIXME - Calling am_memtracker_reset is really bad since it destroyed all buffers allocated by the HCC runtime as well + //such as the printf buffer. Re-initialze the printf buffer as a workaround for now. +#if (__hcc_workweek__ >= 17423) + Kalmar::getContext()->initPrintfBuffer(); +#endif }; #define ErrorCheck(x) error_check(x, __LINE__, __FILE__) diff --git a/projects/hip/tests/hipify-clang/axpy.cu b/projects/hip/tests/hipify-clang/axpy.cu index 8472e60209..8c6b0e0d8d 100644 --- a/projects/hip/tests/hipify-clang/axpy.cu +++ b/projects/hip/tests/hipify-clang/axpy.cu @@ -1,9 +1,9 @@ -// RUN: hipify "%s" -o=%t -- +// RUN: %run_test hipify "%s" "%t" %cuda_args #include __global__ void axpy(float a, float* x, float* y) { - // RUN: sh -c "test `grep -c -F 'y[hipThreadIdx_x] = a * x[hipThreadIdx_x];' %t` -eq 2" + // CHECK: y[hipThreadIdx_x] = a * x[hipThreadIdx_x]; y[threadIdx.x] = a * x[threadIdx.x]; } @@ -17,21 +17,25 @@ int main(int argc, char* argv[]) { // Copy input data to device. float* device_x; float* device_y; - // RUN: sh -c "test `grep -c -F 'hipMalloc(&device_x, kDataLen * sizeof(float));' %t` -eq 2" + + // CHECK: hipMalloc(&device_x, kDataLen * sizeof(float)); cudaMalloc(&device_x, kDataLen * sizeof(float)); - // RUN: sh -c "test `grep -c -F 'hipMalloc(&device_y, kDataLen * sizeof(float));' %t` -eq 2" + + // CHECK: hipMalloc(&device_y, kDataLen * sizeof(float)); cudaMalloc(&device_y, kDataLen * sizeof(float)); - // RUN: sh -c "test `grep -c -F 'hipMemcpy(device_x, host_x, kDataLen * sizeof(float), hipMemcpyHostToDevice);' %t` -eq 2" + + // CHECK: hipMemcpy(device_x, host_x, kDataLen * sizeof(float), hipMemcpyHostToDevice); cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), cudaMemcpyHostToDevice); // Launch the kernel. - // RUN: sh -c "test `grep -c -F 'hipLaunchKernel(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);' %t` -eq 2" + // CHECK: hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); axpy<<<1, kDataLen>>>(a, device_x, device_y); // Copy output data to host. - // RUN: sh -c "test `grep -c -F 'hipDeviceSynchronize();' %t` -eq 2" + // CHECK: hipDeviceSynchronize(); cudaDeviceSynchronize(); - // RUN: sh -c "test `grep -c -F 'hipMemcpy(host_y, device_y, kDataLen * sizeof(float), hipMemcpyDeviceToHost);' %t` -eq 2" + + // CHECK: hipMemcpy(host_y, device_y, kDataLen * sizeof(float), hipMemcpyDeviceToHost); cudaMemcpy(host_y, device_y, kDataLen * sizeof(float), cudaMemcpyDeviceToHost); // Print the results. @@ -39,7 +43,7 @@ int main(int argc, char* argv[]) { std::cout << "y[" << i << "] = " << host_y[i] << "\n"; } - // RUN: sh -c "test `grep -c -F 'hipDeviceReset();' %t` -eq 2" + // CHECK: hipDeviceReset(); cudaDeviceReset(); return 0; } diff --git a/projects/hip/tests/hipify-clang/concurentKernels.cu b/projects/hip/tests/hipify-clang/concurentKernels.cu new file mode 100644 index 0000000000..e369baaf3e --- /dev/null +++ b/projects/hip/tests/hipify-clang/concurentKernels.cu @@ -0,0 +1,240 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args +/* + * Copyright 1993-2015 NVIDIA Corporation. All rights reserved. + * + * Please refer to the NVIDIA end user license agreement (EULA) associated + * with this source code for terms and conditions that govern your use of + * this software. Any use, reproduction, disclosure, or distribution of + * this software and related documentation outside the terms of the EULA + * is strictly prohibited. + * + */ + +// +// This sample demonstrates the use of streams for concurrent execution. It also illustrates how to +// introduce dependencies between CUDA streams with the new cudaStreamWaitEvent function introduced +// in CUDA 3.2. +// +// Devices of compute capability 1.x will run the kernels one after another +// Devices of compute capability 2.0 or higher can overlap the kernels +// +#include +#include +#include + +// This is a kernel that does no real work but runs at least for a specified number of clocks +__global__ void clock_block(clock_t *d_o, clock_t clock_count) +{ + unsigned int start_clock = (unsigned int) clock(); + + clock_t clock_offset = 0; + + while (clock_offset < clock_count) + { + unsigned int end_clock = (unsigned int) clock(); + + // The code below should work like + // this (thanks to modular arithmetics): + // + // clock_offset = (clock_t) (end_clock > start_clock ? + // end_clock - start_clock : + // end_clock + (0xffffffffu - start_clock)); + // + // Indeed, let m = 2^32 then + // end - start = end + m - start (mod m). + + clock_offset = (clock_t)(end_clock - start_clock); + } + + d_o[0] = clock_offset; +} + + +// Single warp reduction kernel +__global__ void sum(clock_t *d_clocks, int N) +{ + __shared__ clock_t s_clocks[32]; + + clock_t my_sum = 0; + + for (int i = threadIdx.x; i < N; i+= blockDim.x) + { + my_sum += d_clocks[i]; + } + + s_clocks[threadIdx.x] = my_sum; + syncthreads(); + + for (int i=16; i>0; i/=2) + { + if (threadIdx.x < i) + { + s_clocks[threadIdx.x] += s_clocks[threadIdx.x + i]; + } + + syncthreads(); + } + + d_clocks[0] = s_clocks[0]; +} + +int main(int argc, char **argv) +{ + int nkernels = 8; // number of concurrent kernels + int nstreams = nkernels + 1; // use one more stream than concurrent kernel + int nbytes = nkernels * sizeof(clock_t); // number of data bytes + float kernel_time = 10; // time the kernel should run in ms + float elapsed_time; // timing variables + int cuda_device = 0; + + printf("[%s] - Starting...\n", argv[0]); + + // get number of kernels if overridden on the command line + if (checkCmdLineFlag(argc, (const char **)argv, "nkernels")) + { + nkernels = getCmdLineArgumentInt(argc, (const char **)argv, "nkernels"); + nstreams = nkernels + 1; + } + + // use command-line specified CUDA device, otherwise use device with highest Gflops/s + cuda_device = findCudaDevice(argc, (const char **)argv); + + cudaDeviceProp deviceProp; + // CHECK: checkCudaErrors(hipGetDevice(&cuda_device)); + checkCudaErrors(cudaGetDevice(&cuda_device)); + + // CHECK: checkCudaErrors(hipGetDeviceProperties(&deviceProp, cuda_device)); + checkCudaErrors(cudaGetDeviceProperties(&deviceProp, cuda_device)); + + if ((deviceProp.concurrentKernels == 0)) + { + printf("> GPU does not support concurrent kernel execution\n"); + printf(" CUDA kernel runs will be serialized\n"); + } + + printf("> Detected Compute SM %d.%d hardware with %d multi-processors\n", + deviceProp.major, deviceProp.minor, deviceProp.multiProcessorCount); + + // allocate host memory + clock_t *a = 0; // pointer to the array data in host memory + // CHECK: checkCudaErrors(hipHostMalloc((void **)&a, nbytes)); + checkCudaErrors(cudaMallocHost((void **)&a, nbytes)); + + // allocate device memory + clock_t *d_a = 0; // pointers to data and init value in the device memory + // CHECK: checkCudaErrors(hipMalloc((void **)&d_a, nbytes)); + checkCudaErrors(cudaMalloc((void **)&d_a, nbytes)); + + // CHECK: hipStream_t *streams = (hipStream_t *) malloc(nstreams * sizeof(hipStream_t)); + // allocate and initialize an array of stream handles + cudaStream_t *streams = (cudaStream_t *) malloc(nstreams * sizeof(cudaStream_t)); + + for (int i = 0; i < nstreams; i++) + { + // CHECK: checkCudaErrors(hipStreamCreate(&(streams[i]))); + checkCudaErrors(cudaStreamCreate(&(streams[i]))); + } + + // create CUDA event handles + cudaEvent_t start_event, stop_event; + + // CHECK: checkCudaErrors(hipEventCreate(&start_event)); + // CHECK: checkCudaErrors(hipEventCreate(&stop_event)); + checkCudaErrors(cudaEventCreate(&start_event)); + checkCudaErrors(cudaEventCreate(&stop_event)); + + // the events are used for synchronization only and hence do not need to record timings + // this also makes events not introduce global sync points when recorded which is critical to get overlap + + // CHECK: hipEvent_t *kernelEvent; + // CHECK: kernelEvent = (hipEvent_t *) malloc(nkernels * sizeof(hipEvent_t)); + cudaEvent_t *kernelEvent; + kernelEvent = (cudaEvent_t *) malloc(nkernels * sizeof(cudaEvent_t)); + + for (int i = 0; i < nkernels; i++) + { + // CHECK: checkCudaErrors(hipEventCreateWithFlags(&(kernelEvent[i]), hipEventDisableTiming)); + checkCudaErrors(cudaEventCreateWithFlags(&(kernelEvent[i]), cudaEventDisableTiming)); + } + + ////////////////////////////////////////////////////////////////////// + // time execution with nkernels streams + clock_t total_clocks = 0; +#if defined(__arm__) || defined(__aarch64__) + // the kernel takes more time than the channel reset time on arm archs, so to prevent hangs reduce time_clocks. + clock_t time_clocks = (clock_t)(kernel_time * (deviceProp.clockRate / 1000)); +#else + clock_t time_clocks = (clock_t)(kernel_time * deviceProp.clockRate); +#endif + + // CHECK: hipEventRecord(start_event, 0); + cudaEventRecord(start_event, 0); + + // queue nkernels in separate streams and record when they are done + for (int i=0; i>>(&d_a[i], time_clocks); + total_clocks += time_clocks; + + // CHECK: checkCudaErrors(hipEventRecord(kernelEvent[i], streams[i])); + checkCudaErrors(cudaEventRecord(kernelEvent[i], streams[i])); + + // make the last stream wait for the kernel event to be recorded + // CHECK: checkCudaErrors(hipStreamWaitEvent(streams[nstreams-1], kernelEvent[i],0)); + checkCudaErrors(cudaStreamWaitEvent(streams[nstreams-1], kernelEvent[i],0)); + } + + // queue a sum kernel and a copy back to host in the last stream. + // the commands in this stream get dispatched as soon as all the kernel events have been recorded + // CHECK: hipLaunchKernelGGL(sum, dim3(1), dim3(32), 0, streams[nstreams-1], d_a, nkernels); + // CHECK: checkCudaErrors(hipMemcpyAsync(a, d_a, sizeof(clock_t), hipMemcpyDeviceToHost, streams[nstreams-1])); + sum<<<1,32,0,streams[nstreams-1]>>>(d_a, nkernels); + checkCudaErrors(cudaMemcpyAsync(a, d_a, sizeof(clock_t), cudaMemcpyDeviceToHost, streams[nstreams-1])); + + // at this point the CPU has dispatched all work for the GPU and can continue processing other tasks in parallel + + // in this sample we just wait until the GPU is done + // CHECK: checkCudaErrors(hipEventRecord(stop_event, 0)); + // CHECK: checkCudaErrors(hipEventSynchronize(stop_event)); + // CHECK: checkCudaErrors(hipEventElapsedTime(&elapsed_time, start_event, stop_event)); + checkCudaErrors(cudaEventRecord(stop_event, 0)); + checkCudaErrors(cudaEventSynchronize(stop_event)); + checkCudaErrors(cudaEventElapsedTime(&elapsed_time, start_event, stop_event)); + + printf("Expected time for serial execution of %d kernels = %.3fs\n", nkernels, nkernels * kernel_time/1000.0f); + printf("Expected time for concurrent execution of %d kernels = %.3fs\n", nkernels, kernel_time/1000.0f); + printf("Measured time for sample = %.3fs\n", elapsed_time/1000.0f); + + bool bTestResult = (a[0] > total_clocks); + + // release resources + for (int i = 0; i < nkernels; i++) + { + // CHECK: hipStreamDestroy(streams[i]); + // CHECK: hipEventDestroy(kernelEvent[i]); + cudaStreamDestroy(streams[i]); + cudaEventDestroy(kernelEvent[i]); + } + + free(streams); + free(kernelEvent); + + // CHECK: hipEventDestroy(start_event); + // CHECK: hipEventDestroy(stop_event); + // CHECK: hipHostFree(a); + // CHECK: hipFree(d_a); + cudaEventDestroy(start_event); + cudaEventDestroy(stop_event); + cudaFreeHost(a); + cudaFree(d_a); + + if (!bTestResult) + { + printf("Test failed!\n"); + exit(EXIT_FAILURE); + } + + printf("Test passed\n"); + exit(EXIT_SUCCESS); +} diff --git a/projects/hip/tests/hipify-clang/cudaRegister.cu b/projects/hip/tests/hipify-clang/cudaRegister.cu new file mode 100644 index 0000000000..80d17f65b9 --- /dev/null +++ b/projects/hip/tests/hipify-clang/cudaRegister.cu @@ -0,0 +1,111 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include +#include +#include +#include + +#define LEN 1024 +#define SIZE LEN * sizeof(float) +#define ITER 1024*1024 + +// CHECK: if(status != hipSuccess) { +#define check(msg, status){ \ +if(status != cudaSuccess) { \ + printf("%s failed. \n", #msg); \ +} \ +} + +__global__ void Inc1(float *Ad, float *Bd){ + // CHECK: int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + int tx = threadIdx.x + blockIdx.x * blockDim.x; + if(tx < 1 ){ + for(int i=0;i>>(Ad, Bd); + sleep(3); + A[0] = -(ITER*1.0f); + std::cout<<"Same cache line before completion: \t"<< A[0]<>>(Ad, Bd); + sleep(3); + A[0] = -(ITER*1.0f); + std::cout<<"Diff cache line before completion: \t"< $TMP_SCRIPT +set -o errexit +set -o xtrace +rm $TMP_SCRIPT +$HIPIFY -o=$TMP_FILE $IN_FILE -- $@ && cat $TMP_FILE | sed -Ee 's|//.+|// |g' | FileCheck $IN_FILE +EOF +chmod a+x $TMP_SCRIPT + +# Run the script via socat, spawning a virtual terminal and propagating exit code, and hence failure. +socat -du EXEC:$TMP_SCRIPT,pty,stderr STDOUT diff --git a/projects/hip/tests/hipify-clang/square.cu b/projects/hip/tests/hipify-clang/square.cu new file mode 100644 index 0000000000..e0c72094a8 --- /dev/null +++ b/projects/hip/tests/hipify-clang/square.cu @@ -0,0 +1,114 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include + +#define CHECK(cmd) \ +{\ + cudaError_t error = cmd;\ + if (error != cudaSuccess) { \ + fprintf(stderr, "error: '%s'(%d) at %s:%d\n", cudaGetErrorString(error), error,__FILE__, __LINE__); \ + exit(EXIT_FAILURE);\ + }\ +} + + +/* + * Square each element in the array A and write to array C. + */ +template +__global__ void +vector_square(T *C_d, const T *A_d, size_t N) +{ + // CHECK: size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); + // CHECK: size_t stride = hipBlockDim_x * hipGridDim_x; + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (size_t i=offset; i>> (C_d, A_d, N); + + printf ("info: copy Device2Host\n"); + // CHECK: CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + CHECK ( cudaMemcpy(C_h, C_d, Nbytes, cudaMemcpyDeviceToHost)); + + printf ("info: check result\n"); + for (size_t i=0; i