SWDEV-299127 - Merge 'develop' into 'amd-staging'
Change-Id: I8e7e97d9bf619743c300f68f51d5bf115fdd534e
This commit is contained in:
Jenkins
@@ -0,0 +1,101 @@
|
||||
/* 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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#ifndef __SINESIM_H__
|
||||
#define __SINESIM_H__
|
||||
|
||||
#include <vector>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#include <stdint.h>
|
||||
#include "linmath.h"
|
||||
|
||||
class SineWaveSimulation
|
||||
{
|
||||
float *m_heightMap;
|
||||
size_t m_width, m_height;
|
||||
int m_blocks, m_threads;
|
||||
public:
|
||||
SineWaveSimulation(size_t width, size_t height);
|
||||
~SineWaveSimulation();
|
||||
void initSimulation(float *heightMap);
|
||||
void stepSimulation(float time, hipStream_t stream = 0);
|
||||
void initCudaLaunchConfig(int device);
|
||||
int initCuda(uint8_t *vkDeviceUUID, size_t UUID_SIZE);
|
||||
|
||||
size_t getWidth() const {
|
||||
return m_width;
|
||||
}
|
||||
size_t getHeight() const {
|
||||
return m_height;
|
||||
}
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
void check(T result, char const* const func, const char* const file,
|
||||
int const line) {
|
||||
if (result) {
|
||||
fprintf(stderr, "HIP error at %s:%d code=%d \"%s\" \n", file, line, static_cast<unsigned int>(result), func);
|
||||
// static_cast<unsigned int>(result), _cudaGetErrorEnum(result), func);
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
|
||||
// This will output the proper CUDA error strings in the event
|
||||
// that a CUDA host call returns an error
|
||||
#define checkHIPErrors(val) check((val), #val, __FILE__, __LINE__)
|
||||
|
||||
// This will output the proper error string when calling cudaGetLastError
|
||||
#define getLastHIPError(msg) __getLastHIPError(msg, __FILE__, __LINE__)
|
||||
|
||||
inline void __getLastHIPError(const char* errorMessage, const char* file,
|
||||
const int line) {
|
||||
hipError_t err = hipGetLastError();
|
||||
|
||||
if (hipSuccess != err) {
|
||||
fprintf(stderr,
|
||||
"%s(%i) : getLastHIPError() HIP error :"
|
||||
" %s : (%d) %s.\n",
|
||||
file, line, errorMessage, static_cast<int>(err),
|
||||
hipGetErrorString(err));
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
|
||||
|
||||
#ifndef MAX
|
||||
#define MAX(a, b) (a > b ? a : b)
|
||||
#endif
|
||||
|
||||
#endif // __SINESIM_H__
|
||||
@@ -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);
|
||||
}
|
||||
تفاوت فایلی نمایش داده نمی شود زیرا این فایل بسیار بزرگ است
Diff را بارگزاری کن
@@ -0,0 +1,146 @@
|
||||
/* 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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#ifndef __VULKANBASEAPP_H__
|
||||
#define __VULKANBASEAPP_H__
|
||||
|
||||
#include <string>
|
||||
#include <vector>
|
||||
#include <vulkan/vulkan.h>
|
||||
#ifdef _WIN64
|
||||
#define NOMINMAX
|
||||
#include <windows.h>
|
||||
#include <vulkan/vulkan_win32.h>
|
||||
#endif /* _WIN64 */
|
||||
|
||||
struct GLFWwindow;
|
||||
|
||||
class VulkanBaseApp
|
||||
{
|
||||
public:
|
||||
VulkanBaseApp(const std::string& appName, bool enableValidation = false);
|
||||
static VkExternalSemaphoreHandleTypeFlagBits getDefaultSemaphoreHandleType();
|
||||
static VkExternalMemoryHandleTypeFlagBits getDefaultMemHandleType();
|
||||
virtual ~VulkanBaseApp();
|
||||
void init();
|
||||
void *getMemHandle(VkDeviceMemory memory, VkExternalMemoryHandleTypeFlagBits handleType);
|
||||
void *getSemaphoreHandle(VkSemaphore semaphore, VkExternalSemaphoreHandleTypeFlagBits handleType);
|
||||
void createExternalSemaphore(VkSemaphore& semaphore, VkExternalSemaphoreHandleTypeFlagBits handleType);
|
||||
void createBuffer(VkDeviceSize size, VkBufferUsageFlags usage, VkMemoryPropertyFlags properties, VkBuffer& buffer, VkDeviceMemory& bufferMemory);
|
||||
void createExternalBuffer(VkDeviceSize size, VkBufferUsageFlags usage, VkMemoryPropertyFlags properties, VkExternalMemoryHandleTypeFlagsKHR extMemHandleType, VkBuffer& buffer, VkDeviceMemory& bufferMemory);
|
||||
void importExternalBuffer(void *handle, VkExternalMemoryHandleTypeFlagBits handleType, size_t size, VkBufferUsageFlags usage, VkMemoryPropertyFlags properties, VkBuffer& buffer, VkDeviceMemory& memory);
|
||||
void copyBuffer(VkBuffer dst, VkBuffer src, VkDeviceSize size);
|
||||
VkCommandBuffer beginSingleTimeCommands();
|
||||
void endSingleTimeCommands(VkCommandBuffer commandBuffer);
|
||||
void mainLoop();
|
||||
protected:
|
||||
const std::string m_appName;
|
||||
const bool m_enableValidation;
|
||||
VkInstance m_instance;
|
||||
VkDebugUtilsMessengerEXT m_debugMessenger;
|
||||
VkSurfaceKHR m_surface;
|
||||
VkPhysicalDevice m_physicalDevice;
|
||||
VkDevice m_device;
|
||||
VkQueue m_graphicsQueue;
|
||||
VkQueue m_presentQueue;
|
||||
VkSwapchainKHR m_swapChain;
|
||||
std::vector<VkImage> m_swapChainImages;
|
||||
VkFormat m_swapChainFormat;
|
||||
VkExtent2D m_swapChainExtent;
|
||||
std::vector<VkImageView> m_swapChainImageViews;
|
||||
std::vector<std::pair<VkShaderStageFlagBits, std::string> > m_shaderFiles;
|
||||
VkRenderPass m_renderPass;
|
||||
VkPipelineLayout m_pipelineLayout;
|
||||
VkPipeline m_graphicsPipeline;
|
||||
std::vector<VkFramebuffer> m_swapChainFramebuffers;
|
||||
VkCommandPool m_commandPool;
|
||||
std::vector<VkCommandBuffer> m_commandBuffers;
|
||||
std::vector<VkSemaphore> m_imageAvailableSemaphores;
|
||||
std::vector<VkSemaphore> m_renderFinishedSemaphores;
|
||||
std::vector<VkFence> m_inFlightFences;
|
||||
std::vector<VkBuffer> m_uniformBuffers;
|
||||
std::vector<VkDeviceMemory> m_uniformMemory;
|
||||
VkDescriptorSetLayout m_descriptorSetLayout;
|
||||
VkDescriptorPool m_descriptorPool;
|
||||
std::vector<VkDescriptorSet> m_descriptorSets;
|
||||
VkImage m_depthImage;
|
||||
VkDeviceMemory m_depthImageMemory;
|
||||
VkImageView m_depthImageView;
|
||||
size_t m_currentFrame;
|
||||
bool m_framebufferResized;
|
||||
uint8_t m_vkDeviceUUID[VK_UUID_SIZE];
|
||||
|
||||
virtual void initVulkanApp() {}
|
||||
virtual void fillRenderingCommandBuffer(VkCommandBuffer& buffer) {}
|
||||
virtual std::vector<const char *> getRequiredExtensions() const;
|
||||
virtual std::vector<const char *> getRequiredDeviceExtensions() const;
|
||||
virtual void getVertexDescriptions(std::vector<VkVertexInputBindingDescription>& bindingDesc, std::vector<VkVertexInputAttributeDescription>& attribDesc);
|
||||
virtual void getAssemblyStateInfo(VkPipelineInputAssemblyStateCreateInfo& info);
|
||||
virtual void getWaitFrameSemaphores(std::vector<VkSemaphore>& wait, std::vector< VkPipelineStageFlags>& waitStages) const;
|
||||
virtual void getSignalFrameSemaphores(std::vector<VkSemaphore>& signal) const;
|
||||
virtual VkDeviceSize getUniformSize() const;
|
||||
virtual void updateUniformBuffer(uint32_t imageIndex);
|
||||
virtual void drawFrame();
|
||||
private:
|
||||
GLFWwindow *m_window;
|
||||
|
||||
void initWindow();
|
||||
void initVulkan();
|
||||
void createInstance();
|
||||
void createSurface();
|
||||
void createDevice();
|
||||
void createSwapChain();
|
||||
void createImageViews();
|
||||
void createRenderPass();
|
||||
void createDescriptorSetLayout();
|
||||
void createGraphicsPipeline();
|
||||
void createFramebuffers();
|
||||
void createCommandPool();
|
||||
void createDepthResources();
|
||||
void createUniformBuffers();
|
||||
void createDescriptorPool();
|
||||
void createDescriptorSets();
|
||||
void createCommandBuffers();
|
||||
void createSyncObjects();
|
||||
|
||||
void cleanupSwapChain();
|
||||
void recreateSwapChain();
|
||||
|
||||
bool isSuitableDevice(VkPhysicalDevice dev) const;
|
||||
static void resizeCallback(GLFWwindow *window, int width, int height);
|
||||
};
|
||||
|
||||
void readFile(std::istream& s, std::vector<char>& data);
|
||||
|
||||
#endif /* __VULKANBASEAPP_H__ */
|
||||
@@ -0,0 +1,11 @@
|
||||
• Install hip and visual studio
|
||||
• Install vulkan sdk from vulkan.lunarg.com
|
||||
• Download GLFW binaries from glfw.org
|
||||
• Convert sinwave.farg and vert files to spv
|
||||
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
|
||||
|
||||
• 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
|
||||
• run a.exe, you should see a 3D sinewave simulation
|
||||
|
||||
@@ -0,0 +1,502 @@
|
||||
/*
|
||||
* Copyright (c) 2015-2016 The Khronos Group Inc.
|
||||
* Copyright (c) 2015-2016 Valve Corporation
|
||||
* Copyright (c) 2015-2016 LunarG, Inc.
|
||||
*
|
||||
* Licensed under the Apache License, Version 2.0 (the "License");
|
||||
* you may not use this file except in compliance with the License.
|
||||
* You may obtain a copy of the License at
|
||||
*
|
||||
* http://www.apache.org/licenses/LICENSE-2.0
|
||||
*
|
||||
* Unless required by applicable law or agreed to in writing, software
|
||||
* distributed under the License is distributed on an "AS IS" BASIS,
|
||||
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
* See the License for the specific language governing permissions and
|
||||
* limitations under the License.
|
||||
*
|
||||
* Relicensed from the WTFPL (http://www.wtfpl.net/faq/).
|
||||
*/
|
||||
|
||||
#ifndef LINMATH_H
|
||||
#define LINMATH_H
|
||||
|
||||
#define _USE_MATH_DEFINES
|
||||
#include <math.h>
|
||||
|
||||
// Converts degrees to radians.
|
||||
#define degreesToRadians(angleDegrees) (angleDegrees * M_PI / 180.0)
|
||||
|
||||
// Converts radians to degrees.
|
||||
#define radiansToDegrees(angleRadians) (angleRadians * 180.0 / M_PI)
|
||||
|
||||
typedef float vec3[3];
|
||||
static inline void vec3_add(vec3 r, vec3 const a, vec3 const b) {
|
||||
int i;
|
||||
for (i = 0; i < 3; ++i) r[i] = a[i] + b[i];
|
||||
}
|
||||
static inline void vec3_sub(vec3 r, vec3 const a, vec3 const b) {
|
||||
int i;
|
||||
for (i = 0; i < 3; ++i) r[i] = a[i] - b[i];
|
||||
}
|
||||
static inline void vec3_scale(vec3 r, vec3 const v, float const s) {
|
||||
int i;
|
||||
for (i = 0; i < 3; ++i) r[i] = v[i] * s;
|
||||
}
|
||||
static inline float vec3_mul_inner(vec3 const a, vec3 const b) {
|
||||
float p = 0.f;
|
||||
int i;
|
||||
for (i = 0; i < 3; ++i) p += b[i] * a[i];
|
||||
return p;
|
||||
}
|
||||
static inline void vec3_mul_cross(vec3 r, vec3 const a, vec3 const b) {
|
||||
r[0] = a[1] * b[2] - a[2] * b[1];
|
||||
r[1] = a[2] * b[0] - a[0] * b[2];
|
||||
r[2] = a[0] * b[1] - a[1] * b[0];
|
||||
}
|
||||
static inline float vec3_len(vec3 const v) { return sqrtf(vec3_mul_inner(v, v)); }
|
||||
static inline void vec3_norm(vec3 r, vec3 const v) {
|
||||
float k = 1.f / vec3_len(v);
|
||||
vec3_scale(r, v, k);
|
||||
}
|
||||
static inline void vec3_reflect(vec3 r, vec3 const v, vec3 const n) {
|
||||
float p = 2.f * vec3_mul_inner(v, n);
|
||||
int i;
|
||||
for (i = 0; i < 3; ++i) r[i] = v[i] - p * n[i];
|
||||
}
|
||||
|
||||
typedef float vec4[4];
|
||||
static inline void vec4_add(vec4 r, vec4 const a, vec4 const b) {
|
||||
int i;
|
||||
for (i = 0; i < 4; ++i) r[i] = a[i] + b[i];
|
||||
}
|
||||
static inline void vec4_sub(vec4 r, vec4 const a, vec4 const b) {
|
||||
int i;
|
||||
for (i = 0; i < 4; ++i) r[i] = a[i] - b[i];
|
||||
}
|
||||
static inline void vec4_scale(vec4 r, vec4 v, float s) {
|
||||
int i;
|
||||
for (i = 0; i < 4; ++i) r[i] = v[i] * s;
|
||||
}
|
||||
static inline float vec4_mul_inner(vec4 a, vec4 b) {
|
||||
float p = 0.f;
|
||||
int i;
|
||||
for (i = 0; i < 4; ++i) p += b[i] * a[i];
|
||||
return p;
|
||||
}
|
||||
static inline void vec4_mul_cross(vec4 r, vec4 a, vec4 b) {
|
||||
r[0] = a[1] * b[2] - a[2] * b[1];
|
||||
r[1] = a[2] * b[0] - a[0] * b[2];
|
||||
r[2] = a[0] * b[1] - a[1] * b[0];
|
||||
r[3] = 1.f;
|
||||
}
|
||||
static inline float vec4_len(vec4 v) { return sqrtf(vec4_mul_inner(v, v)); }
|
||||
static inline void vec4_norm(vec4 r, vec4 v) {
|
||||
float k = 1.f / vec4_len(v);
|
||||
vec4_scale(r, v, k);
|
||||
}
|
||||
static inline void vec4_reflect(vec4 r, vec4 v, vec4 n) {
|
||||
float p = 2.f * vec4_mul_inner(v, n);
|
||||
int i;
|
||||
for (i = 0; i < 4; ++i) r[i] = v[i] - p * n[i];
|
||||
}
|
||||
|
||||
typedef vec4 mat4x4[4];
|
||||
static inline void mat4x4_identity(mat4x4 M) {
|
||||
int i, j;
|
||||
for (i = 0; i < 4; ++i)
|
||||
for (j = 0; j < 4; ++j) M[i][j] = i == j ? 1.f : 0.f;
|
||||
}
|
||||
static inline void mat4x4_dup(mat4x4 M, mat4x4 N) {
|
||||
int i, j;
|
||||
for (i = 0; i < 4; ++i)
|
||||
for (j = 0; j < 4; ++j) M[i][j] = N[i][j];
|
||||
}
|
||||
static inline void mat4x4_row(vec4 r, mat4x4 M, int i) {
|
||||
int k;
|
||||
for (k = 0; k < 4; ++k) r[k] = M[k][i];
|
||||
}
|
||||
static inline void mat4x4_col(vec4 r, mat4x4 M, int i) {
|
||||
int k;
|
||||
for (k = 0; k < 4; ++k) r[k] = M[i][k];
|
||||
}
|
||||
static inline void mat4x4_transpose(mat4x4 M, mat4x4 N) {
|
||||
int i, j;
|
||||
for (j = 0; j < 4; ++j)
|
||||
for (i = 0; i < 4; ++i) M[i][j] = N[j][i];
|
||||
}
|
||||
static inline void mat4x4_add(mat4x4 M, mat4x4 a, mat4x4 b) {
|
||||
int i;
|
||||
for (i = 0; i < 4; ++i) vec4_add(M[i], a[i], b[i]);
|
||||
}
|
||||
static inline void mat4x4_sub(mat4x4 M, mat4x4 a, mat4x4 b) {
|
||||
int i;
|
||||
for (i = 0; i < 4; ++i) vec4_sub(M[i], a[i], b[i]);
|
||||
}
|
||||
static inline void mat4x4_scale(mat4x4 M, mat4x4 a, float k) {
|
||||
int i;
|
||||
for (i = 0; i < 4; ++i) vec4_scale(M[i], a[i], k);
|
||||
}
|
||||
static inline void mat4x4_scale_aniso(mat4x4 M, mat4x4 a, float x, float y, float z) {
|
||||
int i;
|
||||
vec4_scale(M[0], a[0], x);
|
||||
vec4_scale(M[1], a[1], y);
|
||||
vec4_scale(M[2], a[2], z);
|
||||
for (i = 0; i < 4; ++i) {
|
||||
M[3][i] = a[3][i];
|
||||
}
|
||||
}
|
||||
static inline void mat4x4_mul(mat4x4 M, mat4x4 a, mat4x4 b) {
|
||||
int k, r, c;
|
||||
for (c = 0; c < 4; ++c)
|
||||
for (r = 0; r < 4; ++r) {
|
||||
M[c][r] = 0.f;
|
||||
for (k = 0; k < 4; ++k) M[c][r] += a[k][r] * b[c][k];
|
||||
}
|
||||
}
|
||||
static inline void mat4x4_mul_vec4(vec4 r, mat4x4 M, vec4 v) {
|
||||
int i, j;
|
||||
for (j = 0; j < 4; ++j) {
|
||||
r[j] = 0.f;
|
||||
for (i = 0; i < 4; ++i) r[j] += M[i][j] * v[i];
|
||||
}
|
||||
}
|
||||
static inline void mat4x4_translate(mat4x4 T, float x, float y, float z) {
|
||||
mat4x4_identity(T);
|
||||
T[3][0] = x;
|
||||
T[3][1] = y;
|
||||
T[3][2] = z;
|
||||
}
|
||||
static inline void mat4x4_translate_in_place(mat4x4 M, float x, float y, float z) {
|
||||
vec4 t = {x, y, z, 0};
|
||||
vec4 r;
|
||||
int i;
|
||||
for (i = 0; i < 4; ++i) {
|
||||
mat4x4_row(r, M, i);
|
||||
M[3][i] += vec4_mul_inner(r, t);
|
||||
}
|
||||
}
|
||||
static inline void mat4x4_from_vec3_mul_outer(mat4x4 M, vec3 a, vec3 b) {
|
||||
int i, j;
|
||||
for (i = 0; i < 4; ++i)
|
||||
for (j = 0; j < 4; ++j) M[i][j] = i < 3 && j < 3 ? a[i] * b[j] : 0.f;
|
||||
}
|
||||
static inline void mat4x4_rotate(mat4x4 R, mat4x4 M, float x, float y, float z, float angle) {
|
||||
float s = sinf(angle);
|
||||
float c = cosf(angle);
|
||||
vec3 u = {x, y, z};
|
||||
|
||||
if (vec3_len(u) > 1e-4) {
|
||||
vec3_norm(u, u);
|
||||
mat4x4 T;
|
||||
mat4x4_from_vec3_mul_outer(T, u, u);
|
||||
|
||||
mat4x4 S = {{0, u[2], -u[1], 0}, {-u[2], 0, u[0], 0}, {u[1], -u[0], 0, 0}, {0, 0, 0, 0}};
|
||||
mat4x4_scale(S, S, s);
|
||||
|
||||
mat4x4 C;
|
||||
mat4x4_identity(C);
|
||||
mat4x4_sub(C, C, T);
|
||||
|
||||
mat4x4_scale(C, C, c);
|
||||
|
||||
mat4x4_add(T, T, C);
|
||||
mat4x4_add(T, T, S);
|
||||
|
||||
T[3][3] = 1.;
|
||||
mat4x4_mul(R, M, T);
|
||||
} else {
|
||||
mat4x4_dup(R, M);
|
||||
}
|
||||
}
|
||||
static inline void mat4x4_rotate_X(mat4x4 Q, mat4x4 M, float angle) {
|
||||
float s = sinf(angle);
|
||||
float c = cosf(angle);
|
||||
mat4x4 R = {{1.f, 0.f, 0.f, 0.f}, {0.f, c, s, 0.f}, {0.f, -s, c, 0.f}, {0.f, 0.f, 0.f, 1.f}};
|
||||
mat4x4_mul(Q, M, R);
|
||||
}
|
||||
static inline void mat4x4_rotate_Y(mat4x4 Q, mat4x4 M, float angle) {
|
||||
float s = sinf(angle);
|
||||
float c = cosf(angle);
|
||||
mat4x4 R = {{c, 0.f, s, 0.f}, {0.f, 1.f, 0.f, 0.f}, {-s, 0.f, c, 0.f}, {0.f, 0.f, 0.f, 1.f}};
|
||||
mat4x4_mul(Q, M, R);
|
||||
}
|
||||
static inline void mat4x4_rotate_Z(mat4x4 Q, mat4x4 M, float angle) {
|
||||
float s = sinf(angle);
|
||||
float c = cosf(angle);
|
||||
mat4x4 R = {{c, s, 0.f, 0.f}, {-s, c, 0.f, 0.f}, {0.f, 0.f, 1.f, 0.f}, {0.f, 0.f, 0.f, 1.f}};
|
||||
mat4x4_mul(Q, M, R);
|
||||
}
|
||||
static inline void mat4x4_invert(mat4x4 T, mat4x4 M) {
|
||||
float s[6];
|
||||
float c[6];
|
||||
s[0] = M[0][0] * M[1][1] - M[1][0] * M[0][1];
|
||||
s[1] = M[0][0] * M[1][2] - M[1][0] * M[0][2];
|
||||
s[2] = M[0][0] * M[1][3] - M[1][0] * M[0][3];
|
||||
s[3] = M[0][1] * M[1][2] - M[1][1] * M[0][2];
|
||||
s[4] = M[0][1] * M[1][3] - M[1][1] * M[0][3];
|
||||
s[5] = M[0][2] * M[1][3] - M[1][2] * M[0][3];
|
||||
|
||||
c[0] = M[2][0] * M[3][1] - M[3][0] * M[2][1];
|
||||
c[1] = M[2][0] * M[3][2] - M[3][0] * M[2][2];
|
||||
c[2] = M[2][0] * M[3][3] - M[3][0] * M[2][3];
|
||||
c[3] = M[2][1] * M[3][2] - M[3][1] * M[2][2];
|
||||
c[4] = M[2][1] * M[3][3] - M[3][1] * M[2][3];
|
||||
c[5] = M[2][2] * M[3][3] - M[3][2] * M[2][3];
|
||||
|
||||
/* Assumes it is invertible */
|
||||
float idet = 1.0f / (s[0] * c[5] - s[1] * c[4] + s[2] * c[3] + s[3] * c[2] - s[4] * c[1] + s[5] * c[0]);
|
||||
|
||||
T[0][0] = (M[1][1] * c[5] - M[1][2] * c[4] + M[1][3] * c[3]) * idet;
|
||||
T[0][1] = (-M[0][1] * c[5] + M[0][2] * c[4] - M[0][3] * c[3]) * idet;
|
||||
T[0][2] = (M[3][1] * s[5] - M[3][2] * s[4] + M[3][3] * s[3]) * idet;
|
||||
T[0][3] = (-M[2][1] * s[5] + M[2][2] * s[4] - M[2][3] * s[3]) * idet;
|
||||
|
||||
T[1][0] = (-M[1][0] * c[5] + M[1][2] * c[2] - M[1][3] * c[1]) * idet;
|
||||
T[1][1] = (M[0][0] * c[5] - M[0][2] * c[2] + M[0][3] * c[1]) * idet;
|
||||
T[1][2] = (-M[3][0] * s[5] + M[3][2] * s[2] - M[3][3] * s[1]) * idet;
|
||||
T[1][3] = (M[2][0] * s[5] - M[2][2] * s[2] + M[2][3] * s[1]) * idet;
|
||||
|
||||
T[2][0] = (M[1][0] * c[4] - M[1][1] * c[2] + M[1][3] * c[0]) * idet;
|
||||
T[2][1] = (-M[0][0] * c[4] + M[0][1] * c[2] - M[0][3] * c[0]) * idet;
|
||||
T[2][2] = (M[3][0] * s[4] - M[3][1] * s[2] + M[3][3] * s[0]) * idet;
|
||||
T[2][3] = (-M[2][0] * s[4] + M[2][1] * s[2] - M[2][3] * s[0]) * idet;
|
||||
|
||||
T[3][0] = (-M[1][0] * c[3] + M[1][1] * c[1] - M[1][2] * c[0]) * idet;
|
||||
T[3][1] = (M[0][0] * c[3] - M[0][1] * c[1] + M[0][2] * c[0]) * idet;
|
||||
T[3][2] = (-M[3][0] * s[3] + M[3][1] * s[1] - M[3][2] * s[0]) * idet;
|
||||
T[3][3] = (M[2][0] * s[3] - M[2][1] * s[1] + M[2][2] * s[0]) * idet;
|
||||
}
|
||||
static inline void mat4x4_orthonormalize(mat4x4 R, mat4x4 M) {
|
||||
mat4x4_dup(R, M);
|
||||
float s = 1.;
|
||||
vec3 h;
|
||||
|
||||
vec3_norm(R[2], R[2]);
|
||||
|
||||
s = vec3_mul_inner(R[1], R[2]);
|
||||
vec3_scale(h, R[2], s);
|
||||
vec3_sub(R[1], R[1], h);
|
||||
vec3_norm(R[2], R[2]);
|
||||
|
||||
s = vec3_mul_inner(R[1], R[2]);
|
||||
vec3_scale(h, R[2], s);
|
||||
vec3_sub(R[1], R[1], h);
|
||||
vec3_norm(R[1], R[1]);
|
||||
|
||||
s = vec3_mul_inner(R[0], R[1]);
|
||||
vec3_scale(h, R[1], s);
|
||||
vec3_sub(R[0], R[0], h);
|
||||
vec3_norm(R[0], R[0]);
|
||||
}
|
||||
|
||||
static inline void mat4x4_frustum(mat4x4 M, float l, float r, float b, float t, float n, float f) {
|
||||
M[0][0] = 2.f * n / (r - l);
|
||||
M[0][1] = M[0][2] = M[0][3] = 0.f;
|
||||
|
||||
M[1][1] = 2.f * n / (t - b);
|
||||
M[1][0] = M[1][2] = M[1][3] = 0.f;
|
||||
|
||||
M[2][0] = (r + l) / (r - l);
|
||||
M[2][1] = (t + b) / (t - b);
|
||||
M[2][2] = -(f + n) / (f - n);
|
||||
M[2][3] = -1.f;
|
||||
|
||||
M[3][2] = -2.f * (f * n) / (f - n);
|
||||
M[3][0] = M[3][1] = M[3][3] = 0.f;
|
||||
}
|
||||
static inline void mat4x4_ortho(mat4x4 M, float l, float r, float b, float t, float n, float f) {
|
||||
M[0][0] = 2.f / (r - l);
|
||||
M[0][1] = M[0][2] = M[0][3] = 0.f;
|
||||
|
||||
M[1][1] = 2.f / (t - b);
|
||||
M[1][0] = M[1][2] = M[1][3] = 0.f;
|
||||
|
||||
M[2][2] = -2.f / (f - n);
|
||||
M[2][0] = M[2][1] = M[2][3] = 0.f;
|
||||
|
||||
M[3][0] = -(r + l) / (r - l);
|
||||
M[3][1] = -(t + b) / (t - b);
|
||||
M[3][2] = -(f + n) / (f - n);
|
||||
M[3][3] = 1.f;
|
||||
}
|
||||
static inline void mat4x4_perspective(mat4x4 m, float y_fov, float aspect, float n, float f) {
|
||||
/* NOTE: Degrees are an unhandy unit to work with.
|
||||
* linmath.h uses radians for everything! */
|
||||
float const a = (float)(1.f / tan(y_fov / 2.f));
|
||||
|
||||
m[0][0] = a / aspect;
|
||||
m[0][1] = 0.f;
|
||||
m[0][2] = 0.f;
|
||||
m[0][3] = 0.f;
|
||||
|
||||
m[1][0] = 0.f;
|
||||
m[1][1] = a;
|
||||
m[1][2] = 0.f;
|
||||
m[1][3] = 0.f;
|
||||
|
||||
m[2][0] = 0.f;
|
||||
m[2][1] = 0.f;
|
||||
m[2][2] = -((f + n) / (f - n));
|
||||
m[2][3] = -1.f;
|
||||
|
||||
m[3][0] = 0.f;
|
||||
m[3][1] = 0.f;
|
||||
m[3][2] = -((2.f * f * n) / (f - n));
|
||||
m[3][3] = 0.f;
|
||||
}
|
||||
static inline void mat4x4_look_at(mat4x4 m, vec3 eye, vec3 center, vec3 up) {
|
||||
/* Adapted from Android's OpenGL Matrix.java. */
|
||||
/* See the OpenGL GLUT documentation for gluLookAt for a description */
|
||||
/* of the algorithm. We implement it in a straightforward way: */
|
||||
|
||||
/* TODO: The negation of of can be spared by swapping the order of
|
||||
* operands in the following cross products in the right way. */
|
||||
vec3 f;
|
||||
vec3_sub(f, center, eye);
|
||||
vec3_norm(f, f);
|
||||
|
||||
vec3 s;
|
||||
vec3_mul_cross(s, f, up);
|
||||
vec3_norm(s, s);
|
||||
|
||||
vec3 t;
|
||||
vec3_mul_cross(t, s, f);
|
||||
|
||||
m[0][0] = s[0];
|
||||
m[0][1] = t[0];
|
||||
m[0][2] = -f[0];
|
||||
m[0][3] = 0.f;
|
||||
|
||||
m[1][0] = s[1];
|
||||
m[1][1] = t[1];
|
||||
m[1][2] = -f[1];
|
||||
m[1][3] = 0.f;
|
||||
|
||||
m[2][0] = s[2];
|
||||
m[2][1] = t[2];
|
||||
m[2][2] = -f[2];
|
||||
m[2][3] = 0.f;
|
||||
|
||||
m[3][0] = 0.f;
|
||||
m[3][1] = 0.f;
|
||||
m[3][2] = 0.f;
|
||||
m[3][3] = 1.f;
|
||||
|
||||
mat4x4_translate_in_place(m, -eye[0], -eye[1], -eye[2]);
|
||||
}
|
||||
|
||||
typedef float quat[4];
|
||||
static inline void quat_identity(quat q) {
|
||||
q[0] = q[1] = q[2] = 0.f;
|
||||
q[3] = 1.f;
|
||||
}
|
||||
static inline void quat_add(quat r, quat a, quat b) {
|
||||
int i;
|
||||
for (i = 0; i < 4; ++i) r[i] = a[i] + b[i];
|
||||
}
|
||||
static inline void quat_sub(quat r, quat a, quat b) {
|
||||
int i;
|
||||
for (i = 0; i < 4; ++i) r[i] = a[i] - b[i];
|
||||
}
|
||||
static inline void quat_mul(quat r, quat p, quat q) {
|
||||
vec3 w;
|
||||
vec3_mul_cross(r, p, q);
|
||||
vec3_scale(w, p, q[3]);
|
||||
vec3_add(r, r, w);
|
||||
vec3_scale(w, q, p[3]);
|
||||
vec3_add(r, r, w);
|
||||
r[3] = p[3] * q[3] - vec3_mul_inner(p, q);
|
||||
}
|
||||
static inline void quat_scale(quat r, quat v, float s) {
|
||||
int i;
|
||||
for (i = 0; i < 4; ++i) r[i] = v[i] * s;
|
||||
}
|
||||
static inline float quat_inner_product(quat a, quat b) {
|
||||
float p = 0.f;
|
||||
int i;
|
||||
for (i = 0; i < 4; ++i) p += b[i] * a[i];
|
||||
return p;
|
||||
}
|
||||
static inline void quat_conj(quat r, quat q) {
|
||||
int i;
|
||||
for (i = 0; i < 3; ++i) r[i] = -q[i];
|
||||
r[3] = q[3];
|
||||
}
|
||||
#define quat_norm vec4_norm
|
||||
static inline void quat_mul_vec3(vec3 r, quat q, vec3 v) {
|
||||
quat v_ = {v[0], v[1], v[2], 0.f};
|
||||
|
||||
quat_conj(r, q);
|
||||
quat_norm(r, r);
|
||||
quat_mul(r, v_, r);
|
||||
quat_mul(r, q, r);
|
||||
}
|
||||
static inline void mat4x4_from_quat(mat4x4 M, quat q) {
|
||||
float a = q[3];
|
||||
float b = q[0];
|
||||
float c = q[1];
|
||||
float d = q[2];
|
||||
float a2 = a * a;
|
||||
float b2 = b * b;
|
||||
float c2 = c * c;
|
||||
float d2 = d * d;
|
||||
|
||||
M[0][0] = a2 + b2 - c2 - d2;
|
||||
M[0][1] = 2.f * (b * c + a * d);
|
||||
M[0][2] = 2.f * (b * d - a * c);
|
||||
M[0][3] = 0.f;
|
||||
|
||||
M[1][0] = 2 * (b * c - a * d);
|
||||
M[1][1] = a2 - b2 + c2 - d2;
|
||||
M[1][2] = 2.f * (c * d + a * b);
|
||||
M[1][3] = 0.f;
|
||||
|
||||
M[2][0] = 2.f * (b * d + a * c);
|
||||
M[2][1] = 2.f * (c * d - a * b);
|
||||
M[2][2] = a2 - b2 - c2 + d2;
|
||||
M[2][3] = 0.f;
|
||||
|
||||
M[3][0] = M[3][1] = M[3][2] = 0.f;
|
||||
M[3][3] = 1.f;
|
||||
}
|
||||
|
||||
static inline void mat4x4o_mul_quat(mat4x4 R, mat4x4 M, quat q) {
|
||||
/* XXX: The way this is written only works for othogonal matrices. */
|
||||
/* TODO: Take care of non-orthogonal case. */
|
||||
quat_mul_vec3(R[0], q, M[0]);
|
||||
quat_mul_vec3(R[1], q, M[1]);
|
||||
quat_mul_vec3(R[2], q, M[2]);
|
||||
|
||||
R[3][0] = R[3][1] = R[3][2] = 0.f;
|
||||
R[3][3] = 1.f;
|
||||
}
|
||||
static inline void quat_from_mat4x4(quat q, mat4x4 M) {
|
||||
float r = 0.f;
|
||||
int i;
|
||||
|
||||
int perm[] = {0, 1, 2, 0, 1};
|
||||
int *p = perm;
|
||||
|
||||
for (i = 0; i < 3; i++) {
|
||||
float m = M[i][i];
|
||||
if (m < r) continue;
|
||||
m = r;
|
||||
p = &perm[i];
|
||||
}
|
||||
|
||||
r = sqrtf(1.f + M[p[0]][p[0]] - M[p[1]][p[1]] - M[p[2]][p[2]]);
|
||||
|
||||
if (r < 1e-6) {
|
||||
q[0] = 1.f;
|
||||
q[1] = q[2] = q[3] = 0.f;
|
||||
return;
|
||||
}
|
||||
|
||||
q[0] = r / 2.f;
|
||||
q[1] = (M[p[0]][p[1]] - M[p[1]][p[0]]) / (2.f * r);
|
||||
q[2] = (M[p[2]][p[0]] - M[p[0]][p[2]]) / (2.f * r);
|
||||
q[3] = (M[p[2]][p[1]] - M[p[1]][p[2]]) / (2.f * r);
|
||||
}
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,454 @@
|
||||
/* 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 "VulkanBaseApp.h"
|
||||
|
||||
#include <iostream>
|
||||
#include <iomanip>
|
||||
#include <chrono>
|
||||
#include <algorithm>
|
||||
#include "linmath.h"
|
||||
#include "hip_runtime.h"
|
||||
#include "SineWaveSimulation.h"
|
||||
|
||||
|
||||
typedef float vec2[2];
|
||||
std::string execution_path;
|
||||
|
||||
#ifndef NDEBUG
|
||||
#define ENABLE_VALIDATION (false)
|
||||
#else
|
||||
#define ENABLE_VALIDATION (true)
|
||||
#endif
|
||||
|
||||
class VulkanCudaSineWave : public VulkanBaseApp
|
||||
{
|
||||
|
||||
typedef struct UniformBufferObject_st {
|
||||
mat4x4 modelViewProj;
|
||||
} UniformBufferObject;
|
||||
|
||||
VkBuffer m_heightBuffer, m_xyBuffer, m_indexBuffer;
|
||||
VkDeviceMemory m_heightMemory, m_xyMemory, m_indexMemory;
|
||||
UniformBufferObject m_ubo;
|
||||
VkSemaphore m_vkWaitSemaphore, m_vkSignalSemaphore;
|
||||
SineWaveSimulation m_sim;
|
||||
hipStream_t m_stream;
|
||||
hipExternalSemaphore_t m_cudaWaitSemaphore, m_cudaSignalSemaphore;
|
||||
hipExternalMemory_t m_cudaVertMem;
|
||||
float *m_cudaHeightMap;
|
||||
using chrono_tp = std::chrono::time_point<std::chrono::high_resolution_clock>;
|
||||
chrono_tp m_lastTime;
|
||||
size_t m_lastFrame;
|
||||
public:
|
||||
VulkanCudaSineWave(size_t width, size_t height) :
|
||||
VulkanBaseApp("vulkanCudaSineWave", ENABLE_VALIDATION),
|
||||
m_heightBuffer(VK_NULL_HANDLE),
|
||||
m_xyBuffer(VK_NULL_HANDLE),
|
||||
m_indexBuffer(VK_NULL_HANDLE),
|
||||
m_heightMemory(VK_NULL_HANDLE),
|
||||
m_xyMemory(VK_NULL_HANDLE),
|
||||
m_indexMemory(VK_NULL_HANDLE),
|
||||
m_ubo(),
|
||||
m_sim(width, height),
|
||||
m_stream(0),
|
||||
m_vkWaitSemaphore(VK_NULL_HANDLE),
|
||||
m_vkSignalSemaphore(VK_NULL_HANDLE),
|
||||
m_cudaWaitSemaphore(),
|
||||
m_cudaSignalSemaphore(),
|
||||
m_cudaVertMem(),
|
||||
m_cudaHeightMap(nullptr),
|
||||
m_lastFrame(0) {
|
||||
// Our index buffer can only index 32-bits of the vertex buffer
|
||||
if ((width * height) > (1ULL << 32ULL)) {
|
||||
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 };
|
||||
GetModuleFileName(NULL, buffer, MAX_PATH);
|
||||
std::string str3 = std::string(buffer);
|
||||
std::string str1 = "vert.spv" ; //sdkFindFilePath("sinewave.vert", execution_path.c_str());
|
||||
char* vertex_shader_path = strdup(str1.c_str());
|
||||
|
||||
std::string str2 = "frag.spv" ; //sdkFindFilePath("sinewave.frag", execution_path.c_str());
|
||||
char* fragment_shader_path = strdup(str2.c_str());
|
||||
|
||||
m_shaderFiles.push_back(std::make_pair(VK_SHADER_STAGE_VERTEX_BIT, vertex_shader_path));
|
||||
m_shaderFiles.push_back(std::make_pair(VK_SHADER_STAGE_FRAGMENT_BIT, fragment_shader_path));
|
||||
|
||||
}
|
||||
~VulkanCudaSineWave() {
|
||||
// Make sure there's no pending work before we start tearing down
|
||||
checkHIPErrors(hipStreamSynchronize(m_stream));
|
||||
|
||||
if (m_vkSignalSemaphore != VK_NULL_HANDLE) {
|
||||
checkHIPErrors(hipDestroyExternalSemaphore(m_cudaSignalSemaphore));
|
||||
vkDestroySemaphore(m_device, m_vkSignalSemaphore, nullptr);
|
||||
}
|
||||
if (m_vkWaitSemaphore != VK_NULL_HANDLE) {
|
||||
checkHIPErrors(hipDestroyExternalSemaphore(m_cudaWaitSemaphore));
|
||||
vkDestroySemaphore(m_device, m_vkWaitSemaphore, nullptr);
|
||||
}
|
||||
|
||||
if (m_xyBuffer != VK_NULL_HANDLE) {
|
||||
vkDestroyBuffer(m_device, m_xyBuffer, nullptr);
|
||||
}
|
||||
if (m_xyMemory != VK_NULL_HANDLE) {
|
||||
vkFreeMemory(m_device, m_xyMemory, nullptr);
|
||||
}
|
||||
|
||||
if (m_heightBuffer != VK_NULL_HANDLE) {
|
||||
vkDestroyBuffer(m_device, m_heightBuffer, nullptr);
|
||||
}
|
||||
if (m_heightMemory != VK_NULL_HANDLE) {
|
||||
vkFreeMemory(m_device, m_heightMemory, nullptr);
|
||||
}
|
||||
if (m_cudaHeightMap) {
|
||||
checkHIPErrors(hipDestroyExternalMemory(m_cudaVertMem));
|
||||
}
|
||||
|
||||
if (m_indexBuffer != VK_NULL_HANDLE) {
|
||||
vkDestroyBuffer(m_device, m_indexBuffer, nullptr);
|
||||
}
|
||||
if (m_indexMemory != VK_NULL_HANDLE) {
|
||||
vkFreeMemory(m_device, m_indexMemory, nullptr);
|
||||
}
|
||||
}
|
||||
|
||||
void fillRenderingCommandBuffer(VkCommandBuffer& commandBuffer) {
|
||||
VkBuffer vertexBuffers[] = { m_heightBuffer, m_xyBuffer };
|
||||
VkDeviceSize offsets[] = { 0, 0 };
|
||||
vkCmdBindVertexBuffers(commandBuffer, 0, sizeof(vertexBuffers) / sizeof(vertexBuffers[0]), vertexBuffers, offsets);
|
||||
vkCmdBindIndexBuffer(commandBuffer, m_indexBuffer, 0, VK_INDEX_TYPE_UINT32);
|
||||
vkCmdDrawIndexed(commandBuffer, (uint32_t)((m_sim.getWidth() - 1) * (m_sim.getHeight() - 1) * 6), 1, 0, 0, 0);
|
||||
}
|
||||
|
||||
void getVertexDescriptions(std::vector<VkVertexInputBindingDescription>& bindingDesc, std::vector<VkVertexInputAttributeDescription>& attribDesc) {
|
||||
bindingDesc.resize(2);
|
||||
attribDesc.resize(2);
|
||||
|
||||
bindingDesc[0].binding = 0;
|
||||
bindingDesc[0].stride = sizeof(float);
|
||||
bindingDesc[0].inputRate = VK_VERTEX_INPUT_RATE_VERTEX;
|
||||
|
||||
bindingDesc[1].binding = 1;
|
||||
bindingDesc[1].stride = sizeof(vec2);
|
||||
bindingDesc[1].inputRate = VK_VERTEX_INPUT_RATE_VERTEX;
|
||||
|
||||
attribDesc[0].binding = 0;
|
||||
attribDesc[0].location = 0;
|
||||
attribDesc[0].format = VK_FORMAT_R32_SFLOAT;
|
||||
attribDesc[0].offset = 0;
|
||||
|
||||
attribDesc[1].binding = 1;
|
||||
attribDesc[1].location = 1;
|
||||
attribDesc[1].format = VK_FORMAT_R32G32_SFLOAT;
|
||||
attribDesc[1].offset = 0;
|
||||
}
|
||||
|
||||
void getAssemblyStateInfo(VkPipelineInputAssemblyStateCreateInfo& info) {
|
||||
info.sType = VK_STRUCTURE_TYPE_PIPELINE_INPUT_ASSEMBLY_STATE_CREATE_INFO;
|
||||
info.topology = VK_PRIMITIVE_TOPOLOGY_TRIANGLE_LIST;
|
||||
info.primitiveRestartEnable = VK_FALSE;
|
||||
}
|
||||
|
||||
void getWaitFrameSemaphores(std::vector<VkSemaphore>& wait, std::vector< VkPipelineStageFlags>& waitStages) const {
|
||||
if (m_currentFrame != 0) {
|
||||
// Have vulkan wait until cuda is done with the vertex buffer before rendering
|
||||
// We don't do this on the first frame, as the wait semaphore hasn't been initialized yet
|
||||
wait.push_back(m_vkWaitSemaphore);
|
||||
// We want to wait until all the pipeline commands are complete before letting cuda work
|
||||
waitStages.push_back(VK_PIPELINE_STAGE_ALL_COMMANDS_BIT);
|
||||
}
|
||||
}
|
||||
|
||||
void getSignalFrameSemaphores(std::vector<VkSemaphore>& signal) const {
|
||||
// Add this semaphore for vulkan to signal once the vertex buffer is ready for cuda to modify
|
||||
signal.push_back(m_vkSignalSemaphore);
|
||||
}
|
||||
|
||||
void initVulkanApp() {
|
||||
int cuda_device = -1;
|
||||
|
||||
// Select cuda device where vulkan is running.
|
||||
cuda_device = m_sim.initCuda(m_vkDeviceUUID, VK_UUID_SIZE);
|
||||
if (cuda_device == -1)
|
||||
{
|
||||
printf("Error: No CUDA-Vulkan interop capable device found\n");
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
|
||||
m_sim.initCudaLaunchConfig(cuda_device);
|
||||
|
||||
// Create the cuda stream we'll be using
|
||||
checkHIPErrors(hipStreamCreateWithFlags(&m_stream, hipStreamNonBlocking));
|
||||
|
||||
const size_t nVerts = m_sim.getWidth() * m_sim.getHeight();
|
||||
const size_t nInds = (m_sim.getWidth() - 1) * (m_sim.getHeight() - 1) * 6;
|
||||
|
||||
// Create the height map cuda will write to
|
||||
createExternalBuffer(nVerts * sizeof(float),
|
||||
VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_VERTEX_BUFFER_BIT,
|
||||
VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT,
|
||||
getDefaultMemHandleType(),
|
||||
m_heightBuffer, m_heightMemory);
|
||||
|
||||
// Create the vertex buffer that will hold the xy coordinates for the grid
|
||||
createBuffer(nVerts * sizeof(vec2),
|
||||
VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_VERTEX_BUFFER_BIT,
|
||||
VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT,
|
||||
m_xyBuffer, m_xyMemory);
|
||||
|
||||
// Create the index buffer that references from both buffers above
|
||||
createBuffer(nInds * sizeof(uint32_t),
|
||||
VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_INDEX_BUFFER_BIT,
|
||||
VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT,
|
||||
m_indexBuffer, m_indexMemory);
|
||||
|
||||
// Import the height map into cuda and retrieve a device pointer to use
|
||||
importHipExternalMemory((void **)&m_cudaHeightMap, m_cudaVertMem, m_heightMemory, nVerts * sizeof(*m_cudaHeightMap), getDefaultMemHandleType());
|
||||
// Set the height map to use in the simulation
|
||||
m_sim.initSimulation(m_cudaHeightMap);
|
||||
|
||||
{
|
||||
// Set up the initial values for the vertex buffers with Vulkan
|
||||
void *stagingBase;
|
||||
VkBuffer stagingBuffer;
|
||||
VkDeviceMemory stagingMemory;
|
||||
VkDeviceSize stagingSz = std::max(nVerts * sizeof(vec2), nInds * sizeof(uint32_t));
|
||||
createBuffer(stagingSz, VK_BUFFER_USAGE_TRANSFER_SRC_BIT,
|
||||
VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT, stagingBuffer, stagingMemory);
|
||||
|
||||
vkMapMemory(m_device, stagingMemory, 0, stagingSz, 0, &stagingBase);
|
||||
|
||||
memset(stagingBase, 0, nVerts * sizeof(float));
|
||||
copyBuffer(m_heightBuffer, stagingBuffer, nVerts * sizeof(float));
|
||||
|
||||
for (size_t y = 0; y < m_sim.getHeight(); y++) {
|
||||
for (size_t x = 0; x < m_sim.getWidth(); x++) {
|
||||
vec2 *stagedVert = (vec2 *)stagingBase;
|
||||
stagedVert[y * m_sim.getWidth() + x][0] = (2.0f * x) / (m_sim.getWidth() - 1) - 1;
|
||||
stagedVert[y * m_sim.getWidth() + x][1] = (2.0f * y) / (m_sim.getHeight() - 1) - 1;
|
||||
}
|
||||
}
|
||||
copyBuffer(m_xyBuffer, stagingBuffer, nVerts * sizeof(vec2));
|
||||
|
||||
{
|
||||
uint32_t *indices = (uint32_t *)stagingBase;
|
||||
for (size_t y = 0; y < m_sim.getHeight() - 1; y++) {
|
||||
for (size_t x = 0; x < m_sim.getWidth() - 1; x++) {
|
||||
indices[0] = (uint32_t)((y + 0) * m_sim.getWidth() + (x + 0));
|
||||
indices[1] = (uint32_t)((y + 1) * m_sim.getWidth() + (x + 0));
|
||||
indices[2] = (uint32_t)((y + 0) * m_sim.getWidth() + (x + 1));
|
||||
indices[3] = (uint32_t)((y + 1) * m_sim.getWidth() + (x + 0));
|
||||
indices[4] = (uint32_t)((y + 1) * m_sim.getWidth() + (x + 1));
|
||||
indices[5] = (uint32_t)((y + 0) * m_sim.getWidth() + (x + 1));
|
||||
indices += 6;
|
||||
}
|
||||
}
|
||||
}
|
||||
copyBuffer(m_indexBuffer, stagingBuffer, nInds * sizeof(uint32_t));
|
||||
|
||||
vkUnmapMemory(m_device, stagingMemory);
|
||||
vkDestroyBuffer(m_device, stagingBuffer, nullptr);
|
||||
vkFreeMemory(m_device, stagingMemory, nullptr);
|
||||
}
|
||||
|
||||
// Create the semaphore vulkan will signal when it's done with the vertex buffer
|
||||
createExternalSemaphore(m_vkSignalSemaphore, getDefaultSemaphoreHandleType());
|
||||
// Create the semaphore vulkan will wait for before using the vertex buffer
|
||||
createExternalSemaphore(m_vkWaitSemaphore, getDefaultSemaphoreHandleType());
|
||||
// Import the semaphore cuda will use -- vulkan's signal will be cuda's wait
|
||||
importCudaExternalSemaphore(m_cudaWaitSemaphore, m_vkSignalSemaphore, getDefaultSemaphoreHandleType());
|
||||
// Import the semaphore cuda will use -- cuda's signal will be vulkan's wait
|
||||
importCudaExternalSemaphore(m_cudaSignalSemaphore, m_vkWaitSemaphore, getDefaultSemaphoreHandleType());
|
||||
|
||||
}
|
||||
|
||||
void importHipExternalMemory(void **cudaPtr, hipExternalMemory_t& cudaMem, VkDeviceMemory& vkMem, VkDeviceSize size, VkExternalMemoryHandleTypeFlagBits handleType) {
|
||||
hipExternalMemoryHandleDesc externalMemoryHandleDesc = {};
|
||||
|
||||
if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT) {
|
||||
externalMemoryHandleDesc.type = hipExternalMemoryHandleTypeOpaqueWin32;
|
||||
}
|
||||
else if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT) {
|
||||
externalMemoryHandleDesc.type = hipExternalMemoryHandleTypeOpaqueWin32Kmt;
|
||||
}
|
||||
else if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD_BIT) {
|
||||
externalMemoryHandleDesc.type = hipExternalMemoryHandleTypeOpaqueFd;
|
||||
}
|
||||
else {
|
||||
throw std::runtime_error("Unknown handle type requested!");
|
||||
}
|
||||
|
||||
externalMemoryHandleDesc.size = size;
|
||||
|
||||
#ifdef _WIN64
|
||||
externalMemoryHandleDesc.handle.win32.handle = (HANDLE)getMemHandle(vkMem, handleType);
|
||||
#else
|
||||
externalMemoryHandleDesc.handle.fd = (int)(uintptr_t)getMemHandle(vkMem, handleType);
|
||||
#endif
|
||||
checkHIPErrors(hipImportExternalMemory(&cudaMem, &externalMemoryHandleDesc));
|
||||
|
||||
hipExternalMemoryBufferDesc externalMemBufferDesc = {};
|
||||
externalMemBufferDesc.offset = 0;
|
||||
externalMemBufferDesc.size = size;
|
||||
externalMemBufferDesc.flags = 0;
|
||||
|
||||
checkHIPErrors(hipExternalMemoryGetMappedBuffer(cudaPtr, cudaMem, &externalMemBufferDesc));
|
||||
}
|
||||
|
||||
void importCudaExternalSemaphore(hipExternalSemaphore_t& cudaSem, VkSemaphore& vkSem, VkExternalSemaphoreHandleTypeFlagBits handleType) {
|
||||
hipExternalSemaphoreHandleDesc externalSemaphoreHandleDesc = {};
|
||||
|
||||
if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT) {
|
||||
externalSemaphoreHandleDesc.type = hipExternalSemaphoreHandleTypeOpaqueWin32;
|
||||
}
|
||||
else if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT) {
|
||||
externalSemaphoreHandleDesc.type = hipExternalSemaphoreHandleTypeOpaqueWin32Kmt;
|
||||
}
|
||||
else if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD_BIT) {
|
||||
externalSemaphoreHandleDesc.type = hipExternalSemaphoreHandleTypeOpaqueFd;
|
||||
}
|
||||
else {
|
||||
throw std::runtime_error("Unknown handle type requested!");
|
||||
}
|
||||
|
||||
#ifdef _WIN64
|
||||
externalSemaphoreHandleDesc.handle.win32.handle = (HANDLE)getSemaphoreHandle(vkSem, handleType);
|
||||
#else
|
||||
externalSemaphoreHandleDesc.handle.fd = (int)(uintptr_t)getSemaphoreHandle(vkSem, handleType);
|
||||
#endif
|
||||
|
||||
externalSemaphoreHandleDesc.flags = 0;
|
||||
|
||||
checkHIPErrors(hipImportExternalSemaphore(&cudaSem, &externalSemaphoreHandleDesc));
|
||||
}
|
||||
|
||||
VkDeviceSize getUniformSize() const {
|
||||
return sizeof(UniformBufferObject);
|
||||
}
|
||||
|
||||
void updateUniformBuffer(uint32_t imageIndex) {
|
||||
{
|
||||
mat4x4 view, proj;
|
||||
vec3 eye = { 1.75f, 1.75f, 1.25f };
|
||||
vec3 center = { 0.0f, 0.0f, -0.25f };
|
||||
vec3 up = { 0.0f, 0.0f, 1.0f };
|
||||
|
||||
mat4x4_perspective(proj, (float)degreesToRadians(45.0f), m_swapChainExtent.width / (float)m_swapChainExtent.height, 0.1f, 10.0f);
|
||||
proj[1][1] *= -1.0f; // Flip y axis
|
||||
|
||||
mat4x4_look_at(view, eye, center, up);
|
||||
mat4x4_mul(m_ubo.modelViewProj, proj, view);
|
||||
}
|
||||
|
||||
void *data;
|
||||
vkMapMemory(m_device, m_uniformMemory[imageIndex], 0, getUniformSize(), 0, &data);
|
||||
memcpy(data, &m_ubo, sizeof(m_ubo));
|
||||
vkUnmapMemory(m_device, m_uniformMemory[imageIndex]);
|
||||
}
|
||||
|
||||
std::vector<const char *> getRequiredExtensions() const {
|
||||
std::vector<const char *> extensions;
|
||||
extensions.push_back(VK_KHR_EXTERNAL_MEMORY_CAPABILITIES_EXTENSION_NAME);
|
||||
extensions.push_back(VK_KHR_EXTERNAL_SEMAPHORE_CAPABILITIES_EXTENSION_NAME);
|
||||
return extensions;
|
||||
}
|
||||
|
||||
std::vector<const char *> getRequiredDeviceExtensions() const {
|
||||
std::vector<const char *> extensions;
|
||||
extensions.push_back(VK_KHR_EXTERNAL_MEMORY_EXTENSION_NAME);
|
||||
extensions.push_back(VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME);
|
||||
#ifdef _WIN64
|
||||
extensions.push_back(VK_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_NAME);
|
||||
extensions.push_back(VK_KHR_EXTERNAL_SEMAPHORE_WIN32_EXTENSION_NAME);
|
||||
#else
|
||||
extensions.push_back(VK_KHR_EXTERNAL_MEMORY_FD_EXTENSION_NAME);
|
||||
extensions.push_back(VK_KHR_EXTERNAL_SEMAPHORE_FD_EXTENSION_NAME);
|
||||
#endif /* _WIN64 */
|
||||
return extensions;
|
||||
}
|
||||
|
||||
void drawFrame() {
|
||||
static chrono_tp startTime = std::chrono::high_resolution_clock::now();
|
||||
|
||||
chrono_tp currentTime = std::chrono::high_resolution_clock::now();
|
||||
float time = std::chrono::duration<float, std::chrono::seconds::period>(currentTime - startTime).count();
|
||||
|
||||
if (m_currentFrame == 0) {
|
||||
m_lastTime = startTime;
|
||||
}
|
||||
|
||||
float frame_time = std::chrono::duration<float, std::chrono::seconds::period>(currentTime - m_lastTime).count();
|
||||
|
||||
hipExternalSemaphoreWaitParams waitParams = {};
|
||||
waitParams.flags = 0;
|
||||
waitParams.params.fence.value = 1;
|
||||
|
||||
hipExternalSemaphoreSignalParams signalParams = {};
|
||||
signalParams.flags = 0;
|
||||
signalParams.params.fence.value = 0;
|
||||
|
||||
|
||||
// Have vulkan draw the current frame...
|
||||
VulkanBaseApp::drawFrame();
|
||||
// Wait for vulkan to complete it's work
|
||||
checkHIPErrors(hipWaitExternalSemaphoresAsync(&m_cudaWaitSemaphore, &waitParams, 1, m_stream));
|
||||
// Now step the simulation
|
||||
m_sim.stepSimulation(time, m_stream);
|
||||
|
||||
// Signal vulkan to continue with the updated buffers
|
||||
checkHIPErrors(hipSignalExternalSemaphoresAsync(&m_cudaSignalSemaphore, &signalParams, 1, m_stream));
|
||||
|
||||
// Output a naive measurement of the frames per second every five seconds
|
||||
if (frame_time > 5) {
|
||||
std::cout << "Average FPS (over "
|
||||
<< std::fixed << std::setprecision(2) << frame_time
|
||||
<< " seconds): "
|
||||
<< std::fixed << std::setprecision(2)
|
||||
<< ((m_currentFrame - m_lastFrame) / frame_time)
|
||||
<< std::endl;
|
||||
m_lastFrame = m_currentFrame;
|
||||
m_lastTime = currentTime;
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
int main(int argc, char **argv)
|
||||
{
|
||||
execution_path = argv[0];
|
||||
VulkanCudaSineWave app((1ULL << 8ULL), (1ULL << 8ULL));
|
||||
app.init();
|
||||
app.mainLoop();
|
||||
return 0;
|
||||
}
|
||||
@@ -0,0 +1,38 @@
|
||||
/* 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.
|
||||
*/
|
||||
|
||||
|
||||
#version 450
|
||||
#extension GL_ARB_separate_shader_objects : enable
|
||||
|
||||
layout(location = 0) in vec3 fragColor;
|
||||
|
||||
layout(location = 0) out vec4 outColor;
|
||||
|
||||
void main() {
|
||||
outColor = vec4(fragColor, 1.0);
|
||||
}
|
||||
@@ -0,0 +1,43 @@
|
||||
/* 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.
|
||||
*/
|
||||
|
||||
#version 450
|
||||
#extension GL_ARB_separate_shader_objects : enable
|
||||
|
||||
layout(binding = 0) uniform UniformBufferObject {
|
||||
mat4 modelViewProj;
|
||||
} ubo;
|
||||
|
||||
layout(location = 0) in float height;
|
||||
layout(location = 1) in vec2 xyPos;
|
||||
|
||||
layout(location = 0) out vec3 fragColor;
|
||||
|
||||
void main() {
|
||||
gl_Position = ubo.modelViewProj * vec4(xyPos.xy, height, 1.0f);
|
||||
fragColor = vec3(0.0f, (height + 0.5f), 0.0f);
|
||||
}
|
||||
@@ -80,6 +80,15 @@ Tests fall in 5 categories and its file name prefix are as follows:
|
||||
- Multi Process tests (Prefix: MultiProc_\*API\*_\*Optional Scenario\*, example: MultiProc_hipIPCMemHandle_GetDataFromProc): These tests are multi process tests and will only run on linux. They are used to test HIP APIs in multi process environment
|
||||
- Performance tests(Prefix: Perf_\*Intent\*_\*Optional Scenario\*, example: Perf_DispatchLatenc y): Performance tests are used to get results of HIP APIs.
|
||||
|
||||
There is a special interface available for process isolation. ```hip::SpawnProc``` in ```hip_test_process.hh```. Using this interface test can spawn of process and place passing conditions on its return value or its output to stdout. This can be useful for testing printf tests.
|
||||
Sample Usage:
|
||||
```cpp
|
||||
hip::SpawnProc proc(<relative path of exe with test folder>, <optional bool value, if output is to be recorded>);
|
||||
REQUIRE(0 == proc.run()); // Test of return value of the proc
|
||||
REQUIRE(exepctedOutput == proc.getOutput()); // Test on expected output of the process
|
||||
```
|
||||
The process can be a standalone exe (see tests/catch/unit/printfExe for more information).
|
||||
|
||||
General Guidelines:
|
||||
- Do not use the catch2 tags. Tags wont be used for filtering
|
||||
- Add as many INFO() as you can in tests which prints state of the t est, this will help the debugger when the test fails (INFO macro only prints when the test fails)
|
||||
|
||||
@@ -36,6 +36,7 @@ target_link_libraries(UnitTests PRIVATE UnitDeviceTests
|
||||
OccupancyTest
|
||||
DeviceTest
|
||||
RTC
|
||||
printfTests
|
||||
TextureTest
|
||||
stdc++fs)
|
||||
|
||||
|
||||
@@ -105,7 +105,7 @@ bool TestContext::skipTest() const {
|
||||
return false;
|
||||
}
|
||||
|
||||
std::string TestContext::currentPath() { return fs::current_path().string(); }
|
||||
std::string TestContext::currentPath() const { return fs::current_path().string(); }
|
||||
|
||||
bool TestContext::parseJsonFile() {
|
||||
// Check if file exists
|
||||
|
||||
@@ -1,3 +1,25 @@
|
||||
/*
|
||||
Copyright (c) 2021 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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#include "hip_test_common.hh"
|
||||
#include <iostream>
|
||||
|
||||
@@ -1,3 +1,25 @@
|
||||
/*
|
||||
Copyright (c) 2021 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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <vector>
|
||||
@@ -72,7 +94,7 @@ class TestContext {
|
||||
bool skipTest() const;
|
||||
|
||||
const std::string& getCurrentTest() const { return current_test; }
|
||||
std::string currentPath();
|
||||
std::string currentPath() const;
|
||||
|
||||
TestContext(const TestContext&) = delete;
|
||||
void operator=(const TestContext&) = delete;
|
||||
|
||||
@@ -0,0 +1,94 @@
|
||||
/*
|
||||
Copyright (c) 2021 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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#include "hip_test_common.hh"
|
||||
|
||||
#include <string>
|
||||
#include <array>
|
||||
#include <cstdlib>
|
||||
#include <random>
|
||||
#include <fstream>
|
||||
#include <streambuf>
|
||||
|
||||
#if __has_include(<filesystem>)
|
||||
#include <filesystem>
|
||||
namespace fs = std::filesystem;
|
||||
#elif __has_include(<experimental/filesystem>)
|
||||
#include <experimental/filesystem>
|
||||
namespace fs = std::experimental::filesystem;
|
||||
#else
|
||||
#error "gg filesystem"
|
||||
#endif
|
||||
|
||||
namespace hip {
|
||||
class SpawnProc {
|
||||
std::string exeName;
|
||||
std::string resultStr;
|
||||
std::string tmpFileName;
|
||||
bool captureOutput;
|
||||
|
||||
std::string getRandomString(size_t len = 6) {
|
||||
std::random_device dev;
|
||||
std::mt19937 rng(dev());
|
||||
std::uniform_int_distribution<std::mt19937::result_type> dist(0, 25);
|
||||
|
||||
std::string res;
|
||||
for (size_t i = 0; i < len; i++) {
|
||||
res += 'a' + dist(rng);
|
||||
}
|
||||
return res;
|
||||
}
|
||||
|
||||
public:
|
||||
SpawnProc(std::string exeName_, bool captureOutput_ = false)
|
||||
: exeName(exeName_), captureOutput(captureOutput_) {
|
||||
auto dir = fs::path(TestContext::get().currentPath()).parent_path();
|
||||
dir /= exeName;
|
||||
exeName = dir.string();
|
||||
if (captureOutput) {
|
||||
auto path = fs::temp_directory_path();
|
||||
path /= getRandomString();
|
||||
tmpFileName = path.string();
|
||||
}
|
||||
}
|
||||
|
||||
int run() {
|
||||
std::string execCmd = exeName;
|
||||
if (captureOutput) {
|
||||
execCmd += " > ";
|
||||
execCmd += tmpFileName;
|
||||
}
|
||||
|
||||
auto res = std::system(execCmd.c_str());
|
||||
|
||||
if (captureOutput) {
|
||||
std::ifstream t(tmpFileName.c_str());
|
||||
resultStr =
|
||||
std::string((std::istreambuf_iterator<char>(t)), std::istreambuf_iterator<char>());
|
||||
}
|
||||
return res;
|
||||
}
|
||||
|
||||
std::string getOutput() { return resultStr; }
|
||||
};
|
||||
} // namespace hip
|
||||
@@ -25,4 +25,6 @@ add_subdirectory(event)
|
||||
add_subdirectory(occupancy)
|
||||
add_subdirectory(device)
|
||||
add_subdirectory(rtc)
|
||||
add_subdirectory(texture)
|
||||
add_subdirectory(printf)
|
||||
add_subdirectory(printfExe)
|
||||
add_subdirectory(texture)
|
||||
|
||||
@@ -127,7 +127,7 @@ TEST_CASE("Unit_unsafeAtomicAdd") {
|
||||
|
||||
HIP_CHECK(hipModuleUnload(module));
|
||||
|
||||
REQUIRE(fabs(res_f - (1000 * f_val)) <= 0.2f);
|
||||
REQUIRE(fabs(res_d - (1000 * d_val)) <= 0.2);
|
||||
REQUIRE(fabs((res_f/1000) - f_val) <= 0.2f);
|
||||
REQUIRE(fabs((res_d/1000) - d_val) <= 0.2);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -45,6 +45,9 @@ set(TEST_SRC
|
||||
hipMemset3DNegative.cc
|
||||
hipMemset3DRegressMultiThread.cc
|
||||
hipMallocManagedFlagsTst.cc
|
||||
hipMemPrefetchAsyncExtTsts.cc
|
||||
hipMemAdviseMmap.cc
|
||||
hipMallocManaged.cc
|
||||
)
|
||||
else()
|
||||
set(TEST_SRC
|
||||
@@ -90,6 +93,9 @@ set(TEST_SRC
|
||||
hipMemset3DNegative.cc
|
||||
hipMemset3DRegressMultiThread.cc
|
||||
hipMallocManagedFlagsTst.cc
|
||||
hipMemPrefetchAsyncExtTsts.cc
|
||||
hipMemAdviseMmap.cc
|
||||
hipMallocManaged.cc
|
||||
)
|
||||
endif()
|
||||
# Create shared lib of all tests
|
||||
|
||||
@@ -0,0 +1,344 @@
|
||||
/*
|
||||
Copyright (c) 2021 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 WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/*
|
||||
List of Test cases:
|
||||
1) Unit_hipMallocManaged_Basic
|
||||
2) Unit_hipMallocManaged_MultiSize
|
||||
3) Unit_hipMallocManaged_MultiKrnlHmmAccess
|
||||
4) Unit_hipMallocManaged_KrnlWth2MemTypes
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
|
||||
|
||||
|
||||
// Kernel functions
|
||||
__global__ void KrnlWth2MemTypes(int *Hmm, int *Dptr, size_t n) {
|
||||
size_t index = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
for (size_t i = index; i < n; i++) {
|
||||
Hmm[i] = Dptr[i] + 10;
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void KernelMul_MngdMem(int *Hmm, int *Dptr, size_t n) {
|
||||
size_t index = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
for (size_t i = index; i < n; i += stride) {
|
||||
Hmm[i] = Dptr[i] * 10;
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void KernelMulAdd_MngdMem(int *Hmm, size_t n) {
|
||||
size_t index = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
for (size_t i = index; i < n; i += stride) {
|
||||
Hmm[i] = Hmm[i] * 2 + 10;
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void KrnlWth2MemTypesC(unsigned char *Hmm, unsigned char *Dptr,
|
||||
size_t n) {
|
||||
size_t index = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
for (size_t i = index; i < n; i += stride) {
|
||||
Hmm[i] = Dptr[i] + 10;
|
||||
}
|
||||
}
|
||||
|
||||
// The following variable will be used to get the result of computation
|
||||
// from multiple threads
|
||||
static bool IfTestPassed = true;
|
||||
|
||||
static int HmmAttrPrint() {
|
||||
int managed = 0;
|
||||
INFO("The following are the attribute values related to HMM for"
|
||||
" device 0:\n");
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
|
||||
INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeConcurrentManagedAccess, 0));
|
||||
INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccess, 0));
|
||||
INFO("hipDeviceAttributePageableMemoryAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
|
||||
INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
|
||||
<< managed);
|
||||
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
|
||||
0));
|
||||
INFO("hipDeviceAttributeManagedMemory: " << managed);
|
||||
return managed;
|
||||
}
|
||||
|
||||
|
||||
static void LaunchKrnl4(size_t NumElms, int InitVal) {
|
||||
int *Hmm = NULL, *Dptr = NULL, blockSize = 64, DataMismatch = 0;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, (sizeof(int) * NumElms)));
|
||||
HIP_CHECK(hipMalloc(&Dptr, (sizeof(int) * NumElms)));
|
||||
int *Hstptr = reinterpret_cast<int*>(new int[NumElms]);
|
||||
for (size_t i = 0; i < NumElms; ++i) {
|
||||
Hstptr[i] = InitVal;
|
||||
}
|
||||
HIP_CHECK(hipMemcpy(Dptr, Hstptr, (NumElms * sizeof(int)),
|
||||
hipMemcpyHostToDevice));
|
||||
dim3 dimBlock(blockSize, 1, 1);
|
||||
dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
|
||||
KrnlWth2MemTypes<<<dimGrid, dimBlock, 0, strm>>>(Hmm, Dptr, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
for (size_t i = 0; i < NumElms; ++i) {
|
||||
if (Hmm[i] != (InitVal + 10)) {
|
||||
DataMismatch++;
|
||||
}
|
||||
}
|
||||
if (DataMismatch != 0) {
|
||||
INFO("Data Mismatch observed after the Kernel: KrnlWth2MemTypes!!\n");
|
||||
REQUIRE(false);
|
||||
}
|
||||
DataMismatch = 0;
|
||||
KernelMul_MngdMem<<<dimGrid, dimBlock, 0, strm>>>(Hmm, Dptr, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
// Verifying the result
|
||||
for (size_t i = 0; i < NumElms; ++i) {
|
||||
if (Hmm[i] != (InitVal * 10)) {
|
||||
DataMismatch++;
|
||||
}
|
||||
}
|
||||
if (DataMismatch != 0) {
|
||||
INFO("Data Mismatch observedafter the Kernel: KernelMul_MngdMem!!\n");
|
||||
REQUIRE(false);
|
||||
}
|
||||
DataMismatch = 0;
|
||||
KernelMulAdd_MngdMem<<<dimGrid, dimBlock, 0, strm>>>(Hmm, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
// Verifying the result
|
||||
|
||||
for (size_t i = 0; i < NumElms; ++i) {
|
||||
if (Hmm[i] != (InitVal * 10 * 2 + 10)) {
|
||||
DataMismatch++;
|
||||
}
|
||||
}
|
||||
if (DataMismatch != 0) {
|
||||
INFO("Data Mismatch observedafter the Kernel: KernelMul_MngdMem!!\n");
|
||||
REQUIRE(false);
|
||||
}
|
||||
delete[] Hstptr;
|
||||
}
|
||||
|
||||
|
||||
|
||||
static size_t N{4 * 1024 * 1024};
|
||||
static unsigned blocksPerCU{6};
|
||||
static unsigned threadsPerBlock{256};
|
||||
|
||||
/*
|
||||
This testcase verifies the hipMallocManaged basic scenario - supported on all devices
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipMallocManaged_Basic") {
|
||||
int numElements = (N < (64 * 1024 * 1024)) ? 64 * 1024 * 1024 : N;
|
||||
float *A, *B, *C;
|
||||
|
||||
HIP_CHECK(hipMallocManaged(&A, numElements*sizeof(float)));
|
||||
HIP_CHECK(hipMallocManaged(&B, numElements*sizeof(float)));
|
||||
HIP_CHECK(hipMallocManaged(&C, numElements*sizeof(float)));
|
||||
}
|
||||
|
||||
/*
|
||||
This testcase verifies the hipMallocManaged basic scenario - supported only on HMM enabled devices
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipMallocManaged_Advanced") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int numElements = (N < (64 * 1024 * 1024)) ? 64 * 1024 * 1024 : N;
|
||||
float *A, *B, *C;
|
||||
|
||||
HIP_CHECK(hipMallocManaged(&A, numElements*sizeof(float)));
|
||||
HIP_CHECK(hipMallocManaged(&B, numElements*sizeof(float)));
|
||||
HIP_CHECK(hipMallocManaged(&C, numElements*sizeof(float)));
|
||||
HipTest::setDefaultData(numElements, A, B, C);
|
||||
|
||||
hipDevice_t device = hipCpuDeviceId;
|
||||
|
||||
HIP_CHECK(hipMemAdvise(A, numElements*sizeof(float),
|
||||
hipMemAdviseSetReadMostly, device));
|
||||
HIP_CHECK(hipMemPrefetchAsync(A, numElements*sizeof(float), 0));
|
||||
HIP_CHECK(hipMemPrefetchAsync(B, numElements*sizeof(float), 0));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&device, sizeof(device),
|
||||
hipMemRangeAttributeLastPrefetchLocation,
|
||||
A, numElements*sizeof(float)));
|
||||
if (device != 0) {
|
||||
INFO("hipMemRangeGetAttribute error, device = " << device);
|
||||
}
|
||||
uint32_t read_only = 0xf;
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&read_only, sizeof(read_only),
|
||||
hipMemRangeAttributeReadMostly,
|
||||
A, numElements*sizeof(float)));
|
||||
if (read_only != 1) {
|
||||
SUCCEED("hipMemRangeGetAttribute error, read_only = " << read_only);
|
||||
}
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock,
|
||||
numElements);
|
||||
hipEvent_t event0, event1;
|
||||
HIP_CHECK(hipEventCreate(&event0));
|
||||
HIP_CHECK(hipEventCreate(&event1));
|
||||
HIP_CHECK(hipEventRecord(event0, 0));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, 0, static_cast<const float*>(A),
|
||||
static_cast<const float*>(B), C, numElements);
|
||||
HIP_CHECK(hipEventRecord(event1, 0));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
float time = 0.0f;
|
||||
HIP_CHECK(hipEventElapsedTime(&time, event0, event1));
|
||||
printf("Time %.3f ms\n", time);
|
||||
float maxError = 0.0f;
|
||||
HIP_CHECK(hipMemPrefetchAsync(B, numElements*sizeof(float),
|
||||
hipCpuDeviceId));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
device = 0;
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&device, sizeof(device),
|
||||
hipMemRangeAttributeLastPrefetchLocation,
|
||||
A, numElements*sizeof(float)));
|
||||
if (device != hipCpuDeviceId) {
|
||||
SUCCEED("hipMemRangeGetAttribute error device = " << device);
|
||||
}
|
||||
|
||||
for (int i = 0; i < numElements; i++) {
|
||||
maxError = fmax(maxError, fabs(B[i]-3.0f));
|
||||
}
|
||||
HIP_CHECK(hipFree(A));
|
||||
HIP_CHECK(hipFree(B));
|
||||
REQUIRE(maxError != 0.0f);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
// The following test case tests the behavior of kernel with a HMM memory and
|
||||
// hipMalloc memory
|
||||
|
||||
TEST_CASE("Unit_hipMallocManaged_KrnlWth2MemTypes") {
|
||||
IfTestPassed = true;
|
||||
int *Hmm = NULL, *Dptr = NULL, InitVal = 123;
|
||||
size_t NumElms = (1024 * 1024);
|
||||
int *Hptr = new int[NumElms], blockSize = 64, DataMismatch = 0;
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, sizeof(int) * NumElms));
|
||||
HIP_CHECK(hipMalloc(&Dptr, sizeof(int) * NumElms));
|
||||
for (size_t i = 0; i < NumElms; ++i) {
|
||||
Hmm[i] = 0;
|
||||
Hptr[i] = InitVal;
|
||||
}
|
||||
HIP_CHECK(hipMemcpy(Dptr, Hptr, sizeof(int) * NumElms,
|
||||
hipMemcpyHostToDevice));
|
||||
dim3 dimBlock(blockSize, 1, 1);
|
||||
dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
|
||||
KrnlWth2MemTypes<<<dimGrid, dimBlock, 0, strm>>>(Hmm, Dptr, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
// Verifying the results
|
||||
for (size_t k = 0; k < NumElms; ++k) {
|
||||
if (Hmm[k] != (InitVal + 10)) {
|
||||
DataMismatch++;
|
||||
}
|
||||
}
|
||||
if (DataMismatch != 0) {
|
||||
WARN("DataMismatch observed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
HIP_CHECK(hipFree(Dptr));
|
||||
delete[] Hptr;
|
||||
REQUIRE(IfTestPassed);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
// The following test case tests when the same Hmm memory is used for
|
||||
// launching multiple different kernels will results in any issue
|
||||
TEST_CASE("Unit_hipMallocManaged_MultiKrnlHmmAccess") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed) {
|
||||
int InitVal = 123, NumElms = (1024 * 1024);
|
||||
LaunchKrnl4(NumElms, InitVal);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
// The following test case allocation, host access, device access of HMM
|
||||
// memory from size 1 to 10KB
|
||||
|
||||
TEST_CASE("Unit_hipMallocManaged_MultiSize") {
|
||||
IfTestPassed = true;
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
unsigned char *Hmm1 = NULL, *Hmm2 = NULL;
|
||||
int InitVal = 100, blockSize = 64, DataMismatch = 0;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
dim3 dimBlock(blockSize, 1, 1);
|
||||
for (int i = 1; i < (1024*1024); ++i) {
|
||||
HIP_CHECK(hipMallocManaged(&Hmm1, i));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm2, i));
|
||||
for (int j = 0; j < i; ++j) {
|
||||
Hmm1[j] = InitVal;
|
||||
}
|
||||
dim3 dimGrid((i + blockSize -1)/blockSize, 1, 1);
|
||||
KrnlWth2MemTypesC<<<dimGrid, dimBlock, 0, strm>>>(Hmm2, Hmm1, i);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
// Verifying the results
|
||||
for (int k = 0; k < i; ++k) {
|
||||
if (Hmm2[k] != (InitVal + 10)) {
|
||||
DataMismatch++;
|
||||
}
|
||||
}
|
||||
if (DataMismatch != 0) {
|
||||
WARN("DataMismatch observed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
DataMismatch = 0;
|
||||
HIP_CHECK(hipFree(Hmm1));
|
||||
HIP_CHECK(hipFree(Hmm2));
|
||||
REQUIRE(IfTestPassed);
|
||||
}
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,87 @@
|
||||
/*
|
||||
Copyright (c) 2021 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 <hip_test_common.hh>
|
||||
#ifdef __linux__
|
||||
#include <fcntl.h>
|
||||
#include <sys/mman.h>
|
||||
#include <sys/stat.h>
|
||||
#include <unistd.h>
|
||||
#endif
|
||||
|
||||
|
||||
/* Test Description: The following test case tests the working of
|
||||
hipMemAdvise() works with mmap() memory
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipMemAdvise_MmapMem") {
|
||||
int managed = 0, NUM_ELMS = 212992, PageableMem = 0;
|
||||
INFO("The following are the attribute values related to HMM for"
|
||||
" device 0:\n");
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
|
||||
INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeConcurrentManagedAccess, 0));
|
||||
INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&PageableMem,
|
||||
hipDeviceAttributePageableMemoryAccess, 0));
|
||||
INFO("hipDeviceAttributePageableMemoryAccess: " << PageableMem);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
|
||||
INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
|
||||
<< managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
|
||||
0));
|
||||
INFO("hipDeviceAttributeManagedMemory: " << managed);
|
||||
if ((managed == 1) && (PageableMem == 1)) {
|
||||
#ifdef __linux__
|
||||
// For now this test is enabled only for linux platforms
|
||||
FILE *fptr;
|
||||
fptr = fopen("ForTest1.txt", "w");
|
||||
for (int m = 0; m < NUM_ELMS; ++m) {
|
||||
putw(m, fptr);
|
||||
}
|
||||
fclose(fptr);
|
||||
int fd = open("./ForTest1.txt", O_RDWR, S_IRUSR | S_IWUSR);
|
||||
struct stat sb;
|
||||
if (fstat(fd, &sb) == -1) {
|
||||
perror("couldn't get file size.\n");
|
||||
close(fd);
|
||||
REQUIRE(false);
|
||||
}
|
||||
void *MmpdFile = nullptr;
|
||||
MmpdFile = mmap(NULL, sb.st_size, PROT_READ | PROT_WRITE, MAP_SHARED, fd,
|
||||
0);
|
||||
if (MmpdFile == nullptr) {
|
||||
INFO("mmap() call failed!\n. Cant proceed with the test.");
|
||||
REQUIRE(false);
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(MmpdFile, sb.st_size, hipMemAdviseSetReadMostly, 0));
|
||||
HIP_CHECK(hipMemAdvise(MmpdFile, sb.st_size,
|
||||
hipMemAdviseSetPreferredLocation, 0));
|
||||
HIP_CHECK(hipMemAdvise(MmpdFile, sb.st_size, hipMemAdviseSetAccessedBy, 0));
|
||||
munmap(MmpdFile, sb.st_size);
|
||||
close(fd);
|
||||
#endif
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,419 @@
|
||||
/*
|
||||
Copyright (c) 2021 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 <hip_test_common.hh>
|
||||
// Kernel function
|
||||
__global__ void MemPrftchAsyncKernel(int* Hmm, const int* Hmm1, size_t N) {
|
||||
size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
|
||||
size_t stride = hipBlockDim_x * hipGridDim_x;
|
||||
for (size_t i = offset; i < N; i += stride) {
|
||||
Hmm[i] = Hmm1[i] * Hmm1[i];
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void MemPrftchAsyncKernel1(int* Hmm, size_t N) {
|
||||
size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
|
||||
size_t stride = hipBlockDim_x * hipGridDim_x;
|
||||
for (size_t i = offset; i < N; i += stride) {
|
||||
Hmm[i] = Hmm[i] * Hmm[i];
|
||||
}
|
||||
}
|
||||
|
||||
static int HmmAttrPrint() {
|
||||
int managed = 0;
|
||||
WARN("The following are the attribute values related to HMM for"
|
||||
" device 0:\n");
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
|
||||
WARN("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeConcurrentManagedAccess, 0));
|
||||
WARN("hipDeviceAttributeConcurrentManagedAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccess, 0));
|
||||
WARN("hipDeviceAttributePageableMemoryAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
|
||||
WARN("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
|
||||
<< managed);
|
||||
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
|
||||
0));
|
||||
WARN("hipDeviceAttributeManagedMemory: " << managed);
|
||||
return managed;
|
||||
}
|
||||
|
||||
static void ReleaseResource(int *Hmm, int *Hmm1, hipStream_t *strm) {
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
HIP_CHECK(hipFree(Hmm1));
|
||||
HIP_CHECK(hipStreamDestroy(*strm));
|
||||
}
|
||||
|
||||
/* The following test allocates a managed memory and prefetch it in
|
||||
one-to-all and all-to-one fahsion followed by kernel launch within available
|
||||
devices*/
|
||||
TEST_CASE("Unit_hipMemPrefetchAsyncOneToAll") {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
int *Hmm = nullptr, *Hmm1 = nullptr, NumDevs, MemSz = (4096 * 4);
|
||||
int InitVal = 123, NumElms = MemSz/4;
|
||||
bool IfTestPassed = true;
|
||||
HIP_CHECK(hipGetDeviceCount(&NumDevs));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MemSz));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm1, MemSz));
|
||||
for (int i = 0; i < NumElms; ++i) {
|
||||
Hmm1[i] = InitVal;
|
||||
}
|
||||
hipStream_t strm;
|
||||
for (int i = -1; i < NumDevs; ++i) {
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm1, MemSz, i, 0));
|
||||
for (int j = -1; j < NumDevs; ++j) {
|
||||
if (i == j) {
|
||||
continue;
|
||||
}
|
||||
if (j != -1) {
|
||||
HIP_CHECK(hipSetDevice(j));
|
||||
}
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
// Prefetching memory from i to j
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm1, MemSz, j, strm));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
MemPrftchAsyncKernel<<<(NumElms/32), 32, 0, strm>>>(Hmm, Hmm1, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
// Verifying the result
|
||||
for (int m = 0; m < NumElms; ++m) {
|
||||
if (Hmm[m] != (InitVal * InitVal)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
if (!IfTestPassed) {
|
||||
ReleaseResource(Hmm, Hmm1, &strm);
|
||||
INFO("Did not find expected value!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
// Resetting the values in Hmm
|
||||
HIP_CHECK(hipMemset(Hmm, 0, MemSz));
|
||||
// Prefetching memory from j to i
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm1, MemSz, i, strm));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
MemPrftchAsyncKernel<<<(NumElms/32), 32, 0, strm>>>(Hmm, Hmm1, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
// Verifying the result
|
||||
for (int m = 0; m < NumElms; ++m) {
|
||||
if (Hmm[m] != (InitVal * InitVal)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
if (!IfTestPassed) {
|
||||
ReleaseResource(Hmm, Hmm1, &strm);
|
||||
INFO("Did not find expected value!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
// Resetting the values in Hmm
|
||||
HIP_CHECK(hipMemset(Hmm, 0, MemSz));
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
}
|
||||
}
|
||||
// Releasing the resources in case all the scenarios passed
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
HIP_CHECK(hipFree(Hmm1));
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
/* Test Case Description: Allocate managed memory --> prefetch to gpu 0
|
||||
call hipMemAdvise() on the memory and apply the flags ReadMostly,
|
||||
AccessedBy, and PreferredLocation for gpus other than gpu 0 and verify
|
||||
the flags using hipMemGetAttribute()*/
|
||||
TEST_CASE("Unit_hipMemPrefetchAsyncAdviseFlgTst") {
|
||||
int NGpus = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&NGpus));
|
||||
if (NGpus >= 2) {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
int *Hmm = nullptr, MemSz = (4096 * 4), InitVal = 123;
|
||||
int Outpt = 9999, NumElms = MemSz/4;
|
||||
bool IfTestPassed = true;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MemSz));
|
||||
// Initializing the memory
|
||||
for (int i = 0; i < NumElms; ++i) {
|
||||
Hmm[i] = InitVal;
|
||||
}
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm, MemSz, 0, strm));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MemSz, hipMemAdviseSetReadMostly, 1));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&Outpt, sizeof(int),
|
||||
hipMemRangeAttributeReadMostly, Hmm, MemSz));
|
||||
if (Outpt != 1) {
|
||||
WARN("hipMemRangeAttributeReadMostly flag did not take effect"
|
||||
" as expected!!");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MemSz, hipMemAdviseSetAccessedBy, 1));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&Outpt, sizeof(int),
|
||||
hipMemRangeAttributeAccessedBy, Hmm, MemSz));
|
||||
if (Outpt != 1) {
|
||||
WARN("hipMemRangeAttributeAccessedBy flag did not take effect"
|
||||
" as expected!!");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MemSz, hipMemAdviseSetPreferredLocation,
|
||||
1));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&Outpt, sizeof(int),
|
||||
hipMemRangeAttributePreferredLocation, Hmm, MemSz));
|
||||
if (Outpt != 1) {
|
||||
WARN("hipMemRangeAttributePreferredLocation flag did not take effect"
|
||||
" as expected!!");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
REQUIRE(IfTestPassed);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
} else {
|
||||
SUCCEED("This test needs atleast 2 gpus, but the system has less than"
|
||||
" 2 gpus hence skipping the test");
|
||||
}
|
||||
}
|
||||
|
||||
/* Test Case description: Allocate managed memory --> set AccessedBy using
|
||||
hipMemAdvise() to gpu1 prefetch the memory to gpu 0 and then query for
|
||||
AccessedBy using hipMemGetAttribute() and validate if AccessedBy is still
|
||||
set to gpu1. Similar tests are done with ReadMostly and PreferredLocation
|
||||
flags */
|
||||
TEST_CASE("Unit_hipMemPrefetchAsyncAccsdByTst") {
|
||||
int NGpus = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&NGpus));
|
||||
if (NGpus >= 2) {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
int *Hmm = nullptr, MemSz = (4096 * 4), InitVal = 123, NumElms = MemSz/4;
|
||||
int Outpt = 9999;
|
||||
bool IfTestPassed = true;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MemSz));
|
||||
// Initializing the memory
|
||||
for (int i = 0; i < NumElms; ++i) {
|
||||
Hmm[i] = InitVal;
|
||||
}
|
||||
SECTION("Test AccessedBy with Prefetch") {
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MemSz, hipMemAdviseSetAccessedBy, 1));
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm, MemSz, 0, strm));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&Outpt, sizeof(int),
|
||||
hipMemRangeAttributeAccessedBy, Hmm, MemSz));
|
||||
if (Outpt != 1) {
|
||||
WARN("hipMemRangeAttributeAccessedBy flag did not take effect"
|
||||
" as expected!!");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
SECTION("Test ReadMostly with Prefetch") {
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MemSz, hipMemAdviseSetReadMostly, 1));
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm, MemSz, 0, strm));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
MemPrftchAsyncKernel1<<<(NumElms/32), 32, 0, strm>>>(Hmm, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&Outpt, sizeof(int),
|
||||
hipMemRangeAttributeReadMostly, Hmm, MemSz));
|
||||
if (Outpt != 1) {
|
||||
WARN("hipMemRangeAttributeReadMostly flag did not take effect"
|
||||
" as expected!!");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
// Verifying the results
|
||||
for (int i = 0; i < NumElms; ++i) {
|
||||
if (Hmm[i] != (InitVal * InitVal)) {
|
||||
WARN("Did not receive expected value!!");
|
||||
IfTestPassed = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
SECTION("Test PreferredLocation with Prefetch") {
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MemSz, hipMemAdviseSetPreferredLocation,
|
||||
1));
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm, MemSz, 0, strm));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
MemPrftchAsyncKernel1<<<(NumElms/32), 32, 0, strm>>>(Hmm, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&Outpt, sizeof(int),
|
||||
hipMemRangeAttributePreferredLocation, Hmm, MemSz));
|
||||
if (Outpt != 1) {
|
||||
WARN("hipMemRangeAttributePreferredLocation flag did not take effect"
|
||||
" as expected!!");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
// Verifying the results
|
||||
for (int i = 0; i < NumElms; ++i) {
|
||||
if (Hmm[i] != (InitVal * InitVal)) {
|
||||
WARN("Did not receive expected value!!");
|
||||
IfTestPassed = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
REQUIRE(IfTestPassed);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
} else {
|
||||
SUCCEED("This test needs atleast 2 gpus, but the system has less than"
|
||||
" 2 gpus hence skipping the test");
|
||||
}
|
||||
}
|
||||
|
||||
/*Test Case description: Negative testing with hipMemPrefetchAsync() api*/
|
||||
TEST_CASE("Unit_hipMemPrefetchAsyncNegativeTst") {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
int *Hmm = nullptr, MemSz = 4096*4, NumElms = MemSz/4, InitVal = 123;
|
||||
bool IfTestPassed = true;
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MemSz));
|
||||
for (int i = 0; i < NumElms; ++i) {
|
||||
Hmm[i] = InitVal;
|
||||
}
|
||||
hipError_t err;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
SECTION("Passing null for dev ptr") {
|
||||
int *Ptr;
|
||||
err = hipMemPrefetchAsync(NULL, MemSz, 0, strm);
|
||||
if (err == hipSuccess) {
|
||||
WARN("hipMemPrefetchAsync() gives hipSuccess when NULL is passed!!");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
err = hipMemPrefetchAsync(Ptr, MemSz, 0, strm);
|
||||
if (err == hipSuccess) {
|
||||
WARN("hipMemPrefetchAsync() gives hipSuccess when uninitialized"
|
||||
" pointer is passed!!");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
|
||||
SECTION("Passing unusual count size(2nd param)") {
|
||||
// Passing count size as zero
|
||||
// expectation: Api should return error
|
||||
err = hipMemPrefetchAsync(Hmm, 0, 0, strm);
|
||||
if (err == hipSuccess) {
|
||||
WARN("hipMemPrefetchAsync() gives hipSuccess when count size is"
|
||||
" passed as zero!!");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
// Passing count size half of actually allocated
|
||||
// expectation: No issue should be observed
|
||||
err = hipMemPrefetchAsync(Hmm, MemSz/2, 0, strm);
|
||||
if (err != hipSuccess) {
|
||||
WARN("hipMemPrefetchAsync() returned error when count size passed is"
|
||||
" half of actually allocated!!");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
// Passing count size double that of actually allocated
|
||||
// expectation: Api should return error
|
||||
err = hipMemPrefetchAsync(Hmm, MemSz*2, 0, strm);
|
||||
if (err == hipSuccess) {
|
||||
WARN("hipMemPrefetchAsync() gives hipSuccess when count size passed is"
|
||||
" double that of actually allocated!!");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
SECTION("Passing invalid device Ordinal") {
|
||||
err = hipMemPrefetchAsync(Hmm, MemSz, 9999, strm);
|
||||
if (err == hipSuccess) {
|
||||
WARN("hipMemPrefetchAsync() gives hipSuccess when Invalid device"
|
||||
" ordinal is passed!!");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
SECTION("Checking behavior with stream object") {
|
||||
// Passing Null stream
|
||||
// expectation: No issue should be observed
|
||||
err = hipMemPrefetchAsync(Hmm, MemSz, 0, 0);
|
||||
if (err != hipSuccess) {
|
||||
WARN("hipMemPrefetchAsync() returns error when Null stream is"
|
||||
"passed!!");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
// Passing stream object belong to destination device
|
||||
// expectation: No issue should be observed
|
||||
int NGpus = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&NGpus));
|
||||
if (NGpus > 1) {
|
||||
hipStream_t strm1;
|
||||
HIP_CHECK(hipSetDevice(1));
|
||||
HIP_CHECK(hipStreamCreate(&strm1));
|
||||
err = hipMemPrefetchAsync(Hmm, MemSz, 1, strm1);
|
||||
if (err != hipSuccess) {
|
||||
WARN("hipMemPrefetchAsync() returns error when stream object"
|
||||
" created in the context of destination gpu is passed!!");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipStreamDestroy(strm1));
|
||||
}
|
||||
}
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
REQUIRE(IfTestPassed);
|
||||
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/* Test Case description: In this test case I am trying to allocate HMM memory
|
||||
which is not multiple of page Size, but still trying to launch kernel and
|
||||
see if we are getting values as expected.*/
|
||||
TEST_CASE("Unit_hipMemPrefetchAsync_NonPageSz") {
|
||||
int *Hmm = nullptr, NumElms = 4096*2, InitVal = 123;
|
||||
hipStream_t strm;
|
||||
bool IfTestPassed = true;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
// Allocating memory = 2*Page Size + 8 bytes
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, (NumElms * sizeof(int) + 8)));
|
||||
for (int i = 0; i < (NumElms + 2); ++i) {
|
||||
Hmm[i] = InitVal;
|
||||
}
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm, (NumElms * sizeof(int) + 8), 0, strm));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
MemPrftchAsyncKernel1<<<((NumElms + 2)/32 + 1), 32>>>(Hmm, (NumElms + 2));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
for (int i = 0; i < (NumElms + 2); ++i) {
|
||||
if (Hmm[i] != (InitVal * InitVal)) {
|
||||
WARN("Didnt receive expected output after kernel launch!!");
|
||||
IfTestPassed = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
REQUIRE(IfTestPassed);
|
||||
}
|
||||
@@ -0,0 +1,16 @@
|
||||
# AMD Tests
|
||||
set(TEST_SRC
|
||||
printfFlags.cc
|
||||
printfSpecifiers.cc
|
||||
)
|
||||
|
||||
# Create shared lib of all tests
|
||||
add_library(printfTests SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
set_property(TARGET printfTests PROPERTY CXX_STANDARD 17)
|
||||
else()
|
||||
target_compile_options(printfTests PUBLIC -std=c++17)
|
||||
endif()
|
||||
|
||||
# Add dependency on build_tests to build it on this custom target
|
||||
add_dependencies(build_tests printfTests)
|
||||
@@ -0,0 +1,57 @@
|
||||
/*
|
||||
Copyright (c) 2021 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 <hip_test_common.hh>
|
||||
#include <hip_test_process.hh>
|
||||
|
||||
__global__ void test_kernel() {
|
||||
printf("%08d\n", 42);
|
||||
printf("%08i\n", -42);
|
||||
printf("%08u\n", 42);
|
||||
printf("%08g\n", 123.456);
|
||||
printf("%0+8d\n", 42);
|
||||
printf("%+d\n", -42);
|
||||
printf("%+08d\n", 42);
|
||||
printf("%-8s\n", "xyzzy");
|
||||
printf("% i\n", -42);
|
||||
printf("%-16.8d\n", 42);
|
||||
printf("%16.8d\n", 42);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_printf_flags") {
|
||||
std::string reference(R"here(00000042
|
||||
-0000042
|
||||
00000042
|
||||
0123.456
|
||||
+0000042
|
||||
-42
|
||||
+0000042
|
||||
xyzzy
|
||||
-42
|
||||
00000042
|
||||
00000042
|
||||
)here");
|
||||
|
||||
hip::SpawnProc proc("unit/printfExe/printfFlags", true);
|
||||
REQUIRE(proc.run() == 0);
|
||||
REQUIRE(proc.getOutput() == reference);
|
||||
}
|
||||
@@ -0,0 +1,95 @@
|
||||
/*
|
||||
Copyright (c) 2021 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 <hip_test_common.hh>
|
||||
#include <hip_test_process.hh>
|
||||
|
||||
TEST_CASE("Unit_printf_specifier") {
|
||||
#ifdef __HIP_PLATFORM_NVIDIA__
|
||||
std::string reference(R"here(xyzzy
|
||||
%
|
||||
hello % world
|
||||
%s
|
||||
%s0xf01dab1eca55e77e
|
||||
%cxyzzy
|
||||
sep
|
||||
-42
|
||||
42
|
||||
123.456000
|
||||
-123.456000
|
||||
-1.234560e+02
|
||||
1.234560E+02
|
||||
123.456
|
||||
-123.456
|
||||
x
|
||||
(null)
|
||||
(nil)
|
||||
3.14159000 hello 0xf01dab1eca55e77e
|
||||
)here");
|
||||
#elif !defined(_WIN32)
|
||||
std::string reference(R"here(xyzzy
|
||||
%
|
||||
hello % world
|
||||
%s
|
||||
%s0xf01dab1eca55e77e
|
||||
%cxyzzy
|
||||
sep
|
||||
-42
|
||||
42
|
||||
123.456000
|
||||
-123.456000
|
||||
-1.234560e+02
|
||||
1.234560E+02
|
||||
123.456
|
||||
-123.456
|
||||
x
|
||||
|
||||
(nil)
|
||||
3.14159000 hello 0xf01dab1eca55e77e
|
||||
)here");
|
||||
#else
|
||||
std::string reference(R"here(xyzzy
|
||||
%
|
||||
hello % world
|
||||
%s
|
||||
%sF01DAB1ECA55E77E
|
||||
%cxyzzy
|
||||
sep
|
||||
-42
|
||||
42
|
||||
123.456000
|
||||
-123.456000
|
||||
-1.234560e+02
|
||||
1.234560E+02
|
||||
123.456
|
||||
-123.456
|
||||
x
|
||||
|
||||
0000000000000000
|
||||
3.14159000 hello F01DAB1ECA55E77E
|
||||
)here");
|
||||
#endif
|
||||
|
||||
hip::SpawnProc proc("unit/printfExe/printfSepcifiers", true);
|
||||
REQUIRE(0 == proc.run());
|
||||
REQUIRE(proc.getOutput() == reference);
|
||||
}
|
||||
@@ -0,0 +1,5 @@
|
||||
add_executable(printfFlags EXCLUDE_FROM_ALL printfFlags.cc)
|
||||
add_executable(printfSepcifiers EXCLUDE_FROM_ALL printfSepcifiers.cc)
|
||||
|
||||
add_dependencies(build_tests printfFlags)
|
||||
add_dependencies(build_tests printfSepcifiers)
|
||||
@@ -0,0 +1,42 @@
|
||||
/*
|
||||
Copyright (c) 2021 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 <hip/hip_runtime.h>
|
||||
|
||||
__global__ void test_kernel() {
|
||||
printf("%08d\n", 42);
|
||||
printf("%08i\n", -42);
|
||||
printf("%08u\n", 42);
|
||||
printf("%08g\n", 123.456);
|
||||
printf("%0+8d\n", 42);
|
||||
printf("%+d\n", -42);
|
||||
printf("%+08d\n", 42);
|
||||
printf("%-8s\n", "xyzzy");
|
||||
printf("% i\n", -42);
|
||||
printf("%-16.8d\n", 42);
|
||||
printf("%16.8d\n", 42);
|
||||
}
|
||||
|
||||
int main() {
|
||||
test_kernel<<<1, 1>>>();
|
||||
hipDeviceSynchronize();
|
||||
}
|
||||
@@ -0,0 +1,65 @@
|
||||
/*
|
||||
Copyright (c) 2021 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 <hip/hip_runtime.h>
|
||||
|
||||
__global__ void test_kernel() {
|
||||
const char* N = nullptr;
|
||||
const char* s = "hello world";
|
||||
printf("xyzzy\n");
|
||||
printf("%%\n");
|
||||
printf("hello %% world\n");
|
||||
printf("%%s\n");
|
||||
// Two special tests to make sure that the compiler pass correctly
|
||||
// skips over a '%%' without affecting the logic for locating
|
||||
// string arguments.
|
||||
printf("%%s%p\n", (void*)0xf01dab1eca55e77e);
|
||||
printf("%%c%s\n", "xyzzy");
|
||||
printf("%c%c%c\n", 's', 'e', 'p');
|
||||
printf("%d\n", -42);
|
||||
printf("%u\n", 42);
|
||||
printf("%f\n", 123.456);
|
||||
#ifdef __HIP_PLATFORM_AMD__
|
||||
printf("%F\n", -123.456);
|
||||
#else
|
||||
printf("%f\n", -123.456);
|
||||
#endif
|
||||
printf("%e\n", -123.456);
|
||||
printf("%E\n", 123.456);
|
||||
printf("%g\n", 123.456);
|
||||
printf("%G\n", -123.456);
|
||||
printf("%c\n", 'x');
|
||||
printf("%s\n", N);
|
||||
printf("%p\n", (void *)N);
|
||||
#ifdef __HIP_PLATFORM_AMD__
|
||||
printf("%.*f %*.*s %p\n", 8, 3.14159, 8, 5, s, (void*)0xf01dab1eca55e77e);
|
||||
#else
|
||||
// In Cuda, printf doesn't support %.*, %*.*
|
||||
printf("%.8f %8.5s %p\n", 3.14159, s, (void*)0xf01dab1eca55e77e);
|
||||
#endif
|
||||
}
|
||||
|
||||
int main() {
|
||||
test_kernel<<<1, 1>>>();
|
||||
hipDeviceSynchronize();
|
||||
return 0;
|
||||
}
|
||||
@@ -238,6 +238,7 @@ int main(int argc, char** argv) {
|
||||
HIPCHECK(hipGetDeviceCount(&device_num));
|
||||
for (int dev = 0; dev < device_num; ++dev) {
|
||||
/*************************************************************************/
|
||||
HIPCHECK(hipSetDevice(dev));
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, dev));
|
||||
|
||||
|
||||
@@ -35,7 +35,7 @@ THE SOFTWARE.
|
||||
* TEST: %t --tests 0x11
|
||||
* TEST: %t --tests 0x12
|
||||
* TEST: %t --tests 0x13
|
||||
* TEST: %t --tests 0x14 EXCLUDE_HIP_PLATFORM amd
|
||||
* TEST: %t --tests 0x14
|
||||
* TEST: %t --tests 0x15
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
مرجع در شماره جدید
Block a user