diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h
index d7c8d3a675..079d681c39 100644
--- a/include/hip/hcc_detail/hip_runtime.h
+++ b/include/hip/hcc_detail/hip_runtime.h
@@ -478,7 +478,7 @@ __host__ __device__ int max(int arg1, int arg2);
 __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);
+
 __device__ float __expf(float x);
 __device__ float __frsqrt_rn(float x);
 __device__ float __fsqrt_rd(float x);
@@ -486,11 +486,13 @@ __device__ float __fsqrt_rn(float x);
 __device__ float __fsqrt_ru(float x);
 __device__ float __fsqrt_rz(float x);
 __device__ float __log10f(float x);
-__device__ float __log2f(float x);
+//__device__ float __log2f(float x);
 __device__ float __logf(float x);
 __device__ float __powf(float base, float exponent);
 __device__ void __sincosf(float x, float *s, float *c) ;
-__device__ float __sinf(float x);
+extern __attribute__((const)) float __sinf(float) __asm("llvm.sin.f32");
+extern __attribute__((const)) float __cosf(float) __asm("llvm.cos.f32");
+extern __attribute__((const)) float __log2f(float) __asm("llvm.log2.f32");
 __device__ float __tanf(float x);
 __device__ float __dsqrt_rd(double x);
 __device__ float __dsqrt_rn(double x);
diff --git a/src/device_util.cpp b/src/device_util.cpp
index deb0db0a34..fc966f3395 100644
--- a/src/device_util.cpp
+++ b/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); };
diff --git a/src/hip_device.cpp b/src/hip_device.cpp
index 51fd5e4f81..371578ca2c 100644
--- a/src/hip_device.cpp
+++ b/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);
diff --git a/tests/src/deviceLib/hip_trig.cpp b/tests/src/deviceLib/hip_trig.cpp
new file mode 100644
index 0000000000..7f9b5d60b0
--- /dev/null
+++ b/tests/src/deviceLib/hip_trig.cpp
@@ -0,0 +1,81 @@
+/*
+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.
+*/
+
+/* HIT_START
+ * BUILD: %t %s 
+ * RUN: %t
+ * HIT_END
+ */
+
+#include<hip/hip_runtime_api.h>
+#include<hip/hip_runtime.h>
+#include<iostream>
+#include"test_common.h"
+
+#define LEN 512
+#define SIZE LEN<<2
+
+__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 = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_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]);
+}
+
+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;
+  In = new float[LEN];
+  sin_h = new float[LEN];
+  cos_h = new float[LEN];
+  tan_h = new float[LEN];
+  sin_ph = new float[LEN];
+  cos_ph = new float[LEN];
+  for(int i=0;i<LEN;i++){
+    In[i] = 1.0f;
+    sin_h[i] = 0.0f;
+    cos_h[i] = 0.0f;
+    tan_h[i] = 0.0f;
+    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);
+  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, sin_pd, cos_pd);
+  hipMemcpy(sin_h, sin_d, SIZE, hipMemcpyDeviceToHost);
+  hipMemcpy(cos_h, cos_d, SIZE, hipMemcpyDeviceToHost);
+  hipMemcpy(tan_h, tan_d, SIZE, hipMemcpyDeviceToHost);
+  hipMemcpy(sin_ph, sin_pd, SIZE, hipMemcpyDeviceToHost);
+  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;
+    }
+  }
+  passed();
+}