Test for CUDA overheads when using Dynamic Parallelism.

gistfile1.rb

#!/usr/bin/ruby

for narg in 0..10 do
  File.open('kernel_dp' + narg.to_s + '.cu', 'w') do |kernel|
    kernel.write('extern "C" __global__ void kernel(');
    for iarg in 0...narg do
      if (iarg != 0) then
        kernel.write(',');
      end
      kernel.write('void *a');
      kernel.write(iarg);
    end
    kernel.write('){}');
    kernel.write('extern "C" __global__ void outer(){');
    kernel.write('void *a=0;');
    kernel.write('for(int i=0;i<1000000;++i){');
    kernel.write('kernel<<<1,1>>>(');
    for iarg in 0...narg do
      if (iarg != 0) then
        kernel.write(',');
      end
      kernel.write('a');
    end
    kernel.write(');');
    kernel.write('}');
    kernel.write('}');
  end
end

for narg in 0..10 do
  File.open('runtime_dp' + narg.to_s + '.cu', 'w') do |runtime_dp|
    runtime_dp.write("#include <cstdio>\n");
    runtime_dp.write("#include <sys/time.h>\n");
    runtime_dp.write("#include \"kernel_dp" + narg.to_s + ".cu\"\n");
    runtime_dp.write('int main(){');
    runtime_dp.write('void *a=0;');
    runtime_dp.write('kernel<<<1,1>>>(');
    for iarg in 0...narg do
      if (iarg != 0) then
        runtime_dp.write(',');
      end
      runtime_dp.write('a');
    end
    runtime_dp.write(');');
    runtime_dp.write('cudaDeviceSynchronize();');
    runtime_dp.write('struct timeval s;gettimeofday(&s,0);');
    runtime_dp.write('outer<<<1,1>>>();');
    runtime_dp.write('cudaDeviceSynchronize();');
    runtime_dp.write('struct timeval e;gettimeofday(&e,0);');
    runtime_dp.write('double sec=e.tv_sec-s.tv_sec+(e.tv_usec-s.tv_usec)*1e-6;');
    runtime_dp.write('printf("%.06e sec\n",sec);');
    runtime_dp.write('return 0;');
    runtime_dp.write('}');
  end
end