SWDEV-253247: add ulong and ulonglong version of__shfl*

Change-Id: I40ab6cfa12175f334e8392b71f567054d8256e2a

tests/src/kernel/hipShflTests.cpp

@@ -57,6 +57,15 @@ void matrixTransposeCPUReference(T* output, T* input, const unsigned int width)
     }
 }
 
+void getFactor(int& fact) { fact = 101; }
+void getFactor(unsigned int& fact) { fact = static_cast<unsigned int>(INT32_MAX)+1; }
+void getFactor(float& fact) { fact = 2.5; }
+void getFactor(double& fact) { fact = 2.5; }
+void getFactor(long& fact) { fact = 202; }
+void getFactor(unsigned long& fact) { fact = static_cast<unsigned long>(__LONG_MAX__)+1; }
+void getFactor(long long& fact) { fact = 303; }
+void getFactor(unsigned long long& fact) { fact = static_cast<unsigned long long>(__LONG_LONG_MAX__)+1; }
+
 template<typename T>
 void runTest() {
     T* Matrix;
@@ -77,8 +86,10 @@ void runTest() {
     cpuTransposeMatrix = (T*)malloc(NUM * sizeof(T));
 
     // initialize the input data
+    T factor;
+    getFactor(factor);
     for (i = 0; i < NUM; i++) {
-        Matrix[i] = (T)i * 10l;
+        Matrix[i] = (T)i + factor;
     }
 
     // allocate the memory on the device side
@@ -124,7 +135,11 @@ void runTest() {
 int main() {
     runTest<int>();
     runTest<float>();
+    runTest<double>();
     runTest<long>();
     runTest<long long>();
+    runTest<unsigned int>();
+    runTest<unsigned long>();
+    runTest<unsigned long long>();
     passed();
 }

tests/src/kernel/hipShflUpDownTest.cpp

@@ -47,13 +47,31 @@ __global__ void shflUpSum(T* a, int size) {
     a[threadIdx.x] = val;
 }
 
+template <typename T>
+__global__ void shflXorSum(T* a, int size) {
+  T val = a[threadIdx.x];
+  for (int i = size/2; i > 0; i /= 2)
+    val += __shfl_xor(val, i, size);
+  a[threadIdx.x] = val;
+}
+
+void getFactor(int& fact) { fact = 101; }
+void getFactor(unsigned int& fact) { fact = static_cast<unsigned int>(INT32_MAX)+1; }
+void getFactor(float& fact) { fact = 2.5; }
+void getFactor(double& fact) { fact = 2.5; }
+void getFactor(long& fact) { fact = 202; }
+void getFactor(unsigned long& fact) { fact = static_cast<unsigned long>(__LONG_MAX__)+1; }
+void getFactor(long long& fact) { fact = 303; }
+void getFactor(unsigned long long& fact) { fact = static_cast<unsigned long long>(__LONG_LONG_MAX__)+1; }
+
 template <typename T>
 void runTestShflUp() {
     const int size = 32;
     T a[size];
     T cpuSum = 0;
+    T factor; getFactor(factor);
     for (int i = 0; i < size; i++) {
-        a[i] = i;
+        a[i] = i + factor;
         cpuSum += a[i];
     }
     T* d_a;
@@ -73,8 +91,9 @@ void runTestShflDown() {
     const int size = 32;
     T a[size];
     T cpuSum = 0;
+    T factor; getFactor(factor);
     for (int i = 0; i < size; i++) {
-        a[i] = i;
+        a[i] = i + factor;
         cpuSum += a[i];
     }
     T* d_a;
@@ -84,19 +103,58 @@ void runTestShflDown() {
     hipMemcpy(&a, d_a, sizeof(T) * size, hipMemcpyDefault);
     if (a[0] != cpuSum) {
         hipFree(d_a);
-        failed("Shfl Up Sum did not match.");
+        failed("Shfl Down Sum did not match.");
+    }
+    hipFree(d_a);
+}
+
+template <typename T>
+void runTestShflXor() {
+    const int size = 32;
+    T a[size];
+    T cpuSum = 0;
+    T factor; getFactor(factor);
+    for (int i = 0; i < size; i++) {
+        a[i] = i + factor;
+        cpuSum += a[i];
+    }
+    T* d_a;
+    hipMalloc(&d_a, sizeof(T) * size);
+    hipMemcpy(d_a, &a, sizeof(T) * size, hipMemcpyDefault);
+    hipLaunchKernelGGL(shflXorSum<T>, 1, size, 0, 0, d_a, size);
+    hipMemcpy(&a, d_a, sizeof(T) * size, hipMemcpyDefault);
+    if (a[0] != cpuSum) {
+        hipFree(d_a);
+        failed("Shfl Xor Sum did not match.");
     }
     hipFree(d_a);
 }
 int main() {
     runTestShflUp<int>();
     runTestShflUp<float>();
+    runTestShflUp<double>();
     runTestShflUp<long>();
     runTestShflUp<long long>();
+    runTestShflUp<unsigned int>();
+    runTestShflUp<unsigned long>();
+    runTestShflUp<unsigned long long>();
 
     runTestShflDown<int>();
     runTestShflDown<float>();
+    runTestShflDown<double>();
     runTestShflDown<long>();
     runTestShflDown<long long>();
+    runTestShflDown<unsigned int>();
+    runTestShflDown<unsigned long>();
+    runTestShflDown<unsigned long long>();
+
+    runTestShflXor<int>();
+    runTestShflXor<float>();
+    runTestShflXor<double>();
+    runTestShflXor<long>();
+    runTestShflXor<long long>();
+    runTestShflXor<unsigned int>();
+    runTestShflXor<unsigned long>();
+    runTestShflXor<unsigned long long>();
     passed();
 }