Skip to content

Commit

Permalink
[SYCLomatic #1361] Add 24 math API tests.
Browse files Browse the repository at this point in the history
Signed-off-by: Tang, Jiajun jiajun.tang@intel.com
  • Loading branch information
tangjj11 committed Nov 14, 2023
1 parent 7843c51 commit 5d70717
Show file tree
Hide file tree
Showing 4 changed files with 1,114 additions and 0 deletions.
278 changes: 278 additions & 0 deletions features/feature_case/math/math-emu-double.cu
Original file line number Diff line number Diff line change
Expand Up @@ -256,6 +256,200 @@ void testRnorm4dCases(const vector<pair<d_tuple4, di_pair>> &TestCases) {
}
}

// Double Precision Intrinsics

__global__ void dadd_rd(float *const Result, float Input1, float Input2) {
*Result = __dadd_rd(Input1, Input2);
}

void testDadd_rdCases(
const vector<pair<pair<float, float>, di_pair>> &TestCases) {
float *Result;
cudaMallocManaged(&Result, sizeof(*Result));
for (const auto &TestCase : TestCases) {
dadd_rd<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second);
cudaDeviceSynchronize();
checkResult("__dadd_rd", {TestCase.first.first, TestCase.first.second},
TestCase.second.first, *Result, TestCase.second.second);
}
}

__global__ void dadd_rn(float *const Result, float Input1, float Input2) {
*Result = __dadd_rn(Input1, Input2);
}

void testDadd_rnCases(
const vector<pair<pair<float, float>, di_pair>> &TestCases) {
float *Result;
cudaMallocManaged(&Result, sizeof(*Result));
for (const auto &TestCase : TestCases) {
dadd_rn<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second);
cudaDeviceSynchronize();
checkResult("__dadd_rn", {TestCase.first.first, TestCase.first.second},
TestCase.second.first, *Result, TestCase.second.second);
}
}

__global__ void dadd_ru(float *const Result, float Input1, float Input2) {
*Result = __dadd_ru(Input1, Input2);
}

void testDadd_ruCases(
const vector<pair<pair<float, float>, di_pair>> &TestCases) {
float *Result;
cudaMallocManaged(&Result, sizeof(*Result));
for (const auto &TestCase : TestCases) {
dadd_ru<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second);
cudaDeviceSynchronize();
checkResult("__dadd_ru", {TestCase.first.first, TestCase.first.second},
TestCase.second.first, *Result, TestCase.second.second);
}
}

__global__ void dadd_rz(float *const Result, float Input1, float Input2) {
*Result = __dadd_rz(Input1, Input2);
}

void testDadd_rzCases(
const vector<pair<pair<float, float>, di_pair>> &TestCases) {
float *Result;
cudaMallocManaged(&Result, sizeof(*Result));
for (const auto &TestCase : TestCases) {
dadd_rz<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second);
cudaDeviceSynchronize();
checkResult("__dadd_rz", {TestCase.first.first, TestCase.first.second},
TestCase.second.first, *Result, TestCase.second.second);
}
}

__global__ void dmul_rd(float *const Result, float Input1, float Input2) {
*Result = __dmul_rd(Input1, Input2);
}

void testDmul_rdCases(
const vector<pair<pair<float, float>, di_pair>> &TestCases) {
float *Result;
cudaMallocManaged(&Result, sizeof(*Result));
for (const auto &TestCase : TestCases) {
dmul_rd<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second);
cudaDeviceSynchronize();
checkResult("__dmul_rd", {TestCase.first.first, TestCase.first.second},
TestCase.second.first, *Result, TestCase.second.second);
}
}

__global__ void dmul_rn(float *const Result, float Input1, float Input2) {
*Result = __dmul_rn(Input1, Input2);
}

void testDmul_rnCases(
const vector<pair<pair<float, float>, di_pair>> &TestCases) {
float *Result;
cudaMallocManaged(&Result, sizeof(*Result));
for (const auto &TestCase : TestCases) {
dmul_rn<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second);
cudaDeviceSynchronize();
checkResult("__dmul_rn", {TestCase.first.first, TestCase.first.second},
TestCase.second.first, *Result, TestCase.second.second);
}
}

