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

[ROCm/hip commit: 1d2f267dc3]

projects/hip/samples/0_Intro/bit_extract/bit_extract.cpp

@@ -30,7 +30,7 @@ THE SOFTWARE.
     if (error != hipSuccess) { \
       fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \
     exit(EXIT_FAILURE);\
-	}\
+    }\
 }
 
 void __global__
@@ -43,59 +43,60 @@ bit_extract_kernel(hipLaunchParm lp, uint32_t *C_d, const uint32_t *A_d, size_t
 #ifdef __HIP_PLATFORM_HCC__
         C_d[i] = hc::__bitextract_u32(A_d[i], 8, 4); 
 #else /* defined __HIP_PLATFORM_NVCC__ or other path */
-		C_d[i] = ((A_d[i] & 0xf00)  >> 8);
+        C_d[i] = ((A_d[i] & 0xf00)  >> 8);
 #endif
-	}
+    }
 }
 
 
 int main(int argc, char *argv[])
 {
-	uint32_t *A_d, *C_d;
+    uint32_t *A_d, *C_d;
-	uint32_t *A_h, *C_h;
+    uint32_t *A_h, *C_h;
-	size_t N = 1000000;
+    size_t N = 1000000;
-	size_t Nbytes = N * sizeof(uint32_t);
+    size_t Nbytes = N * sizeof(uint32_t);
 
     int deviceId;
     CHECK (hipGetDevice(&deviceId));
-	hipDeviceProp_t props;
+    hipDeviceProp_t props;
-	CHECK(hipGetDeviceProperties(&props, deviceId));
+    CHECK(hipGetDeviceProperties(&props, deviceId));
-	printf ("info: running on device #%d %s\n", deviceId, props.name);
+    printf ("info: running on device #%d %s\n", deviceId, props.name);
 
 
-	printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
+    printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
-	A_h = (uint32_t*)malloc(Nbytes);
+    A_h = (uint32_t*)malloc(Nbytes);
-	CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
+    CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
-	C_h = (uint32_t*)malloc(Nbytes);
+    C_h = (uint32_t*)malloc(Nbytes);
-	CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
+    CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
 
     for (size_t i=0; i<N; i++) 
-	{
+    {
-		A_h[i] = i;
+        A_h[i] = i;
-	}
+    }
 
-	printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
+    printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
-	CHECK(hipMalloc(&A_d, Nbytes));
+    CHECK(hipMalloc(&A_d, Nbytes));
-	CHECK(hipMalloc(&C_d, Nbytes));
+    CHECK(hipMalloc(&C_d, Nbytes));
 
-	printf ("info: copy Host2Device\n");
+    printf ("info: copy Host2Device\n");
     CHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
 
-	printf ("info: launch 'bit_extract_kernel' \n");
+    printf ("info: launch 'bit_extract_kernel' \n");
-	const unsigned blocks = 512;
+    const unsigned blocks = 512;
-	const unsigned threadsPerBlock = 256;
+    const unsigned threadsPerBlock = 256;
-	hipLaunchKernel(bit_extract_kernel, dim3(blocks), dim3(threadsPerBlock), 0, 0,   C_d, A_d, N);
+    hipLaunchKernel(bit_extract_kernel, dim3(blocks), dim3(threadsPerBlock), 0, 0,   C_d, A_d, N);
 
-	printf ("info: copy Device2Host\n");
+    printf ("info: copy Device2Host\n");
     CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
 
-	printf ("info: check result\n");
+    printf ("info: check result\n");
     for (size_t i=0; i<N; i++)  {
         unsigned Agold = ((A_h[i] & 0xf00) >> 8);
-		if (C_h[i] != Agold) {
+        if (C_h[i] != Agold) {
             fprintf (stderr, "mismatch detected.\n");
             printf ("%zu: %08x =? %08x (Ain=%08x)\n", i, C_h[i], Agold, A_h[i]);
-			CHECK(hipErrorUnknown);
+            CHECK(hipErrorUnknown);
-		}
+        }
-	}
+    }
+    printf ("PASSED!\n");
 }

projects/hip/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp

@@ -16,6 +16,7 @@
 int main(int argc, char *argv[])
 {
     int sizeElements = 1000000;
+    bool pass = true;
 
     // Allocate auto-managed host/device views of data:
     concurrency::array_view<float> A(sizeElements);
@@ -43,6 +44,8 @@ int main(int argc, char *argv[])
         // Because C is an array_view, the HCC runtime will copy C back to host at first access here:
         if (C[i] != ref) {
             printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C[i], ref);
+            pass = false;
         }
     };
+    if (pass) printf ("PASSED!\n");
 }

projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp

@@ -11,6 +11,7 @@ int main(int argc, char *argv[])
 {
     int sizeElements = 1000000;
     size_t sizeBytes = sizeElements * sizeof(float);
+    bool pass = true;
 
     // Allocate host memory
     float *A_h = (float*)malloc(sizeBytes);
@@ -54,6 +55,8 @@ int main(int argc, char *argv[])
         float ref= 1.618f * i + 3.142f * i;
         if (C_h[i] != ref) {
             printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C_h[i], ref);
+            pass = false;
         }
     };
+    if (pass) printf ("PASSED!\n");
 }

projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp

@@ -11,6 +11,7 @@ int main(int argc, char *argv[])
 {
     int sizeElements = 1000000;
     size_t sizeBytes = sizeElements * sizeof(float);
+    bool pass = true;
 
     // Allocate host memory
     float *A_h = (float*)malloc(sizeBytes);
@@ -48,6 +49,8 @@ int main(int argc, char *argv[])
         float ref= 1.618f * i + 3.142f * i;
         if (C_h[i] != ref) {
             printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C_h[i], ref);
+            pass = false;
         }
     };
+    if (pass) printf ("PASSED!\n");
 }

projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_array.hc

@@ -3,6 +3,7 @@
 int main(int argc, char *argv[])
 {
     int size = 1000000;
+    bool pass = true;
 
     // Allocate auto-managed host/device views of data:
     hc::array_view<float> A(size);
@@ -28,6 +29,8 @@ int main(int argc, char *argv[])
         float ref= 1.618f * i + 3.142f * i;
         if (C[i] != ref) {
             printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C[i], ref);
+            pass = false;
         }
     };
+    if (pass) printf ("PASSED!\n");
 }

projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp

@@ -16,6 +16,7 @@
 int main(int argc, char *argv[])
 {
     int sizeElements = 1000000;
+    bool pass = true;
 
     // Allocate auto-managed host/device views of data:
     hc::array_view<float> A(sizeElements);
@@ -43,6 +44,8 @@ int main(int argc, char *argv[])
         // Because C is an array_view, the HCC runtime will copy C back to host at first access here:
         if (C[i] != ref) {
             printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C[i], ref);
+            pass = false;
         }
     };
+    if (pass) printf ("PASSED!\n");
 }

projects/hip/samples/0_Intro/hcc_dialects/vadd_hip.cpp

@@ -14,6 +14,7 @@ int main(int argc, char *argv[])
 {
     int sizeElements = 1000000;
     size_t sizeBytes = sizeElements * sizeof(float);
+    bool pass = true;
 
     // Allocate host memory
     float *A_h = (float*)malloc(sizeBytes);
@@ -46,6 +47,8 @@ int main(int argc, char *argv[])
         float ref= 1.618f * i + 3.142f * i;
         if (C_h[i] != ref) {
             printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C_h[i], ref);
+            pass = false;
         }
     };
+    if (pass) printf ("PASSED!\n");
 }