added fast math APIs

1. Added fast math apis for sin, cos, tan, sincos
2. Added test for trig math functions
3. Added logarithm fast math
4. Changed how hipGetDevice, hipDeviceGetCacheConfig emit errors

Change-Id: Ie6ab594ddd5853cbe85e39a2f6d3479a807fa323

src/device_util.cpp

@@ -156,7 +156,7 @@ __device__ char4 __hip_hc_sub8pk(char4 in1, char4 in2) {
     one1 = in1.val & MASK2;
     one2 = in2.val & MASK2;
     out.val = out.val | ((one1 - one2) & MASK2);
-    return out; 
+    return out;
 }
 
 __device__ char4 __hip_hc_mul8pk(char4 in1, char4 in2) {
@@ -2045,7 +2045,7 @@ __device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr()
 
 
 //TODO - add a couple fast math operations here, the set here will grow :
-__device__  float __cosf(float x) {return hc::fast_math::cosf(x); };
+//__device__  float __cosf(float x) {return hc::fast_math::cosf(x); };
 __device__  float __expf(float x) {return hc::fast_math::expf(x); };
 __device__  float __frsqrt_rn(float x) {return hc::fast_math::rsqrt(x); };
 __device__  float __fsqrt_rd(float x) {return hc::fast_math::sqrt(x); };
@@ -2053,12 +2053,12 @@ __device__  float __fsqrt_rn(float x) {return hc::fast_math::sqrt(x); };
 __device__  float __fsqrt_ru(float x) {return hc::fast_math::sqrt(x); };
 __device__  float __fsqrt_rz(float x) {return hc::fast_math::sqrt(x); };
 __device__  float __log10f(float x) {return hc::fast_math::log10f(x); };
-__device__  float __log2f(float x) {return hc::fast_math::log2f(x); };
+//__device__  float __log2f(float x) {return hc::fast_math::log2f(x); };
 __device__  float __logf(float x) {return hc::fast_math::logf(x); };
 __device__  float __powf(float base, float exponent) {return hc::fast_math::powf(base, exponent); };
-__device__  void __sincosf(float x, float *s, float *c) {return hc::fast_math::sincosf(x, s, c); };
-__device__  float __sinf(float x) {return hc::fast_math::sinf(x); };
-__device__  float __tanf(float x) {return hc::fast_math::tanf(x); };
+__device__  void __sincosf(float x, float *s, float *c) { *s = __sinf(x); *c = __cosf(x); };
+//__device__  float __sinf(float x) {return hc::fast_math::sinf(x); };
+__device__  float __tanf(float x) {return __sinf(x)/__cosf(x); };
 __device__  float __dsqrt_rd(double x) {return hc::fast_math::sqrt(x); };
 __device__  float __dsqrt_rn(double x) {return hc::fast_math::sqrt(x); };
 __device__  float __dsqrt_ru(double x) {return hc::fast_math::sqrt(x); };

src/hip_device.cpp

@@ -44,7 +44,7 @@ hipError_t hipGetDevice(int *deviceId)
             *deviceId = ctx->getDevice()->_deviceId;
         }
     }else{
-        e = hipErrorInvalidDevice;
+        e = hipErrorInvalidValue;
     }
 
     return ihipLogStatus(e);
@@ -66,7 +66,7 @@ hipError_t hipGetDeviceCount(int *count)
             e = ihipLogStatus(hipErrorNoDevice);
         }
     } else {
-        e = ihipLogStatus(hipErrorNoDevice);
+        e = ihipLogStatus(hipErrorInvalidValue);
     }
     return e;
 }
@@ -84,6 +84,10 @@ hipError_t hipDeviceGetCacheConfig(hipFuncCache_t *cacheConfig)
 {
     HIP_INIT_API(cacheConfig);
 
+    if(cacheConfig == nullptr) {
+        return ihipLogStatus(hipErrorInvalidValue);
+    }
+
     *cacheConfig = hipFuncCachePreferNone;
 
     return ihipLogStatus(hipSuccess);