run4flat/Minimal.c

## Minimal.c
/*
 * This file was generated automatically by ExtUtils::ParseXS version 3.18 from the
 * contents of Minimal.xs. Do not edit this file, edit Minimal.xs instead.
 *
 *    ANY CHANGES MADE HERE WILL BE LOST!
 *
 */

#line 1 "lib/CUDA/Minimal.xs"
#include "EXTERN.h"
#include "perl.h"
#include "XSUB.h"

#include "ppport.h"

#line 17 "lib/CUDA/Minimal.c"
#ifndef PERL_UNUSED_VAR
#  define PERL_UNUSED_VAR(var) if (0) var = var
#endif

#ifndef dVAR
#  define dVAR		dNOOP
#endif


/* This stuff is not part of the API! You have been warned. */
#ifndef PERL_VERSION_DECIMAL
#  define PERL_VERSION_DECIMAL(r,v,s) (r*1000000 + v*1000 + s)
#endif
#ifndef PERL_DECIMAL_VERSION
#  define PERL_DECIMAL_VERSION \
	  PERL_VERSION_DECIMAL(PERL_REVISION,PERL_VERSION,PERL_SUBVERSION)
#endif
#ifndef PERL_VERSION_GE
#  define PERL_VERSION_GE(r,v,s) \
	  (PERL_DECIMAL_VERSION >= PERL_VERSION_DECIMAL(r,v,s))
#endif
#ifndef PERL_VERSION_LE
#  define PERL_VERSION_LE(r,v,s) \
	  (PERL_DECIMAL_VERSION <= PERL_VERSION_DECIMAL(r,v,s))
#endif

/* XS_INTERNAL is the explicit static-linkage variant of the default
 * XS macro.
 *
 * XS_EXTERNAL is the same as XS_INTERNAL except it does not include
 * "STATIC", ie. it exports XSUB symbols. You probably don't want that
 * for anything but the BOOT XSUB.
 *
 * See XSUB.h in core!
 */


/* TODO: This might be compatible further back than 5.10.0. */
#if PERL_VERSION_GE(5, 10, 0) && PERL_VERSION_LE(5, 15, 1)
#  undef XS_EXTERNAL
#  undef XS_INTERNAL
#  if defined(__CYGWIN__) && defined(USE_DYNAMIC_LOADING)
#    define XS_EXTERNAL(name) __declspec(dllexport) XSPROTO(name)
#    define XS_INTERNAL(name) STATIC XSPROTO(name)
#  endif
#  if defined(__SYMBIAN32__)
#    define XS_EXTERNAL(name) EXPORT_C XSPROTO(name)
#    define XS_INTERNAL(name) EXPORT_C STATIC XSPROTO(name)
#  endif
#  ifndef XS_EXTERNAL
#    if defined(HASATTRIBUTE_UNUSED) && !defined(__cplusplus)
#      define XS_EXTERNAL(name) void name(pTHX_ CV* cv __attribute__unused__)
#      define XS_INTERNAL(name) STATIC void name(pTHX_ CV* cv __attribute__unused__)
#    else
#      ifdef __cplusplus
#        define XS_EXTERNAL(name) extern "C" XSPROTO(name)
#        define XS_INTERNAL(name) static XSPROTO(name)
#      else
#        define XS_EXTERNAL(name) XSPROTO(name)
#        define XS_INTERNAL(name) STATIC XSPROTO(name)
#      endif
#    endif
#  endif
#endif

/* perl >= 5.10.0 && perl <= 5.15.1 */


/* The XS_EXTERNAL macro is used for functions that must not be static
 * like the boot XSUB of a module. If perl didn't have an XS_EXTERNAL
 * macro defined, the best we can do is assume XS is the same.
 * Dito for XS_INTERNAL.
 */
#ifndef XS_EXTERNAL
#  define XS_EXTERNAL(name) XS(name)
#endif
#ifndef XS_INTERNAL
#  define XS_INTERNAL(name) XS(name)
#endif

/* Now, finally, after all this mess, we want an ExtUtils::ParseXS
 * internal macro that we're free to redefine for varying linkage due
 * to the EXPORT_XSUB_SYMBOLS XS keyword. This is internal, use
 * XS_EXTERNAL(name) or XS_INTERNAL(name) in your code if you need to!
 */

#undef XS_EUPXS
#if defined(PERL_EUPXS_ALWAYS_EXPORT)
#  define XS_EUPXS(name) XS_EXTERNAL(name)
#else
   /* default to internal */
#  define XS_EUPXS(name) XS_INTERNAL(name)
#endif

#ifndef PERL_ARGS_ASSERT_CROAK_XS_USAGE
#define PERL_ARGS_ASSERT_CROAK_XS_USAGE assert(cv); assert(params)

/* prototype to pass -Wmissing-prototypes */
STATIC void
S_croak_xs_usage(pTHX_ const CV *const cv, const char *const params);

STATIC void
S_croak_xs_usage(pTHX_ const CV *const cv, const char *const params)
{
    const GV *const gv = CvGV(cv);

    PERL_ARGS_ASSERT_CROAK_XS_USAGE;

    if (gv) {
        const char *const gvname = GvNAME(gv);
        const HV *const stash = GvSTASH(gv);
        const char *const hvname = stash ? HvNAME(stash) : NULL;

        if (hvname)
            Perl_croak(aTHX_ "Usage: %s::%s(%s)", hvname, gvname, params);
        else
            Perl_croak(aTHX_ "Usage: %s(%s)", gvname, params);
    } else {
        /* Pants. I don't think that it should be possible to get here. */
        Perl_croak(aTHX_ "Usage: CODE(0x%"UVxf")(%s)", PTR2UV(cv), params);
    }
}
#undef  PERL_ARGS_ASSERT_CROAK_XS_USAGE

