#!/usr/bin/perl

# by bliako @ PerlMonks.org
# date: 01-Jul-2021
# see https://perlmonks.org/?node_id=11134582
# lame example for utilising GPGPU via Inline::C
# TODO: extend to taking params and returning back results

use strict;
use warnings;

use FindBin;

use Inline C => Config =>
	cc => $FindBin::Bin.'/nvcc-compile.pl',
	ld => $FindBin::Bin.'/nvcc-link.pl',
;

use Inline C => <<'EOC';
// from https://developer.nvidia.com/blog/easy-introduction-cuda-c-and-c/
#include <stdio.h>

__global__
void saxpy(int n, float a, float *x, float *y)
{
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if (i < n) y[i] = a*x[i] + y[i];
}

int main()
{
  int N = 1<<20;
  float *x, *y, *d_x, *d_y;
  x = (float*)malloc(N*sizeof(float));
  y = (float*)malloc(N*sizeof(float));

  cudaMalloc(&d_x, N*sizeof(float)); 
  cudaMalloc(&d_y, N*sizeof(float));

  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice);

  // Perform SAXPY on 1M elements
  saxpy<<<(N+255)/256, 256>>>(N, 2.0f, d_x, d_y);

  cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);

  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = max(maxError, abs(y[i]-4.0f));
  printf("Max error: %f\n", maxError);

  cudaFree(d_x);
  cudaFree(d_y);
  free(x);
  free(y);
  return 0; // << late edit!
}
EOC
main;

##</code><code>##

#!/usr/bin/perl

# nvcc-compile.pl
# by bliako @ PerlMonks.org
# date: 01-Jul-2021
# see https://perlmonks.org/?node_id=11134582
# tools for running cuda code on the GPU via Perl and Inline::C
# script to be provided to Inline::C as its 'cc' parameter, like
#  use Inline C => Config =>
#	cc => $FindBin::Bin.'/nvcc-compile.pl',
#	ld => $FindBin::Bin.'/nvcc-link.pl',
#;
# Below, set $EXE and $CC to point to nvcc and gcc commands
# Note that nvcc requires specific gcc versions ONLY

# WARNING: if you make changes here, it is unlikely that Inline::C
# will notice. It's better to delete the temp _Inline directory and start afresh

use strict;
use warnings;

use Cwd;

my $verbose = 0;
my $EXE = '/usr/local/cuda/bin/nvcc';
my $CC = '/usr/local/gcc84/bin/gcc84';

#########################
# nothing to change below
#########################
my $PWD = Cwd::cwd;

my @remove = (
	qr/\-Werror=format\-security(?=\s|$)/,
	qr/\-m64(?=\s|$)/,
	qr/\-mtune=generic(?=\s|$)/,
	qr/\-iquote[^ ]*(?=\s|$)/,
	qr/\-grecord\-gcc\-switches(?=\s|$)/,
	qr/\-pipe(?=\s|$)/,
	qr/\-Wall(?=\s|$)/,
	qr/\-Wp,\-D_FORTIFY_SOURCE=[0-9]+(?=\s|$)/,
	qr/\-Wp,\-D_GLIBCXX_ASSERTIONS(?=\s|$)/,
	qr/\-specs=[^ ]+(?=\s|$)/,
	qr/\-DVERSION=[^ ]+(?=\s|$)/,
	qr/\-DXS_VERSION=[^ ]+(?=\s|$)/,
);

my @replace_compiler_options = (
	qr/(\-flto=auto)(?=\s|$)/,
	qr/(\-ffat\-lto\-objects)(?=\s|$)/,
	qr/(\-fexceptions)(?=\s|$)/,
	qr/(\-fstack\-protector\-strong)(?=\s|$)/,
	qr/(\-fasynchronous\-unwind\-tables)(?=\s|$)/,
	qr/(\-fstack\-clash\-protection)(?=\s|$)/,
	qr/(\-fcf\-protection)(?=\s|$)/,
	qr/(\-fwrapv)(?=\s|$)/,
	qr/(\-fPIC)(?=\s|$)/,
	qr/(\-fno\-strict\-aliasing)(?=\s|$)/,
	qr/(\-Wl,\-\-as-needed)(?=\s|$)/,
	qr/(\-Wl,\-z,now)(?=\s|$)/,
);

