Deleted asm signed bit field extract -- it did not work. Changed xtract32 to use amdgcn_builtin with proper #if defined.

Author: gwoltman

src/cl/carryutil.cl

@@ -15,11 +15,10 @@ void updateStats(global uint *bufROE, u32 posROE, float roundMax) {
 }
 #endif
 
-#if 0 && HAS_ASM
-i32  lowBits(i32 u, u32 bits) { i32 tmp; __asm("v_bfe_i32 %0, %1, 0, %2" : "=v" (tmp) : "v" (u), "v" (bits)); return tmp; }
-i32 xtract32(i64 x, u32 bits) { i32 tmp; __asm("v_alignbit_b32 %0, %1, %2, %3" : "=v"(tmp) : "v"(as_int2(x).y), "v"(as_int2(x).x), "v"(bits)); return tmp; }
+i32 lowBits(i32 u, u32 bits) { return ((u << (32 - bits)) >> (32 - bits)); }
+#if defined(__has_builtin) && __has_builtin(__builtin_amdgcn_alignbit)
+i32 xtract32(i64 x, u32 bits) { return __builtin_amdgcn_alignbit(as_int2(x).y, as_int2(x).x, bits); }
 #else
-i32  lowBits(i32 u, u32 bits) { return ((u << (32 - bits)) >> (32 - bits)); }
 i32 xtract32(i64 x, u32 bits) { return x >> bits; }
 #endif