__global__ void dmul_ru(float *const Result, float Input1, float Input2) {
*Result = __dmul_ru(Input1, Input2);
}

void testDmul_ruCases(
const vector<pair<pair<float, float>, di_pair>> &TestCases) {
float *Result;
cudaMallocManaged(&Result, sizeof(*Result));
for (const auto &TestCase : TestCases) {
dmul_ru<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second);
cudaDeviceSynchronize();
checkResult("__dmul_ru", {TestCase.first.first, TestCase.first.second},
TestCase.second.first, *Result, TestCase.second.second);
}
}

__global__ void dmul_rz(float *const Result, float Input1, float Input2) {
*Result = __dmul_rz(Input1, Input2);
}

void testDmul_rzCases(
const vector<pair<pair<float, float>, di_pair>> &TestCases) {
float *Result;
cudaMallocManaged(&Result, sizeof(*Result));
for (const auto &TestCase : TestCases) {
dmul_rz<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second);
cudaDeviceSynchronize();
checkResult("__dmul_rz", {TestCase.first.first, TestCase.first.second},
TestCase.second.first, *Result, TestCase.second.second);
}
}

__global__ void dsub_rd(float *const Result, float Input1, float Input2) {
*Result = __dsub_rd(Input1, Input2);
}

void testDsub_rdCases(
const vector<pair<pair<float, float>, di_pair>> &TestCases) {
float *Result;
cudaMallocManaged(&Result, sizeof(*Result));
for (const auto &TestCase : TestCases) {
dsub_rd<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second);
cudaDeviceSynchronize();
checkResult("__dsub_rd", {TestCase.first.first, TestCase.first.second},
TestCase.second.first, *Result, TestCase.second.second);
}
}

__global__ void dsub_rn(float *const Result, float Input1, float Input2) {
*Result = __dsub_rn(Input1, Input2);
}

void testDsub_rnCases(
const vector<pair<pair<float, float>, di_pair>> &TestCases) {
float *Result;
cudaMallocManaged(&Result, sizeof(*Result));
for (const auto &TestCase : TestCases) {
dsub_rn<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second);
cudaDeviceSynchronize();
checkResult("__dsub_rn", {TestCase.first.first, TestCase.first.second},
TestCase.second.first, *Result, TestCase.second.second);
}
}

__global__ void dsub_ru(float *const Result, float Input1, float Input2) {
*Result = __dsub_ru(Input1, Input2);
}

void testDsub_ruCases(
const vector<pair<pair<float, float>, di_pair>> &TestCases) {
float *Result;
cudaMallocManaged(&Result, sizeof(*Result));
for (const auto &TestCase : TestCases) {
dsub_ru<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second);
cudaDeviceSynchronize();
checkResult("__dsub_ru", {TestCase.first.first, TestCase.first.second},
TestCase.second.first, *Result, TestCase.second.second);
}
}

__global__ void dsub_rz(float *const Result, float Input1, float Input2) {
*Result = __dsub_rz(Input1, Input2);
}

void testDsub_rzCases(
const vector<pair<pair<float, float>, di_pair>> &TestCases) {
float *Result;
cudaMallocManaged(&Result, sizeof(*Result));
for (const auto &TestCase : TestCases) {
dsub_rz<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second);
cudaDeviceSynchronize();
checkResult("__dsub_rz", {TestCase.first.first, TestCase.first.second},
TestCase.second.first, *Result, TestCase.second.second);
}
}

