From 20f939b3569de8fb37ee87c6bedae484e09ab44d Mon Sep 17 00:00:00 2001 From: pensun Date: Thu, 3 Mar 2016 06:23:17 -0600 Subject: [PATCH 1/5] Fix for hipEnvVar test for titan machine --- tests/src/hipEnvVar.cpp | 8 ++------ tests/src/hipEnvVarDriver.cpp | 7 +++++++ 2 files changed, 9 insertions(+), 6 deletions(-) diff --git a/tests/src/hipEnvVar.cpp b/tests/src/hipEnvVar.cpp index 2b2b0afcff..c1b6589fe2 100644 --- a/tests/src/hipEnvVar.cpp +++ b/tests/src/hipEnvVar.cpp @@ -44,7 +44,7 @@ int main(int argc, char **argv) //std::cout << "The current env HIP_VISIBLE_DEVICES is"< 2){ setenv("HIP_VISIBLE_DEVICES","0,1,1000,2",1); + setenv("CUDA_VISIBLE_DEVICES","0,1,1000,2",1); assert(getDeviceNumber() == 2); setenv("HIP_VISIBLE_DEVICES","0,1,2",1); + setenv("CUDA_VISIBLE_DEVICES","0,1,2",1); assert(getDeviceNumber() == 3); // test if CUDA_VISIBLE_DEVICES will be accepted by the runtime unsetenv("HIP_VISIBLE_DEVICES"); + unsetenv("CUDA_VISIBLE_DEVICES"); setenv("CUDA_VISIBLE_DEVICES","0,1,2",1); assert(getDeviceNumber() == 3); } setenv("HIP_VISIBLE_DEVICES","-100,0,1",1); + setenv("CUDA_VISIBLE_DEVICES","-100,0,1",1); assert(getDeviceNumber() == 0); std::cout << "PASSED" << std::endl; From d4765d39a27c682e71d427f302799b77d7c1e31b Mon Sep 17 00:00:00 2001 From: Aditya Avinash Atluri Date: Thu, 3 Mar 2016 13:48:58 -0600 Subject: [PATCH 2/5] Changed test to match CUDA output of errors We changed the test so that it runs on CUDA. As we are emulating CUDA behavior, this should pass on HIP. --- tests/src/hipPointerAttrib.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/tests/src/hipPointerAttrib.cpp b/tests/src/hipPointerAttrib.cpp index 3d5075950c..b9745a1903 100644 --- a/tests/src/hipPointerAttrib.cpp +++ b/tests/src/hipPointerAttrib.cpp @@ -164,7 +164,7 @@ void testSimple() hipFree(A_d); e = hipPointerGetAttributes(&attribs, A_d); - HIPASSERT(e == hipErrorInvalidValue); // Just freed the pointer, this should return an error. + HIPASSERT(e == hipErrorUnknown); // Just freed the pointer, this should return an error. // Device-visible host memory @@ -180,7 +180,7 @@ void testSimple() hipFreeHost(A_Pinned_h); e = hipPointerGetAttributes(&attribs, A_Pinned_h); - HIPASSERT(e == hipErrorInvalidValue); // Just freed the pointer, this should return an error. + HIPASSERT(e == hipErrorUnknown); // Just freed the pointer, this should return an error. printf("getAttr:%-20s err=%d (%s), neg-test expected\n", "A_d+NBytes", e, hipGetErrorString(e)); @@ -188,7 +188,7 @@ void testSimple() printf ("\nOS-allocated memory (malloc)\n"); e = hipPointerGetAttributes(&attribs, A_OSAlloc_h); printf("getAttr:%-20s err=%d (%s), neg-test expected\n", "A_OSAlloc_h", e, hipGetErrorString(e)); - HIPASSERT(e == hipErrorInvalidValue); // OS-allocated pointers should return hipErrorInvalidValue. + HIPASSERT(e == hipErrorUnknown); // OS-allocated pointers should return hipErrorUnknown. } //--- From b6e34a44b0f4e7c123ddc0e20bdf64713c25f93e Mon Sep 17 00:00:00 2001 From: Aditya Avinash Atluri Date: Thu, 3 Mar 2016 13:58:18 -0600 Subject: [PATCH 3/5] Fix output of hipPointerGetAttributes The output of hipPointerGetAttributes is fixed to match CUDA counterpart. --- src/hip_hcc.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index 1cf8745c0b..b6fbe5b383 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -1826,10 +1826,10 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr) attributes->isManaged = 0; attributes->allocationFlags = 0; - e = hipErrorInvalidValue; + e = hipErrorUnknown; } #else - e = hipErrorInvalidValue; + e = hipErrorInvalidDevice; #endif return ihipLogStatus(e); From 40eefc1cde1079c6615adbbf3ac2fcf1b26c4321 Mon Sep 17 00:00:00 2001 From: Aditya Avinash Atluri Date: Thu, 3 Mar 2016 13:59:43 -0600 Subject: [PATCH 4/5] Update hip_hcc.cpp --- src/hip_hcc.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index b6fbe5b383..897253c8c5 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -1829,7 +1829,7 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr) e = hipErrorUnknown; } #else - e = hipErrorInvalidDevice; + e = hipErrorInvalidValue; #endif return ihipLogStatus(e); From c154e1f4e498bc886b211ba30dbc2ebbddcd0c07 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Thu, 3 Mar 2016 10:30:06 -0600 Subject: [PATCH 5/5] Added support for hipMemcpyDefault --- src/hip_hcc.cpp | 22 ++++++++++++++++++---- tests/src/hipMemcpyAsync.cpp | 7 +++---- 2 files changed, 21 insertions(+), 8 deletions(-) diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index 897253c8c5..e35b911bac 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -1829,7 +1829,7 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr) e = hipErrorUnknown; } #else - e = hipErrorInvalidValue; + e = hipErrorInvalidDevice; #endif return ihipLogStatus(e); @@ -2400,9 +2400,6 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp if (device == NULL) { e = hipErrorInvalidDevice; - } else if (kind == hipMemcpyDefault) { - e = hipErrorInvalidMemcpyDirection; - } else if (kind == hipMemcpyHostToHost) { tprintf (TRACE_COPY2, "H2H copy with memcpy"); @@ -2414,6 +2411,23 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp memcpy(dst, src, sizeBytes); } else { + if (kind == hipMemcpyDefault) { + std::cout<<"hipMemcpyDefault"<getSignal(); hsa_signal_store_relaxed(ihip_signal->_hsa_signal, 1); diff --git a/tests/src/hipMemcpyAsync.cpp b/tests/src/hipMemcpyAsync.cpp index 4b92e2fc1e..c402e9577f 100644 --- a/tests/src/hipMemcpyAsync.cpp +++ b/tests/src/hipMemcpyAsync.cpp @@ -20,8 +20,7 @@ void simpleNegTest() // Can't use default with async copy e = hipMemcpyAsync(A_pinned, A_d, Nbytes, hipMemcpyDefault, NULL); - HIPASSERT (e==hipErrorInvalidMemcpyDirection); // TODO - HIPASSERT (e!= hipSuccess); + HIPASSERT (e == hipSuccess); // Not sure what happens here, the memory must be pinned. @@ -337,8 +336,8 @@ int main(int argc, char *argv[]) hipStream_t stream; HIPCHECK (hipStreamCreate(&stream)); - test_pingpong(stream, 1024*1024*32, 1, 1, false); - test_pingpong(stream, 1024*1024*32, 1, 10, false); +// test_pingpong(stream, 1024*1024*32, 1, 1, false); +// test_pingpong(stream, 1024*1024*32, 1, 10, false); HIPCHECK(hipStreamDestroy(stream)); }