diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux index 7b3d444c4d..24ec58d55e 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux @@ -1319,7 +1319,6 @@ "Unit_hipGraphAddNodeTypeMemset_Positive_Basic - uint8_t", "Unit_hipGraphAddNodeTypeMemset_Positive_Basic - uint16_t", "Unit_hipGraphAddNodeTypeMemset_Positive_Basic - uint32_t", - "Unit_hipMemAdvise_TstAlignedAllocMem_XNACK", "Unit_hipCreateSurfaceObject_Negative_Parameters", "Unit_hipDestroySurfaceObject_Negative_Parameters", "Unit_Device___float2half_rd_Accuracy_Limited_Positive", diff --git a/projects/hip-tests/catch/unit/memory/hipMemAdvise_AlignedAllocMem_Exe.cc b/projects/hip-tests/catch/unit/memory/hipMemAdvise_AlignedAllocMem_Exe.cc index 0fe16a273b..baaa100594 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemAdvise_AlignedAllocMem_Exe.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemAdvise_AlignedAllocMem_Exe.cc @@ -23,6 +23,7 @@ THE SOFTWARE. #include #include #include "hip/hip_runtime_api.h" +#include "hip/hip_runtime.h" #define HIP_CHECK(error) \ { \ @@ -35,8 +36,10 @@ THE SOFTWARE. // Kernel __global__ void MemAdvise_Exe(int *Hmm, int n) { - for (int i = 0; i < n; i ++) { - Hmm[i] = Hmm[i] + 10; + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + for (int i = offset; i < n; i+=stride) { + Hmm[i] += 10; } } @@ -51,8 +54,7 @@ static int hipMemAdvise_AlignedAllocMem_Exe() { 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(aligned_alloc(4096, 4096*4)); + Mllc = reinterpret_cast(aligned_alloc(4096, MemSz)); for (int i = 0; i < NumElms; ++i) { Mllc[i] = InitVal; @@ -64,8 +66,7 @@ static int hipMemAdvise_AlignedAllocMem_Exe() { // The following hipMemAdvise() call is made to know if advise on // aligned_alloc() is causing any issue HIP_CHECK(hipMemAdvise(Mllc, MemSz, hipMemAdviseSetPreferredLocation, 0)); - hipError_t err = hipMemPrefetchAsync(Mllc, MemSz, 0, strm); - if (err != hipErrorInvalidValue) return -1; + HIP_CHECK(hipMemPrefetchAsync(Mllc, MemSz, 0, strm)); HIP_CHECK(hipStreamSynchronize(strm)); MemAdvise_Exe<<<(NumElms/32), 32, 0, strm>>>(Mllc, NumElms); HIP_CHECK(hipStreamSynchronize(strm));