0
\$\begingroup\$

I implemented virtual array using OpenCL and graphics cards as backing store, with caching some of data on RAM using a LRU algorithm and accessing data through some simple setters/getters. Accession of an array element goes through a page lock (using array of mutexes for all independent pages), then LRU cache, then pcie data transmission in case of eviction and returning with data (or editing page data if its a write).

#include "GraphicsCardSupplyDepot.h"
#include "VirtualMultiArray.h"
#include "PcieBandwidthBenchmarker.h"
#include "CpuBenchmarker.h"

// testing
#include <random>
#include <iostream>
#include "omp.h"

constexpr bool TEST_BANDWIDTH=true;
constexpr bool TEST_LATENCY=false;
constexpr bool testType = TEST_BANDWIDTH;

// a test object for virtual array
// testing bandwidth: 512kB size
// testing latency: 8byte size
class Object
{
public:
    Object():id(-1){}
    Object(int p):id(p){}
    const int getId() const {return id;}
private:
    char data[testType?(1024*512 - 4):(4)];
    int id;
};

int main()
{
    // number of elements per cache page
    const long long pageSize = 1;

    // number of elements of array
    const long long n = pageSize*(testType?1000:100000);

    // number of benchmark runs
    const int numTestsPerThread = 25;

    // virtual array of objects in video memory
    VirtualMultiArray<Object> test(n,GraphicsCardSupplyDepot().requestGpus(),pageSize,3,PcieBandwidthBenchmarker().bestBandwidth(10));

    // heating cpu to get precise benchmark results
    #pragma omp parallel for
    for(long long j=0;j<n;j++)
    {
            test.set(j,Object(j));
    }

    // test for single thread, 2 threads, .. 64 threads
    for(int i=1;i<=64;i++)
    {
        // benchmark set method
        {
            CpuBenchmarker bench(i*numTestsPerThread*sizeof(Object),std::string("scalar set, ")+std::to_string(i)+std::string("threads"),i*numTestsPerThread);
            #pragma omp parallel for num_threads(i)
            for(long long j=0;j<i;j++)
            {
                // random-access to data
                std::random_device rd;
                std::mt19937 rng(rd());
                std::uniform_real_distribution<float> rnd(0,n-1);

                for(int k=0;k<numTestsPerThread;k++)
                {
                    int rndv = rnd(rng);
                    test.set(rndv,Object(rndv));
                }
            }
        }

        // benchmark get method
        {
            CpuBenchmarker bench(i*numTestsPerThread*sizeof(Object),std::string("scalar get, ")+std::to_string(i)+std::string("threads"),i*numTestsPerThread);
            #pragma omp parallel for num_threads(i)
            for(long long j=0;j<i;j++)
            {
                // random-access to data
                std::random_device rd;
                std::mt19937 rng(rd());
                std::uniform_real_distribution<float> rnd(0,n-1);

                for(int k=0;k<numTestsPerThread;k++)
                {
                    int rndv = rnd(rng);
                    const auto obj = test.get(rndv);
                    if(obj.getId()!=rndv)
                    {
                        throw std::invalid_argument("Error: set/get");
                    }
                }
            }
        }

        std::cout<<"==================================================================="<<std::endl;
    }
    return 0;
}

When I test it on Ubuntu 18.04, it uses up to 3/4 of the combined pcie bandwidth of system but on Windows 10, it uses mostly 1/4 (which is only 2GB/s) of peak.

My system has this specs:

  • fx8150
  • 4 GB single channel ddr3 RAM (10.6 GB/s hardware peak)
  • 3 low end graphics cards (gt1030, 2xk420) that support OpenCL (8GB/s hardware peak)
    • TCC driver mode enabled for 2 of cards, but first card cant do TCC mode
  • g++-10 for Ubuntu C++17
  • MSVC 2019 for Windows C++17

I tried to make it as platform-independent as possible but at one part I had to branch some code for aligned-allocations. What am I doing wrong with Windows-side? (besides the aligned allocations)

For example, in Ubuntu, I get this output:

