replace __inline__ with __forceinline__

KlausT · KlausT · commit 35e6a9c0239e · 2019-02-03T16:10:49.000+01:00
diff --git a/cuda_helper.h b/cuda_helper.h
@@ -668,7 +668,7 @@ static __device__ __forceinline__ uint2 operator* (uint2 a, uint2 b)
 
 // uint2 method
 #if  __CUDA_ARCH__ >= 320 && !defined NOASM
-static __device__ __inline__ uint2 ROR2(const uint2 a, const int offset)
+static __device__ __forceinline__ uint2 ROR2(const uint2 a, const int offset)
 {
 	uint2 result;
 	if (offset < 32) {
@@ -682,7 +682,7 @@ static __device__ __inline__ uint2 ROR2(const uint2 a, const int offset)
 	return result;
 }
 #else
-static __device__ __inline__ uint2 ROR2(const uint2 v, const int n)
+static __device__ __forceinline__ uint2 ROR2(const uint2 v, const int n)
 {
 	uint2 result;
 	if (n <= 32) 
@@ -699,26 +699,26 @@ static __device__ __inline__ uint2 ROR2(const uint2 v, const int n)
 }
 #endif
 
-static __device__ __inline__ uint32_t ROL8(const uint32_t x)
+static __device__ __forceinline__ uint32_t ROL8(const uint32_t x)
 {
 #ifdef __CUDA_ARCH__
 	return __byte_perm(x, x, 0x2103);
 #endif
 }
-static __device__ __inline__ uint32_t ROL16(const uint32_t x)
+static __device__ __forceinline__ uint32_t ROL16(const uint32_t x)
 {
 #ifdef __CUDA_ARCH__
 	return __byte_perm(x, x, 0x1032);
 #endif
 }
-static __device__ __inline__ uint32_t ROL24(const uint32_t x)
+static __device__ __forceinline__ uint32_t ROL24(const uint32_t x)
 {
 #ifdef __CUDA_ARCH__
 	return __byte_perm(x, x, 0x0321);
 #endif
 }
 
-static __device__ __inline__ uint2 ROR8(const uint2 a)
+static __device__ __forceinline__ uint2 ROR8(const uint2 a)
 {
 #ifdef __CUDA_ARCH__
 	uint2 result;
@@ -729,7 +729,7 @@ static __device__ __inline__ uint2 ROR8(const uint2 a)
 #endif
 }
 
-static __device__ __inline__ uint2 ROR16(const uint2 a)
+static __device__ __forceinline__ uint2 ROR16(const uint2 a)
 {
 #ifdef __CUDA_ARCH__
 	uint2 result;
@@ -740,7 +740,7 @@ static __device__ __inline__ uint2 ROR16(const uint2 a)
 #endif
 }
 
-static __device__ __inline__ uint2 ROR24(const uint2 a)
+static __device__ __forceinline__ uint2 ROR24(const uint2 a)
 {
 #ifdef __CUDA_ARCH__
 	uint2 result;
@@ -751,7 +751,7 @@ static __device__ __inline__ uint2 ROR24(const uint2 a)
 #endif
 }
 
-static __device__ __inline__ uint2 ROL8(const uint2 a)
+static __device__ __forceinline__ uint2 ROL8(const uint2 a)
 {
 #ifdef __CUDA_ARCH__
 	uint2 result;
@@ -762,7 +762,7 @@ static __device__ __inline__ uint2 ROL8(const uint2 a)
 #endif
 }
 
-static __device__ __inline__ uint2 ROL16(const uint2 a)
+static __device__ __forceinline__ uint2 ROL16(const uint2 a)
 {
 #ifdef __CUDA_ARCH__
 	uint2 result;
@@ -773,7 +773,7 @@ static __device__ __inline__ uint2 ROL16(const uint2 a)
 #endif
 }
 
-static __device__ __inline__ uint2 ROL24(const uint2 a)
+static __device__ __forceinline__ uint2 ROL24(const uint2 a)
 {
 #ifdef __CUDA_ARCH__
 	uint2 result;
@@ -787,7 +787,7 @@ static __device__ __inline__ uint2 ROL24(const uint2 a)
 #if  __CUDA_ARCH__ >= 320 && !defined NOASM
 
 
-__inline__ static __device__ uint2 ROL2(const uint2 a, const int offset) {
+__forceinline__ static __device__ uint2 ROL2(const uint2 a, const int offset) {
 	uint2 result;
 	if (offset >= 32) {
 		asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(a.x), "r"(a.y), "r"(offset));
@@ -800,7 +800,7 @@ __inline__ static __device__ uint2 ROL2(const uint2 a, const int offset) {
 	return result;
 }
 #else
-__inline__ static __device__ uint2 ROL2(const uint2 v, const int n)
+__forceinline__ static __device__ uint2 ROL2(const uint2 v, const int n)
 {
 		uint2 result;
 		if (n <= 32) 

Original file line number	Diff line number	Diff line change
`@@ -668,7 +668,7 @@ static __device__ __forceinline__ uint2 operator* (uint2 a, uint2 b)`
`668`	`668`
`669`	`669`	`// uint2 method`
`670`	`670`	`#if __CUDA_ARCH__ >= 320 && !defined NOASM`
`671`		`-static __device__ __inline__ uint2 ROR2(const uint2 a, const int offset)`
	`671`	`+static __device__ __forceinline__ uint2 ROR2(const uint2 a, const int offset)`
`672`	`672`	`{`
`673`	`673`	`uint2 result;`
`674`	`674`	`if (offset < 32) {`
`@@ -682,7 +682,7 @@ static __device__ __inline__ uint2 ROR2(const uint2 a, const int offset)`
`682`	`682`	`return result;`
`683`	`683`	`}`
`684`	`684`	`#else`
`685`		`-static __device__ __inline__ uint2 ROR2(const uint2 v, const int n)`
	`685`	`+static __device__ __forceinline__ uint2 ROR2(const uint2 v, const int n)`
`686`	`686`	`{`
`687`	`687`	`uint2 result;`
`688`	`688`	`if (n <= 32)`
`@@ -699,26 +699,26 @@ static __device__ __inline__ uint2 ROR2(const uint2 v, const int n)`
`699`	`699`	`}`
`700`	`700`	`#endif`
`701`	`701`
`702`		`-static __device__ __inline__ uint32_t ROL8(const uint32_t x)`
	`702`	`+static __device__ __forceinline__ uint32_t ROL8(const uint32_t x)`
`703`	`703`	`{`
`704`	`704`	`#ifdef __CUDA_ARCH__`
`705`	`705`	`return __byte_perm(x, x, 0x2103);`
`706`	`706`	`#endif`
`707`	`707`	`}`
`708`		`-static __device__ __inline__ uint32_t ROL16(const uint32_t x)`
	`708`	`+static __device__ __forceinline__ uint32_t ROL16(const uint32_t x)`
`709`	`709`	`{`
`710`	`710`	`#ifdef __CUDA_ARCH__`
`711`	`711`	`return __byte_perm(x, x, 0x1032);`
`712`	`712`	`#endif`
`713`	`713`	`}`
`714`		`-static __device__ __inline__ uint32_t ROL24(const uint32_t x)`
	`714`	`+static __device__ __forceinline__ uint32_t ROL24(const uint32_t x)`
`715`	`715`	`{`
`716`	`716`	`#ifdef __CUDA_ARCH__`
`717`	`717`	`return __byte_perm(x, x, 0x0321);`
`718`	`718`	`#endif`
`719`	`719`	`}`
`720`	`720`
`721`		`-static __device__ __inline__ uint2 ROR8(const uint2 a)`
	`721`	`+static __device__ __forceinline__ uint2 ROR8(const uint2 a)`
`722`	`722`	`{`
`723`	`723`	`#ifdef __CUDA_ARCH__`
`724`	`724`	`uint2 result;`
`@@ -729,7 +729,7 @@ static __device__ __inline__ uint2 ROR8(const uint2 a)`
`729`	`729`	`#endif`
`730`	`730`	`}`
`731`	`731`
`732`		`-static __device__ __inline__ uint2 ROR16(const uint2 a)`
	`732`	`+static __device__ __forceinline__ uint2 ROR16(const uint2 a)`
`733`	`733`	`{`
`734`	`734`	`#ifdef __CUDA_ARCH__`
`735`	`735`	`uint2 result;`
`@@ -740,7 +740,7 @@ static __device__ __inline__ uint2 ROR16(const uint2 a)`
`740`	`740`	`#endif`
`741`	`741`	`}`
`742`	`742`
`743`		`-static __device__ __inline__ uint2 ROR24(const uint2 a)`
	`743`	`+static __device__ __forceinline__ uint2 ROR24(const uint2 a)`
`744`	`744`	`{`
`745`	`745`	`#ifdef __CUDA_ARCH__`
`746`	`746`	`uint2 result;`
`@@ -751,7 +751,7 @@ static __device__ __inline__ uint2 ROR24(const uint2 a)`
`751`	`751`	`#endif`
`752`	`752`	`}`
`753`	`753`
`754`		`-static __device__ __inline__ uint2 ROL8(const uint2 a)`
	`754`	`+static __device__ __forceinline__ uint2 ROL8(const uint2 a)`
`755`	`755`	`{`
`756`	`756`	`#ifdef __CUDA_ARCH__`
`757`	`757`	`uint2 result;`
`@@ -762,7 +762,7 @@ static __device__ __inline__ uint2 ROL8(const uint2 a)`
`762`	`762`	`#endif`
`763`	`763`	`}`
`764`	`764`
`765`		`-static __device__ __inline__ uint2 ROL16(const uint2 a)`
	`765`	`+static __device__ __forceinline__ uint2 ROL16(const uint2 a)`
`766`	`766`	`{`
`767`	`767`	`#ifdef __CUDA_ARCH__`
`768`	`768`	`uint2 result;`
`@@ -773,7 +773,7 @@ static __device__ __inline__ uint2 ROL16(const uint2 a)`
`773`	`773`	`#endif`
`774`	`774`	`}`
`775`	`775`
`776`		`-static __device__ __inline__ uint2 ROL24(const uint2 a)`
	`776`	`+static __device__ __forceinline__ uint2 ROL24(const uint2 a)`
`777`	`777`	`{`
`778`	`778`	`#ifdef __CUDA_ARCH__`
`779`	`779`	`uint2 result;`
`@@ -787,7 +787,7 @@ static __device__ __inline__ uint2 ROL24(const uint2 a)`
`787`	`787`	`#if __CUDA_ARCH__ >= 320 && !defined NOASM`
`788`	`788`
`789`	`789`
`790`		`-__inline__ static __device__ uint2 ROL2(const uint2 a, const int offset) {`
	`790`	`+__forceinline__ static __device__ uint2 ROL2(const uint2 a, const int offset) {`
`791`	`791`	`uint2 result;`
`792`	`792`	`if (offset >= 32) {`
`793`	`793`	`asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(a.x), "r"(a.y), "r"(offset));`
`@@ -800,7 +800,7 @@ __inline__ static __device__ uint2 ROL2(const uint2 a, const int offset) {`
`800`	`800`	`return result;`
`801`	`801`	`}`
`802`	`802`	`#else`
`803`		`-__inline__ static __device__ uint2 ROL2(const uint2 v, const int n)`
	`803`	`+__forceinline__ static __device__ uint2 ROL2(const uint2 v, const int n)`
`804`	`804`	`{`
`805`	`805`	`uint2 result;`
`806`	`806`	`if (n <= 32)`