#ifdef PERL_IMPLICIT_CONTEXT
#define croak_xs_usage(a,b)    S_croak_xs_usage(aTHX_ a,b)
#else
#define croak_xs_usage        S_croak_xs_usage
#endif

#endif

/* NOTE: the prototype of newXSproto() is different in versions of perls,
 * so we define a portable version of newXSproto()
 */
#ifdef newXS_flags
#define newXSproto_portable(name, c_impl, file, proto) newXS_flags(name, c_impl, file, proto, 0)
#else
#define newXSproto_portable(name, c_impl, file, proto) (PL_Sv=(SV*)newXS(name, c_impl, file), sv_setpv(PL_Sv, proto), (CV*)PL_Sv)
#endif /* !defined(newXS_flags) */

#line 159 "lib/CUDA/Minimal.c"

XS_EUPXS(XS_CUDA__Minimal__free); /* prototype to pass -Wmissing-prototypes */
XS_EUPXS(XS_CUDA__Minimal__free)
{
    dVAR; dXSARGS;
    if (items != 1)
       croak_xs_usage(cv,  "dev_ptr_SV");
    {
	SV *	dev_ptr_SV = ST(0)
;
#line 12 "lib/CUDA/Minimal.xs"
		// Only free the memory if the pointer is not null:
		if (SvIV(dev_ptr_SV) != 0) {
			// Cast the SV to a device pointer:
			void * dev_ptr = INT2PTR(void*, SvIV(dev_ptr_SV));
			// Free the memory:
			cudaError_t err = cudaFree(dev_ptr);
			// Croak on failure.
			if (err != cudaSuccess)
				croak("Unable to free memory on the device: %s"
							, cudaGetErrorString(err));
			// Set SV to have a value of zero to prevent accidental double frees:
			sv_setiv(dev_ptr_SV, 0);
		}
#line 184 "lib/CUDA/Minimal.c"
    }
    XSRETURN_EMPTY;
}


XS_EUPXS(XS_CUDA__Minimal__malloc); /* prototype to pass -Wmissing-prototypes */
XS_EUPXS(XS_CUDA__Minimal__malloc)
{
    dVAR; dXSARGS;
    if (items != 1)
       croak_xs_usage(cv,  "data_SV");
    {
	SV *	RETVAL;
	SV *	data_SV = ST(0)
;
#line 29 "lib/CUDA/Minimal.xs"
		// First thing's first: guard against calls in void context:
		if (GIMME_V == G_VOID)
			croak("Cannot call Malloc in void context");
		void * dev_ptr = 0;
		size_t data_len = 0;
		// Check the input arguments:
		if (SvTYPE(data_SV) == SVt_PV) {
			// If the host scalar is a PV, use its length:
			data_len = (size_t)SvCUR(data_SV);
		}
		else {
			// Otherwise interpret the scalar as an integer
			// and use it as the length:
			data_len = (size_t)SvIV(data_SV);
		}
		// Allocate the memory:
		cudaError_t err = cudaMalloc(&dev_ptr, data_len);
		// Check for errors:
		if (err != cudaSuccess)
			croak("Unable to allocate %lu bytes on the device: %s"
						, (long unsigned)data_len, cudaGetErrorString(err));
		// Set the return:
		RETVAL = newSViv(PTR2IV(dev_ptr));
#line 224 "lib/CUDA/Minimal.c"
	ST(0) = RETVAL;
	sv_2mortal(ST(0));
    }
    XSRETURN(1);
}


