AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-299127 - Merge 'develop' into 'amd-staging'

Browse Source 
Change-Id: Ia52c6c4db36af615ffb50a61a54c569a03e81a54
This commit is contained in:
Jenkins
2022-03-10 19:12:39 -05:00
parent 8699a20c2b c52a81e96f
commit f5c415e0d3
10 changed files with 150 additions and 63 deletions
Download Patch File Download Diff File Expand all files Collapse all files
bin/hipcc.pl
+1 -1
View File
@@ -707,7 +707,7 @@ if ($HIP_PLATFORM eq "amd") {
if ($linkType eq 0) {
$toolArgs = " -L$HIP_LIB_PATH -lamdhip64 -L$ROCM_PATH/lib -lhsa-runtime64 -ldl -lnuma " . ${toolArgs};
} else {
$toolArgs = " -Wl,--enable-new-dtags -Wl,-rpath=$HIP_LIB_PATH:$ROCM_PATH/lib -lamdhip64 " . ${toolArgs};
$toolArgs = ${toolArgs} . " -Wl,--enable-new-dtags -Wl,-rpath=$HIP_LIB_PATH:$ROCM_PATH/lib -lamdhip64 ";
}
# To support __fp16 and _Float16, explicitly link with compiler-rt
$HIP_CLANG_BUILTIN_LIB="$HIP_CLANG_PATH/../lib/clang/$HIP_CLANG_VERSION/lib/$HIP_CLANG_TARGET/libclang_rt.builtins.a";
include/hip/hip_runtime_api.h
+64
View File
@@ -73,6 +73,9 @@ typedef struct {
unsigned hasDynamicParallelism : 1; ///< Dynamic parallelism.
} hipDeviceArch_t;
typedef struct hipUUID_t {
char bytes[16];
} hipUUID;
//---
// Common headers for both NVCC and HCC paths:
@@ -162,6 +165,31 @@ typedef enum hipMemoryType {
hipMemoryTypeUnified ///< Not used currently
} hipMemoryType;
/**
* @brief hipKernelNodeAttrID
* @enum
*
*/
typedef enum hipKernelNodeAttrID {
hipKernelNodeAttributeAccessPolicyWindow = 1,
hipKernelNodeAttributeCooperative = 2,
} hipKernelNodeAttrID;
typedef enum hipAccessProperty {
hipAccessPropertyNormal = 0,
hipAccessPropertyStreaming = 1,
hipAccessPropertyPersisting = 2,
} hipAccessProperty;
typedef struct hipAccessPolicyWindow {
void* base_ptr;
hipAccessProperty hitProp;
float hitRatio;
hipAccessProperty missProp;
size_t num_bytes;
} hipAccessPolicyWindow;
typedef union hipKernelNodeAttrValue {
hipAccessPolicyWindow accessPolicyWindow;
int cooperative;
} hipKernelNodeAttrValue;
/**
* Pointer attributes
@@ -998,6 +1026,18 @@ hipError_t hipDeviceComputeCapability(int* major, int* minor, hipDevice_t device
* @returns #hipSuccess, #hipErrorInavlidDevice
*/
hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device);
/**
* @brief Returns an UUID for the device.[BETA]
* @param [out] uuid
* @param [in] device
*
* @beta This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*
* @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue, #hipErrorNotInitialized,
* #hipErrorDeInitialized
*/
hipError_t hipDeviceGetUuid(hipUUID* uuid, hipDevice_t device);
/**
* @brief Returns a value for attr of link between two devices
* @param [out] value
@@ -4872,6 +4912,30 @@ hipError_t hipGraphMemcpyNodeGetParams(hipGraphNode_t node, hipMemcpy3DParms* pN
*/
hipError_t hipGraphMemcpyNodeSetParams(hipGraphNode_t node, const hipMemcpy3DParms* pNodeParams);
/**
* @brief Sets a node attribute.
*
* @param [in] hNode - instance of the node to set parameters to.
* @param [in] attr - the attribute node is set to.
* @param [in] value - const pointer to the parameters.
* @returns #hipSuccess, #hipErrorInvalidValue
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipGraphKernelNodeSetAttribute(hipGraphNode_t hNode, hipKernelNodeAttrID attr,
const hipKernelNodeAttrValue* value);
/**
* @brief Gets a node attribute.
*
* @param [in] hNode - instance of the node to set parameters to.
* @param [in] attr - the attribute node is set to.
* @param [in] value - const pointer to the parameters.
* @returns #hipSuccess, #hipErrorInvalidValue
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipGraphKernelNodeGetAttribute(hipGraphNode_t hNode, hipKernelNodeAttrID attr,
hipKernelNodeAttrValue* value);
/**
* @brief Sets the parameters for a memcpy node in the given graphExec.
*
tests/catch/multiproc/hipIpcMemAccessTest.cc
+5 -5
View File
@@ -118,12 +118,12 @@ TEST_CASE("Unit_hipIpcMemAccess_Semaphores") {
}
for (int i = 0; i < Num_devices; ++i) {
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipMalloc(&C_d, Nbytes));
HIP_CHECK(hipIpcOpenMemHandle(reinterpret_cast<void **>(&B_d),
shrd_mem->memHandle,
hipIpcMemLazyEnablePeerAccess));
HIP_CHECK(hipDeviceCanAccessPeer(&CanAccessPeer, i, shrd_mem->device));
if (CanAccessPeer == 1) {
HIP_CHECK(hipMalloc(&C_d, Nbytes));
HIP_CHECK(hipIpcOpenMemHandle(reinterpret_cast<void **>(&B_d),
shrd_mem->memHandle,
hipIpcMemLazyEnablePeerAccess));
HIP_CHECK(hipMemcpy(C_d, B_d, Nbytes, hipMemcpyDeviceToDevice));
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HipTest::checkTest<int>(A_h, C_h, N);
@@ -131,9 +131,9 @@ TEST_CASE("Unit_hipIpcMemAccess_Semaphores") {
// Checking if the data obtained from Ipc shared memory is consistent
HIP_CHECK(hipMemcpy(C_h, B_d, Nbytes, hipMemcpyDeviceToHost));
HipTest::checkTest<int>(A_h, C_h, N);
}
HIP_CHECK(hipIpcCloseMemHandle(reinterpret_cast<void*>(B_d)));
HIP_CHECK(hipFree(C_d));
}
}
if ((sem_post(sem_ob2)) == -1) {
shrd_mem->IfTestPassed = false;
tests/catch/multiproc/hipMemCoherencyTstMProc.cc
+5 -5
View File
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2021 - 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
@@ -128,7 +128,7 @@ TEST_CASE("Unit_malloc_CoherentTst") {
// Read string from child and close reading end.
read(fd1[0], GpuId, 2 * sizeof(int));
close(fd1[0]);
if ((GpuId[0] == 1) || (GpuId[0] == 1)) {
if ((GpuId[0] == 1) || (GpuId[1] == 1)) {
WARN("This test is not applicable on MI60 & MI100."
"Skipping the test!!");
exit(0);
@@ -223,7 +223,7 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") {
// Read string from child and close reading end.
read(fd1[0], GpuId, 2 * sizeof(int));
close(fd1[0]);
if ((GpuId[0] == 1) || (GpuId[0] == 1)) {
if ((GpuId[0] == 1) || (GpuId[1] == 1)) {
WARN("This test is not applicable on MI60 & MI100."
"Skipping the test!!");
exit(0);
@@ -320,7 +320,7 @@ TEST_CASE("Unit_mmap_CoherentTst") {
// Read string from child and close reading end.
read(fd1[0], GpuId, 2 * sizeof(int));
close(fd1[0]);
if ((GpuId[0] == 1) || (GpuId[0] == 1)) {
if ((GpuId[0] == 1) || (GpuId[1] == 1)) {
WARN("This test is not applicable on MI60 & MI100."
"Skipping the test!!");
exit(0);
@@ -418,7 +418,7 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
// Read string from child and close reading end.
read(fd1[0], GpuId, 2 * sizeof(int));
close(fd1[0]);
if ((GpuId[0] == 1) || (GpuId[0] == 1)) {
if ((GpuId[0] == 1) || (GpuId[1] == 1)) {
WARN("This test is not applicable on MI60 & MI100."
"Skipping the test!!");
exit(0);
tests/catch/unit/device/hipChooseDevice.cc
+1 -4
View File
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
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
@@ -41,8 +41,6 @@ TEST_CASE("Unit_hipChooseDevice_ValidateDevId") {
* Scenario1: Validates if dev = nullptr returns error code
* Scenario2: Validates if prop = nullptr returns error code
*/
#if HT_AMD
// These test scenarios fail on NVIDIA.
TEST_CASE("Unit_hipChooseDevice_NegTst") {
hipDeviceProp_t prop;
int dev = -1;
@@ -57,4 +55,3 @@ TEST_CASE("Unit_hipChooseDevice_NegTst") {
REQUIRE_FALSE(hipSuccess == hipChooseDevice(&dev, nullptr));
}
}
#endif
tests/catch/unit/device/hipGetDeviceProperties.cc
+1 -4
View File
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
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
@@ -182,14 +182,11 @@ TEST_CASE("Unit_hipGetDeviceProperties_ArchPropertiesTst") {
TEST_CASE("Unit_hipGetDeviceProperties_NegTst") {
hipDeviceProp_t prop;
#if HT_AMD
SECTION("props is nullptr") {
int device;
HIP_CHECK(hipGetDevice(&device));
// this test case results in segmentation fault on NVCC
REQUIRE_FALSE(hipSuccess == hipGetDeviceProperties(nullptr, device));
}
#endif
SECTION("device is -1") {
REQUIRE_FALSE(hipSuccess == hipGetDeviceProperties(&prop, -1));
tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc
+26 -2
View File
@@ -248,7 +248,19 @@ by creating dependency graph and instantiate, launching and verifying
the result
*/
TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraph") {
GraphInstantiateWithFlags_DependencyGraph();
int numDevices = 0;
int canAccessPeer = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
GraphInstantiateWithFlags_DependencyGraph();
} else {
SUCCEED("Machine does not seem to have P2P");
}
} else {
SUCCEED("skipped the testcase as no of devices is less than 2");
}
}
/*
This testcase verifies hipGraphInstantiateWithFlags API
@@ -278,7 +290,19 @@ by creating capture graph and instantiate, launching and verifying
the result
*/
TEST_CASE("Unit_hipGraphInstantiateWithFlags_StreamCapture") {
GraphInstantiateWithFlags_StreamCapture();
int numDevices = 0;
int canAccessPeer = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
GraphInstantiateWithFlags_StreamCapture();
} else {
SUCCEED("Machine does not seem to have P2P");
}
} else {
SUCCEED("skipped the testcase as no of devices is less than 2");
}
}
/*
tests/catch/unit/memory/hipMemcpy3DAsync.cc
+45 -39
View File
@@ -481,58 +481,64 @@ template <typename T>
void Memcpy3DAsync<T>::D2D_SameDeviceMem_StreamDiffDevice() {
HIP_CHECK(hipSetDevice(0));
// Allocating the Memory
AllocateMemory();
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipStreamCreate(&stream));
memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
SetDefaultData();
int peerAccess = 0;
HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 0, 1));
if (peerAccess) {
AllocateMemory();
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipStreamCreate(&stream));
memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
SetDefaultData();
// Host to Device
myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
myparms.dstArray = arr;
// Host to Device
myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
myparms.dstArray = arr;
#ifdef __HIP_PLATFORM_NVCC__
myparms.kind = cudaMemcpyHostToDevice;
myparms.kind = cudaMemcpyHostToDevice;
#else
myparms.kind = hipMemcpyHostToDevice;
myparms.kind = hipMemcpyHostToDevice;
#endif
REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
HIP_CHECK(hipStreamSynchronize(stream));
REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
HIP_CHECK(hipStreamSynchronize(stream));
// Array to Array
memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
SetDefaultData();
myparms.srcArray = arr;
myparms.dstArray = arr1;
// Array to Array
memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
SetDefaultData();
myparms.srcArray = arr;
myparms.dstArray = arr1;
#ifdef __HIP_PLATFORM_NVCC__
myparms.kind = cudaMemcpyDeviceToDevice;
myparms.kind = cudaMemcpyDeviceToDevice;
#else
myparms.kind = hipMemcpyDeviceToDevice;
myparms.kind = hipMemcpyDeviceToDevice;
#endif
REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
HIP_CHECK(hipStreamSynchronize(stream));
T *hOutputData = reinterpret_cast<T*>(malloc(size));
memset(hOutputData, 0, size);
REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
HIP_CHECK(hipStreamSynchronize(stream));
T *hOutputData = reinterpret_cast<T*>(malloc(size));
memset(hOutputData, 0, size);
// Device to host
memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
SetDefaultData();
myparms.dstPtr = make_hipPitchedPtr(hOutputData,
width * sizeof(T), width, height);
myparms.srcArray = arr1;
// Device to host
memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
SetDefaultData();
myparms.dstPtr = make_hipPitchedPtr(hOutputData,
width * sizeof(T), width, height);
myparms.srcArray = arr1;
#ifdef __HIP_PLATFORM_NVCC__
myparms.kind = cudaMemcpyDeviceToHost;
myparms.kind = cudaMemcpyDeviceToHost;
#else
myparms.kind = hipMemcpyDeviceToHost;
myparms.kind = hipMemcpyDeviceToHost;
#endif
REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
HIP_CHECK(hipStreamSynchronize(stream));
REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
HIP_CHECK(hipStreamSynchronize(stream));
// Validating the result
HipTest::checkArray(hData, hOutputData, width, height, depth);
// Validating the result
HipTest::checkArray(hData, hOutputData, width, height, depth);
// Deallocating the memory
free(hOutputData);
DeAllocateMemory();
// Deallocating the memory
free(hOutputData);
DeAllocateMemory();
} else {
SUCCEED("Skipped the test as there is no peer access");
}
}
/*
@@ -630,7 +636,7 @@ void Memcpy3DAsync<T>::simple_Memcpy3DAsync() {
memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
SetDefaultData();
myparms.dstPtr = make_hipPitchedPtr(hOutputData,
width * sizeof(T), width, height);
width * sizeof(T), width, height);
myparms.srcArray = arr1;
#ifdef __HIP_PLATFORM_NVCC__
myparms.kind = cudaMemcpyDeviceToHost;
tests/src/ipc/hipMultiProcIpcMem.cpp
+1 -1
View File
@@ -19,7 +19,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* TEST: %t
* TEST: %t --N 4
* HIT_END
*/
tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp
+1 -2
View File
@@ -1,5 +1,5 @@
/*
Copyright (c) 2020 - 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2020 - 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
@@ -185,7 +185,6 @@ int main(int argc, char* argv[]) {
bool testStatus = true;
HipTest::parseStandardArguments(argc, argv, false);
if (p_tests == 1) {
testStatus &= ConcurencyCheck_GlobalVar(1);
testStatus &= ConcurencyCheck_GlobalVar(0);
} else if (p_tests == 2) {
testStatus &= KernelTimeExecution();