View | Details | Raw Unified | Return to bug 244850 | Differences between
and this patch

Collapse All | Expand All

(-)sc/source/core/opencl/formulagroupcl.cxx (-369 / +377 lines)
Lines 1026-1034 Link Here
1026
/// Handling a Double Vector that is used as a sliding window input
1026
/// Handling a Double Vector that is used as a sliding window input
1027
/// to either a sliding window average or sum-of-products
1027
/// to either a sliding window average or sum-of-products
1028
/// Generate a sequential loop for reductions
1028
/// Generate a sequential loop for reductions
1029
class OpAverage;
1030
class OpCount;
1031
1032
template<class Base>
1029
template<class Base>
1033
class DynamicKernelSlidingArgument : public Base
1030
class DynamicKernelSlidingArgument : public Base
1034
{
1031
{
Lines 1335-1520 Link Here
1335
    }
1332
    }
1336
1333
1337
    /// Emit the definition for the auxiliary reduction kernel
1334
    /// Emit the definition for the auxiliary reduction kernel
1338
    virtual void GenSlidingWindowFunction( std::stringstream& ss )
1335
    virtual void GenSlidingWindowFunction( std::stringstream& ss );
1339
    {
1340
        if (!dynamic_cast<OpAverage*>(mpCodeGen.get()))
1341
        {
1342
            std::string name = Base::GetName();
1343
            ss << "__kernel void " << name;
1344
            ss << "_reduction(__global double* A, "
1345
                "__global double *result,int arrayLength,int windowSize){\n";
1346
            ss << "    double tmp, current_result =" <<
1347
                mpCodeGen->GetBottom();
1348
            ss << ";\n";
1349
            ss << "    int writePos = get_group_id(1);\n";
1350
            ss << "    int lidx = get_local_id(0);\n";
1351
            ss << "    __local double shm_buf[256];\n";
1352
            if (mpDVR->IsStartFixed())
1353
                ss << "    int offset = 0;\n";
1354
            else // if (!mpDVR->IsStartFixed())
1355
                ss << "    int offset = get_group_id(1);\n";
1356
            if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
1357
                ss << "    int end = windowSize;\n";
1358
            else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
1359
                ss << "    int end = offset + windowSize;\n";
1360
            else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
1361
                ss << "    int end = windowSize + get_group_id(1);\n";
1362
            else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
1363
                ss << "    int end = windowSize;\n";
1364
            ss << "    end = min(end, arrayLength);\n";
1365
1336
1366
            ss << "    barrier(CLK_LOCAL_MEM_FENCE);\n";
1367
            ss << "    int loop = arrayLength/512 + 1;\n";
1368
            ss << "    for (int l=0; l<loop; l++){\n";
1369
            ss << "    tmp = " << mpCodeGen->GetBottom() << ";\n";
1370
            ss << "    int loopOffset = l*512;\n";
1371
            ss << "    if((loopOffset + lidx + offset + 256) < end) {\n";
1372
            ss << "        tmp = legalize(" << mpCodeGen->Gen2(
1373
                "A[loopOffset + lidx + offset]", "tmp") << ", tmp);\n";
1374
            ss << "        tmp = legalize(" << mpCodeGen->Gen2(
1375
                "A[loopOffset + lidx + offset + 256]", "tmp") << ", tmp);\n";
1376
            ss << "    } else if ((loopOffset + lidx + offset) < end)\n";
1377
            ss << "        tmp = legalize(" << mpCodeGen->Gen2(
1378
                "A[loopOffset + lidx + offset]", "tmp") << ", tmp);\n";
1379
            ss << "    shm_buf[lidx] = tmp;\n";
1380
            ss << "    barrier(CLK_LOCAL_MEM_FENCE);\n";
1381
            ss << "    for (int i = 128; i >0; i/=2) {\n";
1382
            ss << "        if (lidx < i)\n";
1383
            ss << "            shm_buf[lidx] = ";
1384
            // Special case count
1385
            if (dynamic_cast<OpCount*>(mpCodeGen.get()))
1386
                ss << "shm_buf[lidx] + shm_buf[lidx + i];\n";
1387
            else
1388
                ss << mpCodeGen->Gen2("shm_buf[lidx]", "shm_buf[lidx + i]") << ";\n";
1389
            ss << "        barrier(CLK_LOCAL_MEM_FENCE);\n";
1390
            ss << "    }\n";
1391
            ss << "        if (lidx == 0)\n";
1392
            ss << "            current_result =";
1393
            if (dynamic_cast<OpCount*>(mpCodeGen.get()))
1394
                ss << "current_result + shm_buf[0]";
1395
            else
1396
                ss << mpCodeGen->Gen2("current_result", "shm_buf[0]");
1397
            ss << ";\n";
1398
            ss << "        barrier(CLK_LOCAL_MEM_FENCE);\n";
1399
            ss << "    }\n";
1400
            ss << "    if (lidx == 0)\n";
1401
            ss << "        result[writePos] = current_result;\n";
1402
            ss << "}\n";
1403
        }
1404
        else
1405
        {
1406
            std::string name = Base::GetName();
1407
            /*sum reduction*/
1408
            ss << "__kernel void " << name << "_sum";
1409
            ss << "_reduction(__global double* A, "
1410
                "__global double *result,int arrayLength,int windowSize){\n";
1411
            ss << "    double tmp, current_result =" <<
1412
                mpCodeGen->GetBottom();
1413
            ss << ";\n";
1414
            ss << "    int writePos = get_group_id(1);\n";
1415
            ss << "    int lidx = get_local_id(0);\n";
1416
            ss << "    __local double shm_buf[256];\n";
1417
            if (mpDVR->IsStartFixed())
1418
                ss << "    int offset = 0;\n";
1419
            else // if (!mpDVR->IsStartFixed())
1420
                ss << "    int offset = get_group_id(1);\n";
1421
            if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
1422
                ss << "    int end = windowSize;\n";
1423
            else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
1424
                ss << "    int end = offset + windowSize;\n";
1425
            else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
1426
                ss << "    int end = windowSize + get_group_id(1);\n";
1427
            else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
1428
                ss << "    int end = windowSize;\n";
1429
            ss << "    end = min(end, arrayLength);\n";
1430
            ss << "    barrier(CLK_LOCAL_MEM_FENCE);\n";
1431
            ss << "    int loop = arrayLength/512 + 1;\n";
1432
            ss << "    for (int l=0; l<loop; l++){\n";
1433
            ss << "    tmp = " << mpCodeGen->GetBottom() << ";\n";
1434
            ss << "    int loopOffset = l*512;\n";
1435
            ss << "    if((loopOffset + lidx + offset + 256) < end) {\n";
1436
            ss << "        tmp = legalize(";
1437
            ss << "(A[loopOffset + lidx + offset]+ tmp)";
1438
            ss << ", tmp);\n";
1439
            ss << "        tmp = legalize((A[loopOffset + lidx + offset + 256]+ tmp)";
1440
            ss << ", tmp);\n";
1441
            ss << "    } else if ((loopOffset + lidx + offset) < end)\n";
1442
            ss << "        tmp = legalize((A[loopOffset + lidx + offset] + tmp)";
1443
            ss << ", tmp);\n";
1444
            ss << "    shm_buf[lidx] = tmp;\n";
1445
            ss << "    barrier(CLK_LOCAL_MEM_FENCE);\n";
1446
            ss << "    for (int i = 128; i >0; i/=2) {\n";
1447
            ss << "        if (lidx < i)\n";
1448
            ss << "            shm_buf[lidx] = ";
1449
            ss << "shm_buf[lidx] + shm_buf[lidx + i];\n";
1450
            ss << "        barrier(CLK_LOCAL_MEM_FENCE);\n";
1451
            ss << "    }\n";
1452
            ss << "        if (lidx == 0)\n";
1453
            ss << "            current_result =";
1454
            ss << "current_result + shm_buf[0]";
1455
            ss << ";\n";
1456
            ss << "        barrier(CLK_LOCAL_MEM_FENCE);\n";
1457
            ss << "    }\n";
1458
            ss << "    if (lidx == 0)\n";
1459
            ss << "        result[writePos] = current_result;\n";
1460
            ss << "}\n";
1461
            /*count reduction*/
1462
            ss << "__kernel void " << name << "_count";
1463
            ss << "_reduction(__global double* A, "
1464
                "__global double *result,int arrayLength,int windowSize){\n";
1465
            ss << "    double tmp, current_result =" <<
1466
                mpCodeGen->GetBottom();
1467
            ss << ";\n";
1468
            ss << "    int writePos = get_group_id(1);\n";
1469
            ss << "    int lidx = get_local_id(0);\n";
1470
            ss << "    __local double shm_buf[256];\n";
1471
            if (mpDVR->IsStartFixed())
1472
                ss << "    int offset = 0;\n";
1473
            else // if (!mpDVR->IsStartFixed())
1474
                ss << "    int offset = get_group_id(1);\n";
1475
            if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
1476
                ss << "    int end = windowSize;\n";
1477
            else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
1478
                ss << "    int end = offset + windowSize;\n";
1479
            else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
1480
                ss << "    int end = windowSize + get_group_id(1);\n";
1481
            else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
1482
                ss << "    int end = windowSize;\n";
1483
            ss << "    end = min(end, arrayLength);\n";
1484
            ss << "    barrier(CLK_LOCAL_MEM_FENCE);\n";
1485
            ss << "    int loop = arrayLength/512 + 1;\n";
1486
            ss << "    for (int l=0; l<loop; l++){\n";
1487
            ss << "    tmp = " << mpCodeGen->GetBottom() << ";\n";
1488
            ss << "    int loopOffset = l*512;\n";
1489
            ss << "    if((loopOffset + lidx + offset + 256) < end) {\n";
1490
            ss << "        tmp = legalize((isnan(A[loopOffset + lidx + offset])?tmp:tmp+1.0)";
1491
            ss << ", tmp);\n";
1492
            ss << "        tmp = legalize((isnan(A[loopOffset + lidx + offset+256])?tmp:tmp+1.0)";
1493
            ss << ", tmp);\n";
1494
            ss << "    } else if ((loopOffset + lidx + offset) < end)\n";
1495
            ss << "        tmp = legalize((isnan(A[loopOffset + lidx + offset])?tmp:tmp+1.0)";
1496
            ss << ", tmp);\n";
1497
            ss << "    shm_buf[lidx] = tmp;\n";
1498
            ss << "    barrier(CLK_LOCAL_MEM_FENCE);\n";
1499
            ss << "    for (int i = 128; i >0; i/=2) {\n";
1500
            ss << "        if (lidx < i)\n";
1501
            ss << "            shm_buf[lidx] = ";
1502
            ss << "shm_buf[lidx] + shm_buf[lidx + i];\n";
1503
            ss << "        barrier(CLK_LOCAL_MEM_FENCE);\n";
1504
            ss << "    }\n";
1505
            ss << "        if (lidx == 0)\n";
1506
            ss << "            current_result =";
1507
            ss << "current_result + shm_buf[0];";
1508
            ss << ";\n";
1509
            ss << "        barrier(CLK_LOCAL_MEM_FENCE);\n";
1510
            ss << "    }\n";
1511
            ss << "    if (lidx == 0)\n";
1512
            ss << "        result[writePos] = current_result;\n";
1513
            ss << "}\n";
1514
        }
1515
1516
    }
