diff --git a/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/CMakeLists.txt b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/CMakeLists.txt new file mode 100644 index 0000000000..288c47c3d2 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.cpp b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.cpp new file mode 100644 index 0000000000..41adae8065 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/buildcmd.txt b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/buildcmd.txt index b847596510..791466ec42 100644 --- a/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/buildcmd.txt +++ b/projects/hip-tests/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/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/main.cpp b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/main.cpp index 0788243e52..93c8db75db 100644 --- a/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/main.cpp +++ b/projects/hip-tests/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());