Test for CUDA overheads with varying numbers of arguments.

gistfile1.rb

#!/usr/bin/ruby

for narg in 0..10 do
  File.open('kernel' + 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('){}');
  end
end

for narg in 0..10 do
  File.open('driver' + narg.to_s + '.cpp', 'w') do |kernel|
    kernel.write("#include <cstdio>\n");
    kernel.write("#include <sys/time.h>\n");
    kernel.write("#include <cuda.h>\n");
    kernel.write('int main(){');
    kernel.write('cuInit(0);');
    kernel.write('CUdevice dev;cuDeviceGet(&dev,0);');
    kernel.write('CUcontext ctx;cuCtxCreate(&ctx,0,dev);');
    kernel.write('CUmodule module;');
    kernel.write('cuModuleLoad(&module,"kernel' + narg.to_s + '.cubin");');
    kernel.write('CUfunction f;cuModuleGetFunction(&f,module,"kernel");');
    kernel.write('void *a=0;');
    kernel.write('{');
    kernel.write('void *params[]={');
    for iarg in 0...narg do
      kernel.write('&a,');
    end
    kernel.write('};');
    kernel.write('cuLaunchKernel(f,1,1,1,1,1,1,0,0,params,0);');
    kernel.write('}');
    kernel.write('cuCtxSynchronize();');
    kernel.write('struct timeval s;gettimeofday(&s,0);');
    kernel.write('for(int i=0;i<1000000;++i){');
    kernel.write('void *params[]={');
    for iarg in 0...narg do
      kernel.write('&a,');
    end
    kernel.write('};');
    kernel.write('cuLaunchKernel(f,1,1,1,1,1,1,0,0,params,0);');
    kernel.write('}');
    kernel.write('cuCtxSynchronize();');
    kernel.write('struct timeval e;gettimeofday(&e,0);');
    kernel.write('double sec=e.tv_sec-s.tv_sec+(e.tv_usec-s.tv_usec)*1e-6;');
    kernel.write('printf("%.06e sec\n",sec);');
    kernel.write('return 0;');
    kernel.write('}');
  end
end

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