Merge github pull request #41 from ekondis/master into amd-develop

Change-Id: I0b32cc7479d5e17895e0a5dffecd23c4b614518e

docs/markdown/hip_kernel_language.md

@@ -451,6 +451,8 @@ Following is the list of supported integer intrinsics. Note that intrinsics are
 | unsigned int __ffsll(long long int x) <br><sub>Find the position of least signigicant bit set to 1 in a 64 bit signed integer.</sub> |
 | unsigned int __popc ( unsigned int x ) <br><sub>Count the number of bits that are set to 1 in a 32 bit integer.</sub> |
 | int __popcll ( unsigned long long int x )<br><sub>Count the number of bits that are set to 1 in a 64 bit integer.</sub> |
+| int __mul24 ( int x, int y )<br><sub>Multiply two 24bit integers.</sub> |
+| unsigned int __umul24 ( unsigned int x, unsigned int y )<br><sub>Multiply two 24bit unsigned integers.</sub> |
 <sub><b id="f3"><sup>[1]</sup></b> 
 The hcc implementation of __ffs() and __ffsll() contains code to add a constant +1 to produce the ffs result format.
 For the cases where this overhead is not acceptable and programmer is willing to specialize for the platform, 

include/hip/hcc_detail/hip_runtime.h

@@ -430,6 +430,9 @@ __device__ unsigned int atomicInc(unsigned int* address,
 __device__ unsigned int atomicDec(unsigned int* address,
                        unsigned int val);
 
+//__mul24 __umul24
+__device__  int __mul24(int arg1, int arg2);
+__device__  unsigned int __umul24(unsigned int arg1, unsigned int arg2);
 
 // integer intrinsic function __poc __clz __ffs __brev
 __device__ unsigned int __popc( unsigned int input);

src/device_util.cpp

@@ -1672,6 +1672,17 @@ __device__  unsigned int atomicDec(unsigned int* address,
 	return hc::__atomic_wrapdec(address,val);
 }
 
+//__mul24 __umul24
+__device__  int __mul24(int arg1,
+                       int arg2)
+{
+	return hc::__mul24(arg1, arg2);
+}
+__device__  unsigned int __umul24(unsigned int arg1,
+                       unsigned int arg2)
+{
+	return hc::__mul24(arg1, arg2);
+}
 
 __device__ unsigned int test__popc(unsigned int input)
 {

tests/src/deviceLib/hipIntegerIntrinsics.cpp

@@ -35,7 +35,7 @@ __device__ void integer_intrinsics()
     __ffs((int)10);
     __ffsll((long long)10);
     //__hadd((int)1, (int)3);
-    //__mul24((int)1, (int)2);
+    __mul24((int)1, (int)2);
     //__mul64hi((long long)1, (long long)2);
     //__mulhi((int)1, (int)2);
     __popc((unsigned int)4);
@@ -45,7 +45,7 @@ __device__ void integer_intrinsics()
     //__rhadd((int)1, (int)2);
     //__sad((int)1, (int)2, 0);
     //__uhadd((unsigned int)1, (unsigned int)3);
-    //__umul24((unsigned int)1, (unsigned int)2);
+    __umul24((unsigned int)1, (unsigned int)2);
     //__umul64hi((unsigned long long)1, (unsigned long long)2);
     //__umulhi((unsigned int)1, (unsigned int)2);
     //__urhadd((unsigned int)1, (unsigned int)2);