ccudaparallel-processinggpgputesla

Concurrent Kernel Launch Example - CUDA


I'm attempting to implement concurrent kernel launches for a very complex CUDA kernel, so I thought I'd start out with a simple example. It just launches a kernel which does a sum reduction. Simple enough. Here it is:

#include <stdlib.h>
#include <stdio.h>
#include <time.h>
#include <cuda.h>

extern __shared__ char dsmem[];
__device__ double *scratch_space;

__device__ double NDreduceSum(double *a, unsigned short length)
{
    const int tid = threadIdx.x;
    unsigned short k = length;
    double *b;

    b = scratch_space;

    for (int i = tid; i < length; i+= blockDim.x)
        b[i] = a[i];

    __syncthreads();

    do {
        k = (k + 1) / 2;

        if (tid < k && tid + k < length)
            b[tid] += b[tid + k];

        length = k;
        __syncthreads();
    } while (k > 1);

    return b[0];
}

__device__ double reduceSum(double *a, unsigned short length)
{
    const int tid = threadIdx.x;
    unsigned short k = length;

    do
    {
        k = (k + 1) / 2;

        if (tid < k && tid + k < length)
            a[tid] += a[tid + k];

        length = k;
        __syncthreads();
    }
    while (k > 1);

    return a[0];
}

__global__ void kernel_thing(double *ad, int size)
{
    double sum_1, sum_2, sum_3;
    time_t begin, end, t1, t2, t3;

    scratch_space = (double *) &dsmem[0];

    for (int j = 0; j < 1000000; j++) {
        begin = clock();
        sum_1 = NDreduceSum(ad, size);
        end = clock();
    }

    __syncthreads();

    t1 = end - begin;

    begin = clock();

    sum_2 = 0;
    if (threadIdx.x == 0) {
        for (int i = 0; i < size; i++) {
            sum_2 += ad[i];
        }
    }

    __syncthreads();

    end = clock();

    t2 = end - begin;

    __syncthreads();
    begin = clock();
    sum_3 = reduceSum(ad, size);
    end = clock();

    __syncthreads();

    t3 = end - begin;

    if (threadIdx.x == 0) {
        printf("Sum found: %lf and %lf and %lf. In %ld and %ld and %ld ticks.\n", sum_1, sum_2, sum_3, t1, t2, t3);
    }
}

int main(int argc, char **argv)
{
    int i;
    const int size = 512;
    double *a, *ad, *b, *bd;
    double sum_a, sum_b;
    cudaStream_t stream_a, stream_b;
    cudaError_t result;
    cudaEvent_t a_start, a_stop, b_start, b_stop;

    a = (double *) malloc(sizeof(double) * size);
    b = (double *) malloc(sizeof(double) * size);

    srand48(time(0));

    for (i = 0; i < size; i++) {
        a[i] = drand48();
    }

    for (i = 0; i < size; i++) {
        b[i] = drand48();
    }

    sum_a = 0;
    for (i = 0; i < size; i++) {
        sum_a += a[i];
    }

    sum_b = 0;
    for (i = 0; i < size; i++) {
        sum_b += b[i];
    }

    printf("Looking for sum_a %lf\n", sum_a);
    printf("Looking for sum_b %lf\n", sum_b);

    cudaEventCreate(&a_start);
    cudaEventCreate(&b_start);
    cudaEventCreate(&a_stop);
    cudaEventCreate(&b_stop);

    cudaMalloc((void **) &ad, sizeof(double) * size);
    cudaMalloc((void **) &bd, sizeof(double) * size);

    result = cudaStreamCreate(&stream_a);
    result = cudaStreamCreate(&stream_b);

    result = cudaMemcpyAsync(ad, a, sizeof(double) * size, cudaMemcpyHostToDevice, stream_a);
    result = cudaMemcpyAsync(bd, b, sizeof(double) * size, cudaMemcpyHostToDevice, stream_b);

    cudaEventRecord(a_start);
    kernel_thing<<<1, 512, 49152, stream_a>>>(ad, size);
    cudaEventRecord(a_stop);
    cudaEventRecord(b_start);
    kernel_thing<<<1, 512, 49152, stream_b>>>(bd, size);
    cudaEventRecord(b_stop);

    result = cudaMemcpyAsync(a, ad, sizeof(double) * size, cudaMemcpyDeviceToHost, stream_a);
    result = cudaMemcpyAsync(b, bd, sizeof(double) * size, cudaMemcpyDeviceToHost, stream_b);

    cudaEventSynchronize(a_stop);
    cudaEventSynchronize(b_stop);

    float a_ms = 0;
    float b_ms = 0;
    cudaEventElapsedTime(&a_ms, a_start, a_stop);
    cudaEventElapsedTime(&b_ms, b_start, b_stop);

    printf("%lf ms for A.\n", a_ms);
    printf("%lf ms for B.\n", b_ms);

    result = cudaStreamDestroy(stream_a);
    result = cudaStreamDestroy(stream_b);

    if (result != cudaSuccess) {
        printf("I should probably do this after each important operation.\n");
    }

    /*
    printf("Matrix after:\n");
    for (i = 0; i < size; i++) {
        printf("%lf ", a[i]);
    }
    printf("\n");
    */

    free(a);
    free(b);
    cudaFree(ad);
    cudaFree(bd);

    return 0;
}

Compiled like so:

CFLAGS = -arch sm_35

CC = nvcc

all: parallel

parallel: parallel.cu
    $(LINK.c) $^ -o $@

clean:
    rm -f *.o core parallel

I'm using a single Tesla K20X.

When I run this simple example, I get the following output:

Looking for sum_a 247.983945
Looking for sum_b 248.033749
Sum found: 247.983945 and 247.983945 and 247.983945. In 3242 and 51600 and 4792 ticks.
Sum found: 248.033749 and 248.033749 and 248.033749. In 3314 and 52000 and 4497 ticks.
4645.079102 ms for A.
4630.725098 ms for B.
Application 577759 resources: utime ~8s, stime ~2s, Rss ~82764, inblocks ~406, outblocks ~967

So, as you can see, each of the kernels gets the correct results and takes around 4.5 s, which is what I got in an earlier one-kernel version. Great! However, as you can see from the aprun output, the wall time is actually around 10 s, which is much more than the one-kernel version. So, it looks like the kernels are either not launching in parallel, or I'm not getting nearly the speed-up (2x) that I was expecting from concurrent kernel launches.

To tl;dr this question:

  1. Am I missing anything in my code example? Are the kernels actually launching in parallel?
  2. What kind of speed-up should I expect with a Tesla K20X? Shouldn't the kernels run exactly in parallel, completing twice the work in the same time? How many kernels can I expect to run efficiently in parallel?

Thanks for you help.


Solution

  • The cudaEventRecord operations in between your kernels are causing serialization.

    Right now the results you are getting:

    4645.079102 ms for A.
    4630.725098 ms for B.
    

    are back-to-back due to this serialization.

    Instead, just time the entire kernel launch sequence:

    cudaEventRecord(a_start);
    kernel_thing<<<1, 512, 49152, stream_a>>>(ad, size);
    kernel_thing<<<1, 512, 49152, stream_b>>>(bd, size);
    cudaEventRecord(a_stop);
    

    And I think you will see an elapsed time for (a_start, a_stop) that is roughly the same as one of your previous kernels (~4600ms) indicating more or less full concurrency. I used CUDA 6 RC, copied data back to the host rather than printf from kernel, and eliminated the cudaEventRecord operations between the kernel calls, and I got an overall execution time of ~4.8s. If I didn't modify the cudaEventRecord arrangement, instead my execution time was ~8.3s

    A few other notes: