diff --git a/.jenkins/Jenkinsfile b/.jenkins/Jenkinsfile index 3cd6a96649..d64775c026 100644 --- a/.jenkins/Jenkinsfile +++ b/.jenkins/Jenkinsfile @@ -101,7 +101,7 @@ def hipBuildTest(String backendLabel) { # Check if backend label contains string "amd" or backend host is a server with amd gpu if [[ $backendLabel =~ amd ]]; then export HT_CONFIG_FILE="$HIP_DIR/tests/catch/hipTestMain/config/config_amd_linux.json" - LLVM_PATH=/opt/rocm/llvm ctest -E 'Unit_hipGraphChildGraphNodeGetGraph_Functional|Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Negative|Unit_hipPtrGetAttribute_Simple' + LLVM_PATH=/opt/rocm/llvm ctest -E 'Unit_hipGraphChildGraphNodeGetGraph_Functional|Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Negative|Unit_hipPtrGetAttribute_Simple|Unit_hipStreamPerThread_DeviceReset_2' sleep 120 else make test diff --git a/bin/hipcc b/bin/hipcc index 32494dd345..86622c7a36 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -22,7 +22,6 @@ # Need perl > 5.10 to use logic-defined or use 5.006; use v5.10.1; -use strict; use warnings; use File::Basename; @@ -30,6 +29,13 @@ use File::Spec::Functions 'catfile'; # TODO: By default select perl script until change incorporated in HIP build script. my $HIPCC_USE_PERL_SCRIPT = 1; +my $isWindows = ($^O eq 'MSWin32' or $^O eq 'msys'); +# escapes args with quotes SWDEV-341955 +foreach $arg (@ARGV) { + if ($isWindows) { + $arg =~ s/[^-a-zA-Z0-9_=+,.:\/\\]/\\$&/g; + } +} my $SCRIPT_DIR=dirname(__FILE__); if ($HIPCC_USE_PERL_SCRIPT) { diff --git a/bin/hipcc.pl b/bin/hipcc.pl index 24edf7c59c..645ae62d03 100755 --- a/bin/hipcc.pl +++ b/bin/hipcc.pl @@ -595,9 +595,12 @@ foreach $arg (@ARGV) # common characters such as alphanumerics. # Do the quoting here because sometimes the $arg is changed in the loop # Important to have all of '-Xlinker' in the set of unquoted characters. - if (not $isWindows and $escapeArg) { # Windows needs different quoting, ignore for now + if (not $isWindows and $escapeArg) { $arg =~ s/[^-a-zA-Z0-9_=+,.\/]/\\$&/g; } + if ($isWindows and $escapeArg) { + $arg =~ s/[^-a-zA-Z0-9_=+,.:\/\\]/\\$&/g; + } $toolArgs .= " $arg" unless $swallowArg; $prevArg = $arg; } diff --git a/docs/markdown/hip_build.md b/docs/markdown/hip_build.md index d61103aed2..17bc1af955 100755 --- a/docs/markdown/hip_build.md +++ b/docs/markdown/hip_build.md @@ -8,10 +8,12 @@ * [Set the environment variables](#set-the-environment-variables) * [Build HIP](#build-hip) * [Default paths and environment variables](#default-paths-and-environment-variables) + * [Build HIP Tests](#build-hip-tests) - [Build HIP on NVIDIA platform](#build-hip-on-NVIDIA-platform) * [Get HIP source code](#get-hip-source-code) * [Set the environment variables](#set-the-environment-variables) * [Build HIP](#build-hip) + * [Build HIP tests](#build-hip-tests) - [Run HIP](#run-hip) @@ -30,7 +32,6 @@ sudo apt install comgr sudo apt-get -y install rocm-dkms ``` - ## NVIDIA platform Install Nvidia driver and pre-build packages (see HIP Installation Guide at https://docs.amd.com/ for the release) @@ -61,7 +62,6 @@ ROCM_PATH is path where ROCM is installed. BY default ROCM_PATH is at /opt/rocm. git clone -b $HIP_BRANCH https://github.com/ROCm-Developer-Tools/hipamd.git git clone -b $HIP_BRANCH https://github.com/ROCm-Developer-Tools/hip.git git clone -b $HIP_BRANCH https://github.com/ROCm-Developer-Tools/ROCclr.git -git clone -b $HIP_BRANCH https://github.com/RadeonOpenCompute/ROCm-OpenCL-Runtime.git ``` ## Set the environment variables @@ -69,8 +69,6 @@ git clone -b $HIP_BRANCH https://github.com/RadeonOpenCompute/ROCm-OpenCL-Runtim ``` export HIPAMD_DIR="$(readlink -f hipamd)" export HIP_DIR="$(readlink -f hip)" -export ROCclr_DIR="$(readlink -f ROCclr)" -export OPENCL_DIR="$(readlink -f ROCm-OpenCL-Runtime)" ``` ROCclr is defined on AMD platform that HIP use Radeon Open Compute Common Language Runtime (ROCclr), which is a virtual device interface that HIP runtimes interact with different backends. @@ -84,7 +82,7 @@ See https://github.com/ROCm-Developer-Tools/hipamd ``` cd "$HIPAMD_DIR" mkdir -p build; cd build -cmake -DHIP_COMMON_DIR=$HIP_DIR -DAMD_OPENCL_PATH=$OPENCL_DIR -DROCCLR_PATH=$ROCCLR_DIR -DCMAKE_PREFIX_PATH="/" -DCMAKE_INSTALL_PREFIX=$PWD/install .. +cmake -DHIP_COMMON_DIR=$HIP_DIR -DCMAKE_PREFIX_PATH="/" -DCMAKE_INSTALL_PREFIX=$PWD/install .. make -j$(nproc) sudo make install ``` @@ -103,6 +101,77 @@ By default, release version of AMDHIP is built. After make install command, make sure HIP_PATH is pointed to $PWD/install/hip. +## Build HIP tests + +### Build HIP directed tests +Developers can build HIP directed tests right after build HIP commands, + +``` +sudo make install +make -j$(nproc) build_tests +``` +By default, all HIP directed tests will be built and generated under the folder $HIPAMD_DIR/build/directed_tests. +Take HIP directed device APIs tests, as an example, all available test applications will have executable files generated under, +$HIPAMD_DIR/build/directed_tests/runtimeApi/device. + +Run all HIP directed_tests, use the command, + +``` +ctest +``` +Or +``` +make test +``` + +Build and run a single directed test, use the follow command as an example, + +``` +make directed_tests.texture.hipTexObjPitch +cd $HIPAMD_DIR/build/directed_tests/texcture +./hipTexObjPitch +``` +Please note, the integrated HIP directed tests, will be deprecated in future release. + + +### Build HIP catch tests + +After build and install HIP commands, catch tests can be built via the following instructions, + +``` +cd "$HIP_DIR" +mkdir -p build; cd build +export HIP_PATH=$HIPAMD_DIR/build +cmake ../tests/catch/ -DHIP_PLATFORM=amd +make -j$(nproc) build_tests +ctest # run tests +``` + +HIP catch tests are built under the folder $HIP_DIR/build. + +To run a single catch test, the following is an example, + +``` +cd $HIP_DIR/build/unit/texture +./TextureTest +``` + +### Build HIP Catch2 standalone test + +HIP Catch2 supports build a standalone test, for example, + +``` +export PATH=$HIP_DIR/bin:$PATH +export HIP_PATH=$HIPAMD_DIR/build/install + +hipcc $HIP_DIR/tests/catch/unit/memory/hipPointerGetAttributes.cc -I ./tests/catch/include ./tests/catch/hipTestMain/standalone_main.cc -I ./tests/catch/external/Catch2 -g -o hipPointerGetAttributes +./hipPointerGetAttributes +... + +All tests passed +``` + +HIP catch tests, especially new architectured Catch2, will be official HIP tests in the repository and can be built alone as with the instructions shown above. # Build HIP on NVIDIA platform @@ -131,6 +200,9 @@ make -j$(nproc) sudo make install ``` +## Build HIP tests +Build HIP tests commands on NVIDIA platform are basically the same as AMD, except set -DHIP_PLATFORM=nvidia. + # Run HIP Compile and run the [square sample](https://github.com/ROCm-Developer-Tools/HIP/tree/rocm-5.0.x/samples/0_Intro/square). diff --git a/hip-lang-config.cmake.in b/hip-lang-config.cmake.in index 1a72643a1f..dbc8e5348d 100644 --- a/hip-lang-config.cmake.in +++ b/hip-lang-config.cmake.in @@ -94,7 +94,7 @@ find_path(HSA_HEADER hsa/hsa.h /opt/rocm/include ) -if (HSA_HEADER-NOTFOUND) +if (NOT HSA_HEADER) message (FATAL_ERROR "HSA header not found! ROCM_PATH environment not set") endif() @@ -136,7 +136,7 @@ set_property(TARGET hip-lang::device APPEND PROPERTY ) # Add support for __fp16 and _Float16, explicitly link with compiler-rt -if(CLANGRT_BUILTINS-NOTFOUND) +if(NOT CLANGRT_BUILTINS) message(FATAL_ERROR "clangrt builtins lib not found") else() set_property(TARGET hip-lang::device APPEND PROPERTY diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index e6a5dc2ab8..2ac304df2f 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -1192,7 +1192,6 @@ typedef enum hipGraphInstantiateFlags { hipGraphInstantiateFlagAutoFreeOnLaunch = 1, ///< Automatically free memory allocated in a graph before relaunching. } hipGraphInstantiateFlags; -#include // Doxygen end group GlobalDefs /** @} */ @@ -1517,9 +1516,9 @@ hipError_t hipGetDeviceProperties(hipDeviceProp_t* prop, int deviceId); */ hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig); /** - * @brief Set Cache configuration for a specific function + * @brief Get Cache configuration for a specific Device * - * @param [in] cacheConfig + * @param [out] cacheConfig * * @returns #hipSuccess, #hipErrorNotInitialized * Note: AMD devices and some Nvidia GPUS do not support reconfigurable cache. This hint is ignored @@ -6813,6 +6812,8 @@ static inline hipError_t hipMallocManaged(T** devPtr, size_t size, #endif #endif +#include + #if USE_PROF_API #include #endif diff --git a/include/hip/hiprtc.h b/include/hip/hiprtc.h index 9d4a598e01..61f7649305 100644 --- a/include/hip/hiprtc.h +++ b/include/hip/hiprtc.h @@ -269,6 +269,29 @@ hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* code); */ hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, size_t* codeSizeRet); +/** + * @brief Gets the pointer of compiled bitcode by the runtime compilation program instance. + * + * @param [in] prog runtime compilation program instance. + * @param [out] code char pointer to bitcode. + * @return HIPRTC_SUCCESS + * + * @see hiprtcResult + */ +hiprtcResult hiprtcGetBitcode(hiprtcProgram prog, char* bitcode); + +/** + * @brief Gets the size of compiled bitcode by the runtime compilation program instance. + * + * + * @param [in] prog runtime compilation program instance. + * @param [out] code the size of bitcode. + * @return HIPRTC_SUCCESS + * + * @see hiprtcResult + */ +hiprtcResult hiprtcGetBitcodeSize(hiprtcProgram prog, size_t* bitcode_size); + /** * @brief Creates the link instance via hiprtc APIs. * diff --git a/samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp b/samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp index 012debc148..fe71a03c5f 100644 --- a/samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp +++ b/samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp @@ -37,8 +37,8 @@ THE SOFTWARE. // Device (Kernel) function, it must be void __global__ void matrixTranspose(float* out, float* in, const int width) { - int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; - int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; out[y * width + x] = in[x * width + y]; } diff --git a/samples/2_Cookbook/1_hipEvent/hipEvent.cpp b/samples/2_Cookbook/1_hipEvent/hipEvent.cpp index a83b99beb6..81013c1dd1 100644 --- a/samples/2_Cookbook/1_hipEvent/hipEvent.cpp +++ b/samples/2_Cookbook/1_hipEvent/hipEvent.cpp @@ -35,8 +35,8 @@ THE SOFTWARE. // Device (Kernel) function, it must be void __global__ void matrixTranspose(float* out, float* in, const int width) { - int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; - int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; out[y * width + x] = in[x * width + y]; } diff --git a/samples/2_Cookbook/20_hip_vulkan/CMakeLists.txt b/samples/2_Cookbook/20_hip_vulkan/CMakeLists.txt new file mode 100644 index 0000000000..288c47c3d2 --- /dev/null +++ b/samples/2_Cookbook/20_hip_vulkan/CMakeLists.txt @@ -0,0 +1,97 @@ + +# Copyright (c) 2020 - 2022 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. + +# hipcc.bat fails to qualify as a valid compiler for CMAKE_CXX_COMPILER_ID = ROCMClang +# so the simple compiler test is skipped and forced to use hipcc.bat as compiler +set(CMAKE_C_COMPILER_WORKS 1) +set(CMAKE_CXX_COMPILER_WORKS 1) +set(CMAKE_CXX_STANDARD 14) +project(hipVulkan) + +cmake_minimum_required(VERSION 3.10) +set(CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/cmake;${CMAKE_MODULE_PATH}") + +if (NOT DEFINED ROCM_PATH ) + set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." ) +endif () + +# Search for rocm in common locations +list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH}) + +# need to set rocm_path for windows +# since clang and hip are two different folders during build/install step +if (WIN32 AND HIPINFO_INTERNAL_BUILD) + set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --rocm-path=${HIP_PATH}") +endif() + + + +# Find hip +find_package(hip REQUIRED) +if (WIN32) +find_package(GLFW3) +if(NOT GLFW_FOUND) + if(EXISTS "${GLFW_PATH}") + message(STATUS "FOUND GLFW SDK: ${GLFW_PATH}") + elseif (EXISTS "$ENV{GLFW_PATH}") + message(STATUS "FOUND GLFW SDK: $ENV{GLFW_PATH}") + set(GLFW_PATH $ENV{GLFW_PATH}) + else() + message("Error: Unable to locate GLFW SDK. please specify GLFW_PATH") + return() + endif() +endif() +endif(WIN32) +find_package(Vulkan) +if(NOT Vulkan_FOUND) + if(EXISTS "${VULKAN_PATH}") + message(STATUS "Vulkan SDK: ${VULKAN_PATH}") + elseif (EXISTS "$ENV{VULKAN_SDK}") + message(STATUS "FOUND VULKAN SDK: $ENV{VULKAN_SDK}") + set(VULKAN_PATH $ENV{VULKAN_SDK}) + else() + message("Error: Unable to locate Vulkan SDK. please specify VULKAN_PATH") + return() + endif() +endif() +set(VULKAN_PATH ${Vulkan_INCLUDE_DIRS}) +STRING(REGEX REPLACE "/[Ii]nclude" "" VULKAN_PATH ${VULKAN_PATH}) +# Include Vulkan header files from Vulkan SDK +include_directories(AFTER ${VULKAN_PATH}/include) +link_directories(${VULKAN_PATH}/bin;${VULKAN_PATH}/lib;) +link_directories(${GLFW_PATH}/lib-vc2019) +# Set compiler and linker +set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE}) +set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE}) +set(CMAKE_BUILD_TYPE Release) + +# Create the excutable +add_executable(hipVulkan VulkanBaseApp.cpp VulkanBaseApp.h main.cpp SineWaveSimulation.cpp SineWaveSimulation.h linmath.h) +include_directories(${HIP_PATH}/include) +include_directories(${GLFW_PATH}/include) + +# Link with HIP +if (WIN32) + target_link_libraries(hipVulkan advapi32 hip::host vulkan-1 glfw3dll) +else (WIN32) + target_link_libraries(hipVulkan hip::host vulkan glfw) +endif (WIN32) + diff --git a/samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.cpp b/samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.cpp new file mode 100644 index 0000000000..41adae8065 --- /dev/null +++ b/samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.cpp @@ -0,0 +1,147 @@ +/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of NVIDIA CORPORATION nor the names of its + * contributors may be used to endorse or promote products derived + * from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY + * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY + * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/* + * Modifications Copyright (C)2021 Advanced + * Micro Devices, Inc. All rights reserved. + */ + +#include "SineWaveSimulation.h" +#include +//#include +#include "hip/hip_runtime.h" + + +__global__ void sinewave(float *heightMap, unsigned int width, unsigned int height, float time) +{ + const float freq = 4.0f; + const size_t stride = gridDim.x * blockDim.x; + + // Iterate through the entire array in a way that is + // independent of the grid configuration + for (size_t tid = blockIdx.x * blockDim.x + threadIdx.x; tid < width * height; tid += stride) { + // Calculate the x, y coordinates + const size_t y = tid / width; + const size_t x = tid - y * width; + // Normalize x, y to [0,1] + const float u = ((2.0f * x) / width) - 1.0f; + const float v = ((2.0f * y) / height) - 1.0f; + // Calculate the new height value + const float w = 0.5f * sinf(u * freq + time) * cosf(v * freq + time); + // Store this new height value + heightMap[tid] = w; + } +} + +SineWaveSimulation::SineWaveSimulation(size_t width, size_t height) + : m_heightMap(nullptr), m_width(width), m_height(height) +{ +} + +void SineWaveSimulation::initCudaLaunchConfig(int device) +{ + hipDeviceProp_t prop = {}; + checkHIPErrors(hipSetDevice(device)); + checkHIPErrors(hipGetDeviceProperties(&prop, device)); + + // We don't need large block sizes, since there's not much inter-thread communication + m_threads = prop.warpSize; + + // Use the occupancy calculator and fill the gpu as best as we can + checkHIPErrors(hipOccupancyMaxActiveBlocksPerMultiprocessor(&m_blocks, sinewave, prop.warpSize, 0)); + m_blocks *= prop.multiProcessorCount; + + // Go ahead and the clamp the blocks to the minimum needed for this height/width + m_blocks = std::min(m_blocks, (int)((m_width * m_height + m_threads - 1) / m_threads)); +} + +int SineWaveSimulation::initCuda(uint8_t *vkDeviceUUID, size_t UUID_SIZE) +{ + int current_device = 0; + int device_count = 0; + int devices_prohibited = 0; + + hipDeviceProp_t deviceProp; + checkHIPErrors(hipGetDeviceCount(&device_count)); + + if (device_count == 0) { + fprintf(stderr, "CUDA error: no devices supporting CUDA.\n"); + exit(EXIT_FAILURE); + } + + // Find the GPU which is selected by Vulkan + while (current_device < device_count) { + hipGetDeviceProperties(&deviceProp, current_device); + + if ((deviceProp.computeMode != hipComputeModeProhibited)) { + // Compare the cuda device UUID with vulkan UUID + // FIXME + int ret = 0; // memcmp((void*)&deviceProp.uuid, vkDeviceUUID, UUID_SIZE); + if (ret == 0) + { + checkHIPErrors(hipSetDevice(current_device)); + checkHIPErrors(hipGetDeviceProperties(&deviceProp, current_device)); + printf("GPU Device %d: \"%s\" with compute capability %d.%d\n\n", + current_device, deviceProp.name, deviceProp.major, + deviceProp.minor); + + return current_device; + } + + } else { + devices_prohibited++; + } + + current_device++; + } + + if (devices_prohibited == device_count) { + fprintf(stderr, + "HIP error:" + " No Vulkan-HIP Interop capable GPU found.\n"); + exit(EXIT_FAILURE); + } + + return -1; +} + +SineWaveSimulation::~SineWaveSimulation() +{ + m_heightMap = NULL; +} + +void SineWaveSimulation::initSimulation(float *heights) +{ + m_heightMap = heights; +} + +void SineWaveSimulation::stepSimulation(float time, hipStream_t stream) +{ + hipLaunchKernelGGL(sinewave, dim3(m_blocks), dim3(m_threads), 0, stream , m_heightMap, m_width, m_height, time); + getLastHIPError("Failed to launch CUDA simulation"); + //hipStreamSynchronize(stream); +} diff --git a/samples/2_Cookbook/20_hip_vulkan/buildcmd.txt b/samples/2_Cookbook/20_hip_vulkan/buildcmd.txt index b847596510..791466ec42 100644 --- a/samples/2_Cookbook/20_hip_vulkan/buildcmd.txt +++ b/samples/2_Cookbook/20_hip_vulkan/buildcmd.txt @@ -5,7 +5,12 @@ o c:\VulkanSDK\1.2.182.0\bin\glslangValidator.exe sinewave.vert -V -o vert.spv o c:\VulkanSDK\1.2.182.0\bin\glslangValidator.exe sinewave.frag -V -o frag.spv +to build without cmake: • set HCC_AMDGPU_TARGET=gfx906:sramecc-:xnack- (for your graphic card, you can get the name from hipinfo ) -$• hipcc -v *.cpp *.hip -Ic:\VulkanSDK\1.2.182.0\include -L c:\VulkanSDK\1.2.182.0\lib -Ic:\glfw-3.3.4.bin.WIN64\include -L c:\glfw-3.3.4.bin.WIN64\lib-vc2019 -Ic:\hip\include\hip -lglfw3dll -lvulkan-1 -ladvapi32 -std=c++14 +• hipcc -v *.cpp *.hip -Ic:\VulkanSDK\1.2.182.0\include -L c:\VulkanSDK\1.2.182.0\lib -Ic:\glfw-3.3.4.bin.WIN64\include -L c:\glfw-3.3.4.bin.WIN64\lib-vc2019 -Ic:\hip\include\hip -lglfw3dll -lvulkan-1 -ladvapi32 -std=c++14 • run a.exe, you should see a 3D sinewave simulation +to build with cmake on windows: +• mkdir build; cd build +• cmake.exe -GNinja -DCMAKE_CXX_COMPILER_ID=ROCMClang -DCMAKE_C_COMPILER_ID=ROCMClang -DCMAKE_PREFIX_PATH=d:\driver2\drivers\drivers\compute\hip_sdk + diff --git a/samples/2_Cookbook/20_hip_vulkan/main.cpp b/samples/2_Cookbook/20_hip_vulkan/main.cpp index 0788243e52..93c8db75db 100644 --- a/samples/2_Cookbook/20_hip_vulkan/main.cpp +++ b/samples/2_Cookbook/20_hip_vulkan/main.cpp @@ -37,7 +37,7 @@ #include #include #include "linmath.h" -#include "hip_runtime.h" +#include "hip/hip_runtime.h" #include "SineWaveSimulation.h" @@ -50,6 +50,18 @@ std::string execution_path; #define ENABLE_VALIDATION (true) #endif +#ifndef _WIN64 +#define MAX_PATH 260 + +int GetModuleFileName(void* hndl, char* name, int size) +{ + FILE* stream = fopen("/proc/self/cmdline", "r"); + fgets(name, size, stream); + fclose(stream); + return strlen(name); +} +#endif + class VulkanCudaSineWave : public VulkanBaseApp { @@ -93,7 +105,7 @@ public: throw std::runtime_error("Requested height and width is too large for this sample!"); } // Add our compiled vulkan shader files - TCHAR buffer[MAX_PATH] = { 0 }; + char buffer[MAX_PATH] = { 0 }; //assuming none unicode GetModuleFileName(NULL, buffer, MAX_PATH); std::string str3 = std::string(buffer); std::string str1 = "vert.spv" ; //sdkFindFilePath("sinewave.vert", execution_path.c_str()); diff --git a/samples/2_Cookbook/21_cmake_hip_cxx_clang/CMakeLists.txt b/samples/2_Cookbook/21_cmake_hip_cxx_clang/CMakeLists.txt new file mode 100644 index 0000000000..3cf4f35336 --- /dev/null +++ b/samples/2_Cookbook/21_cmake_hip_cxx_clang/CMakeLists.txt @@ -0,0 +1,34 @@ +# Copyright (C) 2022 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. + +cmake_minimum_required(VERSION 3.21.3) + +project(cmake_cxx_amdclang++_test + DESCRIPTION "Verifies CXX Language Support with amdclang++" + LANGUAGES CXX) + +# Find HIP +find_package(hip REQUIRED) + +# Create the excutable +add_executable(square square.cpp) + +# Link with HIP +target_link_libraries(square hip::device) diff --git a/samples/2_Cookbook/21_cmake_hip_cxx_clang/README.md b/samples/2_Cookbook/21_cmake_hip_cxx_clang/README.md new file mode 100644 index 0000000000..3d9765308c --- /dev/null +++ b/samples/2_Cookbook/21_cmake_hip_cxx_clang/README.md @@ -0,0 +1,19 @@ +### This sample tests CXX Language support with amdclang++ +I. Build +mkdir -p build; cd build +rm -rf *; +CXX=`hipconfig -l`/amdclang++ cmake .. (or) +cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/amdclang++ .. (or) +cmake -DCMAKE_CXX_COMPILER=/opt/rocm-X.Y.Z/llvm/bin/amdclang++ .. +make + +II. Test +$ ./square +info: running on device Vega 20 [Radeon Pro Vega 20] +info: allocate host mem ( 7.63 MB) +info: allocate device mem ( 7.63 MB) +info: copy Host2Device +info: launch 'vector_square' kernel +info: copy Device2Host +info: check result +PASSED! diff --git a/samples/2_Cookbook/21_cmake_hip_cxx_clang/square.cpp b/samples/2_Cookbook/21_cmake_hip_cxx_clang/square.cpp new file mode 100644 index 0000000000..c3ed799edc --- /dev/null +++ b/samples/2_Cookbook/21_cmake_hip_cxx_clang/square.cpp @@ -0,0 +1,98 @@ +/* +Copyright (c) 2015-2022 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) \ +{\ + hipError_t error = cmd;\ + if (error != hipSuccess) { \ + fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(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, T *A_d, size_t N) +{ + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x ; + + for (size_t i=offset; i +#include + +#define CHECK(cmd) \ +{\ + hipError_t error = cmd;\ + if (error != hipSuccess) { \ + fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(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, T *A_d, size_t N) +{ + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x ; + + for (size_t i=offset; i