scalar set, 1threads: 12182103 nanoseconds     (bandwidth = 1075.94 MB/s)      (throughput = 487284.12 nanoseconds per iteration) 
scalar get, 1threads: 10874521 nanoseconds     (bandwidth = 1205.31 MB/s)      (throughput = 434980.84 nanoseconds per iteration) 
===================================================================
scalar set, 2threads: 15642385 nanoseconds     (bandwidth = 1675.86 MB/s)      (throughput = 312847.70 nanoseconds per iteration) 
scalar get, 2threads: 16902450 nanoseconds     (bandwidth = 1550.92 MB/s)      (throughput = 338049.00 nanoseconds per iteration) 
===================================================================
scalar set, 3threads: 15004388 nanoseconds     (bandwidth = 2620.67 MB/s)      (throughput = 200058.51 nanoseconds per iteration) 
scalar get, 3threads: 16687201 nanoseconds     (bandwidth = 2356.39 MB/s)      (throughput = 222496.01 nanoseconds per iteration) 
...
...
===================================================================
scalar set, 63threads: 212283324 nanoseconds     (bandwidth = 3889.87 MB/s)      (throughput = 134783.06 nanoseconds per iteration) 
scalar get, 63threads: 136367146 nanoseconds     (bandwidth = 6055.37 MB/s)      (throughput = 86582.31 nanoseconds per iteration) 
===================================================================
scalar set, 64threads: 229655008 nanoseconds     (bandwidth = 3652.70 MB/s)      (throughput = 143534.38 nanoseconds per iteration) 
scalar get, 64threads: 149184573 nanoseconds     (bandwidth = 5622.97 MB/s)      (throughput = 93240.36 nanoseconds per iteration) 
===================================================================

but in Windows it tops out at 2GB/s with much less number of threads. Tested it without LRU caching (simply direct-mapping of pages to vram) too but same performance difference happened.

Edit:

According to performance profiler of Visual Studio, the random-number generator is not bottleneck:

enter image description here

When I follow hot path, it ends at an OpenCL command that queries an event:

 clGetEventInfo(evt, evtInf,sizeof(cl_int), &evtStatus0, nullptr)

in the PageCache header, line 248.

enter image description here

This is out of project, maybe related to my system being old? Or maybe its about some wrong api usage elsewhere, such as some wrong flags in opencl buffer construction?

Edit 2:

For Windows, I removed clGetEventInfo() and used blocking-version of read/write opencl commands:

clEnqueueReadBuffer(q->getQueue(), gpu->getMem(), CL_TRUE

this gave +25% performance (2.5GB/s instead of 2.0GB/s) for the benchmark on Windows but there is another bottleneck now:

enter image description here

kernel-related overhead is nearly as high as pcie-i/o overhead! When I follow hot path, again, the i/o part is the most bottlenecking part:

enter image description here

Does this mean, these OpenCL buffer read/write operations also go through kernel of Windows? I guess they have busy-wait inside with some mutex lock which is slow? When I click "include external code" it adds "nvopencl64.dll" to list of hot path with "kernel" tag on its line. Does this mean, nvidia driver uses some spin-wait in OpenCL synchronization which reduces threading i/o scalability because of Windows-related lock-contention? In Ubuntu, it is 2x fast with same nvidia gpus and extra 25% faster with explicit idle-wait loop using events. So, (I guess) its more related to Windows-based issue rather than Nvidia. But not sure.

When I don't use all 3 graphics cards at the same time, I get these results:

  • GT1030: 1GB/s
  • K420 at 8x slot: 2GB/s
  • K420 at 4x slot: 1GB/s

but when I use all 3 with many threads, I get 2.5GB/s total. Thats not same with Ubuntu which reaches 6GB/s total. If there was a way to tell OpenCL to use idle-wait loop instead of busy-wait loop, then threads would actually overlap i/o efficiently. But in Windows they not only stop scaling at 8 threads, but also decrease after more threads. In Ubuntu, it continuously increases until 64 threads. Maybe page-locking contention inside OpenCL driver causes this scalability issue?

Edit-3:

When I test a simple mutex example:

int main()
{

    std::mutex m;
    for (int i = 0; i < 100; i++)
    {
        CpuBenchmarker bench;
        for (int j = 0; j < 1000; j++)
        {
            std::unique_lock<std::mutex> l(m);
        }
    }
    return 0;
}

profiler says 97% of overhead is from kernel space and benchmarker outputs ~70 nanoseconds per lock duration. CPU is FX8150 at 2.1 GHz.

Lastly, does the code structure look ok?

\$\endgroup\$
2
  • \$\begingroup\$ I am not familiar with omp so not sure how it parallelizes it loops. But thestd::random_device is very expensive to create. Could you not do this outside the loop (in an array)? Then your code in the loop will be the only thing you test (and you can access the random device via the array and each loop iteration can look up its own random device. \$\endgroup\$ Mar 11 at 20:12
  • \$\begingroup\$ If I measure only random number generation, it is much less overhead than 25 iterations of copying 512kB objects. Every omp thread creates only 1 random device. Then 25 iterations are run. \$\endgroup\$ Mar 11 at 22:14

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service, privacy policy and cookie policy

Browse other questions tagged or ask your own question.