int main() {
testNormCases({
{{-0.3, -0.34, -0.98}, {1.079814798935447, 15}},
Expand Down Expand Up @@ -306,6 +500,90 @@ int main() {
{{0.5, 456, 23, 1}, {0.002190191670280358, 18}},
{{23, 432, 23, 1}, {0.002308274913317669, 18}},
});
testDadd_rdCases({
{{-0.3, -0.4}, {-0.7000000476837158, 16}},
{{0.3, -0.4}, {-0.09999999403953552, 17}},
{{0.3, 0.4}, {0.7000000476837158, 16}},
{{0.3, 0.8}, {1.100000023841858, 15}},
{{3, 4}, {7, 15}},
});
testDadd_rnCases({
{{-0.3, -0.4}, {-0.7000000476837158, 16}},
{{0.3, -0.4}, {-0.09999999403953552, 17}},
{{0.3, 0.4}, {0.7000000476837158, 16}},
{{0.3, 0.8}, {1.100000023841858, 15}},
{{3, 4}, {7, 15}},
});
testDadd_ruCases({
{{-0.3, -0.4}, {-0.7000000476837158, 16}},
{{0.3, -0.4}, {-0.09999999403953552, 17}},
{{0.3, 0.4}, {0.7000000476837158, 16}},
{{0.3, 0.8}, {1.100000023841858, 15}},
{{3, 4}, {7, 15}},
});
testDadd_rzCases({
{{-0.3, -0.4}, {-0.7000000476837158, 16}},
{{0.3, -0.4}, {-0.09999999403953552, 17}},
{{0.3, 0.4}, {0.7000000476837158, 16}},
{{0.3, 0.8}, {1.100000023841858, 15}},
{{3, 4}, {7, 15}},
});
testDmul_rdCases({
{{-0.3, -0.4}, {0.12000000476837158, 17}},
{{0.3, -0.4}, {-0.12000000476837158, 17}},
{{0.3, 0.4}, {0.12000000476837158, 17}},
{{0.3, 0.8}, {0.2400000095367432, 16}},
{{3, 4}, {12, 15}},
});
testDmul_rnCases({
{{-0.3, -0.4}, {0.12000000476837158, 17}},
{{0.3, -0.4}, {-0.12000000476837158, 17}},
{{0.3, 0.4}, {0.12000000476837158, 17}},
{{0.3, 0.8}, {0.2400000095367432, 16}},
{{3, 4}, {12, 15}},
});
testDmul_ruCases({
{{-0.3, -0.4}, {0.12000000476837158, 17}},
{{0.3, -0.4}, {-0.12000000476837158, 17}},
{{0.3, 0.4}, {0.12000000476837158, 17}},
{{0.3, 0.8}, {0.2400000095367432, 16}},
{{3, 4}, {12, 15}},
});
testDmul_rzCases({
{{-0.3, -0.4}, {0.12000000476837158, 17}},
{{0.3, -0.4}, {-0.12000000476837158, 17}},
{{0.3, 0.4}, {0.12000000476837158, 17}},
{{0.3, 0.8}, {0.2400000095367432, 16}},
{{3, 4}, {12, 15}},
});
testDsub_rdCases({
{{-0.3, -0.4}, {0.09999999403953552, 17}},
{{0.3, -0.4}, {0.7000000476837158, 16}},
{{0.3, 0.4}, {-0.09999999403953552, 17}},
{{0.3, 0.8}, {-0.5, 16}},
{{3, 4}, {-1, 15}},
});
testDsub_rnCases({
{{-0.3, -0.4}, {0.09999999403953552, 17}},
{{0.3, -0.4}, {0.7000000476837158, 16}},
{{0.3, 0.4}, {-0.09999999403953552, 17}},
{{0.3, 0.8}, {-0.5, 16}},
{{3, 4}, {-1, 15}},
});
testDsub_ruCases({
{{-0.3, -0.4}, {0.09999999403953552, 17}},
{{0.3, -0.4}, {0.7000000476837158, 16}},
{{0.3, 0.4}, {-0.09999999403953552, 17}},
{{0.3, 0.8}, {-0.5, 16}},
{{3, 4}, {-1, 15}},
});
testDsub_rzCases({
{{-0.3, -0.4}, {0.09999999403953552, 17}},
{{0.3, -0.4}, {0.7000000476837158, 16}},
{{0.3, 0.4}, {-0.09999999403953552, 17}},
{{0.3, 0.8}, {-0.5, 16}},
{{3, 4}, {-1, 15}},
});
cout << "passed " << passed << "/" << passed + failed << " cases!" << endl;
if (failed) {
cout << "failed!" << endl;
Expand Down
Loading

0 comments on commit 5d70717

Please sign in to comment.