XS_EUPXS(XS_CUDA__Minimal__transfer); /* prototype to pass -Wmissing-prototypes */
XS_EUPXS(XS_CUDA__Minimal__transfer)
{
    dVAR; dXSARGS;
    if (items < 2)
       croak_xs_usage(cv,  "src_SV, dst_SV, ...");
    {
	SV *	src_SV = ST(0)
;
	SV *	dst_SV = ST(1)
;
#line 60 "lib/CUDA/Minimal.xs"
		void * dst_ptr = 0;
		void * src_ptr = 0;
		size_t length = 0;
		size_t host_offset = 0;
		size_t host_length = 0;
		enum cudaMemcpyKind kind;

		// Get the specified length and host offset if they passed it in:
		if (items > 2) length = (size_t)SvIV(ST(2));
		if (items > 3) host_offset = (size_t)SvIV(ST(2));

		// Determine if either of the two SVs are the host memory:
		if (SvTYPE(dst_SV) == SVt_PV && SvTYPE(src_SV) == SVt_PV) {
			// We can't have both of them looking like host memory:
			croak("Transfer requires one or more of %s\n%s"
							, "the arguments to be a device pointer"
							, "but it looks like both are host arrays");
		}
		else if (SvTYPE(dst_SV) == SVt_PV) {
			// Looks like the destination is host memory.
			kind = cudaMemcpyDeviceToHost;
			host_length = (size_t)SvCUR(dst_SV) - host_offset;
			src_ptr = INT2PTR(void*, SvIV(src_SV));
			dst_ptr = sv_2pvbyte_nolen(dst_SV) + host_offset;
			// Make sure the offset is shorter than the host length:
			if (host_length <= 0)
				croak("Host offset must be less than the host's length");
		}
		else if (SvTYPE(src_SV) == SVt_PV) {
			// Looks like the source is host memory.
			kind = cudaMemcpyHostToDevice;
			host_length = (size_t)SvCUR(src_SV) - host_offset;
			src_ptr = sv_2pvbyte_nolen(src_SV) + host_offset;
			dst_ptr = INT2PTR(void*, SvIV(dst_SV));
			// Make sure the offset is shorter than the host length:
			if (host_length <= 0)
				croak("Host offset must be less than the host's length");
		}
		else {
			// Looks like both the source and destination are device pointers.
			kind = cudaMemcpyDeviceToDevice;
			src_ptr = INT2PTR(void*, SvIV(src_SV));
			dst_ptr = INT2PTR(void*, SvIV(dst_SV));
			if (host_offset > 0) {
				croak("Host offsets are not allowed for %s"
						, "device-to-device transfers");
			}
		}

		// Make sure that they provided a length of some sort
		if (length == 0 && host_length == 0)
			croak("You must provide the number of bytes %s"
						, "for device-to-device transfers");

		// Make sure the requested length does not exceed the host's length
		if (host_length > 0 && length > host_length)
			croak("Attempting to transfer more data %s"
						, "than the host can accomodate");

		// Use the host length if no length was explicitly given:
		if (length == 0) length = host_length;

		// Perform the copy and check for errors:
		cudaError_t err = cudaMemcpy(dst_ptr, src_ptr, length, kind);
		if (err != cudaSuccess)
			croak("Unable to copy memory: %s"
						, cudaGetErrorString(err));
#line 311 "lib/CUDA/Minimal.c"
    }
    XSRETURN_EMPTY;
}


XS_EUPXS(XS_CUDA__Minimal_ThreadSynchronize); /* prototype to pass -Wmissing-prototypes */
XS_EUPXS(XS_CUDA__Minimal_ThreadSynchronize)
{
    dVAR; dXSARGS;
    if (items != 0)
       croak_xs_usage(cv,  "");
    {
#line 131 "lib/CUDA/Minimal.xs"
		cudaThreadSynchronize();
#line 326 "lib/CUDA/Minimal.c"
    }
    XSRETURN_EMPTY;
}


XS_EUPXS(XS_CUDA__Minimal_GetLastError); /* prototype to pass -Wmissing-prototypes */
XS_EUPXS(XS_CUDA__Minimal_GetLastError)
{
    dVAR; dXSARGS;
    if (items != 0)
       croak_xs_usage(cv,  "");
    {
	SV *	RETVAL;
#line 136 "lib/CUDA/Minimal.xs"
		cudaError_t err = cudaGetLastError();
		RETVAL = newSVpv(cudaGetErrorString(err), 0);
#line 343 "lib/CUDA/Minimal.c"
	ST(0) = RETVAL;
	sv_2mortal(ST(0));
    }
    XSRETURN(1);
}


XS_EUPXS(XS_CUDA__Minimal_PeekAtLastError); /* prototype to pass -Wmissing-prototypes */
XS_EUPXS(XS_CUDA__Minimal_PeekAtLastError)
{
    dVAR; dXSARGS;
    if (items != 0)
       croak_xs_usage(cv,  "");
    {
	SV *	RETVAL;
#line 144 "lib/CUDA/Minimal.xs"
		cudaError_t err = cudaPeekAtLastError();
		RETVAL = newSVpv(cudaGetErrorString(err), 0);
#line 362 "lib/CUDA/Minimal.c"
	ST(0) = RETVAL;
	sv_2mortal(ST(0));
    }
    XSRETURN(1);
}

#ifdef __cplusplus
extern "C"
#endif
XS_EXTERNAL(boot_CUDA__Minimal); /* prototype to pass -Wmissing-prototypes */
XS_EXTERNAL(boot_CUDA__Minimal)
{
    dVAR; dXSARGS;
#if (PERL_REVISION == 5 && PERL_VERSION < 9)
    char* file = __FILE__;
#else
    const char* file = __FILE__;
#endif

    PERL_UNUSED_VAR(cv); /* -W */
    PERL_UNUSED_VAR(items); /* -W */
#ifdef XS_APIVERSION_BOOTCHECK
    XS_APIVERSION_BOOTCHECK;
#endif
    XS_VERSION_BOOTCHECK;

        newXS("CUDA::Minimal::_free", XS_CUDA__Minimal__free, file);
        newXS("CUDA::Minimal::_malloc", XS_CUDA__Minimal__malloc, file);
        (void)newXSproto_portable("CUDA::Minimal::_transfer", XS_CUDA__Minimal__transfer, file, "$$;$$");
        newXS("CUDA::Minimal::ThreadSynchronize", XS_CUDA__Minimal_ThreadSynchronize, file);
        newXS("CUDA::Minimal::GetLastError", XS_CUDA__Minimal_GetLastError, file);
        newXS("CUDA::Minimal::PeekAtLastError", XS_CUDA__Minimal_PeekAtLastError, file);

    /* Initialisation Section */

#line 150 "lib/CUDA/Minimal.xs"
#undef PERL_VERSION
#define PERL_VERSION 0

#line 402 "lib/CUDA/Minimal.c"

    /* End of Initialisation Section */

#if (PERL_REVISION == 5 && PERL_VERSION >= 9)
  if (PL_unitcheckav)
       call_list(PL_scopestack_ix, PL_unitcheckav);
#endif
    XSRETURN_YES;
}


