From 8408bc55363e5692c7b6280ee8f4ea42581e839f Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Fri, 21 Sep 2018 15:28:10 -0400 Subject: [PATCH 1/3] Improve hip_trig test case Hip_trig is failing but returning a passing value on HIP-Clang. Also update with debug output and free arrays used. [ROCm/hip commit: 357b6844fa9958a98bf9b5eba16fd4efa00765a2] --- projects/hip/tests/src/deviceLib/hip_trig.cpp | 43 +++++++++++++++---- 1 file changed, 34 insertions(+), 9 deletions(-) diff --git a/projects/hip/tests/src/deviceLib/hip_trig.cpp b/projects/hip/tests/src/deviceLib/hip_trig.cpp index 3f200fcf8f..b7542c25a4 100644 --- a/projects/hip/tests/src/deviceLib/hip_trig.cpp +++ b/projects/hip/tests/src/deviceLib/hip_trig.cpp @@ -35,6 +35,8 @@ THE SOFTWARE. #define LEN 512 #define SIZE LEN << 2 +#define TEST_DEBUG (0) + __global__ void kernel_trig(hipLaunchParm lp, float* In, float* sin_d, float* cos_d, float* tan_d, float* sin_pd, float* cos_pd) { int tid = threadIdx.x + blockIdx.x * blockDim.x; @@ -47,6 +49,7 @@ __global__ void kernel_trig(hipLaunchParm lp, float* In, float* sin_d, float* co int main() { float *In, *sin_h, *cos_h, *tan_h, *sin_ph, *cos_ph; float *In_d, *sin_d, *cos_d, *tan_d, *sin_pd, *cos_pd; + int errors = 0; In = new float[LEN]; sin_h = new float[LEN]; cos_h = new float[LEN]; @@ -61,14 +64,16 @@ int main() { sin_ph[i] = 0.0f; cos_ph[i] = 0.0f; } - hipMalloc((void**)&In_d, SIZE); - hipMalloc((void**)&sin_d, SIZE); - hipMalloc((void**)&cos_d, SIZE); - hipMalloc((void**)&tan_d, SIZE); - hipMalloc((void**)&sin_pd, SIZE); - hipMalloc((void**)&cos_pd, SIZE); + HIP_ASSERT(hipMalloc((void**)&In_d, SIZE)); + HIP_ASSERT(hipMalloc((void**)&sin_d, SIZE)); + HIP_ASSERT(hipMalloc((void**)&cos_d, SIZE)); + HIP_ASSERT(hipMalloc((void**)&tan_d, SIZE)); + HIP_ASSERT(hipMalloc((void**)&sin_pd, SIZE)); + HIP_ASSERT(hipMalloc((void**)&cos_pd, SIZE)); + hipMemcpy(In_d, In, SIZE, hipMemcpyHostToDevice); - hipLaunchKernel(kernel_trig, dim3(LEN, 1, 1), dim3(1, 1, 1), 0, 0, In_d, sin_d, cos_d, tan_d, + hipLaunchKernel(kernel_trig, dim3(LEN, 1, 1), dim3(1, 1, 1), 0, 0, + In_d, sin_d, cos_d, tan_d, sin_pd, cos_pd); hipMemcpy(sin_h, sin_d, SIZE, hipMemcpyDeviceToHost); hipMemcpy(cos_h, cos_d, SIZE, hipMemcpyDeviceToHost); @@ -77,8 +82,28 @@ int main() { hipMemcpy(cos_ph, cos_pd, SIZE, hipMemcpyDeviceToHost); for (int i = 0; i < LEN; i++) { if (sin_h[i] != sin_ph[i] || cos_h[i] != cos_ph[i] || tan_h[i] * cos_h[i] != sin_h[i]) { - std::cout << "Failed!" << std::endl; + errors++; +#if TEST_DEBUG + std::cout << "Check Failed!" << std::endl; + std::cout << " sin_h: " << sin_h[i] << " sin_ph: " << sin_ph[i] << "\n" + << " cos_h: " << cos_h[i] << " cos_ph:" << cos_ph[i] << "\n" + << " tan_h * cos_h: " << tan_h[i] * cos_h[i] << " sin_h[i]: " << sin_h[i] << "\n"; +#endif } } - passed(); + + HIP_ASSERT(hipFree(In_d)); + HIP_ASSERT(hipFree(sin_d)); + HIP_ASSERT(hipFree(cos_d)); + HIP_ASSERT(hipFree(tan_d)); + HIP_ASSERT(hipFree(sin_pd)); + HIP_ASSERT(hipFree(cos_pd)); + + if (errors != 0) { + std::cout << "hip_trig FAILED!" << std::endl; + return -1; + } else { + std::cout << "hip_trig PASSED!" << std::endl; + } + return errors; } From 9762215e54426d3b842e480c98d8dc74f22a9b38 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Mon, 24 Sep 2018 15:01:24 +0000 Subject: [PATCH 2/3] Fix missing HIP_ASSERT in hip_trig [ROCm/hip commit: 3d3820272a45a9be58d0c77ee4de18c4c620eb24] --- projects/hip/tests/src/deviceLib/hip_trig.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/hip/tests/src/deviceLib/hip_trig.cpp b/projects/hip/tests/src/deviceLib/hip_trig.cpp index b7542c25a4..29c24cf5b4 100644 --- a/projects/hip/tests/src/deviceLib/hip_trig.cpp +++ b/projects/hip/tests/src/deviceLib/hip_trig.cpp @@ -32,6 +32,8 @@ THE SOFTWARE. #include "test_common.h" #include +#define HIP_ASSERT(x) (assert((x) == hipSuccess)) + #define LEN 512 #define SIZE LEN << 2 From 93a84365abecedb0636372a2ba8b57a00a65e97f Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 25 Sep 2018 15:58:36 +0000 Subject: [PATCH 3/3] Use trig functions from ocml instead [ROCm/hip commit: 5179a72cda7dfebc45c167755e5d439974f7da92] --- projects/hip/tests/src/deviceLib/hip_trig.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/hip/tests/src/deviceLib/hip_trig.cpp b/projects/hip/tests/src/deviceLib/hip_trig.cpp index 29c24cf5b4..7b076065fa 100644 --- a/projects/hip/tests/src/deviceLib/hip_trig.cpp +++ b/projects/hip/tests/src/deviceLib/hip_trig.cpp @@ -42,10 +42,10 @@ THE SOFTWARE. __global__ void kernel_trig(hipLaunchParm lp, float* In, float* sin_d, float* cos_d, float* tan_d, float* sin_pd, float* cos_pd) { int tid = threadIdx.x + blockIdx.x * blockDim.x; - sin_d[tid] = __sinf(In[tid]); - cos_d[tid] = __cosf(In[tid]); - tan_d[tid] = __tanf(In[tid]); - __sincosf(In[tid], &sin_pd[tid], &cos_pd[tid]); + sin_d[tid] = sinf(In[tid]); + cos_d[tid] = cosf(In[tid]); + tan_d[tid] = tanf(In[tid]); + sincosf(In[tid], &sin_pd[tid], &cos_pd[tid]); } int main() {