__device__

#include <stdio.h>
#include <stdlib.h>

/**
 * @brief This function sets the int passed to 12345.
 * @details This function is only callable from a device because of the
 * __device__ function specifier.
 **/
__device__ void initialiseA(int* a)
{
    *a = 12345;
}

/**
 * @brief This function does the interface between the host and the __device__
 * function because its __global__ function specifier means it will be called
 * from the host, and since this function is on the device, it can call
 * __device__ functions such as initialiseA.
 **/
__global__ void interface(int* a)
{
    initialiseA(a);
}

/**
 * @brief Illustrates how to use a __device__ function.
 * @details This application consists in initialising a variable via a
 * __device__ function. To call such a function, a __global__ function serves
 * as the man-in-the-middle; both callable from host and able to call __device__
 * functions.
 **/
int main(int argc, char* argv[])
{
    // Declare variable on host.
    int a_host = 0;

    // Declare pointer that will point to the memory allocated on the device.
    int* a_device;

    // Allocate memory on the device
    cudaMalloc(&a_device, sizeof(int));

    // Launch the kernel on the device
    printf("Before launching the kernel, a = %d.\n", a_host);
    interface<<<1, 1>>>(a_device);

    // Copy the variable initialised back from the device to the host and print its value
    cudaMemcpy(&a_host, a_device, sizeof(int), cudaMemcpyDeviceToHost);
    printf("After launching the kernel, a = %d.\n", a_host);

    // Free resources
    cudaFree(a_device);

    return EXIT_SUCCESS;
}