device-libs: Use ballot(true) instead of calling read_exec builtin

The read_exec builtins are implemented with the ballot intrinsic anyway.
In the wave32 case, these will optimize down to just use the low 32-bits.
This converts a few uses, but others remain.

Apparently you can just use exec_hi as a GPR in wave32 though, so I'm not sure
we should be treating the raw exec read as assumed 0.

Change-Id: Id5621bf31b0bb7fa27456938942138f3dea85a0a

Author: arsenm

amd/device-libs/asanrtl/src/dm.cl

@@ -113,22 +113,10 @@ first(__global void * v)
 }
 
 // The number of active lanes at this point
-static uint
-active_lane_count_w64(void)
-{
-    return __builtin_popcountl(__builtin_amdgcn_read_exec());
-}
-
-static uint
-active_lane_count_w32(void)
-{
-    return __builtin_popcount(__builtin_amdgcn_read_exec_lo());
-}
-
 static uint
 active_lane_count(void)
 {
-    return __oclc_wavefrontsize64 ? active_lane_count_w64() : active_lane_count_w32();
+    return __builtin_popcountl(__builtin_amdgcn_ballot_w64(true));
 }
 
 static ulong

amd/device-libs/opencl/src/pipes/wresvnp.cl

@@ -12,11 +12,7 @@
 static uint
 active_lane_count(void)
 {
-    if (__oclc_wavefrontsize64) {
-        return __builtin_popcountl(__builtin_amdgcn_read_exec());
-    } else {
-        return __builtin_popcount(__builtin_amdgcn_read_exec_lo());
-    }
+    return __builtin_popcountl(__builtin_amdgcn_ballot_w64(true));
 }
 
 size_t