SWDEV-336378 - cmake support for vulkan example (#2725)

Change-Id: I5ac28d2f57c1732e243674aa904cabe23be34797 [ROCm/hip-tests commit: e9202ded5b]
2022-06-13 15:35:33 +05:30
commit 3913b8282b
@@ -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)
+
@@ -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 <algorithm>
+//#include <helper_cuda.h>
+#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);
+}
@@ -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
+
@@ -37,7 +37,7 @@
 #include <chrono>
 #include <algorithm>
 #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());