1517
1518
    virtual std::string GenSlidingWindowDeclRef( bool ) const
1337
    virtual std::string GenSlidingWindowDeclRef( bool ) const
1519
    {
1338
    {
1520
        std::stringstream ss;
1339
        std::stringstream ss;
Lines 1527-1721 Link Here
1527
1346
1528
    /// Controls how the elements in the DoubleVectorRef are traversed
1347
    /// Controls how the elements in the DoubleVectorRef are traversed
1529
    size_t GenReductionLoopHeader(
1348
    size_t GenReductionLoopHeader(
1530
        std::stringstream& ss, int nResultSize, bool& needBody )
1349
        std::stringstream& ss, int nResultSize, bool& needBody );
1531
    {
1532
        assert(mpDVR);
1533
        size_t nCurWindowSize = mpDVR->GetRefRowSize();
1534
        std::string temp = Base::GetName() + "[gid0]";
1535
        ss << "tmp = ";
1536
        // Special case count
1537
        if (dynamic_cast<OpAverage*>(mpCodeGen.get()))
1538
        {
1539
            ss << mpCodeGen->Gen2(temp, "tmp") << ";\n";
1540
            ss << "nCount = nCount-1;\n";
1541
            ss << "nCount = nCount +"; /*re-assign nCount from count reduction*/
1542
            ss << Base::GetName() << "[gid0+" << nResultSize << "]" << ";\n";
1543
        }
1544
        else if (dynamic_cast<OpCount*>(mpCodeGen.get()))
1545
            ss << temp << "+ tmp";
1546
        else
1547
            ss << mpCodeGen->Gen2(temp, "tmp");
1548
        ss << ";\n\t";
1549
        needBody = false;
1550
        return nCurWindowSize;
1551
    }
1552
1350
1553
    virtual size_t Marshal( cl_kernel k, int argno, int w, cl_program mpProgram )
1351
    virtual size_t Marshal( cl_kernel k, int argno, int w, cl_program mpProgram );
1554
    {
1555
        assert(Base::mpClmem == nullptr);
1556
1352
1557
        openclwrapper::KernelEnv kEnv;
1558
        openclwrapper::setKernelEnv(&kEnv);
1559
        cl_int err;
1560
        size_t nInput = mpDVR->GetArrayLength();
1561
        size_t nCurWindowSize = mpDVR->GetRefRowSize();
1562
        // create clmem buffer
1563
        if (mpDVR->GetArrays()[Base::mnIndex].mpNumericArray == nullptr)
1564
            throw Unhandled(__FILE__, __LINE__);
1565
        double* pHostBuffer = const_cast<double*>(
1566
            mpDVR->GetArrays()[Base::mnIndex].mpNumericArray);
1567
        size_t szHostBuffer = nInput * sizeof(double);
1568
        Base::mpClmem = clCreateBuffer(kEnv.mpkContext,
1569
            cl_mem_flags(CL_MEM_READ_ONLY) | CL_MEM_USE_HOST_PTR,
1570
            szHostBuffer,
1571
            pHostBuffer, &err);
1572
        SAL_INFO("sc.opencl", "Created buffer " << Base::mpClmem << " size " << nInput << "*" << sizeof(double) << "=" << szHostBuffer << " using host buffer " << pHostBuffer);
1573
1574
        mpClmem2 = clCreateBuffer(kEnv.mpkContext,
1575
            CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
1576
            sizeof(double) * w, nullptr, nullptr);
1577
        if (CL_SUCCESS != err)
1578
            throw OpenCLError("clCreateBuffer", err, __FILE__, __LINE__);
1579
        SAL_INFO("sc.opencl", "Created buffer " << mpClmem2 << " size " << sizeof(double) << "*" << w << "=" << (sizeof(double)*w));
1580
1581
        // reproduce the reduction function name
1582
        std::string kernelName;
1583
        if (!dynamic_cast<OpAverage*>(mpCodeGen.get()))
1584
            kernelName = Base::GetName() + "_reduction";
1585
        else
1586
            kernelName = Base::GetName() + "_sum_reduction";
1587
        cl_kernel redKernel = clCreateKernel(mpProgram, kernelName.c_str(), &err);
1588
        if (err != CL_SUCCESS)
1589
            throw OpenCLError("clCreateKernel", err, __FILE__, __LINE__);
1590
        SAL_INFO("sc.opencl", "Created kernel " << redKernel << " with name " << kernelName << " in program " << mpProgram);
1591
1592
        // set kernel arg of reduction kernel
1593
        // TODO(Wei Wei): use unique name for kernel
1594
        cl_mem buf = Base::GetCLBuffer();
1595
        SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 0 << ": cl_mem: " << buf);
1596
        err = clSetKernelArg(redKernel, 0, sizeof(cl_mem),
1597
            static_cast<void*>(&buf));
1598
        if (CL_SUCCESS != err)
1599
            throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
1600
1601
        SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 1 << ": cl_mem: " << mpClmem2);
1602
        err = clSetKernelArg(redKernel, 1, sizeof(cl_mem), &mpClmem2);
1603
        if (CL_SUCCESS != err)
1604
            throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
1605
1606
        SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 2 << ": cl_int: " << nInput);
