|
1 |
| -// ====------------ math-emu-double.cu---------- *- CUDA -* -------------===//// |
| 1 | +// ===-------------- math-emu-float.cu---------- *- CUDA -* ---------------===// |
2 | 2 | //
|
3 | 3 | // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
|
4 | 4 | // See https://llvm.org/LICENSE.txt for license information.
|
@@ -39,13 +39,28 @@ void checkResult(const string &FuncName, const vector<T> &Inputs,
|
39 | 39 | for (size_t i = 1; i < Inputs.size(); ++i) {
|
40 | 40 | cout << ", " << Inputs[i];
|
41 | 41 | }
|
42 |
| - cout << ") = " << fixed << setprecision(precision) << DeviceResult |
43 |
| - << " (expect " << Expect - pow(10, -precision) << " ~ " |
| 42 | + cout << ") = " << fixed << setprecision(precision < 0 ? 0 : precision) |
| 43 | + << DeviceResult << " (expect " << Expect - pow(10, -precision) << " ~ " |
44 | 44 | << Expect + pow(10, -precision) << ")";
|
45 | 45 | cout.unsetf(ios::fixed);
|
46 | 46 | check(abs(DeviceResult - Expect) < pow(10, -precision));
|
47 | 47 | }
|
48 | 48 |
|
| 49 | +__global__ void expf(float *const Result, float Input1) { |
| 50 | + *Result = expf(Input1); |
| 51 | +} |
| 52 | + |
| 53 | +void testExpfCases(const vector<pair<float, fi_pair>> &TestCases) { |
| 54 | + float *Result; |
| 55 | + cudaMallocManaged(&Result, sizeof(*Result)); |
| 56 | + for (const auto &TestCase : TestCases) { |
| 57 | + expf<<<1, 1>>>(Result, TestCase.first); |
| 58 | + cudaDeviceSynchronize(); |
| 59 | + checkResult("expf", {TestCase.first}, TestCase.second.first, *Result, |
| 60 | + TestCase.second.second); |
| 61 | + } |
| 62 | +} |
| 63 | + |
49 | 64 | __global__ void _norm3df(float *const DeviceResult, float Input1, float Input2,
|
50 | 65 | float Input3) {
|
51 | 66 | *DeviceResult = norm3df(Input1, Input2, Input3);
|
@@ -256,7 +271,28 @@ void testRnormfCases(const vector<pair<f_vector, fi_pair>> &TestCases) {
|
256 | 271 | }
|
257 | 272 | }
|
258 | 273 |
|
| 274 | +__global__ void _expf(float *const Result, float Input1) { |
| 275 | + *Result = __expf(Input1); |
| 276 | +} |
| 277 | + |
| 278 | +void test_ExpfCases(const vector<pair<float, fi_pair>> &TestCases) { |
| 279 | + float *Result; |
| 280 | + cudaMallocManaged(&Result, sizeof(*Result)); |
| 281 | + for (const auto &TestCase : TestCases) { |
| 282 | + _expf<<<1, 1>>>(Result, TestCase.first); |
| 283 | + cudaDeviceSynchronize(); |
| 284 | + checkResult("__expf", {TestCase.first}, TestCase.second.first, *Result, |
| 285 | + TestCase.second.second); |
| 286 | + } |
| 287 | +} |
| 288 | + |
259 | 289 | int main() {
|
| 290 | + testExpfCases({ |
| 291 | + {-0.3, {0.7408, 4}}, |
| 292 | + {0.34, {1.405, 3}}, |
| 293 | + {23, {9745000000, -6}}, |
| 294 | + {-12, {0.000006144, 9}}, |
| 295 | + }); |
260 | 296 | testNorm3dfCases({
|
261 | 297 | {{-0.3, -0.34, -0.98}, {1.079814791679382, 15}},
|
262 | 298 | {{0.3, 0.34, 0.98}, {1.079814791679382, 15}},
|
@@ -306,6 +342,12 @@ int main() {
|
306 | 342 | {{0.5}, {2, 3}},
|
307 | 343 | {{23, 432, 23, 456, 23}, {0.0015888, 7}},
|
308 | 344 | });
|
| 345 | + test_ExpfCases({ |
| 346 | + {-0.3, {0.7408, 4}}, |
| 347 | + {0.34, {1.405, 3}}, |
| 348 | + {23, {9745000000, -6}}, |
| 349 | + {-12, {0.000006144, 9}}, |
| 350 | + }); |
309 | 351 | cout << "passed " << passed << "/" << passed + failed << " cases!" << endl;
|
310 | 352 | if (failed) {
|
311 | 353 | cout << "failed!" << endl;
|
|
0 commit comments