From 2bc541e15f3acf2a0a86a4f1529bf728e685be51 Mon Sep 17 00:00:00 2001 From: Elias Konstantinidis Date: Sun, 2 Oct 2016 10:07:31 +0300 Subject: [PATCH 1/3] Added support for __mul24 and __umul24 --- include/hip/hcc_detail/hip_runtime.h | 3 +++ src/device_util.cpp | 11 +++++++++++ 2 files changed, 14 insertions(+) mode change 100644 => 100755 include/hip/hcc_detail/hip_runtime.h mode change 100644 => 100755 src/device_util.cpp diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h old mode 100644 new mode 100755 index c89faaae11..47b75f282d --- a/include/hip/hcc_detail/hip_runtime.h +++ b/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); diff --git a/src/device_util.cpp b/src/device_util.cpp old mode 100644 new mode 100755 index 21a58d7070..6c608d891e --- a/src/device_util.cpp +++ b/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) { From 0b4b1b7df812439436621c5bb8bfab42eef0cfb2 Mon Sep 17 00:00:00 2001 From: Elias Konstantinidis Date: Sun, 2 Oct 2016 10:08:07 +0300 Subject: [PATCH 2/3] Added __mul24 & __umul24 documentation entry --- docs/markdown/hip_kernel_language.md | 2 ++ 1 file changed, 2 insertions(+) mode change 100644 => 100755 docs/markdown/hip_kernel_language.md diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md old mode 100644 new mode 100755 index ed9dc2fbe9..4b8762025e --- a/docs/markdown/hip_kernel_language.md +++ b/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)
Find the position of least signigicant bit set to 1 in a 64 bit signed integer. | | unsigned int __popc ( unsigned int x )
Count the number of bits that are set to 1 in a 32 bit integer. | | int __popcll ( unsigned long long int x )
Count the number of bits that are set to 1 in a 64 bit integer. | +| int __mul24 ( int x, int y )
Multiply two 24bit integers. | +| unsigned int __umul24 ( unsigned int x, unsigned int y )
Multiply two 24bit unsigned integers. | [1] 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, From bfd38972bdcb35f91cb35aeab17d3b4b28fb42cb Mon Sep 17 00:00:00 2001 From: Elias Konstantinidis Date: Sun, 2 Oct 2016 08:16:10 +0000 Subject: [PATCH 3/3] Enabled tests for __mul24 & __umul24 --- tests/src/deviceLib/hipIntegerIntrinsics.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/src/deviceLib/hipIntegerIntrinsics.cpp b/tests/src/deviceLib/hipIntegerIntrinsics.cpp index 250dc1d949..abe0b66e90 100644 --- a/tests/src/deviceLib/hipIntegerIntrinsics.cpp +++ b/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);