1607
        err = clSetKernelArg(redKernel, 2, sizeof(cl_int), static_cast<void*>(&nInput));
1608
        if (CL_SUCCESS != err)
1609
            throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
1610
1611
        SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 3 << ": cl_int: " << nCurWindowSize);
1612
        err = clSetKernelArg(redKernel, 3, sizeof(cl_int), static_cast<void*>(&nCurWindowSize));
1613
        if (CL_SUCCESS != err)
1614
            throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
1615
1616
        // set work group size and execute
1617
        size_t global_work_size[] = { 256, static_cast<size_t>(w) };
1618
        size_t const local_work_size[] = { 256, 1 };
1619
        SAL_INFO("sc.opencl", "Enqueing kernel " << redKernel);
1620
        err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, nullptr,
1621
            global_work_size, local_work_size, 0, nullptr, nullptr);
1622
        if (CL_SUCCESS != err)
1623
            throw OpenCLError("clEnqueueNDRangeKernel", err, __FILE__, __LINE__);
1624
        err = clFinish(kEnv.mpkCmdQueue);
1625
        if (CL_SUCCESS != err)
1626
            throw OpenCLError("clFinish", err, __FILE__, __LINE__);
1627
        if (dynamic_cast<OpAverage*>(mpCodeGen.get()))
