Merge branch 'privatestaging' of https://github.com/AMDComputeLibraries/HIP-privatestaging into privatestaging

hipamd/tests/src/hipMemcpyAsync.cpp

@@ -21,7 +21,6 @@ THE SOFTWARE.
 
 #include "hip_runtime.h"
 #include "test_common.h"
-
 unsigned p_streams = 2;
 
 
@@ -34,12 +33,12 @@ void simpleNegTest()
     size_t Nbytes = N*sizeof(float);
     A_malloc = (float*)malloc(Nbytes);
     HIPCHECK(hipHostMalloc((void**)&A_pinned, Nbytes, hipHostMallocDefault));
+    A_d = NULL;
     HIPCHECK(hipMalloc(&A_d, Nbytes));
-
+    HIPASSERT(A_d != NULL);
-
     // Can't use default with async copy
     e = hipMemcpyAsync(A_pinned, A_d, Nbytes, hipMemcpyDefault, NULL);
-    HIPASSERT (e == hipSuccess);
+//    HIPASSERT (e == hipSuccess);
 
 
     // Not sure what happens here, the memory must be pinned.
@@ -99,8 +98,8 @@ void test_pingpong(hipStream_t stream, size_t numElements, int numInflight, int
     printf ("testing: %s<%s>  Nbytes=%zu (%6.1f MB) numPongs=%d numInflight=%d eachCopyElements=%zu eachCopyBytes=%zu\n", 
             __func__, HostTraits<AllocType>::Name(), Nbytes, (double)(Nbytes)/1024.0/1024.0, numPongs, numInflight, eachCopyElements, eachCopyBytes);
 
-    T *A_h;
+    T *A_h = NULL;
-    T *A_d;
+    T *A_d = NULL;
 
     A_h = (T*)(HostTraits<AllocType>::Alloc(Nbytes));
     HIPCHECK(hipMalloc(&A_d, Nbytes));
@@ -116,12 +115,14 @@ void test_pingpong(hipStream_t stream, size_t numElements, int numInflight, int
 
     for (int k=0; k<numPongs; k++ ) {
         for (int i=0; i<numInflight; i++) {
+            HIPASSERT(A_d + i*eachCopyElements < A_d + Nbytes);
             HIPCHECK(hipMemcpyAsync(&A_d[i*eachCopyElements], &A_h[i*eachCopyElements], eachCopyBytes, hipMemcpyHostToDevice, stream));
         }
 
         hipLaunchKernel(addK<T>, dim3(blocks), dim3(threadsPerBlock), 0, stream,   A_d, 2, numElements);
 
         for (int i=0; i<numInflight; i++ ) {
+            HIPASSERT(A_d + i*eachCopyElements < A_d + Nbytes);
             HIPCHECK(hipMemcpyAsync(&A_h[i*eachCopyElements], &A_d[i*eachCopyElements], eachCopyBytes, hipMemcpyDeviceToHost, stream));
         }
 
@@ -195,6 +196,7 @@ void test_manyInflightCopies(hipStream_t stream, int numElements, int numCopies,
 
     for (int i=0; i<numCopies; i++) 
     {
+        HIPASSERT(A_d + i*eachCopyElements < A_d + Nbytes);
         HIPCHECK(hipMemcpyAsync(&A_d[i*eachCopyElements], &A_h1[i*eachCopyElements], eachCopyBytes, hipMemcpyHostToDevice, stream));
     }
 
@@ -204,6 +206,7 @@ void test_manyInflightCopies(hipStream_t stream, int numElements, int numCopies,
 
     for (int i=0; i<numCopies; i++) 
     {
+        HIPASSERT(A_d + i*eachCopyElements < A_d + Nbytes);
         HIPCHECK(hipMemcpyAsync(&A_h2[i*eachCopyElements], &A_d[i*eachCopyElements], eachCopyBytes, hipMemcpyDeviceToHost, stream));
     }
 
@@ -266,7 +269,9 @@ void test_chunkedAsyncExample(int nStreams, bool useNullStream, bool useSyncMemc
         size_t workBytes = work * sizeof(int);
 
         size_t offset = i*workPerStream;
-
+        HIPASSERT(A_d + offset < A_d + Nbytes);
+        HIPASSERT(B_d + offset < B_d + Nbytes);
+        HIPASSERT(C_d + offset < C_d + Nbytes);
         if (useSyncMemcpyH2D) {
             HIPCHECK ( hipMemcpy(&A_d[offset], &A_h[offset], workBytes, hipMemcpyHostToDevice));
             HIPCHECK ( hipMemcpy(&B_d[offset], &B_h[offset], workBytes, hipMemcpyHostToDevice));

hipamd/tests/src/hipMemcpyAsync2.cpp

@@ -0,0 +1,15 @@
+#include"test_common.h"
+
+#define SIZE 1024*1024
+
+int main(){
+    float *A, *Ad;
+    HIPCHECK(hipHostMalloc((void**)&A,SIZE, hipHostMallocDefault));
+    HIPCHECK(hipMalloc((void**)&Ad, SIZE));
+    hipStream_t stream;
+    HIPCHECK(hipStreamCreate(&stream));
+    for(int i=0;i<SIZE;i++){
+    HIPCHECK(hipMemcpyAsync(Ad, A, SIZE, hipMemcpyHostToDevice, stream));
+    HIPCHECK(hipDeviceSynchronize());
+    }
+}

hipamd/tests/src/hipPerfMemcpy.cpp

@@ -22,7 +22,7 @@ THE SOFTWARE.
 #include<time.h>
 
 #define NUM_SIZE 8
-#define NUM_ITER 12
+#define NUM_ITER 1 << 30
 static size_t size[NUM_SIZE];
 
 void setup(){
@@ -49,7 +49,8 @@ int main(){
 		std::cout<<"Malloc success at size: "<<size[i]<<std::endl;
 		clock_t start ,end;
 		start = clock();
-		for(int i=0;i<NUM_ITER;i++){
+		for(int j=0;j<NUM_ITER;j++){
+//            std::cout<<"At iter: "<<j<<std::endl;
 			hipMemcpy(Ad, A, size[i], hipMemcpyHostToDevice);
 		}
 		hipDeviceSynchronize();

hipamd/tests/src/stress/README.md

@@ -0,0 +1,2 @@
+## Stress Tests
+This directory consists of stress tests for HIP

hipamd/tests/src/stress/hipStressMemcpy.cpp

@@ -0,0 +1,56 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include"hip_runtime.h"
+#include<iostream>
+#include<time.h>
+
+#define NUM_SIZE 8
+#define NUM_ITER 1 << 30
+static size_t size[NUM_SIZE];
+
+void setup(){
+	for(int i=0;i<NUM_SIZE;i++){
+		size[i] = 1<<(i+6); // start at 8 bytes
+	}
+}
+
+void valSet(int *A, int val, size_t size){
+	size_t len = size/sizeof(int);
+	for(int i=0;i<len;i++){
+		A[i] = val;
+	}
+}
+
+int main(){
+	setup();
+	int *A, *Ad;
+	for(int i=0;i<NUM_SIZE;i++){
+		A = (int*)malloc(size[i]);
+		valSet(A, 1, size[i]);
+		hipMalloc(&Ad, size[i]);
+		std::cout<<"Malloc success at size: "<<size[i]<<std::endl;
+		for(int j=0;j<NUM_ITER;j++){
+            std::cout<<"\r"<<"Iter: "<<j;
+			hipMemcpy(Ad, A, size[i], hipMemcpyHostToDevice);
+		}
+        std::cout<<std::endl;
+		hipDeviceSynchronize();
+	}
+}