I'm using moderngpu CUDA library (https://github.com/moderngpu/moderngpu)
There's a strange speed up of interval_gather function in moderngpu if there are some bytes allocated by cudaMalloc before calling it.
I call a method mgpu::interval_gather(m0.data(), 1, m1.data(), 1, m2.data(), m3.data(), context); Each of m0, m1, m2, m3 is a mgpu::mem_t array of size of 1 000 000 filled with 0s.
I do it 10 000 times and it takes about 3 seconds.
However, if I allocate some dummy memory before it:
cudaMalloc((void **)(&tmpPtr), sizeof(int));
It speeds up around 10x, taking 0.3 seconds. Why would that be?
I tried allocating the mem_t memory before, or after allocating the dummy byte, same result. I experimented with different size of allocated dummy memory, calling function multiple times, or disabling -O2 compiling flag - nothing seems to change.
However, when I tried to reduce the size of arrays I pass to interval_gather, the effect stops. It then takes 0.3 seconds regardless of whether there was a dummy alloc or not.
I ran all of this on GTX 980.
Here's a full code snippet. If magic variable is set to true, it will speed up.
#include <iostream>
#include <chrono>
#include <string>
#include <moderngpu/kernel_intervalmove.hxx>
void print_ms(std::clock_t start, std::clock_t end, std::string desc)
{
double ms = 1000.0 * (end - start) / CLOCKS_PER_SEC;
std::cout << desc << ": " << ms << " ms." << std::endl;
}
void fun(bool magic, mgpu::context_t &context)
{
int n = 10000000;
mgpu::mem_t<int> m0(n, context);
mgpu::mem_t<int> m1(n, context);
mgpu::mem_t<int> m2(n, context);
mgpu::mem_t<int> m3(n, context);
void *tmpPtr;
cudaMalloc((void **)(&tmpPtr), sizeof(int));
if (!magic)
cudaFree(tmpPtr);
for (int aa = 0; aa < 10000; aa++)
mgpu::interval_gather(m0.data(), 1,
m1.data(), 1,
m2.data(),
m3.data(),
context);
if (magic)
cudaFree(tmpPtr);
}
int main(int argc, char *argv[])
{
mgpu::standard_context_t context(false);
std::clock_t c0 = std::clock();
fun(false, context);
context.synchronize();
std::clock_t c1 = std::clock();
print_ms(c0, c1, "1st");
fun(false, context);
context.synchronize();
std::clock_t c2 = std::clock();
print_ms(c1, c2, "2nd");
fun(true, context);
context.synchronize();
std::clock_t c3 = std::clock();
print_ms(c2, c3, "3rd");
return 0;
}