1628
        {
1629
            /*average need more reduction kernel for count computing*/
1630
            std::unique_ptr<double[]> pAllBuffer(new double[2 * w]);
1631
            double* resbuf = static_cast<double*>(clEnqueueMapBuffer(kEnv.mpkCmdQueue,
1632
                mpClmem2,
1633
                CL_TRUE, CL_MAP_READ, 0,
1634
                sizeof(double) * w, 0, nullptr, nullptr,
1635
                &err));
1636
            if (err != CL_SUCCESS)
1637
                throw OpenCLError("clEnqueueMapBuffer", err, __FILE__, __LINE__);
1638
1639
            for (int i = 0; i < w; i++)
1640
                pAllBuffer[i] = resbuf[i];
1641
            err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem2, resbuf, 0, nullptr, nullptr);
1642
            if (err != CL_SUCCESS)
1643
                throw OpenCLError("clEnqueueUnmapMemObject", err, __FILE__, __LINE__);
1644
1645
            kernelName = Base::GetName() + "_count_reduction";
1646
            redKernel = clCreateKernel(mpProgram, kernelName.c_str(), &err);
1647
            if (err != CL_SUCCESS)
1648
                throw OpenCLError("clCreateKernel", err, __FILE__, __LINE__);
1649
            SAL_INFO("sc.opencl", "Created kernel " << redKernel << " with name " << kernelName << " in program " << mpProgram);
1650
1651
            // set kernel arg of reduction kernel
1652
            buf = Base::GetCLBuffer();
1653
            SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 0 << ": cl_mem: " << buf);
1654
            err = clSetKernelArg(redKernel, 0, sizeof(cl_mem),
1655
                static_cast<void*>(&buf));
1656
            if (CL_SUCCESS != err)
1657
                throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
1658
1659
            SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 1 << ": cl_mem: " << mpClmem2);
1660
            err = clSetKernelArg(redKernel, 1, sizeof(cl_mem), &mpClmem2);
1661
            if (CL_SUCCESS != err)
1662
                throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
1663
1664
            SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 2 << ": cl_int: " << nInput);
1665
            err = clSetKernelArg(redKernel, 2, sizeof(cl_int), static_cast<void*>(&nInput));
1666
            if (CL_SUCCESS != err)
1667
                throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
1668
1669
            SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 3 << ": cl_int: " << nCurWindowSize);
1670
            err = clSetKernelArg(redKernel, 3, sizeof(cl_int), static_cast<void*>(&nCurWindowSize));
1671
            if (CL_SUCCESS != err)
1672
                throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
1673
1674
            // set work group size and execute
1675
            size_t global_work_size1[] = { 256, static_cast<size_t>(w) };
1676
            size_t const local_work_size1[] = { 256, 1 };
1677
            SAL_INFO("sc.opencl", "Enqueing kernel " << redKernel);
1678
            err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, nullptr,
1679
                global_work_size1, local_work_size1, 0, nullptr, nullptr);
1680
            if (CL_SUCCESS != err)
1681
                throw OpenCLError("clEnqueueNDRangeKernel", err, __FILE__, __LINE__);
1682
            err = clFinish(kEnv.mpkCmdQueue);
1683
            if (CL_SUCCESS != err)
1684
                throw OpenCLError("clFinish", err, __FILE__, __LINE__);
1685
            resbuf = static_cast<double*>(clEnqueueMapBuffer(kEnv.mpkCmdQueue,
1686
                mpClmem2,
1687
                CL_TRUE, CL_MAP_READ, 0,
1688
                sizeof(double) * w, 0, nullptr, nullptr,
1689
                &err));
1690
            if (err != CL_SUCCESS)
1691
                throw OpenCLError("clEnqueueMapBuffer", err, __FILE__, __LINE__);
1692
            for (int i = 0; i < w; i++)
1693
                pAllBuffer[i + w] = resbuf[i];
1694
            err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem2, resbuf, 0, nullptr, nullptr);
1695
            // FIXME: Is it intentional to not throw an OpenCLError even if the clEnqueueUnmapMemObject() fails?
1696
            if (CL_SUCCESS != err)
1697
                SAL_WARN("sc.opencl", "clEnqueueUnmapMemObject failed: " << openclwrapper::errorString(err));
1698
            if (mpClmem2)
1699
            {
1700
                err = clReleaseMemObject(mpClmem2);
1701
                SAL_WARN_IF(err != CL_SUCCESS, "sc.opencl", "clReleaseMemObject failed: " << openclwrapper::errorString(err));
1702
                mpClmem2 = nullptr;
1703
            }
1704
            mpClmem2 = clCreateBuffer(kEnv.mpkContext,
1705
                cl_mem_flags(CL_MEM_READ_WRITE) | CL_MEM_COPY_HOST_PTR,
1706
                w * sizeof(double) * 2, pAllBuffer.get(), &err);
1707
            if (CL_SUCCESS != err)
1708
                throw OpenCLError("clCreateBuffer", err, __FILE__, __LINE__);
