From d253806e64c2ec6c362cde04e479ca463748bd98 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 3 May 2016 14:19:25 +0530 Subject: [PATCH 1/2] bit_extract reports PASSED when passed [ROCm/hip commit: cb6a5d94214fade37021f2ba11598c63152a391b] --- .../0_Intro/bit_extract/bit_extract.cpp | 65 ++++++++++--------- 1 file changed, 33 insertions(+), 32 deletions(-) diff --git a/projects/hip/samples/0_Intro/bit_extract/bit_extract.cpp b/projects/hip/samples/0_Intro/bit_extract/bit_extract.cpp index bdc8182c38..53302b7228 100644 --- a/projects/hip/samples/0_Intro/bit_extract/bit_extract.cpp +++ b/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_h, *C_h; - size_t N = 1000000; - size_t Nbytes = N * sizeof(uint32_t); + uint32_t *A_d, *C_d; + uint32_t *A_h, *C_h; + size_t N = 1000000; + size_t Nbytes = N * sizeof(uint32_t); int deviceId; CHECK (hipGetDevice(&deviceId)); - hipDeviceProp_t props; - CHECK(hipGetDeviceProperties(&props, deviceId)); - printf ("info: running on device #%d %s\n", deviceId, props.name); + hipDeviceProp_t props; + CHECK(hipGetDeviceProperties(&props, deviceId)); + 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); - A_h = (uint32_t*)malloc(Nbytes); - CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); - C_h = (uint32_t*)malloc(Nbytes); - CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); + printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); + A_h = (uint32_t*)malloc(Nbytes); + CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); + C_h = (uint32_t*)malloc(Nbytes); + CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); for (size_t i=0; i> 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"); } From fc73044bf2ff7c7e561e3bf264557d153ddd1263 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 3 May 2016 14:32:59 +0530 Subject: [PATCH 2/2] hcc_dialects report PASSED when passed [ROCm/hip commit: 6181988232dcd03443a81365ea42bae177b5db18] --- .../hip/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp | 3 +++ projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp | 3 +++ projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp | 3 +++ projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_array.hc | 3 +++ .../hip/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp | 3 +++ projects/hip/samples/0_Intro/hcc_dialects/vadd_hip.cpp | 3 +++ 6 files changed, 18 insertions(+) diff --git a/projects/hip/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp b/projects/hip/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp index 6fdea5d830..485b64f68d 100644 --- a/projects/hip/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp +++ b/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 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"); } diff --git a/projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp b/projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp index df6a831fd2..5cb2e8c98f 100644 --- a/projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp +++ b/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"); } diff --git a/projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp b/projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp index 63f455a615..bda3adf376 100644 --- a/projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp +++ b/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"); } diff --git a/projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_array.hc b/projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_array.hc index 8a9f8e71e7..d57b9a7e14 100644 --- a/projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_array.hc +++ b/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 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"); } diff --git a/projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp b/projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp index 41d9124ddc..2585f47001 100644 --- a/projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp +++ b/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 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"); } diff --git a/projects/hip/samples/0_Intro/hcc_dialects/vadd_hip.cpp b/projects/hip/samples/0_Intro/hcc_dialects/vadd_hip.cpp index 31751b3419..9d223ba276 100644 --- a/projects/hip/samples/0_Intro/hcc_dialects/vadd_hip.cpp +++ b/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"); }