Do kernel-launched child kernels have the same warp size as host-launched kernels?

Question

When a kernel block is launched from the host, it has a warp size of 32. Is it the same for child kernels launched via dynamic parallelism? My guess would be yes, but I haven't seen it in the docs.

The larger question, of course, is: is it worth it?

__global__ void kernel(const int * vec, float * outvec){
    int idx = threadIdx.x;
    float random_sum=0;
    for(int j=0; j<vec[idx]; j++){
        random_sum+=threadsafe_rand_uniform();
    }
    outvec[idx] = random_sum;
}

Ok, this example is kind of contrived. The point, though, is that if you have a loop of different length from thread to thread, it's tempting to try and dynamically parallelize it. However, if the warp is still 32, you're going to end up wasting a lot of processors on warps of uneven sizes. In this particular example, you may want to sort the data first and then dispatch the dynamically parallelizable indexes in one kernel and the poorly shaped indexes in a different one.

you already answered your second/larger question in the "question", right? — m.s., Jul 06 '15 at 08:42
Assuming there is no flexibility in warp size, for this particular case, I guess so. — mmdanziger, Jul 06 '15 at 08:46

score 1 · Accepted Answer · answered Jul 06 '15 at 10:45

They do indeed have the same warp size. But that is because the warp size is fixed for the graphics card. All kernels running on the same graphics card is going to have the same warp size.

Today virtually all GPUs are using a warp size of 32, but it may change in the future.

Were you possibly thinking about the number of threads in the kernel, and not warp size? If that is the case then no, they are not necessarily the same. You launch a new kernel with dynamic parallelism the same way as you launch it from host:

<<<blocks, threads>>>threadsafe_rand_uniform();

Beware that this is not the same thing as just calling a device function, which is what you are currently doing.

To your question on whether it is worth it? Well, it is hard to tell without considering the alternative. If the alternative is to return data to host, so that the host can launch a new appropriate kernel, then it may well be worth it. But it all depends on the context.

Thanks. I meant warp. Because if you have fewer than 32 iterations (or a non-integer multiple of 32) you're going to have difficulty filling the warp and therefore losing efficiency. — mmdanziger, Jul 06 '15 at 10:49

Do kernel-launched child kernels have the same warp size as host-launched kernels?

1 Answers1