1709
            SAL_INFO("sc.opencl", "Created buffer " << mpClmem2 << " size " << w << "*" << sizeof(double) << "=" << (w*sizeof(double)) << " copying host buffer " << pAllBuffer.get());
1710
        }
1711
        // set kernel arg
1712
        SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": cl_mem: " << mpClmem2);
1713
        err = clSetKernelArg(k, argno, sizeof(cl_mem), &mpClmem2);
1714
        if (CL_SUCCESS != err)
1715
            throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
1716
        return 1;
1717
    }
1718
1719
    ~ParallelReductionVectorRef()
1353
    ~ParallelReductionVectorRef()
1720
    {
1354
    {
1721
        if (mpClmem2)
1355
        if (mpClmem2)
Lines 2314-2319 Link Here
2314
    }
1948
    }
2315
    virtual std::string BinFuncName() const override { return "fsop"; }
1949
    virtual std::string BinFuncName() const override { return "fsop"; }
2316
};
1950
};
1951
1952
template<class Base>
1953
void ParallelReductionVectorRef<Base>::GenSlidingWindowFunction( std::stringstream& ss )
1954
{
1955
    if (!dynamic_cast<OpAverage*>(mpCodeGen.get()))
1956
    {
1957
        std::string name = Base::GetName();
1958
        ss << "__kernel void " << name;
1959
        ss << "_reduction(__global double* A, "
1960
            "__global double *result,int arrayLength,int windowSize){\n";
1961
        ss << "    double tmp, current_result =" <<
1962
            mpCodeGen->GetBottom();
1963
        ss << ";\n";
1964
        ss << "    int writePos = get_group_id(1);\n";
1965
        ss << "    int lidx = get_local_id(0);\n";
1966
        ss << "    __local double shm_buf[256];\n";
1967
        if (mpDVR->IsStartFixed())
1968
            ss << "    int offset = 0;\n";
1969
        else // if (!mpDVR->IsStartFixed())
1970
            ss << "    int offset = get_group_id(1);\n";
1971
        if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
1972
            ss << "    int end = windowSize;\n";
1973
        else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
1974
            ss << "    int end = offset + windowSize;\n";
1975
        else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
1976
            ss << "    int end = windowSize + get_group_id(1);\n";
1977
        else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
1978
            ss << "    int end = windowSize;\n";
1979
        ss << "    end = min(end, arrayLength);\n";
1980
1981
        ss << "    barrier(CLK_LOCAL_MEM_FENCE);\n";
1982
        ss << "    int loop = arrayLength/512 + 1;\n";
1983
        ss << "    for (int l=0; l<loop; l++){\n";
1984
        ss << "    tmp = " << mpCodeGen->GetBottom() << ";\n";
1985
        ss << "    int loopOffset = l*512;\n";
1986
        ss << "    if((loopOffset + lidx + offset + 256) < end) {\n";
1987
        ss << "        tmp = legalize(" << mpCodeGen->Gen2(
1988
            "A[loopOffset + lidx + offset]", "tmp") << ", tmp);\n";
1989
        ss << "        tmp = legalize(" << mpCodeGen->Gen2(
1990
            "A[loopOffset + lidx + offset + 256]", "tmp") << ", tmp);\n";
1991
        ss << "    } else if ((loopOffset + lidx + offset) < end)\n";
1992
        ss << "        tmp = legalize(" << mpCodeGen->Gen2(
1993
            "A[loopOffset + lidx + offset]", "tmp") << ", tmp);\n";
1994
        ss << "    shm_buf[lidx] = tmp;\n";
1995
        ss << "    barrier(CLK_LOCAL_MEM_FENCE);\n";
1996
        ss << "    for (int i = 128; i >0; i/=2) {\n";
1997
        ss << "        if (lidx < i)\n";
1998
        ss << "            shm_buf[lidx] = ";
1999
        // Special case count
2000
        if (dynamic_cast<OpCount*>(mpCodeGen.get()))
2001
            ss << "shm_buf[lidx] + shm_buf[lidx + i];\n";
2002
        else
2003
            ss << mpCodeGen->Gen2("shm_buf[lidx]", "shm_buf[lidx + i]") << ";\n";
2004
        ss << "        barrier(CLK_LOCAL_MEM_FENCE);\n";
2005
        ss << "    }\n";
2006
        ss << "        if (lidx == 0)\n";
2007
        ss << "            current_result =";
2008
        if (dynamic_cast<OpCount*>(mpCodeGen.get()))
2009
            ss << "current_result + shm_buf[0]";
2010
        else
2011
            ss << mpCodeGen->Gen2("current_result", "shm_buf[0]");
2012
        ss << ";\n";
2013
        ss << "        barrier(CLK_LOCAL_MEM_FENCE);\n";
2014
        ss << "    }\n";
2015
        ss << "    if (lidx == 0)\n";
2016
        ss << "        result[writePos] = current_result;\n";
2017
        ss << "}\n";
2018
    }
2019
    else
2020
    {
2021
        std::string name = Base::GetName();
2022
        /*sum reduction*/
2023
        ss << "__kernel void " << name << "_sum";
2024
        ss << "_reduction(__global double* A, "
2025
            "__global double *result,int arrayLength,int windowSize){\n";
2026
        ss << "    double tmp, current_result =" <<
2027
            mpCodeGen->GetBottom();
2028
        ss << ";\n";
2029
        ss << "    int writePos = get_group_id(1);\n";
2030
        ss << "    int lidx = get_local_id(0);\n";
2031
        ss << "    __local double shm_buf[256];\n";
2032
        if (mpDVR->IsStartFixed())
2033
            ss << "    int offset = 0;\n";
2034
        else // if (!mpDVR->IsStartFixed())
2035
            ss << "    int offset = get_group_id(1);\n";
2036
        if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
2037
            ss << "    int end = windowSize;\n";
2038
        else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
2039
            ss << "    int end = offset + windowSize;\n";
2040
        else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
2041
            ss << "    int end = windowSize + get_group_id(1);\n";
2042
        else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
2043
            ss << "    int end = windowSize;\n";
2044
        ss << "    end = min(end, arrayLength);\n";
2045
        ss << "    barrier(CLK_LOCAL_MEM_FENCE);\n";
2046
        ss << "    int loop = arrayLength/512 + 1;\n";
2047
        ss << "    for (int l=0; l<loop; l++){\n";
2048
        ss << "    tmp = " << mpCodeGen->GetBottom() << ";\n";
2049
        ss << "    int loopOffset = l*512;\n";
2050
        ss << "    if((loopOffset + lidx + offset + 256) < end) {\n";
2051
        ss << "        tmp = legalize(";
2052
        ss << "(A[loopOffset + lidx + offset]+ tmp)";
2053
        ss << ", tmp);\n";
2054
        ss << "        tmp = legalize((A[loopOffset + lidx + offset + 256]+ tmp)";
2055
        ss << ", tmp);\n";
2056
        ss << "    } else if ((loopOffset + lidx + offset) < end)\n";
2057
        ss << "        tmp = legalize((A[loopOffset + lidx + offset] + tmp)";
2058
        ss << ", tmp);\n";
2059
        ss << "    shm_buf[lidx] = tmp;\n";
2060
        ss << "    barrier(CLK_LOCAL_MEM_FENCE);\n";
2061
        ss << "    for (int i = 128; i >0; i/=2) {\n";
2062
        ss << "        if (lidx < i)\n";
2063
        ss << "            shm_buf[lidx] = ";
2064
        ss << "shm_buf[lidx] + shm_buf[lidx + i];\n";
2065
        ss << "        barrier(CLK_LOCAL_MEM_FENCE);\n";
2066
        ss << "    }\n";
2067
        ss << "        if (lidx == 0)\n";
2068
        ss << "            current_result =";
2069
        ss << "current_result + shm_buf[0]";
2070
        ss << ";\n";
2071
        ss << "        barrier(CLK_LOCAL_MEM_FENCE);\n";
2072
        ss << "    }\n";
2073
        ss << "    if (lidx == 0)\n";
2074
        ss << "        result[writePos] = current_result;\n";
2075
        ss << "}\n";
2076
        /*count reduction*/
2077
        ss << "__kernel void " << name << "_count";
2078
        ss << "_reduction(__global double* A, "
2079
            "__global double *result,int arrayLength,int windowSize){\n";
2080
        ss << "    double tmp, current_result =" <<
2081
            mpCodeGen->GetBottom();
2082
        ss << ";\n";
2083
        ss << "    int writePos = get_group_id(1);\n";
2084
        ss << "    int lidx = get_local_id(0);\n";
2085
        ss << "    __local double shm_buf[256];\n";
2086
        if (mpDVR->IsStartFixed())
2087
            ss << "    int offset = 0;\n";
2088
        else // if (!mpDVR->IsStartFixed())
2089
            ss << "    int offset = get_group_id(1);\n";
2090
        if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
2091
            ss << "    int end = windowSize;\n";
2092
        else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
2093
            ss << "    int end = offset + windowSize;\n";
2094
        else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed())
2095
            ss << "    int end = windowSize + get_group_id(1);\n";
2096
        else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed())
