SWDEV-348820 - Fix hipPointerAttribute_t incompatibility (#2856)
- Mark hipPointerAttribute_t.memoryType deprecated
- Instead add hipPointerAttribute_t.type
- Update testcases & documentation to use hipPointerAttribute_t.type
Change-Id: I3ba58cd2defe82102f9415bbb90cf151e08499dd
[ROCm/hip commit: 96b3b9eaf5]
This commit is contained in:
ROCm CI Service Account
committed by
GitHub
GitHub
parent
e5de7bbf55
commit
996399531c
@@ -468,31 +468,34 @@ int main()
|
||||
|
||||
## CU_POINTER_ATTRIBUTE_MEMORY_TYPE
|
||||
|
||||
To get pointer's memory type in HIP/HIP-Clang, developers should use hipPointerGetAttributes API. First parameter of the API is hipPointerAttribute_t which has 'memoryType' as member variable. 'memoryType' indicates input pointer is allocated on device or host.
|
||||
To get pointer's memory type in HIP/HIP-Clang, developers should use hipPointerGetAttributes API. First parameter of the API is hipPointerAttribute_t which has 'type' as member variable. 'type' indicates input pointer is allocated on device or host.
|
||||
|
||||
For example:
|
||||
```
|
||||
double * ptr;
|
||||
hipMalloc(reinterpret_cast<void**>(&ptr), sizeof(double));
|
||||
hipPointerAttribute_t attr;
|
||||
hipPointerGetAttributes(&attr, ptr); /*attr.memoryType will have value as hipMemoryTypeDevice*/
|
||||
hipPointerGetAttributes(&attr, ptr); /*attr.type will have value as hipMemoryTypeDevice*/
|
||||
|
||||
double* ptrHost;
|
||||
hipHostMalloc(&ptrHost, sizeof(double));
|
||||
hipPointerAttribute_t attr;
|
||||
hipPointerGetAttributes(&attr, ptrHost); /*attr.memoryType will have value as hipMemoryTypeHost*/
|
||||
hipPointerGetAttributes(&attr, ptrHost); /*attr.type will have value as hipMemoryTypeHost*/
|
||||
```
|
||||
Please note, hipMemoryType enum values are different from cudaMemoryType enum values.
|
||||
|
||||
For example, on AMD platform, memoryType is defined in hip_runtime_api.h,
|
||||
For example, on AMD platform, hipMemoryType is defined in hip_runtime_api.h,
|
||||
```
|
||||
typedef enum hipMemoryType {
|
||||
hipMemoryTypeHost, ///< Memory is physically located on host
|
||||
hipMemoryTypeDevice, ///< Memory is physically located on device.
|
||||
hipMemoryTypeArray, ///< Array memory, physically located on device.
|
||||
hipMemoryTypeUnified ///< Not used currently
|
||||
hipMemoryTypeHost = 0, ///< Memory is physically located on host
|
||||
hipMemoryTypeDevice = 1, ///< Memory is physically located on device. (see deviceId for specific device)
|
||||
hipMemoryTypeArray = 2, ///< Array memory, physically located on device. (see deviceId for specific device)
|
||||
hipMemoryTypeUnified = 3, ///< Not used currently
|
||||
hipMemoryTypeManaged = 4 ///< Managed memory, automaticallly managed by the unified memory system
|
||||
} hipMemoryType;
|
||||
|
||||
Looking into CUDA toolkit, it defines memoryType as following,
|
||||
```
|
||||
Looking into CUDA toolkit, it defines cudaMemoryType as following,
|
||||
```
|
||||
enum cudaMemoryType
|
||||
{
|
||||
cudaMemoryTypeUnregistered = 0, // Unregistered memory.
|
||||
@@ -500,8 +503,8 @@ enum cudaMemoryType
|
||||
cudaMemoryTypeDevice = 2, // Device memory.
|
||||
cudaMemoryTypeManaged = 3, // Managed memory
|
||||
}
|
||||
|
||||
In this case, memoryType translation for hipPointerGetAttributes needs to be handled properly on nvidia platform to get the correct memory type in CUDA, which is done in the file nvidia_hip_runtime_api.h.
|
||||
```
|
||||
In this case, memory type translation for hipPointerGetAttributes needs to be handled properly on nvidia platform to get the correct memory type in CUDA, which is done in the file nvidia_hip_runtime_api.h.
|
||||
|
||||
So in any HIP applications which use HIP APIs involving memory types, developers should use #ifdef in order to assign the correct enum values depending on Nvidia or AMD platform.
|
||||
|
||||
|
||||
@@ -153,24 +153,31 @@ typedef struct hipDeviceProp_t {
|
||||
} hipDeviceProp_t;
|
||||
|
||||
|
||||
/**
|
||||
* Memory type (for pointer attributes)
|
||||
/*
|
||||
* @brief HIP Memory type (for pointer attributes)
|
||||
* @enum
|
||||
* @ingroup Enumerations
|
||||
*/
|
||||
typedef enum hipMemoryType {
|
||||
hipMemoryTypeHost, ///< Memory is physically located on host
|
||||
hipMemoryTypeDevice, ///< Memory is physically located on device. (see deviceId for specific
|
||||
///< device)
|
||||
hipMemoryTypeArray, ///< Array memory, physically located on device. (see deviceId for specific
|
||||
///< device)
|
||||
hipMemoryTypeUnified, ///< Not used currently
|
||||
hipMemoryTypeManaged ///< Managed memory, automaticallly managed by the unified memory system
|
||||
hipMemoryTypeHost = 0, ///< Memory is physically located on host
|
||||
hipMemoryTypeDevice = 1, ///< Memory is physically located on device. (see deviceId for
|
||||
///< specific device)
|
||||
hipMemoryTypeArray = 2, ///< Array memory, physically located on device. (see deviceId for
|
||||
///< specific device)
|
||||
hipMemoryTypeUnified = 3, ///< Not used currently
|
||||
hipMemoryTypeManaged = 4 ///< Managed memory, automaticallly managed by the unified
|
||||
///< memory system
|
||||
} hipMemoryType;
|
||||
|
||||
/**
|
||||
* Pointer attributes
|
||||
*/
|
||||
typedef struct hipPointerAttribute_t {
|
||||
enum hipMemoryType memoryType;
|
||||
union {
|
||||
// Deprecated, use instead type
|
||||
enum hipMemoryType memoryType;
|
||||
enum hipMemoryType type;
|
||||
};
|
||||
int device;
|
||||
void* devicePointer;
|
||||
void* hostPointer;
|
||||
|
||||
@@ -43,7 +43,7 @@ constexpr size_t N{1000000};
|
||||
|
||||
bool operator==(const hipPointerAttribute_t& lhs, const hipPointerAttribute_t& rhs) {
|
||||
return ((lhs.hostPointer == rhs.hostPointer) && (lhs.devicePointer == rhs.devicePointer) &&
|
||||
(lhs.memoryType == rhs.memoryType) && (lhs.device == rhs.device) &&
|
||||
(lhs.type == rhs.type) && (lhs.device == rhs.device) &&
|
||||
(lhs.allocationFlags == rhs.allocationFlags));
|
||||
}
|
||||
|
||||
@@ -68,7 +68,7 @@ const char* memoryTypeToString(hipMemoryType memoryType) {
|
||||
void resetAttribs(hipPointerAttribute_t* attribs) {
|
||||
attribs->hostPointer = reinterpret_cast<void*>(-1);
|
||||
attribs->devicePointer = reinterpret_cast<void*>(-1);
|
||||
attribs->memoryType = hipMemoryTypeHost;
|
||||
attribs->type = hipMemoryTypeHost;
|
||||
attribs->device = -2;
|
||||
attribs->isManaged = -1;
|
||||
attribs->allocationFlags = 0xffff;
|
||||
@@ -77,9 +77,9 @@ void resetAttribs(hipPointerAttribute_t* attribs) {
|
||||
|
||||
void printAttribs(const hipPointerAttribute_t* attribs) {
|
||||
printf(
|
||||
"hostPointer:%p devicePointer:%p memType:%s deviceId:%d isManaged:%d "
|
||||
"hostPointer:%p devicePointer:%p type:%s deviceId:%d isManaged:%d "
|
||||
"allocationFlags:%u\n",
|
||||
attribs->hostPointer, attribs->devicePointer, memoryTypeToString(attribs->memoryType),
|
||||
attribs->hostPointer, attribs->devicePointer, memoryTypeToString(attribs->type),
|
||||
attribs->device, attribs->isManaged, attribs->allocationFlags);
|
||||
}
|
||||
|
||||
@@ -151,14 +151,14 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) {
|
||||
if (isDevice) {
|
||||
totalDeviceAllocated[reference[i]._attrib.device] += reference[i]._sizeBytes;
|
||||
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&ptr), reference[i]._sizeBytes));
|
||||
reference[i]._attrib.memoryType = hipMemoryTypeDevice;
|
||||
reference[i]._attrib.type = hipMemoryTypeDevice;
|
||||
reference[i]._attrib.devicePointer = ptr;
|
||||
reference[i]._attrib.hostPointer = NULL;
|
||||
reference[i]._attrib.allocationFlags = 0;
|
||||
} else {
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&ptr), reference[i]._sizeBytes,
|
||||
hipHostMallocDefault));
|
||||
reference[i]._attrib.memoryType = hipMemoryTypeHost;
|
||||
reference[i]._attrib.type = hipMemoryTypeHost;
|
||||
reference[i]._attrib.devicePointer = ptr;
|
||||
reference[i]._attrib.hostPointer = ptr;
|
||||
reference[i]._attrib.allocationFlags = 0;
|
||||
@@ -190,7 +190,7 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) {
|
||||
checkPointer(ref, i, 2, (ptr + ref._sizeBytes - 1));
|
||||
}
|
||||
|
||||
if (ref._attrib.memoryType == hipMemoryTypeDevice) {
|
||||
if (ref._attrib.type == hipMemoryTypeDevice) {
|
||||
HIP_CHECK(hipFree(ref._pointer));
|
||||
} else {
|
||||
HIP_CHECK(hipHostFree(ref._pointer));
|
||||
@@ -369,11 +369,11 @@ TEST_CASE("Unit_hipPointerGetAttributes_GpuIter") {
|
||||
|
||||
// Memory address and type check
|
||||
if (MemoryType == MemoryTypes::DeviceMemory) {
|
||||
REQUIRE(attributes.memoryType == hipMemoryTypeDevice);
|
||||
REQUIRE(attributes.type == hipMemoryTypeDevice);
|
||||
REQUIRE(attributes.devicePointer == ptr);
|
||||
REQUIRE(attributes.hostPointer == nullptr);
|
||||
} else if (MemoryType == MemoryTypes::HostMemory) {
|
||||
REQUIRE(attributes.memoryType == hipMemoryTypeHost);
|
||||
REQUIRE(attributes.type == hipMemoryTypeHost);
|
||||
REQUIRE(attributes.hostPointer == ptr);
|
||||
} else if (MemoryType == MemoryTypes::MappedMemory) {
|
||||
int* devicePtr{nullptr};
|
||||
|
||||
@@ -40,7 +40,7 @@ size_t Nbytes = 0;
|
||||
|
||||
bool operator==(const hipPointerAttribute_t& lhs, const hipPointerAttribute_t& rhs) {
|
||||
return ((lhs.hostPointer == rhs.hostPointer) && (lhs.devicePointer == rhs.devicePointer) &&
|
||||
(lhs.memoryType == rhs.memoryType) && (lhs.device == rhs.device) &&
|
||||
(lhs.type == rhs.type) && (lhs.device == rhs.device) &&
|
||||
(lhs.allocationFlags == rhs.allocationFlags));
|
||||
};
|
||||
|
||||
@@ -65,7 +65,7 @@ const char* memoryTypeToString(hipMemoryType memoryType) {
|
||||
void resetAttribs(hipPointerAttribute_t* attribs) {
|
||||
attribs->hostPointer = (void*)(-1);
|
||||
attribs->devicePointer = (void*)(-1);
|
||||
attribs->memoryType = hipMemoryTypeHost;
|
||||
attribs->type = hipMemoryTypeHost;
|
||||
attribs->device = -2;
|
||||
attribs->isManaged = -1;
|
||||
attribs->allocationFlags = 0xffff;
|
||||
@@ -74,9 +74,9 @@ void resetAttribs(hipPointerAttribute_t* attribs) {
|
||||
|
||||
void printAttribs(const hipPointerAttribute_t* attribs) {
|
||||
printf(
|
||||
"hostPointer:%p devicePointer:%p memoryType:%s deviceId:%d isManaged:%d "
|
||||
"hostPointer:%p devicePointer:%p type:%s deviceId:%d isManaged:%d "
|
||||
"allocationFlags:%u\n",
|
||||
attribs->hostPointer, attribs->devicePointer, memoryTypeToString(attribs->memoryType),
|
||||
attribs->hostPointer, attribs->devicePointer, memoryTypeToString(attribs->type),
|
||||
attribs->device, attribs->isManaged, attribs->allocationFlags);
|
||||
};
|
||||
|
||||
@@ -229,13 +229,13 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) {
|
||||
if (isDevice) {
|
||||
totalDeviceAllocated[reference[i]._attrib.device] += reference[i]._sizeBytes;
|
||||
HIPCHECK(hipMalloc((void**)&ptr, reference[i]._sizeBytes));
|
||||
reference[i]._attrib.memoryType = hipMemoryTypeDevice;
|
||||
reference[i]._attrib.type = hipMemoryTypeDevice;
|
||||
reference[i]._attrib.devicePointer = ptr;
|
||||
reference[i]._attrib.hostPointer = NULL;
|
||||
reference[i]._attrib.allocationFlags = 0; // TODO-randomize these.
|
||||
} else {
|
||||
HIPCHECK(hipHostMalloc((void**)&ptr, reference[i]._sizeBytes, hipHostMallocDefault));
|
||||
reference[i]._attrib.memoryType = hipMemoryTypeHost;
|
||||
reference[i]._attrib.type = hipMemoryTypeHost;
|
||||
reference[i]._attrib.devicePointer = ptr;
|
||||
reference[i]._attrib.hostPointer = ptr;
|
||||
reference[i]._attrib.allocationFlags = 0; // TODO-randomize these.
|
||||
@@ -265,7 +265,7 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) {
|
||||
checkPointer(ref, i, 2, (char*)ref._pointer + ref._sizeBytes - 1);
|
||||
}
|
||||
|
||||
if (ref._attrib.memoryType == hipMemoryTypeDevice) {
|
||||
if (ref._attrib.type == hipMemoryTypeDevice) {
|
||||
hipFree(ref._pointer);
|
||||
} else {
|
||||
hipHostFree(ref._pointer);
|
||||
|
||||
Reference in New Issue
Block a user