## prepoc-Minimal.c
static void XS_CUDA__Minimal__free(register PerlInterpreter* my_perl __attribute__((unused)), CV* cv __attribute__((unused)));
static void XS_CUDA__Minimal__free(register PerlInterpreter* my_perl __attribute__((unused)), CV* cv __attribute__((unused)))
{
    extern int Perl___notused __attribute__((unused)); SV **sp = (((PerlInterpreter *)pthread_getspecific(PL_thr_key))->Istack_sp); I32 ax = (*(((PerlInterpreter *)pthread_getspecific(PL_thr_key))->Imarkstack_ptr)--); register SV **mark = (((PerlInterpreter *)pthread_getspecific(PL_thr_key))->Istack_base) + ax++; I32 items = (I32)(sp - mark);
    if (items != 1)
       Perl_croak_xs_usage(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), cv,"dev_ptr_SV");
    {
 SV * dev_ptr_SV = (((PerlInterpreter *)pthread_getspecific(PL_thr_key))->Istack_base)[ax + (0)]
;
# 12 "lib/CUDA/Minimal.xs"

  if ((((dev_ptr_SV)->sv_flags & 0x00000100) ? ((XPVIV*) (dev_ptr_SV)->sv_any)->xiv_u.xivu_iv : Perl_sv_2iv_flags(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), dev_ptr_SV,2)) != 0) {

   void * dev_ptr = (void*)(unsigned long)((((dev_ptr_SV)->sv_flags & 0x00000100) ? ((XPVIV*) (dev_ptr_SV)->sv_any)->xiv_u.xivu_iv : Perl_sv_2iv_flags(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), dev_ptr_SV,2)));

   cudaError_t err = cudaFree(dev_ptr);

   if (err != cudaSuccess)
    Perl_croak_nocontext("Unable to free memory on the device: %s"
       , cudaGetErrorString(err));

   Perl_sv_setiv(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), dev_ptr_SV,0);
  }
# 184 "lib/CUDA/Minimal.c"
    }
    (void)( { (void)( { const IV tmpXSoff = (0); (((PerlInterpreter *)pthread_getspecific(PL_thr_key))->Istack_sp) = (((PerlInterpreter *)pthread_getspecific(PL_thr_key))->Istack_base) + ax + (tmpXSoff - 1); return; } ); } );
}


void boot_CUDA__Minimal(register PerlInterpreter* my_perl __attribute__((unused)), CV* cv __attribute__((unused)));
void boot_CUDA__Minimal(register PerlInterpreter* my_perl __attribute__((unused)), CV* cv __attribute__((unused)))
{
    extern int Perl___notused __attribute__((unused)); SV **sp = (((PerlInterpreter *)pthread_getspecific(PL_thr_key))->Istack_sp); I32 ax = (*(((PerlInterpreter *)pthread_getspecific(PL_thr_key))->Imarkstack_ptr)--); register SV **mark = (((PerlInterpreter *)pthread_getspecific(PL_thr_key))->Istack_base) + ax++; I32 items = (I32)(sp - mark);


    const char* file = "lib/CUDA/Minimal.c";


    ((void)cv);
    ((void)items);

    Perl_xs_apiversion_bootcheck(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), (((PerlInterpreter *)pthread_getspecific(PL_thr_key))->Istack_base)[ax + (0)], ("" "v" "5" "." "14" "." "0" ""), (sizeof("v" "5" "." "14" "." "0")-1));

    Perl_xs_version_bootcheck(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), items, ax, ("" 0.01 ""), (sizeof(0.01)-1));

        Perl_newXS(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), "CUDA::Minimal::_free",XS_CUDA__Minimal__free,file);
        Perl_newXS(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), "CUDA::Minimal::_malloc",XS_CUDA__Minimal__malloc,file);
        (void)Perl_newXS_flags(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), "CUDA::Minimal::_transfer",XS_CUDA__Minimal__transfer,file,"$$;$$",0);
        Perl_newXS(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), "CUDA::Minimal::ThreadSynchronize",XS_CUDA__Minimal_ThreadSynchronize,file);
        Perl_newXS(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), "CUDA::Minimal::GetLastError",XS_CUDA__Minimal_GetLastError,file);
        Perl_newXS(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), "CUDA::Minimal::PeekAtLastError",XS_CUDA__Minimal_PeekAtLastError,file);