2097
            ss << "    int end = windowSize;\n";
2098
        ss << "    end = min(end, arrayLength);\n";
2099
        ss << "    barrier(CLK_LOCAL_MEM_FENCE);\n";
2100
        ss << "    int loop = arrayLength/512 + 1;\n";
2101
        ss << "    for (int l=0; l<loop; l++){\n";
2102
        ss << "    tmp = " << mpCodeGen->GetBottom() << ";\n";
2103
        ss << "    int loopOffset = l*512;\n";
2104
        ss << "    if((loopOffset + lidx + offset + 256) < end) {\n";
2105
        ss << "        tmp = legalize((isnan(A[loopOffset + lidx + offset])?tmp:tmp+1.0)";
2106
        ss << ", tmp);\n";
2107
        ss << "        tmp = legalize((isnan(A[loopOffset + lidx + offset+256])?tmp:tmp+1.0)";
2108
        ss << ", tmp);\n";
2109
        ss << "    } else if ((loopOffset + lidx + offset) < end)\n";
2110
        ss << "        tmp = legalize((isnan(A[loopOffset + lidx + offset])?tmp:tmp+1.0)";
2111
        ss << ", tmp);\n";
2112
        ss << "    shm_buf[lidx] = tmp;\n";
2113
        ss << "    barrier(CLK_LOCAL_MEM_FENCE);\n";
2114
        ss << "    for (int i = 128; i >0; i/=2) {\n";
2115
        ss << "        if (lidx < i)\n";
2116
        ss << "            shm_buf[lidx] = ";
2117
        ss << "shm_buf[lidx] + shm_buf[lidx + i];\n";
2118
        ss << "        barrier(CLK_LOCAL_MEM_FENCE);\n";
2119
        ss << "    }\n";
2120
        ss << "        if (lidx == 0)\n";
2121
        ss << "            current_result =";
2122
        ss << "current_result + shm_buf[0];";
2123
        ss << ";\n";
2124
        ss << "        barrier(CLK_LOCAL_MEM_FENCE);\n";
2125
        ss << "    }\n";
2126
        ss << "    if (lidx == 0)\n";
2127
        ss << "        result[writePos] = current_result;\n";
2128
        ss << "}\n";
2129
    }
