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

SWDEV-552537 - Fix nvidia build failures (#1125)

Browse Source 
* SWDEV-552537 - Fix nvidia build failures

* Add string header to fix hip-tests

---------

Co-authored-by: Branislav Brzak <branislav.brzak@amd.com>
Co-authored-by: Danylo Lytovchenko <danylo.lytovchenko@amd.com>
This commit is contained in:
vstojilj
2025-10-13 09:20:17 +02:00
committed by GitHub
parent b86b676514
commit bfedf63575
18 changed files with 150 additions and 132 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/unit/device/chkUUIDInGrandChild_Exe.cc
+1
View File
@@ -19,6 +19,7 @@ THE SOFTWARE.
#include <hip/hip_runtime.h>
#include <cstring>
#include <string>
int main(int argc, char** argv) {
if (argc < 0) {
return -1;
projects/hip-tests/catch/unit/device/multipleUUID_Exe.cc
+1
View File
@@ -18,6 +18,7 @@ THE SOFTWARE.
*/
#include <hip/hip_runtime.h>
#include <cstring>
#include <string>
#include <vector>
int main(int argc, char** argv) {
if (argc < 0) {
projects/hip-tests/catch/unit/device/setEnvInChildProc_Exe.cc
+1 -1
View File
@@ -17,7 +17,7 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip/hip_runtime.h>
#include <cstring>
#include <string>
#ifdef _WIN64
#define setenv(x, y, z) _putenv_s(x, y)
#define unsetenv(x) _putenv(x)
projects/hip-tests/catch/unit/graph/hipLaunchHostFunc.cc
+2 -1
View File
@@ -211,7 +211,8 @@ TEST_CASE("Unit_hipLaunchHostFunc_H2D_Kernel_D2H_Capture") {
HIP_CHECK(hipStreamBeginCapture(captureStream, hipStreamCaptureModeGlobal));
// Device alloc (async so it belongs to the captured stream)
HIP_CHECK(hipMallocAsync(&d_arrayA, arraySize * sizeof(double), captureStream));
HIP_CHECK(hipMallocAsync(reinterpret_cast<void**>(&d_arrayA), arraySize * sizeof(double),
captureStream));
// Initialize host data via a host function in the stream
set_vector_args args{h_array, initValue};
projects/hip-tests/catch/unit/texture/hipTexRefSetGetFilterMode.cc
+2
View File
@@ -18,6 +18,7 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#if defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000
texture<float, 2, hipReadModeElementType> tex;
// Test for hipTexRefSetFilterMode and hipTexRefGetFilterMode, including error handling
TEST_CASE("Unit_hipTexRefSetGetFilterMode") {
@@ -63,3 +64,4 @@ TEST_CASE("Unit_hipTexRefSetGetFilterMode") {
REQUIRE(errGetMode == hipErrorInvalidValue);
}
}
#endif
projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmapFilterMode.cc
+2
View File
@@ -18,6 +18,7 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#if defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000
texture<float, 2, hipReadModeElementType> tex;
// Test for hipTexRefSetMipmapFilterMode and hipTexRefGetMipmapFilterMode, including error handling
TEST_CASE("Unit_hipTexRefSetGetMipmapFilterMode") {
@@ -56,3 +57,4 @@ TEST_CASE("Unit_hipTexRefSetGetMipmapFilterMode") {
REQUIRE(err == hipErrorInvalidValue);
}
}
#endif
projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmapLevelBias.cc
+2
View File
@@ -18,6 +18,7 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#if defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000
texture<float, 2, hipReadModeElementType> tex;
// Test for hipTexRefSetMipmapLevelBias and hipTexRefGetMipmapLevelBias, including error handling
TEST_CASE("Unit_hipTexRefSetGetMipmapLevelBias") {
@@ -50,3 +51,4 @@ TEST_CASE("Unit_hipTexRefSetGetMipmapLevelBias") {
REQUIRE(err == hipErrorInvalidValue);
}
}
#endif
projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmapLevelClamp.cc
+2
View File
@@ -19,6 +19,7 @@ THE SOFTWARE.
#include <hip_test_common.hh>
#include <float.h>
#if defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000
texture<float, 2, hipReadModeElementType> tex;
// Test for hipTexRefSetMipmapLevelClamp and hipTexRefGetMipmapLevelClamp, including error handling
TEST_CASE("Unit_texRefSetGetMipmapLevelClamp") {
@@ -55,3 +56,4 @@ TEST_CASE("Unit_texRefSetGetMipmapLevelClamp") {
REQUIRE(err == hipErrorInvalidValue);
}
}
#endif
projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmappedArray.cc
+2
View File
@@ -17,6 +17,7 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#if defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000
texture<float, 2, hipReadModeElementType> tex;
// Test for hipTexRefSetMipmappedArray and hipTexRefGetMipmappedArray, including error handling
@@ -74,3 +75,4 @@ TEST_CASE("Unit_hipTexRefSetGetMipmappedArray") {
REQUIRE(err == hipErrorInvalidValue);
}
}
#endif
projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc
+2 -2
View File
@@ -62,11 +62,11 @@ TEST_CASE("Unit_hipMemAddressFree_negative") {
REQUIRE(granularity > 0);
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
// Allocate virtual address range
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
SECTION("nullptr to devptr") {
REQUIRE(hipMemAddressFree((hipDeviceptr_t) nullptr, size_mem) == hipErrorInvalidValue);
REQUIRE(hipMemAddressFree(nullptr, size_mem) == hipErrorInvalidValue);
}
SECTION("pass zero to size") { REQUIRE(hipMemAddressFree(ptrA, 0) == hipErrorInvalidValue); }
projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc
+4 -4
View File
@@ -68,7 +68,7 @@ TEST_CASE("Unit_hipMemAddressReserve_AlignmentTest") {
REQUIRE(granularity > 0);
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
// Allocate virtual address range
hipDeviceptr_t ptrA;
void* ptrA;
size_t alignmnt = 1;
hipMemGenericAllocationHandle_t handle;
// Allocate physical memory
@@ -93,8 +93,8 @@ TEST_CASE("Unit_hipMemAddressReserve_AlignmentTest") {
accessDesc.flags = hipMemAccessFlagsProtReadWrite;
// Make the address accessible to GPU 0
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
@@ -131,7 +131,7 @@ TEST_CASE("Unit_hipMemAddressReserve_Negative") {
REQUIRE(granularity > 0);
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
// Allocate virtual address range
hipDeviceptr_t ptrA;
void* ptrA;
SECTION("Nullptr to ptr") {
REQUIRE(hipMemAddressReserve(nullptr, size_mem, 0, 0, 0) == hipErrorInvalidValue);
projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemCreate.cc
+19 -17
View File
@@ -138,7 +138,7 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPostUnmap") {
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
// Set access
@@ -154,8 +154,8 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPostUnmap") {
for (size_t idx = 0; idx < N; idx++) {
A_h[idx] = idx;
}
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
@@ -207,7 +207,7 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPreUse") {
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
HIP_CHECK(hipMemRelease(handle));
@@ -224,8 +224,8 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPreUse") {
for (size_t idx = 0; idx < N; idx++) {
A_h[idx] = idx;
}
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
@@ -276,7 +276,7 @@ TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") {
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
HIP_CHECK(hipMemRelease(handle));
@@ -293,11 +293,11 @@ TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") {
A_h[idx] = idx;
C_h[idx] = idx * idx;
}
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
// Invoke kernel
hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK), dim3(THREADS_PER_BLOCK), 0, 0,
reinterpret_cast<int*>(ptrA));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
@@ -352,12 +352,12 @@ TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") {
HIP_CHECK(hipMemCreate(&handle[count], size_mem, &prop, 0));
}
// Allocate virtual address range for all the memory chunks
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, (numOfBuffers * size_mem), 0, 0, 0));
for (int idx = 0; idx < numOfBuffers; idx++) {
unsigned long long uiptr = reinterpret_cast<unsigned long long>(ptrA);
uiptr = uiptr + idx * size_mem;
HIP_CHECK(hipMemMap(reinterpret_cast<hipDeviceptr_t>(uiptr), size_mem, 0, handle[idx], 0));
HIP_CHECK(hipMemMap(reinterpret_cast<void*>(uiptr), size_mem, 0, handle[idx], 0));
HIP_CHECK(hipMemRelease(handle[idx]));
}
hipMemAccessDesc accessDesc = {};
@@ -373,18 +373,20 @@ TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") {
A_h[idx] = idx;
C_h[idx] = idx * idx;
}
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), numOfBuffers * buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(),
numOfBuffers * buffer_size));
// Launch square kernel
hipLaunchKernelGGL(square_kernel, dim3((N * numOfBuffers) / THREADS_PER_BLOCK),
dim3(THREADS_PER_BLOCK), 0, 0, reinterpret_cast<int*>(ptrA));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, numOfBuffers * buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA),
numOfBuffers * buffer_size));
HIP_CHECK(hipDeviceSynchronize());
// Validate Results
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
for (int idx = 0; idx < numOfBuffers; idx++) {
unsigned long long uiptr = reinterpret_cast<unsigned long long>(ptrA);
uiptr = uiptr + idx * size_mem;
HIP_CHECK(hipMemUnmap(reinterpret_cast<hipDeviceptr_t>(uiptr), size_mem));
HIP_CHECK(hipMemUnmap(reinterpret_cast<void*>(uiptr), size_mem));
}
HIP_CHECK(hipMemAddressFree(ptrA, (numOfBuffers * size_mem)));
CTX_DESTROY();
@@ -434,7 +436,7 @@ TEST_CASE("Unit_hipMemCreate_ChkWithMemset") {
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
// Set access
@@ -445,8 +447,8 @@ TEST_CASE("Unit_hipMemCreate_ChkWithMemset") {
// Make the address accessible to GPU 0
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
std::vector<int> A_h(N);
HIP_CHECK(hipMemset(reinterpret_cast<void*>(ptrA), init_val, buffer_size));
HIP_CHECK(hipMemcpyDtoH(A_h.data(), ptrA, buffer_size));
HIP_CHECK(hipMemset(ptrA, init_val, buffer_size));
HIP_CHECK(hipMemcpyDtoH(A_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
for (int idx = 0; idx < N; idx++) {
REQUIRE(A_h[idx] == init_val);
}
projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMap.cc
+30 -28
View File
@@ -99,7 +99,7 @@ TEST_CASE("Unit_hipMemMap_SameMemoryReuse") {
// Allocate a physical memory chunk
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate num_buf virtual address ranges
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
hipMemAccessDesc accessDesc = {};
accessDesc.location.type = hipMemLocationTypeDevice;
@@ -110,12 +110,12 @@ TEST_CASE("Unit_hipMemMap_SameMemoryReuse") {
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
// Set access to GPU 0
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
square_kernel<<<dim3(N / threadsPerBlk), dim3(threadsPerBlk), 0, 0>>>(
reinterpret_cast<int*>(ptrA));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
HIP_CHECK(hipStreamSynchronize(0));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
@@ -175,7 +175,7 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU") {
// Allocate a physical memory chunk
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate num_buf virtual address ranges
hipDeviceptr_t ptrA[num_buf];
void* ptrA[num_buf];
for (int buf = 0; buf < num_buf; buf++) {
HIP_CHECK(hipMemAddressReserve(&ptrA[buf], size_mem, 0, 0, 0));
}
@@ -188,12 +188,12 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU") {
HIP_CHECK(hipMemMap(ptrA[buf], size_mem, 0, handle, 0));
// Set access to GPU 0
HIP_CHECK(hipMemSetAccess(ptrA[buf], size_mem, &accessDesc, 1));
HIP_CHECK(hipMemcpyHtoD(ptrA[buf], A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA[buf], buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA[buf]), A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA[buf]), buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
square_kernel<<<dim3(N / threadsPerBlk), dim3(threadsPerBlk), 0, 0>>>(
reinterpret_cast<int*>(ptrA[buf]));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA[buf], buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA[buf]), buffer_size));
HIP_CHECK(hipStreamSynchronize(0));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
HIP_CHECK(hipMemUnmap(ptrA[buf], size_mem));
@@ -255,7 +255,7 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemory_Map2MultVMMs") {
// Allocate a physical memory chunk
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate num_buf virtual address ranges
hipDeviceptr_t ptrA[num_buf];
void* ptrA[num_buf];
for (int buf = 0; buf < num_buf; buf++) {
HIP_CHECK(hipMemAddressReserve(&ptrA[buf], size_mem, 0, 0, 0));
}
@@ -271,12 +271,12 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemory_Map2MultVMMs") {
HIP_CHECK(hipMemSetAccess(ptrA[buf], size_mem, &accessDesc, 1));
}
// Copy data to VMM via ptrA[0]
HIP_CHECK(hipMemcpyHtoD(ptrA[0], A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA[0]), A_h.data(), buffer_size));
// Validate the data contained in VMM using ptrA[0], ptrA[1],
// ......, ptrA[num_buf-1]
for (int buf = 0; buf < num_buf; buf++) {
std::fill(B_h.begin(), B_h.end(), initializer);
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA[buf], buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA[buf]), buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
}
@@ -321,7 +321,7 @@ void physicalMemoryReuse_MultiDev (hipMemAllocationProp prop) {
// Allocate a physical memory chunk
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate devicecount virtual address ranges
std::vector<hipDeviceptr_t> ptrA(devicecount);
std::vector<void*> ptrA(devicecount);
for (int devY = 0; devY < devicecount; devY++) {
HIP_CHECK(hipMemAddressReserve(&ptrA[devY], size_mem, 0, 0, 0));
}
@@ -337,8 +337,10 @@ void physicalMemoryReuse_MultiDev (hipMemAllocationProp prop) {
HIP_CHECK(hipMemMap(ptrA[devY], size_mem, 0, handle, 0));
// Set access to GPU 0
HIP_CHECK(hipMemSetAccess(ptrA[devY], size_mem, &accessDesc, 1));
HIP_CHECK(hipMemcpyHtoD(ptrA[devY], A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA[devY], buffer_size));
HIP_CHECK(
hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA[devY]), A_h.data(), buffer_size));
HIP_CHECK(
hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA[devY]), buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
HIP_CHECK(hipMemUnmap(ptrA[devY], size_mem));
}
@@ -431,7 +433,7 @@ TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_SingleGPU") {
HIP_CHECK(hipMemCreate(&handle[buf], size_mem, &prop, 0));
}
// Allocate num_buf virtual address ranges
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
hipMemAccessDesc accessDesc = {};
accessDesc.location.type = hipMemLocationTypeDevice;
@@ -443,13 +445,13 @@ TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_SingleGPU") {
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle[buf], 0));
// Set access to GPU 0
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
#if HT_NVIDIA
square_kernel<<<dim3(N / threadsPerBlk), dim3(threadsPerBlk), 0, 0>>>(
reinterpret_cast<int*>(ptrA));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
HIP_CHECK(hipStreamSynchronize(0));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
#endif
@@ -497,7 +499,7 @@ void vMMMemoryReuse_MultiGPU (hipMemAllocationProp prop) {
HIP_CHECK(hipMemCreate(&handle[dev], size_mem, &prop, 0));
}
// Allocate devicecount virtual address ranges
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
// Map ptrA to physical chunk
SECTION("Set Access of VMM to Different GPU") {
@@ -512,8 +514,8 @@ void vMMMemoryReuse_MultiGPU (hipMemAllocationProp prop) {
std::fill(B_h.begin(), B_h.end(), initializer);
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle[dev], 0));
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
}
@@ -527,8 +529,8 @@ void vMMMemoryReuse_MultiGPU (hipMemAllocationProp prop) {
std::fill(B_h.begin(), B_h.end(), initializer);
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle[dev], 0));
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
}
@@ -617,7 +619,7 @@ TEST_CASE("Unit_hipMemMap_MapPartialVMMMem") {
// Allocate a bigger physical memory chunk of size_mem
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range of size twice size_mem
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, 2 * size_mem, 0, 0, 0));
hipMemAccessDesc accessDesc = {};
accessDesc.location.type = hipMemLocationTypeDevice;
@@ -626,8 +628,8 @@ TEST_CASE("Unit_hipMemMap_MapPartialVMMMem") {
std::fill(B_h.begin(), B_h.end(), initializer);
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
// Release resources
@@ -663,14 +665,14 @@ TEST_CASE("Unit_hipMemMap_negative") {
REQUIRE(granularity > 0);
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
hipDeviceptr_t ptrA;
void* ptrA;
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
SECTION("nullptr to ptrA") {
REQUIRE(hipMemMap((hipDeviceptr_t) nullptr, size_mem, 0, handle, 0) == hipErrorInvalidValue);
REQUIRE(hipMemMap(nullptr, size_mem, 0, handle, 0) == hipErrorInvalidValue);
}
SECTION("pass zero to size") {
projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc
+5 -5
View File
@@ -65,12 +65,12 @@ TEST_CASE("Unit_hipMemRetainAllocationHandle_SetGet") {
REQUIRE(granularity > 0);
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
hipDeviceptr_t ptrA;
void* ptrA;
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
HIP_CHECK(hipMemAddressReserve(reinterpret_cast<void**>(&ptrA), size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap((void*)ptrA, size_mem, 0, handle, 0));
// Test hipMemRetainAllocationHandle
hipMemGenericAllocationHandle_t gethandle;
// Check beginning of VMM ptr
@@ -109,7 +109,7 @@ TEST_CASE("Unit_hipMemRetainAllocationHandle_NegTst") {
REQUIRE(granularity > 0);
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
hipDeviceptr_t ptrA;
void* ptrA;
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
@@ -125,7 +125,7 @@ TEST_CASE("Unit_hipMemRetainAllocationHandle_NegTst") {
REQUIRE(hipMemRetainAllocationHandle(&gethandle, nullptr) == hipErrorInvalidValue);
}
SECTION("not mapped address") {
hipDeviceptr_t ptrB;
void* ptrB;
HIP_CHECK(hipMemAddressReserve(&ptrB, size_mem, 0, 0, 0));
REQUIRE(hipMemRetainAllocationHandle(&gethandle, reinterpret_cast<void*>(ptrB)) ==
hipErrorInvalidValue);
projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc
+52 -49
View File
@@ -100,7 +100,7 @@ TEST_CASE("Unit_hipMemSetAccess_SetGet") {
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
HIP_CHECK(hipMemRelease(handle));
@@ -176,7 +176,7 @@ TEST_CASE("Unit_hipMemSetAccess_MultDevSetGet") {
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
HIP_CHECK(hipMemRelease(handle));
@@ -237,7 +237,7 @@ TEST_CASE("Unit_hipMemSetAccess_EntireVMMRangeSetGet") {
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
// Allocate physical memory
hipMemGenericAllocationHandle_t handle;
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
@@ -259,7 +259,7 @@ TEST_CASE("Unit_hipMemSetAccess_EntireVMMRangeSetGet") {
REQUIRE(flags == hipMemAccessFlagsProtReadWrite);
unsigned long long uiptr = reinterpret_cast<unsigned long long>(ptrA);
uiptr += (size_mem - 1);
HIP_CHECK(hipMemGetAccess(&flags, &location, reinterpret_cast<hipDeviceptr_t>(uiptr)));
HIP_CHECK(hipMemGetAccess(&flags, &location, reinterpret_cast<void*>(uiptr)));
REQUIRE(flags == hipMemAccessFlagsProtReadWrite);
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
@@ -295,7 +295,7 @@ TEST_CASE("Unit_hipMemGetAccess_NegTst") {
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
// Allocate physical memory
hipMemGenericAllocationHandle_t handle;
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
@@ -320,7 +320,7 @@ TEST_CASE("Unit_hipMemGetAccess_NegTst") {
REQUIRE(status == hipErrorInvalidValue);
unsigned long long uiptr = reinterpret_cast<unsigned long long>(ptrA);
uiptr += size_mem;
status = hipMemGetAccess(&flags, &location, reinterpret_cast<hipDeviceptr_t>(uiptr));
status = hipMemGetAccess(&flags, &location, reinterpret_cast<void*>(uiptr));
REQUIRE(status == hipErrorInvalidValue);
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
@@ -363,7 +363,7 @@ TEST_CASE("Unit_hipMemSetAccess_FuncTstOnMultDev") {
REQUIRE(granularity > 0);
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
// Allocate physical memory
hipDeviceptr_t ptrA;
void* ptrA;
hipMemGenericAllocationHandle_t handle;
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
@@ -382,7 +382,7 @@ TEST_CASE("Unit_hipMemSetAccess_FuncTstOnMultDev") {
for (int idx = 0; idx < N; idx++) {
A_h[idx] = idx;
}
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
// Set the A_h to verify with square kernel.
for (int idx = 0; idx < N; idx++) {
A_h[idx] = idx * idx;
@@ -390,7 +390,7 @@ TEST_CASE("Unit_hipMemSetAccess_FuncTstOnMultDev") {
// Launch square kernel
hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK), dim3(THREADS_PER_BLOCK), 0, 0,
reinterpret_cast<int*>(ptrA));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
@@ -438,7 +438,7 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") {
// Allocate a physical memory chunk
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate num_buf virtual address ranges
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
hipMemAccessDesc accessDesc = {};
@@ -451,8 +451,8 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") {
// Change property of virtual memory range to read only
accessDesc.flags = hipMemAccessFlagsProtReadWrite;
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
}
@@ -462,8 +462,8 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") {
// Change property of virtual memory range to read only
accessDesc.flags = hipMemAccessFlagsProtReadWrite;
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
}
@@ -471,13 +471,15 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") {
SECTION("Check error while writing on Read-Only memory") {
accessDesc.flags = hipMemAccessFlagsProtRead;
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
REQUIRE(hipErrorInvalidValue == hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
REQUIRE(hipErrorInvalidValue ==
hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
}
SECTION("Check error while writing on inaccessible memory") {
accessDesc.flags = hipMemAccessFlagsProtNone;
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
REQUIRE(hipErrorInvalidValue == hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
REQUIRE(hipErrorInvalidValue ==
hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
}
#endif
@@ -524,7 +526,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2UnifiedMemCpy") {
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
// Allocate physical memory
hipMemGenericAllocationHandle_t handle;
hipDeviceptr_t ptrA, ptrB;
void *ptrA, *ptrB;
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
@@ -543,10 +545,11 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2UnifiedMemCpy") {
for (int idx = 0; idx < N; idx++) {
ptrA_h[idx] = idx;
}
HIP_CHECK(hipMemcpyHtoD(ptrA, ptrA_h, buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), ptrA_h, buffer_size));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&ptrB), buffer_size));
HIP_CHECK(hipMemcpyDtoD(ptrB, ptrA, buffer_size));
HIP_CHECK(hipMemcpyDtoH(ptrB_h, ptrB, buffer_size));
HIP_CHECK(hipMemcpyDtoD(reinterpret_cast<hipDeviceptr_t>(ptrB),
reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
HIP_CHECK(hipMemcpyDtoH(ptrB_h, reinterpret_cast<hipDeviceptr_t>(ptrB), buffer_size));
bool bPassed = true;
for (int idx = 0; idx < N; idx++) {
if (ptrB_h[idx] != idx) {
@@ -594,7 +597,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2DevMemCpy") {
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
// Allocate physical memory
hipMemGenericAllocationHandle_t handle;
hipDeviceptr_t ptrA, ptrB;
void *ptrA, *ptrB;
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
@@ -611,10 +614,11 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2DevMemCpy") {
for (int idx = 0; idx < N; idx++) {
A_h[idx] = idx;
}
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&ptrB), buffer_size));
HIP_CHECK(hipMemcpyDtoD(ptrB, ptrA, buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrB, buffer_size));
HIP_CHECK(hipMemcpyDtoD(reinterpret_cast<hipDeviceptr_t>(ptrB),
reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrB), buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
HIP_CHECK(hipFree(reinterpret_cast<void*>(ptrB)));
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
@@ -659,7 +663,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerDevMemCpy") {
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
// Allocate physical memory
hipMemGenericAllocationHandle_t handle;
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
@@ -676,7 +680,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerDevMemCpy") {
for (int idx = 0; idx < N; idx++) {
A_h[idx] = idx;
}
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
// Check Peer Access
for (deviceId = 1; deviceId < devicecount; deviceId++) {
int canAccessPeer = 0;
@@ -703,7 +707,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerDevMemCpy") {
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &access, 1));
hipDeviceptr_t dptr_peer;
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&dptr_peer), buffer_size));
HIP_CHECK(hipMemcpyDtoD(dptr_peer, ptrA, buffer_size));
HIP_CHECK(hipMemcpyDtoD(dptr_peer, reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), dptr_peer, buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
HIP_CHECK(hipFree(reinterpret_cast<void*>(dptr_peer)));
@@ -750,7 +754,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerPeerMemCpy") {
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
// Allocate physical memory
hipMemGenericAllocationHandle_t handle;
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
@@ -767,7 +771,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerPeerMemCpy") {
for (int idx = 0; idx < N; idx++) {
A_h[idx] = idx;
}
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
// Check Peer Access
for (deviceId = 1; deviceId < devicecount; deviceId++) {
std::fill(B_h.begin(), B_h.end(), initializer);
@@ -843,7 +847,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMMemCpy") {
HIP_CHECK(hipMemCreate(&handle1, size_mem, &prop, 0));
HIP_CHECK(hipMemCreate(&handle2, size_mem, &prop, 0));
// Allocate virtual address range
hipDeviceptr_t ptrA, ptrB;
void *ptrA, *ptrB;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemAddressReserve(&ptrB, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle1, 0));
@@ -862,9 +866,10 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMMemCpy") {
for (int idx = 0; idx < N; idx++) {
A_h[idx] = idx;
}
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoD(ptrB, ptrA, buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrB, buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoD(reinterpret_cast<hipDeviceptr_t>(ptrB),
reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrB), buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
HIP_CHECK(hipMemUnmap(ptrB, size_mem));
@@ -913,7 +918,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") {
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
HIP_CHECK(hipMemRelease(handle));
@@ -928,7 +933,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") {
for (int idx = 0; idx < N; idx++) {
A_h[idx] = idx;
}
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
for (deviceId = 1; deviceId < devicecount; deviceId++) {
int canAccessPeer = 0;
hipDevice_t device_other;
@@ -958,7 +963,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") {
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle_loc, size_mem_loc, &prop_loc, 0));
// Allocate virtual address range
hipDeviceptr_t ptrB;
void* ptrB;
HIP_CHECK(hipMemAddressReserve(&ptrB, size_mem_loc, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrB, size_mem_loc, 0, handle_loc, 0));
HIP_CHECK(hipMemRelease(handle_loc));
@@ -969,9 +974,8 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") {
accessDesc_loc.flags = hipMemAccessFlagsProtReadWrite;
// Make the address accessible to GPU 0
HIP_CHECK(hipMemSetAccess(ptrB, size_mem_loc, &accessDesc_loc, 1));
HIP_CHECK(hipMemcpyPeer(reinterpret_cast<void*>(ptrB), deviceId, reinterpret_cast<void*>(ptrA),
0, buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrB, buffer_size));
HIP_CHECK(hipMemcpyPeer(ptrB, deviceId, ptrA, 0, buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrB), buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
HIP_CHECK(hipMemUnmap(ptrB, size_mem_loc));
HIP_CHECK(hipMemAddressFree(ptrB, size_mem_loc));
@@ -984,7 +988,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") {
class vmm_resize_class {
size_t current_size_tot;
size_t current_size_rounded_tot;
hipDeviceptr_t ptrVmm;
void* ptrVmm;
std::vector<hipMemGenericAllocationHandle_t> vhandle;
std::vector<size_t> vsize;
// allocate initial VMM memory chunk
@@ -1014,7 +1018,7 @@ class vmm_resize_class {
accessDesc.flags = hipMemAccessFlagsProtReadWrite;
// Make the address accessible to GPU device
HIP_CHECK(hipMemSetAccess(ptrVmm, size_rounded, &accessDesc, 1));
*ptr = ptrVmm;
*ptr = reinterpret_cast<hipDeviceptr_t>(ptrVmm);
current_size_tot += size;
current_size_rounded_tot += size_rounded;
return 0;
@@ -1067,7 +1071,7 @@ class vmm_resize_class {
} else {
unsigned long long uiptr = reinterpret_cast<unsigned long long>(ptrVmm);
uiptr = uiptr + vsize[idx - 1];
HIP_CHECK(hipMemMap(reinterpret_cast<hipDeviceptr_t>(uiptr), vsize[idx], 0, myhandle, 0));
HIP_CHECK(hipMemMap(reinterpret_cast<void*>(uiptr), vsize[idx], 0, myhandle, 0));
}
idx++;
}
@@ -1078,7 +1082,7 @@ class vmm_resize_class {
accessDesc.flags = hipMemAccessFlagsProtReadWrite;
// Make the address accessible to GPU 0
HIP_CHECK(hipMemSetAccess(ptrVmm, (size_rounded + current_size_rounded_tot), &accessDesc, 1));
*ptr = ptrVmm;
*ptr = reinterpret_cast<hipDeviceptr_t>(ptrVmm);
current_size_tot += size;
current_size_rounded_tot += size_rounded;
return 0;
@@ -1253,7 +1257,7 @@ TEST_CASE("Unit_hipMemSetAccess_negative") {
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
hipDeviceptr_t ptrA;
void* ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
// Set access
@@ -1263,8 +1267,7 @@ TEST_CASE("Unit_hipMemSetAccess_negative") {
accessDesc.flags = hipMemAccessFlagsProtReadWrite;
SECTION("nullptr to ptrA") {
REQUIRE(hipMemSetAccess((hipDeviceptr_t) nullptr, size_mem, &accessDesc, 1) ==
hipErrorInvalidValue);
REQUIRE(hipMemSetAccess(nullptr, size_mem, &accessDesc, 1) == hipErrorInvalidValue);
}
SECTION("pass zero to size") {
@@ -1310,7 +1313,7 @@ TEST_CASE("Unit_hipMemSetAccess_negative") {
}
SECTION("uninitialized virtual memory") {
hipDeviceptr_t ptrB;
void* ptrB;
HIP_CHECK(hipMemAddressReserve(&ptrB, size_mem, 0, 0, 0));
REQUIRE(hipMemSetAccess(ptrB, size_mem, &accessDesc, 1) == hipErrorInvalidValue);
HIP_CHECK(hipMemAddressFree(ptrB, size_mem));
@@ -1335,7 +1338,7 @@ TEST_CASE("Unit_hipMemSetAccessHostDevice_hostalloc") {
prop.type = hipMemAllocationTypePinned; // pinned system RAM
prop.location.type = hipMemLocationTypeHost; // generic host
prop.location.id = 0; // host id must be 0
prop.requestedHandleType = hipMemHandleTypeNone;
prop.requestedHandleTypes = hipMemHandleTypeNone;
constexpr size_t N = 1024;
constexpr size_t bytes = N * sizeof(int);
@@ -1407,7 +1410,7 @@ TEST_CASE("Unit_hipMemSetAccessHost_devicealloc") {
prop.type = hipMemAllocationTypePinned; // pinned system RAM
prop.location.type = hipMemLocationTypeDevice; // generic host
prop.location.id = 0; // host id must be 0
prop.requestedHandleType = hipMemHandleTypeNone;
prop.requestedHandleTypes = hipMemHandleTypeNone;
constexpr size_t N = 1024;
constexpr size_t bytes = N * sizeof(int);
projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemUnmap.cc
+2 -4
View File
@@ -66,16 +66,14 @@ TEST_CASE("Unit_hipMemUnmap_negative") {
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
hipDeviceptr_t ptrA;
void* ptrA;
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
SECTION("nullptr to ptrA") {
REQUIRE(hipMemUnmap((hipDeviceptr_t) nullptr, size_mem) == hipErrorInvalidValue);
}
SECTION("nullptr to ptrA") { REQUIRE(hipMemUnmap(nullptr, size_mem) == hipErrorInvalidValue); }
SECTION("pass zero to size") { REQUIRE(hipMemUnmap(ptrA, 0) == hipErrorInvalidValue); }
projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc
+5 -6
View File
@@ -59,13 +59,12 @@ TEST_CASE("Unit_hipMemVmm_Basic") {
size_t size = granularity;
void* reservedAddress{nullptr};
HIP_CHECK(hipMemAddressReserve(reinterpret_cast<hipDeviceptr_t*>(&reservedAddress), size, 0,
(hipDeviceptr_t) nullptr, 0));
HIP_CHECK(hipMemAddressReserve(&reservedAddress, size, 0, nullptr, 0));
hipMemGenericAllocationHandle_t gaHandle;
HIP_CHECK(hipMemCreate(&gaHandle, size, &memAllocationProp, 0));
HIP_CHECK(hipMemMap(reinterpret_cast<hipDeviceptr_t>(reservedAddress), size, 0, gaHandle, 0));
HIP_CHECK(hipMemMap(reservedAddress, size, 0, gaHandle, 0));
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, 0));
@@ -76,7 +75,7 @@ TEST_CASE("Unit_hipMemVmm_Basic") {
std::vector<char> values(size);
const char value = 1;
HIP_CHECK(hipMemSetAccess(reinterpret_cast<hipDeviceptr_t>(reservedAddress), size, &desc, 1));
HIP_CHECK(hipMemSetAccess(reservedAddress, size, &desc, 1));
HIP_CHECK(hipMemset(reservedAddress, value, size));
HIP_CHECK(hipMemcpy(&values[0], reservedAddress, size, hipMemcpyDeviceToHost));
@@ -84,9 +83,9 @@ TEST_CASE("Unit_hipMemVmm_Basic") {
REQUIRE(values[i] == value);
}
HIP_CHECK(hipMemUnmap(reinterpret_cast<hipDeviceptr_t>(reservedAddress), size));
HIP_CHECK(hipMemUnmap(reservedAddress, size));
HIP_CHECK(hipMemRelease(gaHandle));
HIP_CHECK(hipMemAddressFree(reinterpret_cast<hipDeviceptr_t>(reservedAddress), size));
HIP_CHECK(hipMemAddressFree(reservedAddress, size));
CTX_DESTROY();
}
projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
+16 -15
View File
@@ -1930,6 +1930,7 @@ typedef enum CUmemAllocationGranularity_flags_enum hipMemAllocationGranularity_f
typedef enum cudaMemLocationType hipMemLocationType;
#define hipMemLocationTypeInvalid cudaMemLocationTypeInvalid
#define hipMemLocationTypeDevice cudaMemLocationTypeDevice
#define hipMemLocationTypeHost cudaMemLocationTypeHost
#define hipMemHandleTypeNone cudaMemHandleTypeNone
#define hipMemHandleTypePosixFileDescriptor cudaMemHandleTypePosixFileDescriptor
#define hipMemHandleTypeWin32 cudaMemHandleTypeWin32
@@ -3057,12 +3058,12 @@ inline static hipError_t hipMemRelease(hipMemGenericAllocationHandle_t handle) {
return hipCUResultTohipError(cuMemRelease(handle));
}
inline static hipError_t hipMemAddressFree(void* devPtr, size_t size) {
return hipCUResultTohipError(cuMemAddressFree(reinterpret_cast<CUdeviceptr>(devPtr), size));
return hipCUResultTohipError(cuMemAddressFree((CUdeviceptr)devPtr, size));
}
inline static hipError_t hipMemAddressReserve(void** ptr, size_t size, size_t alignment, void* addr,
unsigned long long flags) {
return hipCUResultTohipError(cuMemAddressReserve(reinterpret_cast<CUdeviceptr*>(ptr), size,
alignment, reinterpret_cast<CUdeviceptr>(addr),
return hipCUResultTohipError(cuMemAddressReserve((CUdeviceptr*)ptr, size,
alignment, (CUdeviceptr)addr,
flags));
}
inline static hipError_t hipMemExportToShareableHandle(void* shareableHandle,
@@ -3075,10 +3076,10 @@ inline static hipError_t hipMemExportToShareableHandle(void* shareableHandle,
inline static hipError_t hipMemGetAccess(unsigned long long* flags, const hipMemLocation* location,
void* ptr) {
if (location == NULL) {
return hipCUResultTohipError(cuMemGetAccess(flags, NULL, reinterpret_cast<CUdeviceptr>(ptr)));
return hipCUResultTohipError(cuMemGetAccess(flags, NULL, (CUdeviceptr)ptr));
} else {
CUmemLocation loc = hipMemLocationToCUmemLocation(location);
return hipCUResultTohipError(cuMemGetAccess(flags, &loc, reinterpret_cast<CUdeviceptr>(ptr)));
return hipCUResultTohipError(cuMemGetAccess(flags, &loc, (CUdeviceptr)ptr));
}
}
inline static hipError_t hipMemGetAllocationPropertiesFromHandle(
@@ -3102,7 +3103,7 @@ inline static hipError_t hipMemMap(void* ptr, size_t size, size_t offset,
hipMemGenericAllocationHandle_t handle,
unsigned long long flags) {
return hipCUResultTohipError(
cuMemMap(reinterpret_cast<CUdeviceptr>(ptr), size, offset, handle, flags));
cuMemMap((CUdeviceptr)ptr, size, offset, handle, flags));
}
inline static hipError_t hipMemMapArrayAsync(hipArrayMapInfo* mapInfoList, unsigned int count,
hipStream_t stream) {
@@ -3116,17 +3117,17 @@ inline static hipError_t hipMemSetAccess(void* ptr, size_t size, const hipMemAcc
size_t count) {
if (desc == NULL) {
return hipCUResultTohipError(
cuMemSetAccess(reinterpret_cast<CUdeviceptr>(ptr), size, NULL, count));
cuMemSetAccess((CUdeviceptr)ptr, size, NULL, count));
} else {
CUmemAccessDesc* cuDesc = hipMemAccessDescToCUmemAccessDesc(desc, count);
auto status = hipCUResultTohipError(
cuMemSetAccess(reinterpret_cast<CUdeviceptr>(ptr), size, cuDesc, count));
cuMemSetAccess((CUdeviceptr)ptr, size, cuDesc, count));
free(cuDesc);
return status;
}
}
inline static hipError_t hipMemUnmap(void* ptr, size_t size) {
return hipCUResultTohipError(cuMemUnmap(reinterpret_cast<CUdeviceptr>(ptr), size));
return hipCUResultTohipError(cuMemUnmap((CUdeviceptr)ptr, size));
}
#endif // CUDA_VERSION >= CUDA_10020
@@ -3464,16 +3465,16 @@ inline static hipError_t hipCtxGetDevice(hipDevice_t* device) {
return hipCUResultTohipError(cuCtxGetDevice(device));
}
inline static hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) {
return hipCUResultTohipError(cuCtxGetApiVersion(ctx, (unsigned int*)apiVersion));
inline static hipError_t hipCtxGetApiVersion(hipCtx_t ctx, unsigned int* apiVersion) {
return hipCUResultTohipError(cuCtxGetApiVersion(ctx, apiVersion));
}
inline static hipError_t hipCtxGetCacheConfig(hipFuncCache* cacheConfig) {
return hipCUResultTohipError(cuCtxGetCacheConfig(cacheConfig));
inline static hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) {
return hipCUResultTohipError(cuCtxGetCacheConfig((CUfunc_cache*)cacheConfig));
}
inline static hipError_t hipCtxSetCacheConfig(hipFuncCache cacheConfig) {
return hipCUResultTohipError(cuCtxSetCacheConfig(cacheConfig));
inline static hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) {
return hipCUResultTohipError(cuCtxSetCacheConfig((CUfunc_cache)cacheConfig));
}
inline static hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) {