cuda kernel execution delayed by cpu code

user_4932447 0
Viewed 138

I cannot get my head around the following admittedly very simple code, which is a boiled down version of a more complicated project where I spent many hours on by now. This code as is will run on my system in about 2000 milliseconds. But when I enable the line to put cpu into sleep for 500 ms, the program alltogether will run that time longer, making it about 2500 ms.

I cannot understand how that fits into the statement that cuda kernels execute asynchronously with respect to the host?

Running cuda 11.1 on Vistual Studio 2019

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <chrono>
#include <iostream>
#include <numeric>
#include <thread>

__global__ void kernel(double* val, int siz) {
    for (int i = 0; i < siz; i++) val[i] = sqrt(val[i]); //calculate square root for every value in array
}

int main() {
    auto t1 = std::chrono::high_resolution_clock::now();

    const int siz = 1'000'000; //array length
    double* val = new double[siz];
    std::iota(val, val + siz, 0.0); //fill array with 0, 1, 2,...
    double* d_val;

    cudaMalloc(&d_val, sizeof(double) * siz);
    cudaMemcpy(d_val, val, sizeof(double) * siz, cudaMemcpyDefault);
    kernel <<<1, 1 >>> (d_val, siz); //start kernel
    //std::this_thread::sleep_for(std::chrono::milliseconds(500)); //---- putting cpu to sleep also delays kernel execution?
    cudaError_t err = cudaDeviceSynchronize();
    auto t2 = std::chrono::high_resolution_clock::now();

    std::cout << "status: " << cudaGetErrorString(err) << std::endl;
    std::chrono::duration<double, std::milli> ms = t2 - t1;
    std::cout << "duration: " << ms.count() << std::endl;

    delete[] val;
}
cuda c++
1 Answers

I cannot understand how that fits into the statement that cuda kernels execute asynchronously with respect to the host?

You're experiencing WDDM Command Batching as described here.

In a nutshell, on windows, when in the WDDM driver model, GPU commands (e.g. anything from the cuda runtime API, plus kernel launches) will get sent to a command queue. Every so often, according to an unpublished heuristic, and with no explicit user-controls provided, the command queue will be "flushed" i.e. sent to the GPU, at which time (if not currently busy) the GPU will begin processing those commands.

So, on a WDDM setup, the dispatch of kernels to the command queue is non-blocking (control is returned to the CPU thread immediately). The dispatch of work from the command queue to the GPU follows some other heuristic. (The kernel execution is asynchronous to the host thread, in any event)

You have at least a few options if this is a problem:

  1. On windows, switch to a GPU that is in the TCC driver model.
  2. On windows, attempt to use one of the "hacks" described in the linked answer.
  3. Switch to linux
user_16959600
Related