2130
2131
}
2132
2133
template<class Base>
2134
size_t ParallelReductionVectorRef<Base>::GenReductionLoopHeader(
2135
    std::stringstream& ss, int nResultSize, bool& needBody )
2136
{
2137
    assert(mpDVR);
2138
    size_t nCurWindowSize = mpDVR->GetRefRowSize();
2139
    std::string temp = Base::GetName() + "[gid0]";
2140
    ss << "tmp = ";
2141
    // Special case count
2142
    if (dynamic_cast<OpAverage*>(mpCodeGen.get()))
2143
    {
2144
        ss << mpCodeGen->Gen2(temp, "tmp") << ";\n";
2145
        ss << "nCount = nCount-1;\n";
2146
        ss << "nCount = nCount +"; /*re-assign nCount from count reduction*/
2147
        ss << Base::GetName() << "[gid0+" << nResultSize << "]" << ";\n";
2148
    }
2149
    else if (dynamic_cast<OpCount*>(mpCodeGen.get()))
2150
        ss << temp << "+ tmp";
2151
    else
2152
        ss << mpCodeGen->Gen2(temp, "tmp");
2153
    ss << ";\n\t";
2154
    needBody = false;
2155
    return nCurWindowSize;
2156
}
2157
2158
template<class Base>
2159
size_t ParallelReductionVectorRef<Base>::Marshal( cl_kernel k, int argno, int w, cl_program mpProgram )
2160
{
2161
    assert(Base::mpClmem == nullptr);
2162
2163
    openclwrapper::KernelEnv kEnv;
2164
    openclwrapper::setKernelEnv(&kEnv);
2165
    cl_int err;
2166
    size_t nInput = mpDVR->GetArrayLength();
2167
    size_t nCurWindowSize = mpDVR->GetRefRowSize();
2168
    // create clmem buffer
2169
    if (mpDVR->GetArrays()[Base::mnIndex].mpNumericArray == nullptr)
2170
        throw Unhandled(__FILE__, __LINE__);
2171
    double* pHostBuffer = const_cast<double*>(
2172
        mpDVR->GetArrays()[Base::mnIndex].mpNumericArray);
2173
    size_t szHostBuffer = nInput * sizeof(double);
2174
    Base::mpClmem = clCreateBuffer(kEnv.mpkContext,
2175
        cl_mem_flags(CL_MEM_READ_ONLY) | CL_MEM_USE_HOST_PTR,
2176
        szHostBuffer,
2177
        pHostBuffer, &err);
2178
    SAL_INFO("sc.opencl", "Created buffer " << Base::mpClmem << " size " << nInput << "*" << sizeof(double) << "=" << szHostBuffer << " using host buffer " << pHostBuffer);
2179
2180
    mpClmem2 = clCreateBuffer(kEnv.mpkContext,
2181
        CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
2182
        sizeof(double) * w, nullptr, nullptr);
2183
    if (CL_SUCCESS != err)
2184
        throw OpenCLError("clCreateBuffer", err, __FILE__, __LINE__);
2185
    SAL_INFO("sc.opencl", "Created buffer " << mpClmem2 << " size " << sizeof(double) << "*" << w << "=" << (sizeof(double)*w));
2186
2187
    // reproduce the reduction function name
2188
    std::string kernelName;
2189
    if (!dynamic_cast<OpAverage*>(mpCodeGen.get()))
2190
        kernelName = Base::GetName() + "_reduction";
2191
    else
2192
        kernelName = Base::GetName() + "_sum_reduction";
2193
    cl_kernel redKernel = clCreateKernel(mpProgram, kernelName.c_str(), &err);
2194
    if (err != CL_SUCCESS)
2195
        throw OpenCLError("clCreateKernel", err, __FILE__, __LINE__);
2196
    SAL_INFO("sc.opencl", "Created kernel " << redKernel << " with name " << kernelName << " in program " << mpProgram);
2197
2198
    // set kernel arg of reduction kernel
2199
    // TODO(Wei Wei): use unique name for kernel
2200
    cl_mem buf = Base::GetCLBuffer();
2201
    SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 0 << ": cl_mem: " << buf);
2202
    err = clSetKernelArg(redKernel, 0, sizeof(cl_mem),
2203
        static_cast<void*>(&buf));
2204
    if (CL_SUCCESS != err)
2205
        throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
2206
2207
    SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 1 << ": cl_mem: " << mpClmem2);
2208
    err = clSetKernelArg(redKernel, 1, sizeof(cl_mem), &mpClmem2);
2209
    if (CL_SUCCESS != err)
2210
        throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
2211
2212
    SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 2 << ": cl_int: " << nInput);
2213
    err = clSetKernelArg(redKernel, 2, sizeof(cl_int), static_cast<void*>(&nInput));
2214
    if (CL_SUCCESS != err)
2215
        throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
2216
2217
    SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 3 << ": cl_int: " << nCurWindowSize);
2218
    err = clSetKernelArg(redKernel, 3, sizeof(cl_int), static_cast<void*>(&nCurWindowSize));
2219
    if (CL_SUCCESS != err)
2220
        throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
2221
2222
    // set work group size and execute
2223
    size_t global_work_size[] = { 256, static_cast<size_t>(w) };
2224
    size_t const local_work_size[] = { 256, 1 };
2225
    SAL_INFO("sc.opencl", "Enqueuing kernel " << redKernel);
2226
    err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, nullptr,
2227
        global_work_size, local_work_size, 0, nullptr, nullptr);
2228
    if (CL_SUCCESS != err)
2229
        throw OpenCLError("clEnqueueNDRangeKernel", err, __FILE__, __LINE__);
2230
    err = clFinish(kEnv.mpkCmdQueue);
2231
    if (CL_SUCCESS != err)
2232
        throw OpenCLError("clFinish", err, __FILE__, __LINE__);