# 150 "lib/CUDA/Minimal.xs"
# 402 "lib/CUDA/Minimal.c"


    (void)( { ((((PerlInterpreter *)pthread_getspecific(PL_thr_key))->Istack_base)[ax + (0)] = &(((PerlInterpreter *)pthread_getspecific(PL_thr_key))->Isv_yes) ); (void)( { const IV tmpXSoff = (1); (((PerlInterpreter *)pthread_getspecific(PL_thr_key))->Istack_sp) = (((PerlInterpreter *)pthread_getspecific(PL_thr_key))->Istack_base) + ax + (tmpXSoff - 1); return; } ); } );
}
	/*
	* This file was generated automatically by ExtUtils::ParseXS version 3.18 from the
	* contents of Minimal.xs. Do not edit this file, edit Minimal.xs instead.
	*
	* ANY CHANGES MADE HERE WILL BE LOST!
	*
	*/

	#line 1 "lib/CUDA/Minimal.xs"
	#include "EXTERN.h"
	#include "perl.h"
	#include "XSUB.h"

	#include "ppport.h"

	#line 17 "lib/CUDA/Minimal.c"
	#ifndef PERL_UNUSED_VAR
	# define PERL_UNUSED_VAR(var) if (0) var = var
	#endif

	#ifndef dVAR
	# define dVAR dNOOP
	#endif


	/* This stuff is not part of the API! You have been warned. */
	#ifndef PERL_VERSION_DECIMAL
	# define PERL_VERSION_DECIMAL(r,v,s) (r1000000 + v1000 + s)
	#endif
	#ifndef PERL_DECIMAL_VERSION
	# define PERL_DECIMAL_VERSION \
	PERL_VERSION_DECIMAL(PERL_REVISION,PERL_VERSION,PERL_SUBVERSION)
	#endif
	#ifndef PERL_VERSION_GE
	# define PERL_VERSION_GE(r,v,s) \
	(PERL_DECIMAL_VERSION >= PERL_VERSION_DECIMAL(r,v,s))
	#endif
	#ifndef PERL_VERSION_LE
	# define PERL_VERSION_LE(r,v,s) \
	(PERL_DECIMAL_VERSION <= PERL_VERSION_DECIMAL(r,v,s))
	#endif

	/* XS_INTERNAL is the explicit static-linkage variant of the default
	* XS macro.
	*
	* XS_EXTERNAL is the same as XS_INTERNAL except it does not include
	* "STATIC", ie. it exports XSUB symbols. You probably don't want that
	* for anything but the BOOT XSUB.
	*
	* See XSUB.h in core!
	*/


	/* TODO: This might be compatible further back than 5.10.0. */
	#if PERL_VERSION_GE(5, 10, 0) && PERL_VERSION_LE(5, 15, 1)
	# undef XS_EXTERNAL
	# undef XS_INTERNAL
	# if defined(__CYGWIN__) && defined(USE_DYNAMIC_LOADING)
	# define XS_EXTERNAL(name) __declspec(dllexport) XSPROTO(name)
	# define XS_INTERNAL(name) STATIC XSPROTO(name)
	# endif
	# if defined(__SYMBIAN32__)
	# define XS_EXTERNAL(name) EXPORT_C XSPROTO(name)
	# define XS_INTERNAL(name) EXPORT_C STATIC XSPROTO(name)
	# endif
	# ifndef XS_EXTERNAL
	# if defined(HASATTRIBUTE_UNUSED) && !defined(__cplusplus)
	# define XS_EXTERNAL(name) void name(pTHX_ CV* cv __attribute__unused__)
	# define XS_INTERNAL(name) STATIC void name(pTHX_ CV* cv __attribute__unused__)
	# else
	# ifdef __cplusplus
	# define XS_EXTERNAL(name) extern "C" XSPROTO(name)
	# define XS_INTERNAL(name) static XSPROTO(name)
	# else
	# define XS_EXTERNAL(name) XSPROTO(name)
	# define XS_INTERNAL(name) STATIC XSPROTO(name)
	# endif
	# endif
	# endif
	#endif

	/* perl >= 5.10.0 && perl <= 5.15.1 */


	/* The XS_EXTERNAL macro is used for functions that must not be static
	* like the boot XSUB of a module. If perl didn't have an XS_EXTERNAL
	* macro defined, the best we can do is assume XS is the same.
	* Dito for XS_INTERNAL.
	*/
	#ifndef XS_EXTERNAL
	# define XS_EXTERNAL(name) XS(name)
	#endif
	#ifndef XS_INTERNAL
	# define XS_INTERNAL(name) XS(name)
	#endif

	/* Now, finally, after all this mess, we want an ExtUtils::ParseXS
	* internal macro that we're free to redefine for varying linkage due
	* to the EXPORT_XSUB_SYMBOLS XS keyword. This is internal, use
	* XS_EXTERNAL(name) or XS_INTERNAL(name) in your code if you need to!
	*/

	#undef XS_EUPXS
	#if defined(PERL_EUPXS_ALWAYS_EXPORT)
	# define XS_EUPXS(name) XS_EXTERNAL(name)
	#else
	/* default to internal */
	# define XS_EUPXS(name) XS_INTERNAL(name)
	#endif

	#ifndef PERL_ARGS_ASSERT_CROAK_XS_USAGE
	#define PERL_ARGS_ASSERT_CROAK_XS_USAGE assert(cv); assert(params)

	/* prototype to pass -Wmissing-prototypes */
	STATIC void
	S_croak_xs_usage(pTHX_ const CV const cv, const char const params);

	STATIC void
	S_croak_xs_usage(pTHX_ const CV const cv, const char const params)
	{
	const GV *const gv = CvGV(cv);

	PERL_ARGS_ASSERT_CROAK_XS_USAGE;

	if (gv) {
	const char *const gvname = GvNAME(gv);
	const HV *const stash = GvSTASH(gv);
	const char *const hvname = stash ? HvNAME(stash) : NULL;

	if (hvname)
	Perl_croak(aTHX_ "Usage: %s::%s(%s)", hvname, gvname, params);
	else
	Perl_croak(aTHX_ "Usage: %s(%s)", gvname, params);
	} else {
	/* Pants. I don't think that it should be possible to get here. */
	Perl_croak(aTHX_ "Usage: CODE(0x%"UVxf")(%s)", PTR2UV(cv), params);
	}
	}
	#undef PERL_ARGS_ASSERT_CROAK_XS_USAGE

	#ifdef PERL_IMPLICIT_CONTEXT
	#define croak_xs_usage(a,b) S_croak_xs_usage(aTHX_ a,b)
	#else
	#define croak_xs_usage S_croak_xs_usage
	#endif

	#endif

	/* NOTE: the prototype of newXSproto() is different in versions of perls,
	* so we define a portable version of newXSproto()
	*/
	#ifdef newXS_flags
	#define newXSproto_portable(name, c_impl, file, proto) newXS_flags(name, c_impl, file, proto, 0)
	#else
	#define newXSproto_portable(name, c_impl, file, proto) (PL_Sv=(SV)newXS(name, c_impl, file), sv_setpv(PL_Sv, proto), (CV)PL_Sv)
	#endif /* !defined(newXS_flags) */

	#line 159 "lib/CUDA/Minimal.c"

	XS_EUPXS(XS_CUDA__Minimal__free); /* prototype to pass -Wmissing-prototypes */
	XS_EUPXS(XS_CUDA__Minimal__free)
	{
	dVAR; dXSARGS;
	if (items != 1)
	croak_xs_usage(cv, "dev_ptr_SV");
	{
	SV * dev_ptr_SV = ST(0)
	;
	#line 12 "lib/CUDA/Minimal.xs"
	// Only free the memory if the pointer is not null:
	if (SvIV(dev_ptr_SV) != 0) {
	// Cast the SV to a device pointer:
	void * dev_ptr = INT2PTR(void*, SvIV(dev_ptr_SV));
	// Free the memory:
	cudaError_t err = cudaFree(dev_ptr);
	// Croak on failure.
	if (err != cudaSuccess)
	croak("Unable to free memory on the device: %s"
	, cudaGetErrorString(err));
	// Set SV to have a value of zero to prevent accidental double frees:
	sv_setiv(dev_ptr_SV, 0);
	}
	#line 184 "lib/CUDA/Minimal.c"
	}
	XSRETURN_EMPTY;
	}


	XS_EUPXS(XS_CUDA__Minimal__malloc); /* prototype to pass -Wmissing-prototypes */
	XS_EUPXS(XS_CUDA__Minimal__malloc)
	{
	dVAR; dXSARGS;
	if (items != 1)
	croak_xs_usage(cv, "data_SV");
	{
	SV * RETVAL;
	SV * data_SV = ST(0)
	;
	#line 29 "lib/CUDA/Minimal.xs"
	// First thing's first: guard against calls in void context:
	if (GIMME_V == G_VOID)
	croak("Cannot call Malloc in void context");
	void * dev_ptr = 0;
	size_t data_len = 0;
	// Check the input arguments:
	if (SvTYPE(data_SV) == SVt_PV) {
	// If the host scalar is a PV, use its length:
	data_len = (size_t)SvCUR(data_SV);
	}
	else {
	// Otherwise interpret the scalar as an integer
	// and use it as the length:
	data_len = (size_t)SvIV(data_SV);
	}
	// Allocate the memory:
	cudaError_t err = cudaMalloc(&dev_ptr, data_len);
	// Check for errors:
	if (err != cudaSuccess)
	croak("Unable to allocate %lu bytes on the device: %s"
	, (long unsigned)data_len, cudaGetErrorString(err));
	// Set the return:
	RETVAL = newSViv(PTR2IV(dev_ptr));
	#line 224 "lib/CUDA/Minimal.c"
	ST(0) = RETVAL;
	sv_2mortal(ST(0));
	}
	XSRETURN(1);
	}


	XS_EUPXS(XS_CUDA__Minimal__transfer); /* prototype to pass -Wmissing-prototypes */
	XS_EUPXS(XS_CUDA__Minimal__transfer)
	{
	dVAR; dXSARGS;
	if (items < 2)
	croak_xs_usage(cv, "src_SV, dst_SV, ...");
	{
	SV * src_SV = ST(0)
	;
	SV * dst_SV = ST(1)
	;
	#line 60 "lib/CUDA/Minimal.xs"
	void * dst_ptr = 0;
	void * src_ptr = 0;
	size_t length = 0;
	size_t host_offset = 0;
	size_t host_length = 0;
	enum cudaMemcpyKind kind;

	// Get the specified length and host offset if they passed it in:
	if (items > 2) length = (size_t)SvIV(ST(2));
	if (items > 3) host_offset = (size_t)SvIV(ST(2));

	// Determine if either of the two SVs are the host memory:
	if (SvTYPE(dst_SV) == SVt_PV && SvTYPE(src_SV) == SVt_PV) {
	// We can't have both of them looking like host memory:
	croak("Transfer requires one or more of %s\n%s"
	, "the arguments to be a device pointer"
	, "but it looks like both are host arrays");
	}
	else if (SvTYPE(dst_SV) == SVt_PV) {
	// Looks like the destination is host memory.
	kind = cudaMemcpyDeviceToHost;
	host_length = (size_t)SvCUR(dst_SV) - host_offset;
	src_ptr = INT2PTR(void*, SvIV(src_SV));
	dst_ptr = sv_2pvbyte_nolen(dst_SV) + host_offset;
	// Make sure the offset is shorter than the host length:
	if (host_length <= 0)
	croak("Host offset must be less than the host's length");
	}
	else if (SvTYPE(src_SV) == SVt_PV) {
	// Looks like the source is host memory.
	kind = cudaMemcpyHostToDevice;
	host_length = (size_t)SvCUR(src_SV) - host_offset;
	src_ptr = sv_2pvbyte_nolen(src_SV) + host_offset;
	dst_ptr = INT2PTR(void*, SvIV(dst_SV));
	// Make sure the offset is shorter than the host length:
	if (host_length <= 0)
	croak("Host offset must be less than the host's length");
	}
	else {
	// Looks like both the source and destination are device pointers.
	kind = cudaMemcpyDeviceToDevice;
	src_ptr = INT2PTR(void*, SvIV(src_SV));
	dst_ptr = INT2PTR(void*, SvIV(dst_SV));
	if (host_offset > 0) {
	croak("Host offsets are not allowed for %s"
	, "device-to-device transfers");
	}
	}

	// Make sure that they provided a length of some sort
	if (length == 0 && host_length == 0)
	croak("You must provide the number of bytes %s"
	, "for device-to-device transfers");

	// Make sure the requested length does not exceed the host's length
	if (host_length > 0 && length > host_length)
	croak("Attempting to transfer more data %s"
	, "than the host can accomodate");

	// Use the host length if no length was explicitly given:
	if (length == 0) length = host_length;

	// Perform the copy and check for errors:
	cudaError_t err = cudaMemcpy(dst_ptr, src_ptr, length, kind);
	if (err != cudaSuccess)
	croak("Unable to copy memory: %s"
	, cudaGetErrorString(err));
	#line 311 "lib/CUDA/Minimal.c"
	}
	XSRETURN_EMPTY;
	}


	XS_EUPXS(XS_CUDA__Minimal_ThreadSynchronize); /* prototype to pass -Wmissing-prototypes */
	XS_EUPXS(XS_CUDA__Minimal_ThreadSynchronize)
	{
	dVAR; dXSARGS;
	if (items != 0)
	croak_xs_usage(cv, "");
	{
	#line 131 "lib/CUDA/Minimal.xs"
	cudaThreadSynchronize();
	#line 326 "lib/CUDA/Minimal.c"
	}
	XSRETURN_EMPTY;
	}


	XS_EUPXS(XS_CUDA__Minimal_GetLastError); /* prototype to pass -Wmissing-prototypes */
	XS_EUPXS(XS_CUDA__Minimal_GetLastError)
	{
	dVAR; dXSARGS;
	if (items != 0)
	croak_xs_usage(cv, "");
	{
	SV * RETVAL;
	#line 136 "lib/CUDA/Minimal.xs"
	cudaError_t err = cudaGetLastError();
	RETVAL = newSVpv(cudaGetErrorString(err), 0);
	#line 343 "lib/CUDA/Minimal.c"
	ST(0) = RETVAL;
	sv_2mortal(ST(0));
	}
	XSRETURN(1);
	}


	XS_EUPXS(XS_CUDA__Minimal_PeekAtLastError); /* prototype to pass -Wmissing-prototypes */
	XS_EUPXS(XS_CUDA__Minimal_PeekAtLastError)
	{
	dVAR; dXSARGS;
	if (items != 0)
	croak_xs_usage(cv, "");
	{
	SV * RETVAL;
	#line 144 "lib/CUDA/Minimal.xs"
	cudaError_t err = cudaPeekAtLastError();
	RETVAL = newSVpv(cudaGetErrorString(err), 0);
	#line 362 "lib/CUDA/Minimal.c"
	ST(0) = RETVAL;
	sv_2mortal(ST(0));
	}
	XSRETURN(1);
	}

	#ifdef __cplusplus
	extern "C"
	#endif
	XS_EXTERNAL(boot_CUDA__Minimal); /* prototype to pass -Wmissing-prototypes */
	XS_EXTERNAL(boot_CUDA__Minimal)
	{
	dVAR; dXSARGS;
	#if (PERL_REVISION == 5 && PERL_VERSION < 9)
	char* file = __FILE__;
	#else
	const char* file = __FILE__;
	#endif

	PERL_UNUSED_VAR(cv); /* -W */
	PERL_UNUSED_VAR(items); /* -W */
	#ifdef XS_APIVERSION_BOOTCHECK
	XS_APIVERSION_BOOTCHECK;
	#endif
	XS_VERSION_BOOTCHECK;

	newXS("CUDA::Minimal::_free", XS_CUDA__Minimal__free, file);
	newXS("CUDA::Minimal::_malloc", XS_CUDA__Minimal__malloc, file);
	(void)newXSproto_portable("CUDA::Minimal::_transfer", XS_CUDA__Minimal__transfer, file, "$$;$$");
	newXS("CUDA::Minimal::ThreadSynchronize", XS_CUDA__Minimal_ThreadSynchronize, file);
	newXS("CUDA::Minimal::GetLastError", XS_CUDA__Minimal_GetLastError, file);
	newXS("CUDA::Minimal::PeekAtLastError", XS_CUDA__Minimal_PeekAtLastError, file);

	/* Initialisation Section */

	#line 150 "lib/CUDA/Minimal.xs"
	#undef PERL_VERSION
	#define PERL_VERSION 0

	#line 402 "lib/CUDA/Minimal.c"

	/* End of Initialisation Section */

	#if (PERL_REVISION == 5 && PERL_VERSION >= 9)
	if (PL_unitcheckav)
	call_list(PL_scopestack_ix, PL_unitcheckav);
	#endif
	XSRETURN_YES;
	}
	static void XS_CUDA__Minimal__free(register PerlInterpreter* my_perl __attribute__((unused)), CV* cv __attribute__((unused)));
	static void XS_CUDA__Minimal__free(register PerlInterpreter* my_perl __attribute__((unused)), CV* cv __attribute__((unused)))
	{
	extern int Perl___notused __attribute__((unused)); SV *sp = (((PerlInterpreter )pthread_getspecific(PL_thr_key))->Istack_sp); I32 ax = ((((PerlInterpreter )pthread_getspecific(PL_thr_key))->Imarkstack_ptr)--); register SV *mark = (((PerlInterpreter )pthread_getspecific(PL_thr_key))->Istack_base) + ax++; I32 items = (I32)(sp - mark);
	if (items != 1)
	Perl_croak_xs_usage(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), cv,"dev_ptr_SV");
	{
	SV * dev_ptr_SV = (((PerlInterpreter *)pthread_getspecific(PL_thr_key))->Istack_base)[ax + (0)]
	;
	# 12 "lib/CUDA/Minimal.xs"

	if ((((dev_ptr_SV)->sv_flags & 0x00000100) ? ((XPVIV) (dev_ptr_SV)->sv_any)->xiv_u.xivu_iv : Perl_sv_2iv_flags(((PerlInterpreter )pthread_getspecific(PL_thr_key)), dev_ptr_SV,2)) != 0) {

	void * dev_ptr = (void)(unsigned long)((((dev_ptr_SV)->sv_flags & 0x00000100) ? ((XPVIV) (dev_ptr_SV)->sv_any)->xiv_u.xivu_iv : Perl_sv_2iv_flags(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), dev_ptr_SV,2)));

	cudaError_t err = cudaFree(dev_ptr);

	if (err != cudaSuccess)
	Perl_croak_nocontext("Unable to free memory on the device: %s"
	, cudaGetErrorString(err));

	Perl_sv_setiv(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), dev_ptr_SV,0);
	}
	# 184 "lib/CUDA/Minimal.c"
	}
	(void)( { (void)( { const IV tmpXSoff = (0); (((PerlInterpreter )pthread_getspecific(PL_thr_key))->Istack_sp) = (((PerlInterpreter )pthread_getspecific(PL_thr_key))->Istack_base) + ax + (tmpXSoff - 1); return; } ); } );
	}




	void boot_CUDA__Minimal(register PerlInterpreter* my_perl __attribute__((unused)), CV* cv __attribute__((unused)));
	void boot_CUDA__Minimal(register PerlInterpreter* my_perl __attribute__((unused)), CV* cv __attribute__((unused)))
	{
	extern int Perl___notused __attribute__((unused)); SV *sp = (((PerlInterpreter )pthread_getspecific(PL_thr_key))->Istack_sp); I32 ax = ((((PerlInterpreter )pthread_getspecific(PL_thr_key))->Imarkstack_ptr)--); register SV *mark = (((PerlInterpreter )pthread_getspecific(PL_thr_key))->Istack_base) + ax++; I32 items = (I32)(sp - mark);



	const char* file = "lib/CUDA/Minimal.c";


	((void)cv);
	((void)items);

	Perl_xs_apiversion_bootcheck(((PerlInterpreter )pthread_getspecific(PL_thr_key)), (((PerlInterpreter )pthread_getspecific(PL_thr_key))->Istack_base)[ax + (0)], ("" "v" "5" "." "14" "." "0" ""), (sizeof("v" "5" "." "14" "." "0")-1));

	Perl_xs_version_bootcheck(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), items, ax, ("" 0.01 ""), (sizeof(0.01)-1));

	Perl_newXS(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), "CUDA::Minimal::_free",XS_CUDA__Minimal__free,file);
	Perl_newXS(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), "CUDA::Minimal::_malloc",XS_CUDA__Minimal__malloc,file);
	(void)Perl_newXS_flags(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), "CUDA::Minimal::_transfer",XS_CUDA__Minimal__transfer,file,"$$;$$",0);
	Perl_newXS(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), "CUDA::Minimal::ThreadSynchronize",XS_CUDA__Minimal_ThreadSynchronize,file);
	Perl_newXS(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), "CUDA::Minimal::GetLastError",XS_CUDA__Minimal_GetLastError,file);
	Perl_newXS(((PerlInterpreter *)pthread_getspecific(PL_thr_key)), "CUDA::Minimal::PeekAtLastError",XS_CUDA__Minimal_PeekAtLastError,file);
	# 150 "lib/CUDA/Minimal.xs"
	# 402 "lib/CUDA/Minimal.c"







	(void)( { ((((PerlInterpreter )pthread_getspecific(PL_thr_key))->Istack_base)[ax + (0)] = &(((PerlInterpreter )pthread_getspecific(PL_thr_key))->Isv_yes) ); (void)( { const IV tmpXSoff = (1); (((PerlInterpreter )pthread_getspecific(PL_thr_key))->Istack_sp) = (((PerlInterpreter )pthread_getspecific(PL_thr_key))->Istack_base) + ax + (tmpXSoff - 1); return; } ); } );
	}