AI/rocm-systems
2
0
Fork 0
Koodi Ongelmat Pull-pyynnöt Actions Paketit Projektit Julkaisut Wiki Toiminta

Merge 'develop' into 'amd-staging'

Selaa lähdekoodia 
Change-Id: I27f7446347f663f3fbf988f13f9c5c106c8e0da5
This commit is contained in:
Jenkins
2023-02-08 12:10:40 +00:00
vanhempi 2b5418dacb 554ab5dc40
commit e744717dae
21 muutettua tiedostoa jossa 2179 lisäystä ja 963 poistoa
Download Patch File Download Diff File Expand all files Collapse all files
.jenkins/jenkinsfile
+4 -4
Näytä tiedosto
@@ -1,6 +1,6 @@
def hipBuildTest(String backendLabel) {
node(backendLabel) {
stage("Source sync ${backendLabel}") {
stage("SYNC - ${backendLabel}") {
// Checkout hip repository with the PR patch
dir("${WORKSPACE}/hip-tests") {
@@ -37,7 +37,7 @@ def hipBuildTest(String backendLabel) {
}
}
stage("Build - Catch2 framework") {
stage("BUILD HIP - ${backendLabel}") {
// Running the build on hipamd workspace
dir("${WORKSPACE}/hipamd") {
sh """#!/usr/bin/env bash
@@ -56,7 +56,7 @@ def hipBuildTest(String backendLabel) {
}
}
stage("Build - HIP TESTS") {
stage("BUILD HIP TESTS - ${backendLabel}") {
// Running the build on HIP TESTS workspace
dir("${WORKSPACE}/hip-tests") {
env.HIP_PATH = "${HIPAMD_DIR}" + "/build/install"
@@ -78,7 +78,7 @@ def hipBuildTest(String backendLabel) {
}
}
stage('HIP Unit Tests - Catch2 framework') {
stage("TEST - ${backendLabel}") {
dir("${WORKSPACE}/hip-tests") {
sh """#!/usr/bin/env bash
set -x
catch/external/Catch2/cmake/Catch2/catch_include.cmake
Vendored
+35 -33
Näytä tiedosto
@@ -3,39 +3,41 @@
get_filename_component(_cmake_path cmake ABSOLUTE)
foreach(EXEC_NAME ${exc_names})
if(WIN32)
set(EXEC_NAME ${EXEC_NAME}.exe)
endif()
if(EXISTS "${EXEC_NAME}")
execute_process(
COMMAND "${_cmake_path}"
-D "TEST_TARGET=${TARGET}"
-D "TEST_EXECUTABLE=${EXEC_NAME}"
-D "TEST_EXECUTOR=${crosscompiling_emulator}"
-D "TEST_WORKING_DIR=${_workdir}"
-D "TEST_SPEC=${_TEST_SPEC}"
-D "TEST_EXTRA_ARGS=${_EXTRA_ARGS}"
-D "TEST_PROPERTIES=${_PROPERTIES}"
-D "TEST_PREFIX=${_TEST_PREFIX}"
-D "TEST_SUFFIX=${_TEST_SUFFIX}"
-D "TEST_LIST=${_TEST_LIST}"
-D "TEST_REPORTER=${_REPORTER}"
-D "TEST_OUTPUT_DIR=${_OUTPUT_DIR}"
-D "TEST_OUTPUT_PREFIX=${_OUTPUT_PREFIX}"
-D "TEST_OUTPUT_SUFFIX=${_OUTPUT_SUFFIX}"
-D "CTEST_FILE=${ctestfilepath}"
-P "${_CATCH_ADD_TEST_SCRIPT}"
OUTPUT_VARIABLE output
RESULT_VARIABLE result
WORKING_DIRECTORY "${TEST_WORKING_DIR}"
)
else()
message("executable not found : ${EXEC_NAME}" )
endif()
endforeach()
if(NOT EXISTS "${ctestfilepath}")
foreach(EXEC_NAME ${exc_names})
if(WIN32)
set(EXEC_NAME ${EXEC_NAME}.exe)
endif()
if(EXISTS "${EXEC_NAME}")
execute_process(
COMMAND "${_cmake_path}"
-D "TEST_TARGET=${TARGET}"
-D "TEST_EXECUTABLE=${EXEC_NAME}"
-D "TEST_EXECUTOR=${crosscompiling_emulator}"
-D "TEST_WORKING_DIR=${_workdir}"
-D "TEST_SPEC=${_TEST_SPEC}"
-D "TEST_EXTRA_ARGS=${_EXTRA_ARGS}"
-D "TEST_PROPERTIES=${_PROPERTIES}"
-D "TEST_PREFIX=${_TEST_PREFIX}"
-D "TEST_SUFFIX=${_TEST_SUFFIX}"
-D "TEST_LIST=${_TEST_LIST}"
-D "TEST_REPORTER=${_REPORTER}"
-D "TEST_OUTPUT_DIR=${_OUTPUT_DIR}"
-D "TEST_OUTPUT_PREFIX=${_OUTPUT_PREFIX}"
-D "TEST_OUTPUT_SUFFIX=${_OUTPUT_SUFFIX}"
-D "CTEST_FILE=${ctestfilepath}"
-P "${_CATCH_ADD_TEST_SCRIPT}"
OUTPUT_VARIABLE output
RESULT_VARIABLE result
WORKING_DIRECTORY "${TEST_WORKING_DIR}"
)
else()
message("executable not found : ${EXEC_NAME}" )
endif()
endforeach()
endif()
if(EXISTS "${ctestfilepath}")
# include the generated ctest file for execution
include(${ctestfilepath})
# include the generated ctest file for execution
include(${ctestfilepath})
endif()
catch/hipTestMain/config/config_amd_linux_common.json
+5 -1
Näytä tiedosto
@@ -15,6 +15,8 @@
"Unit_hipMemset_Negative_OutOfBoundsPtr",
"Unit_hipDeviceReset_Positive_Basic",
"Unit_hipDeviceReset_Positive_Threaded",
"Unit_hipMemAdvise_AccessedBy_All_Devices",
"Unit_hipMemAdvise_No_Flag_Interference",
"Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep",
"Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph",
"Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode",
@@ -27,6 +29,8 @@
"Unit_hipStreamAttachMemAsync_Negative_Parameters",
"Unit_hipMemGetAddressRange_Positive",
"Unit_hipGraphAddMemcpyNode1D_Negative_Basic",
"Unit_hipStreamGetCaptureInfo_Nullstream_CaptureInfo"
"Unit_hipStreamGetCaptureInfo_Nullstream_CaptureInfo",
"intermittent issue: corrupted double-linked list",
"Unit_hipGraphRetainUserObject_Functional_2"
]
}
catch/hipTestMain/config/config_amd_windows_common.json
+2
Näytä tiedosto
@@ -94,6 +94,8 @@
"Unit_hipStreamSynchronize_NullStreamAndStreamPerThread",
"Note: intermittent Seg fault failure ",
"Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags",
"Unit_hipMemAdvise_AccessedBy_All_Devices",
"Unit_hipMemAdvise_No_Flag_Interference",
"Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep",
"Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph",
"Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode",
catch/unit/deviceLib/Atomic_func.cc
+119
Näytä tiedosto
@@ -0,0 +1,119 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_checkers.hh>
// Test case to validate atomicInc and atomicDec functions.
// if TestToRun=1, then atomicInc function will be tested and validated
// if TestToRun=2, then atomicDec function will be tested and validated.
// kernel function for atomicInc
static __global__ void AtomicCheckInc(int* g_ptr) {
atomicInc(reinterpret_cast<unsigned int*>(&g_ptr[0]), 17);
}
// kernel function for atomicDec
static __global__ void AtomicCheckDec(int* g_ptr) {
atomicDec(reinterpret_cast<unsigned int*>(&g_ptr[0]), 25);
}
// verify results for atomicInc
static int verifyResultInc(int value) {
int limit = 17;
value = (value >= limit) ? 0 : value + 1;
return value;
}
// verify results for atomicDec
static int verifyResultDec(int value) {
int limit = 25;
value = ((value == 0) || (value > limit)) ? limit : value - 1;
return value;
}
// common fuction to launch atomic functions kernel.
static void launchAtomicFunction(int *Hptr, int val, int TestToRun) {
unsigned int memSize = sizeof(int) * 1;
int *dptr{nullptr};
// allocate device memory
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&dptr), memSize));
// copy host memory to device
HIP_CHECK(hipMemcpy(dptr, Hptr, memSize, hipMemcpyHostToDevice));
// launch kernel function
if (TestToRun == 1) {
AtomicCheckInc<<<1, 1>>>(dptr);
} else if (TestToRun == 2) {
AtomicCheckDec<<<1, 1>>>(dptr);
}
// copy back from device to host
HIP_CHECK(hipMemcpy(Hptr, dptr, memSize, hipMemcpyDeviceToHost));
// verify the results.
if (TestToRun == 1) {
int result = verifyResultInc(val);
REQUIRE(result == Hptr[0]);
} else if (TestToRun == 2) {
int result = verifyResultDec(val);
REQUIRE(result == Hptr[0]);
}
// Cleanup memory
HIP_CHECK(hipFree(dptr));
}
TEST_CASE("Unit_AtomicFunctions_Inc") {
int *Hptr{nullptr};
int val;
// Allocate Host memory
Hptr = reinterpret_cast<int*>(malloc(sizeof(int)));
SECTION("Test case when value is lesser than limit") {
val = Hptr[0] = 10;
launchAtomicFunction(Hptr, val, 1);
}
SECTION("Test case when value is greater than limit") {
val = Hptr[0] = 20;
launchAtomicFunction(Hptr, val, 1);
}
SECTION("Test case when value is equal to the limit") {
val = Hptr[0] = 17;
launchAtomicFunction(Hptr, val, 1);
}
free(Hptr);
}
TEST_CASE("Unit_AtomicFunctions_Dec") {
int *Hptr{nullptr};
int val;
// Allocate Host memory
Hptr = reinterpret_cast<int*>(malloc(sizeof(int)));
SECTION("Test case when value is less than limit") {
val = Hptr[0] = 4;
launchAtomicFunction(Hptr, val, 2);
}
SECTION("Test case when value is greater than limit") {
val = Hptr[0] = 31;
launchAtomicFunction(Hptr, val, 2);
}
SECTION("Test case when value is equal to the limit") {
val = Hptr[0] = 25;
launchAtomicFunction(Hptr, val, 2);
}
free(Hptr);
}
catch/unit/deviceLib/CMakeLists.txt
+1
Näytä tiedosto
@@ -13,6 +13,7 @@ set(TEST_SRC
syncthreadsand.cc
syncthreadscount.cc
syncthreadsor.cc
Atomic_func.cc
)
if(UNIX)
catch/unit/graph/hipGraphGetRootNodes.cc
+119 -1
Näytä tiedosto
@@ -1,5 +1,5 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2023 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
@@ -28,6 +28,11 @@ Functional ::
nodes list will be set to NULL, and the number of nodes actually obtained will be returned in NumRootNodes.
4) Create a graph with stream capture done on multiple dependent streams.
Verify root nodes of created graph are matching the operations pushed which doesn't have dependencies.
5) Functional Test to validate number of root nodes when dependencies in the graph are dynamically varied.
6) Functional Test to validate number of root nodes when dependencies in the graph are dynamically varied
in a cloned graph.
7) Functional Test to validate number of root nodes when a graph with N independent nodes is added as a
child node to another graph.
Argument Validation ::
1) Pass graph as nullptr and verify api returns error code.
@@ -41,6 +46,12 @@ Argument Validation ::
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
#define NUM_OF_DUMMY_NODES 8
static __global__ void dummyKernel() {
return;
}
/**
* Functional Test for API fetching root node list
*/
@@ -346,3 +357,110 @@ TEST_CASE("Unit_hipGraphGetRootNodes_ParamValidation") {
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(C_d));
}
/**
* Functional Test to validate number of root nodes when dependencies
* in the graph are dynamically varied.
*/
TEST_CASE("Unit_hipGraphGetRootNodes_Complx_NumRootNodes") {
hipGraph_t graph;
hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES];
hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES];
HIP_CHECK(hipGraphCreate(&graph, 0));
// Create graph with no dependencies
for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) {
void* kernelArgs[] = {nullptr};
kernelNodeParams[i].func = reinterpret_cast<void *>(dummyKernel);
kernelNodeParams[i].gridDim = dim3(1);
kernelNodeParams[i].blockDim = dim3(1);
kernelNodeParams[i].sharedMemBytes = 0;
kernelNodeParams[i].kernelParams = reinterpret_cast<void**>(kernelArgs);
kernelNodeParams[i].extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr,
0, &kernelNodeParams[i]));
}
size_t numRootNodes{};
HIP_CHECK(hipGraphGetRootNodes(graph, nullptr, &numRootNodes));
REQUIRE(numRootNodes == NUM_OF_DUMMY_NODES);
// Start creating dependencies in a chain
for (size_t i = 0; i < (NUM_OF_DUMMY_NODES - 1); i++) {
numRootNodes = 0;
HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode[i],
&kernelnode[i+1], 1));
HIP_CHECK(hipGraphGetRootNodes(graph, nullptr, &numRootNodes));
REQUIRE(numRootNodes == (NUM_OF_DUMMY_NODES - i - 1));
}
HIP_CHECK(hipGraphDestroy(graph));
}
/**
* Functional Test to validate number of root nodes when dependencies
* in the graph are dynamically varied in a cloned graph.
*/
TEST_CASE("Unit_hipGraphGetRootNodes_Complx_NumRootNodes_ClonedGrph") {
hipGraph_t graph, clonedgraph;
hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES];
hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES];
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphCreate(&clonedgraph, 0));
// Create graph with no dependencies
for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) {
void* kernelArgs[] = {nullptr};
kernelNodeParams[i].func = reinterpret_cast<void *>(dummyKernel);
kernelNodeParams[i].gridDim = dim3(1);
kernelNodeParams[i].blockDim = dim3(1);
kernelNodeParams[i].sharedMemBytes = 0;
kernelNodeParams[i].kernelParams = reinterpret_cast<void**>(kernelArgs);
kernelNodeParams[i].extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr,
0, &kernelNodeParams[i]));
}
size_t numRootNodes{};
HIP_CHECK(hipGraphClone(&clonedgraph, graph));
HIP_CHECK(hipGraphGetRootNodes(clonedgraph, nullptr, &numRootNodes));
REQUIRE(numRootNodes == NUM_OF_DUMMY_NODES);
// Start creating dependencies in a chain
for (size_t i = 0; i < (NUM_OF_DUMMY_NODES - 1); i++) {
numRootNodes = 0;
hipGraphNode_t node1, node2;
HIP_CHECK(hipGraphNodeFindInClone(&node1, kernelnode[i], clonedgraph));
HIP_CHECK(hipGraphNodeFindInClone(&node2, kernelnode[i+1], clonedgraph));
HIP_CHECK(hipGraphAddDependencies(clonedgraph, &node1, &node2, 1));
HIP_CHECK(hipGraphGetRootNodes(clonedgraph, nullptr, &numRootNodes));
REQUIRE(numRootNodes == (NUM_OF_DUMMY_NODES - i - 1));
}
HIP_CHECK(hipGraphDestroy(clonedgraph));
HIP_CHECK(hipGraphDestroy(graph));
}
/**
* Functional Test to validate number of root nodes when a graph with N
* independent nodes is added as a child node to another graph.
*/
TEST_CASE("Unit_hipGraphGetRootNodes_Complx_NRootNodesAsChildGraph") {
hipGraph_t graph, graph1;
hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES];
hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES];
hipGraphNode_t child_node;
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphCreate(&graph1, 0));
// Create graph with no dependencies
for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) {
void* kernelArgs[] = {nullptr};
kernelNodeParams[i].func = reinterpret_cast<void *>(dummyKernel);
kernelNodeParams[i].gridDim = dim3(1);
kernelNodeParams[i].blockDim = dim3(1);
kernelNodeParams[i].sharedMemBytes = 0;
kernelNodeParams[i].kernelParams = reinterpret_cast<void**>(kernelArgs);
kernelNodeParams[i].extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr,
0, &kernelNodeParams[i]));
}
HIP_CHECK(hipGraphAddChildGraphNode(&child_node, graph1,
nullptr, 0, graph));
size_t numRootNodes{};
HIP_CHECK(hipGraphGetRootNodes(graph1, nullptr, &numRootNodes));
REQUIRE(numRootNodes == 1);
HIP_CHECK(hipGraphDestroy(graph1));
HIP_CHECK(hipGraphDestroy(graph));
}
catch/unit/memory/CMakeLists.txt
+2
Näytä tiedosto
@@ -108,6 +108,7 @@ set(TEST_SRC
hipMemcpySync.cc
hipMemsetSync.cc
hipMemsetAsync.cc
hipMemAdvise_old.cc
hipMemAdvise.cc
hipMemRangeGetAttributes.cc
hipStreamAttachMemAsync.cc
@@ -194,6 +195,7 @@ set(TEST_SRC
hipMemcpySync.cc
hipMemsetSync.cc
hipMemsetAsync.cc
hipMemAdvise_old.cc
hipMemAdvise.cc
hipMemRangeGetAttributes.cc
hipMemRangeGetAttributes_old.cc
catch/unit/memory/hipMemAdvise.cc
+283 -904
Näytä tiedosto
File diff suppressed because it is too large Load Diff
catch/unit/memory/hipMemAdvise_old.cc
+938
Näytä tiedosto
@@ -0,0 +1,938 @@
/*
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.
*/
/* Test Case Description:
Scenario-1: The following Function Tests the working of flags which can be
assigned to HMM memory using hipMemAdvise() api
Scenario-2: Negative tests on hipMemAdvise() api
Scenario-3: The following function tests various scenarios around the flag
'hipMemAdviseSetPreferredLocation' using HMM memory and hipMemAdvise() api
Scenario-4: The following function tests various scenarios around the flag
'hipMemAdviseSetReadMostly' using HMM memory and hipMemAdvise() api
Scenario-5: The following function verifies if assigning of a flag
invalidates the earlier flag which was assigned to the same memory region
using hipMemAdvise()
Scenario-6: The following function tests if peers can set
hipMemAdviseSetAccessedBy flag
on HMM memory prefetched on each of the other gpus
Scenario-7: Set AccessedBy flag and check value returned by
hipMemRangeGetAttribute() It should be -2(same is observed on cuda)
Scenario-8: Set AccessedBy flag to device 0 on Hmm memory and prefetch the
memory to device 1, then probe for AccessedBy flag using
hipMemRangeGetAttribute() we should still see the said flag is set for
device 0
Scenario-9: 1) Set AccessedBy to device 0 followed by PreferredLocation to
device 1 check for AccessedBy flag using hipMemRangeGetAttribute() it should
return 0
2) Unset AccessedBy to 0 and set it to device 1 followed by
PreferredLocation to device 1, check for AccessedBy flag using
hipMemRangeGetAttribute() it should return 1
Scenario-10: Set AccessedBy flag to HMM memory launch a kernel and then unset
AccessedBy, launch kernel. We should not have any access issues
Scenario-11: Allocate memory using aligned_alloc(), assign PreferredLocation
flag to the allocated memory and launch a kernel. Kernel should get executed
successfully without hang or segfault
Scenario-12: Allocate Hmm memory, set advise to PreferredLocation and then
get attribute using the api hipMemRangeGetAttribute() for
hipMemRangeAttributeLastPrefetchLocation the value returned should be -2
Scenario-13: Allocate HMM memory, set PreferredLocation to device 0, Prfetch
the mem to device1, probe for hipMemRangeAttributeLastPrefetchLocation using
hipMemRangeGetAttribute(), we should get 1
Scenario-14: Allocate HMM memory, set ReadMostly followed by
PreferredLocation, probe for hipMemRangeAttributeReadMostly and
hipMemRangeAttributePreferredLocation
using hipMemRangeGetAttribute() we should observe 1 and 0 correspondingly.
In other words setting of hipMemRangeAttributePreferredLocation should not
impact hipMemRangeAttributeReadMostly advise to the memory
Scenario-15: Allocate Hmm memory, advise it to ReadMostly for gpu: 0 and
launch kernel on all other gpus except 0. This test case may discover any
effect or access denial case arising due to setting ReadMostly only to a
particular gpu
*/
#include <hip_test_common.hh>
#if __linux__
#include <unistd.h>
#include <sys/mman.h>
#include <sys/wait.h>
#endif
// Kernel function
__global__ void MemAdvseKernel(int n, int *x) {
int index = blockIdx.x * blockDim.x + threadIdx.x;
if (index < n)
x[index] = x[index] * x[index];
}
// Kernel
__global__ void MemAdvise2(int *Hmm, int n) {
int index = blockIdx.x * blockDim.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;
for (int i = index; i < n; i += stride) {
Hmm[i] = Hmm[i] + 10;
}
}
// Kernel
__global__ void MemAdvise3(int *Hmm, int *Hmm1, int n) {
int index = blockIdx.x * blockDim.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;
for (int i = index; i < n; i += stride) {
Hmm1[i] = Hmm[i] + 10;
}
}
static bool CheckError(hipError_t err, int LineNo) {
if (err == hipSuccess) {
WARN("Error expected but received hipSuccess at line no.:" << LineNo);
return false;
} else {
return true;
}
}
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;
}
// The following Function Tests the working of flags which can be assigned
// to HMM memory using hipMemAdvise() api
TEST_CASE("Unit_hipMemAdvise_TstFlags") {
int MangdMem = HmmAttrPrint();
if (MangdMem == 1) {
bool IfTestPassed = true;
int NumDevs = 0, *Outpt = nullptr;
int MEM_SIZE = 4*1024, A_CONST = 9999;
float *Hmm = nullptr;
int AttrVal = 0;
HIP_CHECK(hipGetDeviceCount(&NumDevs));
Outpt = new int[NumDevs];
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 2, hipMemAttachGlobal));
// With the following for loop we iterate through each of the Gpus in the
// system set and unset the flags and check the behavior.
for (int i = 0; i < NumDevs; ++i) {
HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseSetReadMostly, i));
HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal),
hipMemRangeAttributeReadMostly, Hmm,
MEM_SIZE * 2));
if (AttrVal != 1) {
WARN("Attempt to set hipMemAdviseSetReadMostly flag failed!\n");
IfTestPassed = false;
}
HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseUnsetReadMostly,
i));
HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal),
hipMemRangeAttributeReadMostly, Hmm,
(MEM_SIZE * 2)));
if (AttrVal != 0) {
WARN("Attempt to Unset hipMemAdviseSetReadMostly flag failed!\n");
IfTestPassed = false;
}
AttrVal = A_CONST;
// Currently hipMemAdviseSetPreferredLocation and
// hipMemAdviseSetAccessedBy
// flags are resulting in issues: SWDEV-267357
HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2,
hipMemAdviseSetPreferredLocation, i));
HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal),
hipMemRangeAttributePreferredLocation,
Hmm, (MEM_SIZE * 2)));
if (AttrVal != i) {
WARN("Attempt to set hipMemAdviseSetPreferredLocation flag failed!\n");
IfTestPassed = false;
}
AttrVal = A_CONST;
HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2,
hipMemAdviseUnsetPreferredLocation, i));
HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal),
hipMemRangeAttributePreferredLocation,
Hmm, (MEM_SIZE * 2)));
if (AttrVal == i) {
WARN("Attempt to Unset hipMemAdviseUnsetPreferredLocation ");
WARN("flag failed!\n");
IfTestPassed = false;
}
for (int m = 0; m < NumDevs; ++m) {
Outpt[m] = A_CONST;
}
HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseSetAccessedBy, i));
HIP_CHECK(hipMemRangeGetAttribute(Outpt, sizeof(Outpt),
hipMemRangeAttributeAccessedBy, Hmm,
(MEM_SIZE * 2)));
if ((Outpt[0]) != i) {
WARN("Attempt to set hipMemAdviseSetAccessedBy flag failed!\n");
IfTestPassed = false;
}
for (int m = 0; m < NumDevs; ++m) {
Outpt[m] = A_CONST;
}
HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseUnsetAccessedBy,
i));
HIP_CHECK(hipMemRangeGetAttribute(Outpt, sizeof(Outpt),
hipMemRangeAttributeAccessedBy, Hmm,
(MEM_SIZE * 2)));
if ((Outpt[0]) >= 0) {
WARN("Attempt to Unset hipMemAdviseUnsetAccessedBy flag failed!\n");
IfTestPassed = false;
}
}
delete [] Outpt;
HIP_CHECK(hipFree(Hmm));
REQUIRE(IfTestPassed);
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
TEST_CASE("Unit_hipMemAdvise_NegtveTsts") {
int MangdMem = HmmAttrPrint();
if (MangdMem == 1) {
bool IfTestPassed = true;
int NumDevs = 0, MEM_SIZE = 4*1024;
float *Hmm = nullptr;
std::string str;
HIP_CHECK(hipGetDeviceCount(&NumDevs));
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 2, hipMemAttachGlobal));
// Passing NULL as first parameter instead of valid pointer to a memory
IfTestPassed &= CheckError(hipMemAdvise(NULL, MEM_SIZE * 2,
hipMemAdviseSetReadMostly, 0), __LINE__);
// Passing 0 for count(2nd param) parameter
IfTestPassed &= CheckError(hipMemAdvise(Hmm, 0, hipMemAdviseSetReadMostly,
0), __LINE__);
// Passing count much more than actually allocated value
IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 6,
hipMemAdviseSetReadMostly, 0), __LINE__);
REQUIRE(IfTestPassed);
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
// The following function tests various scenarios around the flag
// 'hipMemAdviseSetPreferredLocation' using HMM memory and hipMemAdvise() api
TEST_CASE("Unit_hipMemAdvise_PrefrdLoc") {
int MangdMem = HmmAttrPrint();
if (MangdMem == 1) {
// Check that when a page fault occurs for the memory region set to devPtr,
// the data is migrated to the destn processor
int MEM_SIZE = 4096, A_CONST = 9999;
int *Hmm = nullptr, NumDevs = 0, dev = A_CONST;
bool IfTestPassed = true;
HIP_CHECK(hipGetDeviceCount(&NumDevs));
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 3, hipMemAttachGlobal));
for (int i = 0; i < ((MEM_SIZE * 3)/4); ++i) {
Hmm[i] = 4;
}
for (int devId = 0; devId < NumDevs; ++devId) {
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE * 3,
hipMemAdviseSetPreferredLocation, devId));
int NumElms = ((MEM_SIZE * 3)/4);
MemAdvseKernel<<<NumElms/32, 32>>>(NumElms, Hmm);
int dev = A_CONST;
HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(dev),
hipMemRangeAttributePreferredLocation,
Hmm, MEM_SIZE * 3));
if (dev != devId) {
WARN("Memory observed to be not available on expected location\n");
WARN("line no: " << __LINE__);
WARN("dev: " << dev);
IfTestPassed = false;
}
}
// Check that when preferred location is set for a memory region,
// data can still be prefetched using hipMemPrefetchAsync
hipStream_t strm;
dev = A_CONST;
for (int devId = 0; devId < NumDevs; ++devId) {
HIP_CHECK(hipSetDevice(devId));
HIP_CHECK(hipStreamCreate(&strm));
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE * 3,
hipMemAdviseSetPreferredLocation, devId));
HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE * 3, devId, strm));
HIP_CHECK(hipStreamSynchronize(strm));
HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(dev),
hipMemRangeAttributeLastPrefetchLocation,
Hmm, MEM_SIZE * 3));
if (dev != devId) {
WARN("Memory reported to be not available at the Prefetched ");
WARN("location with device id: " << devId);
WARN("line no: " << __LINE__);
WARN("dev: " << dev);
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");
}
}
// The following function tests various scenarios around the flag
// 'hipMemAdviseSetReadMostly' using HMM memory and hipMemAdvise() api
TEST_CASE("Unit_hipMemAdvise_ReadMostly") {
int MangdMem = HmmAttrPrint();
if (MangdMem == 1) {
bool IfTestPassed = true;
int MEM_SIZE = 4096, A_CONST = 9999;
float *Hmm = nullptr;
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE));
for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) {
Hmm[i] = A_CONST;
}
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, 0));
// Checking if the data can be read after setting hipMemAdviseSetReadMostly
for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) {
if (Hmm[i] != A_CONST) {
WARN("Didn't find expected value in Hmm memory after setting");
WARN(" hipMemAdviseSetReadMostly flag line no.: " << __LINE__);
IfTestPassed = false;
}
}
// Checking if the memory region can be modified
for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) {
Hmm[i] = A_CONST;
}
for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) {
if (Hmm[i] != A_CONST) {
WARN("Didn't find expected value in Hmm memory after Modification\n");
WARN("line no.: " << __LINE__);
IfTestPassed = false;
}
}
int out = A_CONST;
HIP_CHECK(hipMemRangeGetAttribute(&out, 4, hipMemRangeAttributeReadMostly,
Hmm, MEM_SIZE));
if (out != 1) {
WARN("out value: " << out);
IfTestPassed = false;
}
// Checking the advise attribute after prefetch
HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, 0, 0));
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemRangeGetAttribute(&out, sizeof(int),
hipMemRangeAttributeReadMostly, Hmm,
MEM_SIZE));
if (out != 1) {
WARN("Attribute assigned to memory changed after calling ");
WARN("hipMemPrefetchAsync(). line no.: " << __LINE__);
WARN("out value: " << out);
IfTestPassed = false;
}
// hipMemAdvise should succeed for SetReadMostly and UnsetReadMostly
// irrespective of the device
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, 99));
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseUnsetReadMostly, -12));
HIP_CHECK(hipFree(Hmm));
REQUIRE(IfTestPassed);
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
// The following function verifies if assigning of a flag invalidates the
// earlier flag which was assigned to the same memory region using
// hipMemAdvise()
TEST_CASE("Unit_hipMemAdvise_TstFlgOverrideEffect") {
int MangdMem = HmmAttrPrint();
if (MangdMem == 1) {
bool IfTestPassed = true;
int MEM_SIZE = 4*4096, A_CONST = 9999;
float *Hmm = nullptr;
int NumDevs = 0, dev = A_CONST;
HIP_CHECK(hipGetDeviceCount(&NumDevs));
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal));
for (int i = 0; i < NumDevs; ++i) {
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, i));
HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int),
hipMemRangeAttributeReadMostly, Hmm,
MEM_SIZE));
if (dev != 1) {
WARN("hipMemAdviseSetReadMostly flag did not take affect despite ");
WARN("setting it using hipMemAdvise(). line no.: " << __LINE__);
IfTestPassed = false;
break;
}
dev = A_CONST;
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetPreferredLocation,
i));
HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int),
hipMemRangeAttributePreferredLocation,
Hmm, MEM_SIZE));
if (dev != i) {
WARN("hipMemAdviseSetPreferredLocation flag did not take affect ");
WARN("despite setting it using hipMemAdvise()\n");
WARN("line no.: " << __LINE__);
IfTestPassed = false;
break;
}
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, i));
dev = A_CONST;
HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int),
hipMemRangeAttributeAccessedBy, Hmm,
MEM_SIZE));
if (dev != i) {
WARN("hipMemAdviseSetAccessedBy flag did not take affect despite ");
WARN("setting it using hipMemAdvise(). line no.: " << __LINE__);
IfTestPassed = false;
break;
}
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseUnsetAccessedBy, i));
}
HIP_CHECK(hipFree(Hmm));
REQUIRE(IfTestPassed);
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
// The following function tests if peers can set hipMemAdviseSetAccessedBy flag
// on HMM memory prefetched on each of the other gpus
#if HT_AMD
TEST_CASE("Unit_hipMemAdvise_TstAccessedByPeer") {
int MangdMem = HmmAttrPrint();
if (MangdMem == 1) {
bool IfTestPassed = true;
int *Hmm = nullptr, MEM_SIZE = 4*4096, A_CONST = 9999;;
int NumDevs = 0, CanAccessPeer = A_CONST, flag = 0;
HIP_CHECK(hipGetDeviceCount(&NumDevs));
if (NumDevs < 2) {
SUCCEED("Test TestSetAccessedByPeer() need atleast two Gpus to test"
" the scenario. This system has GPUs less than 2");
}
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal));
for (int i = 0; i < NumDevs; ++i) {
HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, i, 0));
for (int j = 0; j < NumDevs; ++j) {
if (i == j)
continue;
HIP_CHECK(hipSetDevice(j));
HIP_CHECK(hipDeviceCanAccessPeer(&CanAccessPeer, j, i));
if (CanAccessPeer) {
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, j));
for (uint64_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) {
Hmm[m] = 4;
}
HIP_CHECK(hipDeviceEnablePeerAccess(i, 0));
MemAdvseKernel<<<(MEM_SIZE/sizeof(int)/32), 32>>>(
(MEM_SIZE/sizeof(int)), Hmm);
HIP_CHECK(hipDeviceSynchronize());
// Verifying the result
for (uint64_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) {
if (Hmm[m] != 16) {
flag = 1;
}
}
if (flag) {
WARN("Didnt get Expected results with device: " << j);
WARN("line no.: " << __LINE__);
IfTestPassed = false;
flag = 0;
}
}
}
}
HIP_CHECK(hipFree(Hmm));
REQUIRE(IfTestPassed);
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
#endif
/* Set AccessedBy flag and check value returned by hipMemRangeGetAttribute()
It should be -2(same is observed on cuda)*/
TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg") {
int managed = HmmAttrPrint();
if (managed == 1) {
int *Hmm = NULL, data = 999;
HIP_CHECK(hipMallocManaged(&Hmm, 2*4096));
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0));
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
hipMemRangeAttributeLastPrefetchLocation,
Hmm, 2*4096));
if (data != -2) {
WARN("Didnt get expected value!!\n");
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
/* Set AccessedBy flag to device 0 on Hmm memory and prefetch the memory to
device 1, then probe for AccessedBy flag using hipMemRangeGetAttribute()
we should still see the said flag is set for device 0*/
TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg2") {
int managed = HmmAttrPrint();
if (managed == 1) {
int *Hmm = NULL, data = 999, Ngpus = 0;
HIP_CHECK(hipGetDeviceCount(&Ngpus));
if (Ngpus >= 2) {
hipStream_t strm;
HIP_CHECK(hipStreamCreate(&strm));
HIP_CHECK(hipMallocManaged(&Hmm, 2*4096));
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0));
HIP_CHECK(hipMemPrefetchAsync(Hmm, 2*4096, 1, strm));
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
hipMemRangeAttributeAccessedBy, Hmm, 2*4096));
if (data != 0) {
WARN("Didnt get expected behavior at line: " << __LINE__);
REQUIRE(false);
}
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetAccessedBy, 0));
HIP_CHECK(hipStreamDestroy(strm));
HIP_CHECK(hipFree(Hmm));
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
/* 1) Set AccessedBy to device 0 followed by PreferredLocation to device 1
check for AccessedBy flag using hipMemRangeGetAttribute() it should
return 0
2) Unset AccessedBy to 0 and set it to device 1 followed by
PreferredLocation to device 1, check for AccessedBy flag using
hipMemRangeGetAttribute() it should return 1*/
TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg3") {
int managed = HmmAttrPrint();
if (managed == 1) {
int *Hmm = NULL, data = 999, Ngpus = 0;
HIP_CHECK(hipGetDeviceCount(&Ngpus));
if (Ngpus >= 2) {
HIP_CHECK(hipMallocManaged(&Hmm, 2*4096));
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0));
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 1));
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
hipMemRangeAttributeAccessedBy, Hmm, 2*4096));
if (data != 0) {
WARN("Didnt get expected behavior at line: " << __LINE__);
REQUIRE(false);
}
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetAccessedBy, 0));
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 1));
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 0));
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
hipMemRangeAttributeAccessedBy, Hmm, 2*4096));
if (data != 1) {
WARN("Didnt get expected behavior at line: " << __LINE__);
REQUIRE(false);
}
HIP_CHECK(hipFree(Hmm));
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
/* Set AccessedBy flag to HMM memory launch a kernel and then unset
AccessedBy, launch kernel. We should not have any access issues*/
TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg4") {
int managed = HmmAttrPrint();
if (managed == 1) {
int *Hmm = NULL, NumElms = (1024 * 1024), InitVal = 123, blockSize = 64;
int DataMismatch = 0;
hipStream_t strm;
HIP_CHECK(hipStreamCreate(&strm));
HIP_CHECK(hipMallocManaged(&Hmm, (NumElms * sizeof(int))));
HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)),
hipMemAdviseSetAccessedBy, 0));
// Initializing memory
for (int i = 0; i < NumElms; ++i) {
Hmm[i] = InitVal;
}
dim3 dimBlock(blockSize, 1, 1);
dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
// launching kernel from each one of the gpus
MemAdvise2<<<dimGrid, dimBlock, 0, strm>>>(Hmm, NumElms);
HIP_CHECK(hipStreamSynchronize(strm));
// verifying the final result
for (int i = 0; i < NumElms; ++i) {
if (Hmm[i] != (InitVal + 10)) {
DataMismatch++;
}
}
if (DataMismatch != 0) {
WARN("DataMismatch is observed at line: " << __LINE__);
REQUIRE(false);
}
HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)),
hipMemAdviseUnsetAccessedBy, 0));
MemAdvise2<<<dimGrid, dimBlock, 0, strm>>>(Hmm, NumElms);
HIP_CHECK(hipStreamSynchronize(strm));
// verifying the final result
for (int i = 0; i < NumElms; ++i) {
if (Hmm[i] != (InitVal + (2*10))) {
DataMismatch++;
}
}
if (DataMismatch != 0) {
WARN("DataMismatch is observed at line: " << __LINE__);
REQUIRE(false);
}
HIP_CHECK(hipFree(Hmm));
HIP_CHECK(hipStreamDestroy(strm));
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
/* Allocate memory using aligned_alloc(), assign PreferredLocation flag to
the allocated memory and launch a kernel. Kernel should get executed
successfully without hang or segfault*/
#if __linux__ && HT_AMD
TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") {
if ((setenv("HSA_XNACK", "1", 1)) != 0) {
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
REQUIRE(false);
}
// The following code block checks for gfx90a so as to skip if the device is not MI200
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
int stat = 0;
if (fork() == 0) {
// The below part should be inside fork
int managedMem = 0, pageMemAccess = 0;
HIP_CHECK(hipDeviceGetAttribute(&pageMemAccess,
hipDeviceAttributePageableMemoryAccess, 0));
WARN("hipDeviceAttributePageableMemoryAccess:" << pageMemAccess);
HIP_CHECK(hipDeviceGetAttribute(&managedMem, hipDeviceAttributeManagedMemory, 0));
WARN("hipDeviceAttributeManagedMemory: " << managedMem);
if ((managedMem == 1) && (pageMemAccess == 1)) {
int *Mllc = nullptr, MemSz = 4096 * 4, NumElms = 4096, InitVal = 123;
// Mllc = reinterpret_cast<(int *)>(aligned_alloc(4096, MemSz));
Mllc = reinterpret_cast<int*>(aligned_alloc(4096, 4096*4));
for (int i = 0; i < NumElms; ++i) {
Mllc[i] = InitVal;
}
hipStream_t strm;
int DataMismatch = 0;
HIP_CHECK(hipStreamCreate(&strm));
// The following hipMemAdvise() call is made to know if advise on
// aligned_alloc() is causing any issue
HIP_CHECK(hipMemAdvise(Mllc, MemSz, hipMemAdviseSetPreferredLocation, 0));
HIP_CHECK(hipMemPrefetchAsync(Mllc, MemSz, 0, strm));
HIP_CHECK(hipStreamSynchronize(strm));
MemAdvise2<<<(NumElms/32), 32, 0, strm>>>(Mllc, NumElms);
HIP_CHECK(hipStreamSynchronize(strm));
for (int i = 0; i < NumElms; ++i) {
if (Mllc[i] != (InitVal + 10)) {
DataMismatch++;
}
}
if (DataMismatch != 0) {
WARN("DataMismatch observed!!");
exit(9); // 9 for failure
} else {
exit(10); // 10 for Pass result
}
} else {
SUCCEED("GPU 0 doesn't support ManagedMemory with hipDeviceAttributePageableMemoryAccess "
"attribute. Hence skipping the testing with Pass result.\n");
exit(Catch::ResultDisposition::ContinueOnFailure);
}
} else {
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result == Catch::ResultDisposition::ContinueOnFailure) {
WARN("GPU 0 doesn't support ManagedMemory with hipDeviceAttributePageableMemoryAccess "
"attribute. Hence skipping the testing with Pass result.\n");
} else {
if (Result != 10) {
REQUIRE(false);
}
}
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
WARN("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
#endif
/* Allocate Hmm memory, set advise to PreferredLocation and then get
attribute using the api hipMemRangeGetAttribute() for
hipMemRangeAttributeLastPrefetchLocation the value returned should be -2*/
TEST_CASE("Unit_hipMemAdvise_TstMemAdvisePrefrdLoc") {
int managed = HmmAttrPrint();
if (managed == 1) {
int *Hmm = NULL, data = 999;
HIP_CHECK(hipMallocManaged(&Hmm, 4096));
HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0));
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
hipMemRangeAttributeLastPrefetchLocation,
Hmm, 4096));
if (data != -2) {
WARN("Didnt receive expected value.");
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
/* Allocate HMM memory, set PreferredLocation to device 0, Prfetch the mem
to device1, probe for hipMemRangeAttributeLastPrefetchLocation using
hipMemRangeGetAttribute(), we should get 1*/
TEST_CASE("Unit_hipMemAdvise_TstMemAdviseLstPreftchLoc") {
int NumDevs = 0;
HIP_CHECK(hipGetDeviceCount(&NumDevs));
if (NumDevs >= 2) {
int managed = HmmAttrPrint();
if (managed == 1) {
int *Hmm = NULL, data = 999;
hipStream_t strm;
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipStreamCreate(&strm));
HIP_CHECK(hipMallocManaged(&Hmm, 4096));
HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0));
HIP_CHECK(hipMemPrefetchAsync(Hmm, 4096, 1, strm));
HIP_CHECK(hipStreamSynchronize(strm));
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
hipMemRangeAttributeLastPrefetchLocation,
Hmm, 4096));
if (data != 1) {
WARN("Didnt receive expected value!!");
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
} else {
SUCCEED("This system has less than 2 gpus hence skipping the test.\n");
}
}
/* Allocate HMM memory, set ReadMostly followed by PreferredLocation, probe
for hipMemRangeAttributeReadMostly and hipMemRangeAttributePreferredLocation
using hipMemRangeGetAttribute() we should observe 1 and 0 correspondingly.
In other words setting of hipMemRangeAttributePreferredLocation should not
impact hipMemRangeAttributeReadMostly advise to the memory*/
TEST_CASE("Unit_hipMemAdvise_TstMemAdviseMultiFlag") {
int managed = HmmAttrPrint();
if (managed == 1) {
int *Hmm = NULL, data = 999;
HIP_CHECK(hipMallocManaged(&Hmm, 4096));
HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetReadMostly, 0));
HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0));
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
hipMemRangeAttributeReadMostly, Hmm,
4096));
if (data != 1) {
WARN("Didnt receive expected value at line: " << data);
REQUIRE(false);
}
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
hipMemRangeAttributePreferredLocation, Hmm,
4096));
if (data != 0) {
WARN("Didnt receive expected value at line: " << data);
REQUIRE(false);
}
HIP_CHECK(hipFree(Hmm));
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
/*Allocate Hmm memory, advise it to ReadMostly for gpu: 0 and launch kernel
on all other gpus except 0. This test case may discover any effect or
access denial case arising due to setting ReadMostly only to a particular
gpu*/
TEST_CASE("Unit_hipMemAdvise_ReadMosltyMgpuTst") {
int managed = HmmAttrPrint();
if (managed == 1) {
int Ngpus = 0;
HIP_CHECK(hipGetDeviceCount(&Ngpus));
if (Ngpus < 2) {
SUCCEED("This test needs atleast two gpus to run."
"Hence skipping the test.\n");
}
int *Hmm = NULL, NumElms = (1024 * 1024), InitVal = 123, blockSize = 64;
int *Hmm1 = NULL, DataMismatch = 0;
hipStream_t strm;
HIP_CHECK(hipStreamCreate(&strm));
HIP_CHECK(hipMallocManaged(&Hmm, (NumElms * sizeof(int))));
// Initializing memory
for (int i = 0; i < NumElms; ++i) {
Hmm[i] = InitVal;
}
HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)),
hipMemAdviseSetReadMostly, 0));
dim3 dimBlock(blockSize, 1, 1);
dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
#if HT_AMD
SECTION("Launch Kernel on all other gpus") {
// launching kernel from each one of the gpus
for (int i = 1; i < Ngpus; ++i) {
DataMismatch = 0;
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipMallocManaged(&Hmm1, (NumElms * sizeof(int))));
MemAdvise3<<<dimGrid, dimBlock, 0, strm>>>(Hmm, Hmm1, NumElms);
HIP_CHECK(hipStreamSynchronize(strm));
// verifying the results
for (int j = 0; j < NumElms; ++j) {
if (Hmm1[j] != (InitVal + 10)) {
DataMismatch++;
}
}
if (DataMismatch != 0) {
WARN("DataMismatch is observed with the gpu: " << i);
REQUIRE(false);
}
HIP_CHECK(hipFree(Hmm1));
}
}
SECTION("Launch Kernel on all other gpus and manipulate the content") {
for (int i = 0; i < Ngpus; ++i) {
DataMismatch = 0;
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)),
hipMemAdviseSetReadMostly, i));
MemAdvise2<<<dimGrid, dimBlock, 0, strm>>>(Hmm, NumElms);
HIP_CHECK(hipStreamSynchronize(strm));
}
// verifying the final result
for (int i = 0; i < NumElms; ++i) {
if (Hmm[i] != (InitVal + Ngpus * 10)) {
DataMismatch++;
}
}
if (DataMismatch != 0) {
WARN("DataMismatch is observed at line: " << __LINE__);
REQUIRE(false);
}
}
#endif
HIP_CHECK(hipFree(Hmm));
HIP_CHECK(hipStreamDestroy(strm));
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
TEST_CASE("Unit_hipMemAdvise_TstSetUnsetPrfrdLoc") {
int managed = HmmAttrPrint();
if (managed == 1) {
int *Hmm = NULL, data = 999;
HIP_CHECK(hipMallocManaged(&Hmm, 2*4096));
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 0));
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
hipMemRangeAttributePreferredLocation, Hmm, 2*4096));
if (data != 0) {
WARN("Didnt receive expected value!!");
REQUIRE(false);
}
HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetPreferredLocation, 0));
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
hipMemRangeAttributePreferredLocation, Hmm, 2*4096));
if (data != -2) {
WARN("Didnt receive expected value!!");
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
catch/unit/memory/hipPointerGetAttributes.cc
-7
Näytä tiedosto
@@ -260,8 +260,6 @@ TEST_CASE("Unit_hipPointerGetAttributes_Basic") {
REQUIRE(attribs.devicePointer != attribs2.devicePointer);
}
HIP_CHECK(hipFree(A_d));
e = hipPointerGetAttributes(&attribs, A_d);
REQUIRE(e == hipErrorInvalidValue);
// Device-visible host memory
printf("\nDevice-visible host memory (hipHostMalloc)\n");
@@ -272,15 +270,10 @@ TEST_CASE("Unit_hipPointerGetAttributes_Basic") {
char* ptr1 = reinterpret_cast<char*>(attribs.hostPointer);
REQUIRE((ptr1 + Nbytes / 2) == reinterpret_cast<char*>(attribs2.hostPointer));
HIP_CHECK(hipHostFree(A_Pinned_h));
e = hipPointerGetAttributes(&attribs, A_Pinned_h);
REQUIRE(e == hipErrorInvalidValue);
// OS memory
printf("\nOS-allocated memory (malloc)\n");
e = hipPointerGetAttributes(&attribs, A_OSAlloc_h);
REQUIRE(e == hipErrorInvalidValue);
}
TEST_CASE("Unit_hipPointerGetAttributes_ClusterAlloc") {
catch/unit/vulkan_interop/CMakeLists.txt
+6
Näytä tiedosto
@@ -1,6 +1,12 @@
set(TEST_SRC
vulkan_test.cc
hipExternalMemoryGetMappedBuffer.cc
hipImportExternalMemory.cc
hipDestroyExternalMemory.cc
hipWaitExternalSemaphoresAsync.cc
hipSignalExternalSemaphoresAsync.cc
hipImportExternalSemaphore.cc
hipDestroyExternalSemaphore.cc
)
find_package(Vulkan)
catch/unit/vulkan_interop/hipDestroyExternalMemory.cc
+45
Näytä tiedosto
@@ -0,0 +1,45 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "vulkan_test.hh"
constexpr bool enable_validation = false;
TEST_CASE("Unit_hipDestroyExternalMemory_Vulkan_Negative_Parameters") {
SECTION("extMem == nullptr") {
HIP_CHECK_ERROR(hipDestroyExternalMemory(nullptr), hipErrorInvalidValue);
}
// Segfaults in CUDA
// Disabled on AMD due to defect - EXSWHTEC-187
#if HT_AMD && 0
SECTION("Double free") {
VulkanTest vkt(enable_validation);
const auto storage = vkt.CreateMappedStorage<int>(1, VK_BUFFER_USAGE_TRANSFER_DST_BIT, true);
auto desc = vkt.BuildMemoryDescriptor(storage.memory, sizeof(*storage.host_ptr));
hipExternalMemory_t ext_memory;
HIP_CHECK(hipImportExternalMemory(&ext_memory, &desc));
HIP_CHECK(hipDestroyExternalMemory(ext_memory));
HIP_CHECK_ERROR(hipDestroyExternalMemory(ext_memory), hipErrorInvalidValue);
}
#endif
}
catch/unit/vulkan_interop/hipDestroyExternalSemaphore.cc
+40
Näytä tiedosto
@@ -0,0 +1,40 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "vulkan_test.hh"
constexpr bool enable_validation = false;
TEST_CASE("Unit_hipDestroyExternalSemaphore_Vulkan_Negative_Parameters") {
SECTION("extSem == nullptr") {
HIP_CHECK_ERROR(hipDestroyExternalSemaphore(nullptr), hipErrorInvalidValue);
}
// Segfaults in CUDA
#if HT_AMD
SECTION("Double free") {
VulkanTest vkt(enable_validation);
const auto ext_semaphore = ImportBinarySemaphore(vkt);
HIP_CHECK(hipDestroyExternalSemaphore(ext_semaphore));
HIP_CHECK_ERROR(hipDestroyExternalSemaphore(ext_semaphore), hipErrorInvalidValue);
}
#endif
}
catch/unit/vulkan_interop/hipImportExternalMemory.cc
+76
Näytä tiedosto
@@ -0,0 +1,76 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "vulkan_test.hh"
constexpr bool enable_validation = false;
TEST_CASE("Unit_hipImportExternalMemory_Vulkan_Negative_Parameters") {
VulkanTest vkt(enable_validation);
const auto storage = vkt.CreateMappedStorage<int>(1, VK_BUFFER_USAGE_TRANSFER_DST_BIT, true);
auto desc = vkt.BuildMemoryDescriptor(storage.memory, sizeof(*storage.host_ptr));
hipExternalMemory_t ext_memory;
// Disabled due to defect - EXSWHTEC-182
#if HT_NVIDIA
SECTION("extMem_out == nullptr") {
HIP_CHECK_ERROR(hipImportExternalMemory(nullptr, &desc), hipErrorInvalidValue);
}
#endif
// Disabled due to defect - EXSWHTEC-183
#if HT_NVIDIA
SECTION("memHandleDesc == nullptr") {
HIP_CHECK_ERROR(hipImportExternalMemory(&ext_memory, nullptr), hipErrorInvalidValue);
}
#endif
// Disabled due to defect - EXSWHTEC-185
#if HT_NVIDIA
SECTION("memHandleDesc.size == 0") {
desc.size = 0;
HIP_CHECK_ERROR(hipImportExternalMemory(&ext_memory, &desc), hipErrorInvalidValue);
}
#endif
// Disabled due to defect - EXSWHTEC-186
#if HT_NVIDIA
SECTION("Invalid memHandleDesc.flags") {
desc.flags = 2;
HIP_CHECK_ERROR(hipImportExternalMemory(&ext_memory, &desc), hipErrorInvalidValue);
}
#endif
// Disabled due to defect - EXSWHTEC-184
#if HT_NVIDIA
SECTION("Invalid memHandleDesc.type") {
desc.type = static_cast<hipExternalMemoryHandleType>(-1);
HIP_CHECK_ERROR(hipImportExternalMemory(&ext_memory, &desc), hipErrorInvalidValue);
}
#endif
#ifdef _WIN32
SECTION("memHandleDesc.handle == NULL") {
desc.handle.win32.handle = NULL;
HIP_CHECK_ERROR(hipImportExternalMemory(&ext_memory, &desc), hipErrorInvalidValue);
}
#endif
}
catch/unit/vulkan_interop/hipImportExternalSemaphore.cc
+56
Näytä tiedosto
@@ -0,0 +1,56 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "vulkan_test.hh"
constexpr bool enable_validation = false;
TEST_CASE("Unit_hipImportExternalSemaphore_Vulkan_Negative_Parameters") {
VulkanTest vkt(enable_validation);
const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY);
auto handle_desc = vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_BINARY);
hipExternalSemaphore_t ext_semaphore;
SECTION("extSem_out == nullptr") {
HIP_CHECK_ERROR(hipImportExternalSemaphore(nullptr, &handle_desc), hipErrorInvalidValue);
}
SECTION("semHandleDesc == nullptr") {
HIP_CHECK_ERROR(hipImportExternalSemaphore(&ext_semaphore, nullptr), hipErrorInvalidValue);
}
SECTION("semHandleDesc.flags != 0") {
handle_desc.flags = 1;
HIP_CHECK_ERROR(hipImportExternalSemaphore(&ext_semaphore, &handle_desc), hipErrorInvalidValue);
}
SECTION("Invalid semHandleDesc.type") {
handle_desc.type = static_cast<hipExternalSemaphoreHandleType>(-1);
HIP_CHECK_ERROR(hipImportExternalSemaphore(&ext_semaphore, &handle_desc), hipErrorInvalidValue);
}
#ifdef _WIN32
SECTION("semHandleDesc.handle == NULL") {
handle_desc.handle.win32.handle = NULL;
HIP_CHECK_ERROR(hipImportExternalSemaphore(&ext_semaphore, &handle_desc), hipErrorInvalidValue);
}
#endif
}
catch/unit/vulkan_interop/hipSignalExternalSemaphoresAsync.cc
+204
Näytä tiedosto
@@ -0,0 +1,204 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "vulkan_test.hh"
constexpr bool enable_validation = false;
TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Binary_Semaphore") {
VulkanTest vkt(enable_validation);
constexpr uint32_t count = 1;
const auto src_storage = vkt.CreateMappedStorage<int>(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT);
const auto dst_storage = vkt.CreateMappedStorage<int>(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT);
const auto command_buffer = vkt.GetCommandBuffer();
VkCommandBufferBeginInfo begin_info = {};
begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO;
begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT;
VK_CHECK_RESULT(vkBeginCommandBuffer(command_buffer, &begin_info));
VkBufferCopy buffer_copy = {};
buffer_copy.size = count * sizeof(*src_storage.host_ptr);
vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy);
VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer));
const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY);
const auto hip_sem_handle_desc =
vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_BINARY);
hipExternalSemaphore_t hip_ext_semaphore;
HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc));
VkSubmitInfo submit_info = {};
submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO;
submit_info.commandBufferCount = 1;
submit_info.pCommandBuffers = &command_buffer;
submit_info.waitSemaphoreCount = 1;
submit_info.pWaitSemaphores = &semaphore;
const auto fence = vkt.CreateFence();
VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence));
REQUIRE(vkGetFenceStatus(vkt.GetDevice(), fence) == VK_NOT_READY);
hipExternalSemaphoreSignalParams signal_params = {};
signal_params.params.fence.value = 0;
HIP_CHECK(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, nullptr));
PollStream(nullptr, hipSuccess);
VK_CHECK_RESULT(
vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/));
HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore));
}
// Timeline semaphores unsupported on AMD
#if HT_NVIDIA
TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Timeline_Semaphore") {
VulkanTest vkt(enable_validation);
constexpr uint64_t signal_value = 2;
const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_TIMELINE);
const auto hip_sem_handle_desc =
vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_TIMELINE);
hipExternalSemaphore_t hip_ext_semaphore;
HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc));
hipExternalSemaphoreSignalParams signal_params = {};
signal_params.params.fence.value = signal_value;
HIP_CHECK(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, nullptr));
PollStream(nullptr, hipSuccess);
uint64_t sem_value = 0u;
VK_CHECK_RESULT(vkGetSemaphoreCounterValue(vkt.GetDevice(), semaphore, &sem_value));
REQUIRE(2 == sem_value);
HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore));
}
#endif
TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Multiple_Semaphores") {
VulkanTest vkt(enable_validation);
constexpr uint32_t count = 1;
const auto src_storage = vkt.CreateMappedStorage<int>(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT);
const auto dst_storage = vkt.CreateMappedStorage<int>(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT);
#if HT_AMD
constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_BINARY;
#else
constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_TIMELINE;
#endif
const auto command_buffer = vkt.GetCommandBuffer();
VkCommandBufferBeginInfo begin_info = {};
begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO;
begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT;
VK_CHECK_RESULT(vkBeginCommandBuffer(command_buffer, &begin_info));
VkBufferCopy buffer_copy = {};
buffer_copy.size = count * sizeof(*src_storage.host_ptr);
vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy);
VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer));
const auto binary_semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY);
const auto hip_binary_sem_handle_desc =
vkt.BuildSemaphoreDescriptor(binary_semaphore, VK_SEMAPHORE_TYPE_BINARY);
hipExternalSemaphore_t hip_binary_ext_semaphore;
HIP_CHECK(hipImportExternalSemaphore(&hip_binary_ext_semaphore, &hip_binary_sem_handle_desc));
const auto timeline_semaphore = vkt.CreateExternalSemaphore(second_semaphore_type);
const auto hip_timeline_sem_handle_desc =
vkt.BuildSemaphoreDescriptor(timeline_semaphore, second_semaphore_type);
hipExternalSemaphore_t hip_timeline_ext_semaphore;
HIP_CHECK(hipImportExternalSemaphore(&hip_timeline_ext_semaphore, &hip_timeline_sem_handle_desc));
uint64_t wait_values[] = {1, 0};
VkTimelineSemaphoreSubmitInfo timeline_submit_info = {};
timeline_submit_info.sType = VK_STRUCTURE_TYPE_TIMELINE_SEMAPHORE_SUBMIT_INFO;
timeline_submit_info.waitSemaphoreValueCount = 2;
timeline_submit_info.pWaitSemaphoreValues = wait_values;
VkSemaphore wait_semaphores[] = {timeline_semaphore, binary_semaphore};
VkSubmitInfo submit_info = {};
submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO;
submit_info.commandBufferCount = 1;
submit_info.pCommandBuffers = &command_buffer;
submit_info.waitSemaphoreCount = 2;
submit_info.pWaitSemaphores = wait_semaphores;
submit_info.pNext =
second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? &timeline_submit_info : nullptr;
const auto fence = vkt.CreateFence();
VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence));
REQUIRE(vkGetFenceStatus(vkt.GetDevice(), fence) == VK_NOT_READY);
hipExternalSemaphoreSignalParams binary_signal_params = {};
binary_signal_params.params.fence.value = 0;
hipExternalSemaphoreSignalParams timeline_signal_params = {};
timeline_signal_params.params.fence.value =
second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? 2 : 0;
hipExternalSemaphore_t ext_semaphores[] = {hip_binary_ext_semaphore, hip_timeline_ext_semaphore};
hipExternalSemaphoreSignalParams signal_params[] = {binary_signal_params, timeline_signal_params};
HIP_CHECK(hipSignalExternalSemaphoresAsync(ext_semaphores, signal_params, 2, nullptr));
VK_CHECK_RESULT(
vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/));
HIP_CHECK(hipDestroyExternalSemaphore(hip_binary_ext_semaphore));
HIP_CHECK(hipDestroyExternalSemaphore(hip_timeline_ext_semaphore));
}
TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Negative_Parameters") {
VulkanTest vkt(enable_validation);
hipExternalSemaphoreSignalParams signal_params = {};
signal_params.params.fence.value = 1;
SECTION("extSemArray == nullptr") {
HIP_CHECK_ERROR(hipSignalExternalSemaphoresAsync(nullptr, &signal_params, 1, nullptr),
hipErrorInvalidValue);
}
SECTION("paramsArray == nullptr") {
const auto hip_ext_semaphore = ImportBinarySemaphore(vkt);
HIP_CHECK_ERROR(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, nullptr, 1, nullptr),
hipErrorInvalidValue);
HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore));
}
SECTION("Wait params flags != 0") {
const auto hip_ext_semaphore = ImportBinarySemaphore(vkt);
signal_params.flags = 1;
HIP_CHECK_ERROR(
hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, nullptr),
hipErrorInvalidValue);
HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore));
}
SECTION("Invalid stream") {
const auto hip_ext_semaphore = ImportBinarySemaphore(vkt);
hipStream_t stream = nullptr;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK_ERROR(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, stream),
hipErrorInvalidValue);
HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore));
}
}
catch/unit/vulkan_interop/hipWaitExternalSemaphoresAsync.cc
+231
Näytä tiedosto
@@ -0,0 +1,231 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "vulkan_test.hh"
constexpr bool enable_validation = false;
TEST_CASE("Unit_hipWaitExternalSemaphoresAsync_Vulkan_Positive_Binary_Semaphore") {
VulkanTest vkt(enable_validation);
constexpr uint32_t count = 1;
const auto src_storage = vkt.CreateMappedStorage<int>(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT);
const auto dst_storage = vkt.CreateMappedStorage<int>(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT);
const auto command_buffer = vkt.GetCommandBuffer();
VkCommandBufferBeginInfo begin_info = {};
begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO;
begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT;
VK_CHECK_RESULT(vkBeginCommandBuffer(command_buffer, &begin_info));
VkBufferCopy buffer_copy = {};
buffer_copy.size = count * sizeof(*src_storage.host_ptr);
vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy);
VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer));
const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY);
const auto hip_sem_handle_desc =
vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_BINARY);
hipExternalSemaphore_t hip_ext_semaphore;
HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc));
hipExternalSemaphoreWaitParams hip_ext_semaphore_wait_params = {};
hip_ext_semaphore_wait_params.flags = 0;
hip_ext_semaphore_wait_params.params.fence.value = 0;
HIP_CHECK(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, &hip_ext_semaphore_wait_params, 1,
nullptr));
PollStream(nullptr, hipErrorNotReady);
VkSubmitInfo submit_info = {};
submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO;
submit_info.commandBufferCount = 1;
submit_info.pCommandBuffers = &command_buffer;
submit_info.signalSemaphoreCount = 1;
submit_info.pSignalSemaphores = &semaphore;
*src_storage.host_ptr = 42;
const auto fence = vkt.CreateFence();
VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence));
VK_CHECK_RESULT(
vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/));
PollStream(nullptr, hipSuccess);
REQUIRE(42 == *dst_storage.host_ptr);
HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore));
}
// Timeline semaphores unsupported on AMD
#if HT_NVIDIA
TEST_CASE("Unit_hipWaitExternalSemaphoresAsync_Vulkan_Positive_Timeline_Semaphore") {
VulkanTest vkt(enable_validation);
const auto [wait_value, signal_value] =
GENERATE(std::make_pair(2, 2), std::make_pair(2, 3), std::make_pair(3, 2));
INFO("Wait value: " << wait_value << ", signal value: " << signal_value);
const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_TIMELINE);
const auto hip_sem_handle_desc =
vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_TIMELINE);
hipExternalSemaphore_t hip_ext_semaphore;
HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc));
hipExternalSemaphoreWaitParams hip_ext_semaphore_wait_params = {};
hip_ext_semaphore_wait_params.flags = 0;
hip_ext_semaphore_wait_params.params.fence.value = wait_value;
HIP_CHECK(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, &hip_ext_semaphore_wait_params, 1,
nullptr));
PollStream(nullptr, hipErrorNotReady);
VkSemaphoreSignalInfo signal_info = {};
signal_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_SIGNAL_INFO;
signal_info.semaphore = semaphore;
signal_info.value = signal_value;
VK_CHECK_RESULT(vkSignalSemaphore(vkt.GetDevice(), &signal_info));
if (wait_value > signal_value) {
PollStream(nullptr, hipErrorNotReady);
signal_info.value = wait_value;
VK_CHECK_RESULT(vkSignalSemaphore(vkt.GetDevice(), &signal_info));
}
PollStream(nullptr, hipSuccess);
HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore));
}
#endif
TEST_CASE("Unit_hipWaitExternalSemaphoresAsync_Vulkan_Positive_Multiple_Semaphores") {
VulkanTest vkt(enable_validation);
#if HT_AMD
constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_BINARY;
#else
constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_TIMELINE;
#endif
constexpr uint32_t count = 1;
const auto src_storage = vkt.CreateMappedStorage<int>(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT);
const auto dst_storage = vkt.CreateMappedStorage<int>(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT);
const auto command_buffer = vkt.GetCommandBuffer();
VkCommandBufferBeginInfo begin_info = {};
begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO;
begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT;
VK_CHECK_RESULT(vkBeginCommandBuffer(command_buffer, &begin_info));
VkBufferCopy buffer_copy = {};
buffer_copy.size = count * sizeof(*src_storage.host_ptr);
vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy);
VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer));
const auto binary_semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY);
const auto hip_binary_sem_handle_desc =
vkt.BuildSemaphoreDescriptor(binary_semaphore, VK_SEMAPHORE_TYPE_BINARY);
hipExternalSemaphore_t hip_binary_ext_semaphore;
HIP_CHECK(hipImportExternalSemaphore(&hip_binary_ext_semaphore, &hip_binary_sem_handle_desc));
const auto timeline_semaphore = vkt.CreateExternalSemaphore(second_semaphore_type);
const auto hip_timeline_sem_handle_desc =
vkt.BuildSemaphoreDescriptor(timeline_semaphore, second_semaphore_type);
hipExternalSemaphore_t hip_timeline_ext_semaphore;
HIP_CHECK(hipImportExternalSemaphore(&hip_timeline_ext_semaphore, &hip_timeline_sem_handle_desc));
hipExternalSemaphoreWaitParams binary_semaphore_wait_params = {};
binary_semaphore_wait_params.params.fence.value = 0;
hipExternalSemaphoreWaitParams timeline_semaphore_wait_params = {};
timeline_semaphore_wait_params.params.fence.value =
second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? 1 : 0;
hipExternalSemaphore_t ext_semaphores[] = {hip_binary_ext_semaphore, hip_timeline_ext_semaphore};
hipExternalSemaphoreWaitParams wait_params[] = {binary_semaphore_wait_params,
timeline_semaphore_wait_params};
HIP_CHECK(hipWaitExternalSemaphoresAsync(ext_semaphores, wait_params, 2, nullptr));
PollStream(nullptr, hipErrorNotReady);
if (second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE) {
VkSemaphoreSignalInfo signal_info = {};
signal_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_SIGNAL_INFO;
signal_info.semaphore = timeline_semaphore;
signal_info.value = 1;
VK_CHECK_RESULT(vkSignalSemaphore(vkt.GetDevice(), &signal_info));
PollStream(nullptr, hipErrorNotReady);
}
VkSubmitInfo submit_info = {};
VkSemaphore signal_semaphores[] = {binary_semaphore, timeline_semaphore};
submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO;
submit_info.commandBufferCount = 1;
submit_info.pCommandBuffers = &command_buffer;
submit_info.signalSemaphoreCount = second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? 1 : 2;
submit_info.pSignalSemaphores =
second_semaphore_type == VK_SEMAPHORE_TYPE_MAX_ENUM ? &binary_semaphore : signal_semaphores;
const auto fence = vkt.CreateFence();
VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence));
VK_CHECK_RESULT(
vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/));
PollStream(nullptr, hipSuccess);
HIP_CHECK(hipDestroyExternalSemaphore(hip_timeline_ext_semaphore));
HIP_CHECK(hipDestroyExternalSemaphore(hip_binary_ext_semaphore));
}
TEST_CASE("Unit_hipWaitExternalSemaphoresAsync_Vulkan_Negative_Parameters") {
VulkanTest vkt(enable_validation);
hipExternalSemaphoreWaitParams wait_params = {};
wait_params.params.fence.value = 1;
SECTION("extSemArray == nullptr") {
HIP_CHECK_ERROR(hipWaitExternalSemaphoresAsync(nullptr, &wait_params, 1, nullptr),
hipErrorInvalidValue);
}
SECTION("paramsArray == nullptr") {
const auto hip_ext_semaphore = ImportBinarySemaphore(vkt);
HIP_CHECK_ERROR(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, nullptr, 1, nullptr),
hipErrorInvalidValue);
HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore));
}
SECTION("Wait params flag != 0") {
const auto hip_ext_semaphore = ImportBinarySemaphore(vkt);
wait_params.flags = 1;
HIP_CHECK_ERROR(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, &wait_params, 1, nullptr),
hipErrorInvalidValue);
HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore));
}
SECTION("Invalid stream") {
const auto hip_ext_semaphore = ImportBinarySemaphore(vkt);
hipStream_t stream = nullptr;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK_ERROR(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, &wait_params, 1, stream),
hipErrorInvalidValue);
HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore));
}
}
samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt
+3 -3
Näytä tiedosto
@@ -25,12 +25,12 @@ endif ()
if(NOT DEFINED HIP_PATH)
if(NOT DEFINED ENV{HIP_PATH})
set(HIP_PATH "${ROCM_PATH}/hip" CACHE PATH "Path to which HIP has been installed")
set(HIP_PATH ${ROCM_PATH} CACHE PATH "Path to which HIP has been installed")
else()
set(HIP_PATH $ENV{HIP_PATH} CACHE PATH "Path to which HIP has been installed")
endif()
endif()
set(CMAKE_MODULE_PATH "${HIP_PATH}/cmake" ${CMAKE_MODULE_PATH})
set(CMAKE_MODULE_PATH "${HIP_PATH}/hip/cmake" ${CMAKE_MODULE_PATH})
set(CMAKE_HIP_ARCHITECTURES OFF)
project(12_cmake)
@@ -53,7 +53,7 @@ set_source_files_properties(${MY_SOURCE_FILES} PROPERTIES HIP_SOURCE_PROPERTY_FO
hip_add_executable(${MY_TARGET_NAME} ${MY_SOURCE_FILES} HIPCC_OPTIONS ${MY_HIPCC_OPTIONS} HCC_OPTIONS ${MY_HCC_OPTIONS} CLANG_OPTIONS ${MY_CLANG_OPTIONS} NVCC_OPTIONS ${MY_NVCC_OPTIONS})
# Search for rocm in common locations
list(APPEND CMAKE_PREFIX_PATH ${HIP_PATH} ${ROCM_PATH})
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
find_package(hip QUIET)
if(TARGET hip::host)
message(STATUS "Found hip::host at ${hip_DIR}")
samples/2_Cookbook/16_assembly_to_executable/Makefile
+4 -4
Näytä tiedosto
@@ -21,15 +21,15 @@ ifeq ($(OS),Windows_NT)
$(error Makefile is not supported on windows platform. Please use cmake instead to build sample.)
endif
ROCM_PATH?= $(wildcard /opt/rocm/)
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
HIP_PATH?= $(ROCM_PATH)
ifeq (,$(HIP_PATH))
HIP_PATH=../../..
endif
HIPCC=$(HIP_PATH)/bin/hipcc
CLANG=$(HIP_PATH)/../llvm/bin/clang
LLVM_MC=$(HIP_PATH)/../llvm/bin/llvm-mc
CLANG_OFFLOAD_BUNDLER=$(HIP_PATH)/../llvm/bin/clang-offload-bundler
CLANG=$(HIP_PATH)/llvm/bin/clang
LLVM_MC=$(HIP_PATH)/llvm/bin/llvm-mc
CLANG_OFFLOAD_BUNDLER=$(HIP_PATH)/llvm/bin/clang-offload-bundler
SRCS=square.cpp
samples/2_Cookbook/17_llvm_ir_to_executable/Makefile
+6 -6
Näytä tiedosto
@@ -21,17 +21,17 @@ ifeq ($(OS),Windows_NT)
$(error Makefile is not supported on windows platform. Please use cmake instead to build sample.)
endif
ROCM_PATH?= $(wildcard /opt/rocm/)
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
HIP_PATH?= $(ROCM_PATH)
ifeq (,$(HIP_PATH))
HIP_PATH=../../..
endif
HIPCC=$(HIP_PATH)/bin/hipcc
CLANG=$(HIP_PATH)/../llvm/bin/clang
LLVM_MC=$(HIP_PATH)/../llvm/bin/llvm-mc
CLANG_OFFLOAD_BUNDLER=$(HIP_PATH)/../llvm/bin/clang-offload-bundler
LLVM_AS=$(HIP_PATH)/../llvm/bin/llvm-as
LLVM_DIS=$(HIP_PATH)/../llvm/bin/llvm-dis
CLANG=$(HIP_PATH)/llvm/bin/clang
LLVM_MC=$(HIP_PATH)/llvm/bin/llvm-mc
CLANG_OFFLOAD_BUNDLER=$(HIP_PATH)/llvm/bin/clang-offload-bundler
LLVM_AS=$(HIP_PATH)/llvm/bin/llvm-as
LLVM_DIS=$(HIP_PATH)/llvm/bin/llvm-dis
SRCS=square.cpp