my @newarg;
for my $anarg (@ARGV){
	print "processing '$anarg'\n" if $verbose;
	for my $q (@remove){
		if( $anarg =~ s/$q//g ){ print "removing $q...\n" if $verbose }
	}
	for my $q (@replace_compiler_options){
		if( $anarg =~ s/$q/-Xcompiler \\"$1\\"/g ){ print "replacing $q...\n" if $verbose }
	}
	if( $anarg !~ /^\s*$/ ){ push @newarg, $anarg }
}

# hack to change the file extension from .c to .cu
# assumes that the file to compile is the last in @ARGV(!!!)
my $cfile = $newarg[-1];
my $cufile = $cfile;
$cufile =~ s/.c$/.cu/;
$newarg[-1] = $cufile;
my $cmdstr = "cp '$cfile' '$cufile'";
die "failed" if mysystem($cmdstr);

$cmdstr = $EXE." --compiler-bindir /usr/local/gcc84/bin/gcc84 ".join(" ", @newarg);
print "$0 : executing:\n$cmdstr\n";
die "failed" if mysystem($cmdstr); 

#system($EXE, @newarg);

sub mysystem {
	my @args = @_;
	system(@args);

	if ($? == -1) {
	    print STDERR "failed to execute: $!\n";
	    return 1;
	}
	elsif ($? & 127) {
	    printf STDERR "child died with signal %d, %s coredump\n",
	        ($? & 127),  ($? & 128) ? 'with' : 'without';
	    return 1;
	}

	my $ex = $? >> 8;
	if( $ex ){
		print STDERR "error, system command failed with exit code $ex";
		return 1;
	}
	printf "success, system command executed.\n";
	return 0;
}

##</code><code>##

#!/usr/bin/perl

# nvcc-link.pl
# by bliako @ PerlMonks.org
# date: 01-Jul-2021
# see https://perlmonks.org/?node_id=11134582
# tools for running cuda code on the GPU via Perl and Inline::C
# script to be provided to Inline::C as its 'ld' parameter, like
#  use Inline C => Config =>
#       cc => $FindBin::Bin.'/nvcc-compile.pl',
#       ld => $FindBin::Bin.'/nvcc-link.pl',
#;
# Below, set $EXE and $CC to point to nvcc and gcc commands
# Note that nvcc requires specific gcc versions ONLY 

# WARNING: if you make changes here, it is unlikely that Inline::C
# will notice. It's better to delete the temp _Inline directory and start afresh

use strict;
use warnings;

use Cwd;

my $verbose = 0;
my $EXE = '/usr/local/cuda/bin/nvcc';
my $CC = '/usr/local/gcc84/bin/gcc84';

###########################
# nothing tho change below
###########################

my $PWD = Cwd::cwd;

my @remove = (
	qr/\-Werror=format\-security(?=\s|$)/,
	qr/\-m64(?=\s|$)/,
	qr/\-mtune=generic(?=\s|$)/,
	qr/\-iquote[^ ]*(?=\s|$)/,
	qr/\-grecord\-gcc\-switches(?=\s|$)/,
	qr/\-pipe(?=\s|$)/,
	qr/\-Wall(?=\s|$)/,
	qr/\-Wp,\-D_FORTIFY_SOURCE=[0-9]+(?=\s|$)/,
	qr/\-Wp,\-D_GLIBCXX_ASSERTIONS(?=\s|$)/,
	qr/\-specs=[^ ]+(?=\s|$)/,
	qr/\-DVERSION=[^ ]+(?=\s|$)/,
	qr/\-DXS_VERSION=[^ ]+(?=\s|$)/,
);

my @replace_compiler_options = (
	qr/(\-flto=auto)(?=\s|$)/,
	qr/(\-ffat\-lto\-objects)(?=\s|$)/,
	qr/(\-fexceptions)(?=\s|$)/,
	qr/(\-fstack\-protector\-strong)(?=\s|$)/,
	qr/(\-fasynchronous\-unwind\-tables)(?=\s|$)/,
	qr/(\-fstack\-clash\-protection)(?=\s|$)/,
	qr/(\-fcf\-protection)(?=\s|$)/,
	qr/(\-fwrapv)(?=\s|$)/,
	qr/(\-fPIC)(?=\s|$)/,
	qr/(\-fno\-strict\-aliasing)(?=\s|$)/,
	qr/(\-Wl,\-z,relro)(?=\s|$)/,
	qr/(\-Wl,\-\-as-needed)(?=\s|$)/,
	qr/(\-Wl,\-z,now)(?=\s|$)/,
);

my @newarg;
for my $anarg (@ARGV){
	print "processing '$anarg'\n" if $verbose;
	for my $q (@remove){
		if( $anarg =~ s/$q//g ){ print "removing $q...\n" if $verbose }
	}
	for my $q (@replace_compiler_options){
		if( $anarg =~ s/$q/-Xcompiler \\"$1\\"/g ){ print "replacing $q...\n" if $verbose }
	}
	if( $anarg !~ /^\s*$/ ){ push @newarg, $anarg }
}

my $cmdstr = $EXE." --compiler-bindir ${CC} ".join(" ", @newarg);
print "$0 : executing:\n$cmdstr\n";
die "failed" if mysystem($cmdstr); 

sub mysystem {
	my @args = @_;
	system(@args);

	if ($? == -1) {
	    print STDERR "failed to execute: $!\n";
	    return 1;
	}
	elsif ($? & 127) {
	    printf STDERR "child died with signal %d, %s coredump\n",
	        ($? & 127),  ($? & 128) ? 'with' : 'without';
	    return 1;
	}

	my $ex = $? >> 8;
	if( $ex ){
		print STDERR "error, system command failed with exit code $ex";
		return 1;
	}
	printf "success, system command executed.\n";
	return 0;
}