AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-253247: add ulong and ulonglong version of__shfl*

Browse Source 
Change-Id: I40ab6cfa12175f334e8392b71f567054d8256e2a
This commit is contained in:
Sarbojit Sarkar
2020-09-23 07:23:09 -04:00
committed by Sarbojit Sarkar
parent b35f5215b2
commit deeb37412c
3 changed files with 213 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/tests/src/kernel/hipShflUpDownTest.cpp
+61 -3
View File
@@ -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();
}