diff --git a/projects/hip/docs/markdown/hip_porting_guide.md b/projects/hip/docs/markdown/hip_porting_guide.md index 179ad64a2f..fe1972e3b3 100644 --- a/projects/hip/docs/markdown/hip_porting_guide.md +++ b/projects/hip/docs/markdown/hip_porting_guide.md @@ -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(&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. diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index 1d8a6cc731..4ffb3f5241 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -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; diff --git a/projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc b/projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc index 8157c8811d..d4ba5960ce 100644 --- a/projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc +++ b/projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc @@ -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(-1); attribs->devicePointer = reinterpret_cast(-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(&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(&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}; diff --git a/projects/hip/tests/src/runtimeApi/memory/hipPointerAttributes.cpp b/projects/hip/tests/src/runtimeApi/memory/hipPointerAttributes.cpp index cd9683ea5b..9666dd4693 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipPointerAttributes.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipPointerAttributes.cpp @@ -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);