2233
    if (dynamic_cast<OpAverage*>(mpCodeGen.get()))
2234
    {
2235
        /*average need more reduction kernel for count computing*/
2236
        std::unique_ptr<double[]> pAllBuffer(new double[2 * w]);
2237
        double* resbuf = static_cast<double*>(clEnqueueMapBuffer(kEnv.mpkCmdQueue,
2238
            mpClmem2,
2239
            CL_TRUE, CL_MAP_READ, 0,
2240
            sizeof(double) * w, 0, nullptr, nullptr,
2241
            &err));
2242
        if (err != CL_SUCCESS)
2243
            throw OpenCLError("clEnqueueMapBuffer", err, __FILE__, __LINE__);
2244
2245
        for (int i = 0; i < w; i++)
2246
            pAllBuffer[i] = resbuf[i];
2247
        err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem2, resbuf, 0, nullptr, nullptr);
2248
        if (err != CL_SUCCESS)
2249
            throw OpenCLError("clEnqueueUnmapMemObject", err, __FILE__, __LINE__);
2250
2251
        kernelName = Base::GetName() + "_count_reduction";
2252
        redKernel = clCreateKernel(mpProgram, kernelName.c_str(), &err);
2253
        if (err != CL_SUCCESS)
2254
            throw OpenCLError("clCreateKernel", err, __FILE__, __LINE__);
2255
        SAL_INFO("sc.opencl", "Created kernel " << redKernel << " with name " << kernelName << " in program " << mpProgram);
2256
2257
        // set kernel arg of reduction kernel
2258
        buf = Base::GetCLBuffer();
2259
        SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 0 << ": cl_mem: " << buf);
2260
        err = clSetKernelArg(redKernel, 0, sizeof(cl_mem),
2261
            static_cast<void*>(&buf));
2262
        if (CL_SUCCESS != err)
2263
            throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
2264
2265
        SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 1 << ": cl_mem: " << mpClmem2);
2266
        err = clSetKernelArg(redKernel, 1, sizeof(cl_mem), &mpClmem2);
2267
        if (CL_SUCCESS != err)
2268
            throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
2269
2270
        SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 2 << ": cl_int: " << nInput);
2271
        err = clSetKernelArg(redKernel, 2, sizeof(cl_int), static_cast<void*>(&nInput));
2272
        if (CL_SUCCESS != err)
2273
            throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
2274
2275
        SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 3 << ": cl_int: " << nCurWindowSize);
2276
        err = clSetKernelArg(redKernel, 3, sizeof(cl_int), static_cast<void*>(&nCurWindowSize));
2277
        if (CL_SUCCESS != err)
2278
            throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
2279
2280
        // set work group size and execute
2281
        size_t global_work_size1[] = { 256, static_cast<size_t>(w) };
2282
        size_t const local_work_size1[] = { 256, 1 };
2283
        SAL_INFO("sc.opencl", "Enqueuing kernel " << redKernel);
2284
        err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, nullptr,
2285
            global_work_size1, local_work_size1, 0, nullptr, nullptr);
2286
        if (CL_SUCCESS != err)
2287
            throw OpenCLError("clEnqueueNDRangeKernel", err, __FILE__, __LINE__);
2288
        err = clFinish(kEnv.mpkCmdQueue);
2289
        if (CL_SUCCESS != err)
2290
            throw OpenCLError("clFinish", err, __FILE__, __LINE__);
2291
        resbuf = static_cast<double*>(clEnqueueMapBuffer(kEnv.mpkCmdQueue,
2292
            mpClmem2,
2293
            CL_TRUE, CL_MAP_READ, 0,
2294
            sizeof(double) * w, 0, nullptr, nullptr,
2295
            &err));
2296
        if (err != CL_SUCCESS)
2297
            throw OpenCLError("clEnqueueMapBuffer", err, __FILE__, __LINE__);
2298
        for (int i = 0; i < w; i++)
2299
            pAllBuffer[i + w] = resbuf[i];
2300
        err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem2, resbuf, 0, nullptr, nullptr);
2301
        // FIXME: Is it intentional to not throw an OpenCLError even if the clEnqueueUnmapMemObject() fails?
2302
        if (CL_SUCCESS != err)
2303
            SAL_WARN("sc.opencl", "clEnqueueUnmapMemObject failed: " << openclwrapper::errorString(err));
2304
        if (mpClmem2)
2305
        {
2306
            err = clReleaseMemObject(mpClmem2);
2307
            SAL_WARN_IF(err != CL_SUCCESS, "sc.opencl", "clReleaseMemObject failed: " << openclwrapper::errorString(err));
2308
            mpClmem2 = nullptr;
2309
        }
2310
        mpClmem2 = clCreateBuffer(kEnv.mpkContext,
2311
            cl_mem_flags(CL_MEM_READ_WRITE) | CL_MEM_COPY_HOST_PTR,
2312
            w * sizeof(double) * 2, pAllBuffer.get(), &err);
2313
        if (CL_SUCCESS != err)
2314
            throw OpenCLError("clCreateBuffer", err, __FILE__, __LINE__);
2315
        SAL_INFO("sc.opencl", "Created buffer " << mpClmem2 << " size " << w << "*" << sizeof(double) << "=" << (w*sizeof(double)) << " copying host buffer " << pAllBuffer.get());
2316
    }
2317
    // set kernel arg
2318
    SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": cl_mem: " << mpClmem2);
2319
    err = clSetKernelArg(k, argno, sizeof(cl_mem), &mpClmem2);
2320
    if (CL_SUCCESS != err)
2321
        throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
2322
    return 1;
2323
}
2324
2317
namespace {
2325
namespace {
2318
struct SumIfsArgs
2326
struct SumIfsArgs
2319
{
2327
{

Return to bug 244850