Example on How to Use the hostcall `__ockl_call_host_function` binding

[matinraayai]

Hi,
I was investigating how to invoke a host function from inside a kernel. I came across the __ockl_call_host_function binding, which seems to go through the same internal function printf goes through. I have some questions regarding this binding's usage:

1. Is it possible to invoke this function from HIP? If so, how do I supply the function's address and the arguments? Why are there 6 args required here? I have a basic working example that I've been trying to compile in HIP. Passing myFunc (casted or not) directly to the first argument doesn't work. Also please provide the correct command of compiling the following HIP file and the device offload link flags.

#include <stdio.h>
#include <algorithm>
#include <stdlib.h>
#include <iostream>
#include "hip/hip_runtime.h"
#include <vector>
#include <random>

#define HIP_ASSERT(x) (assert((x)==hipSuccess))

#define NUM       1024

#define THREADS_PER_BLOCK 256

template<typename T>
void initialize_array(std::vector<T> vec) {
    std::random_device rd;
    std::mt19937 mt(rd());
    std::uniform_real_distribution<double> dist(0.0, 100.0);

    for (auto el: vec) {
        el = 0;
    }
}


extern "C" __device__ long2
__ockl_call_host_function(ulong fptr, ulong arg0, ulong arg1, ulong arg2,
                          ulong arg3, ulong arg4, ulong arg5, ulong arg6);

template<typename T>
int vector_verify(const std::vector<T> &a, const std::vector<T> &b, std::vector<T> &c) {
    int disp = 0;
    for (int i = 0; i < a.size(); i++) {
        if (c[i] != a[i] + b[i])
            disp++;
    }
    return disp;
}

void myFunc() {
  std::cout << "Called by device to run on host\n";
}


template<typename T>
__global__ void k_vector_mult(T* __restrict__ a) {
//      int i = blockDim.x * blockIdx.x + threadIdx.x;
//      if ( i < vec_size)
        // printf("%lx\n", a);
  __ockl_call_host_function(0, 0, 0, 0, 0, 0, 0, 0);
        a[0] = a[0] + 42;
}

template<typename T>
void hipVectorCopy(T* &d_data, std::vector<T> &h_data, hipMemcpyKind memcpyKind) {
    size_t vec_size = sizeof(h_data.data());
    if (memcpyKind == hipMemcpyHostToDevice) {
        HIP_ASSERT(hipMalloc((void**)&d_data, vec_size));
        HIP_ASSERT(hipMemcpy(d_data, h_data.data(), vec_size, hipMemcpyHostToDevice));
    }
    else if (memcpyKind == hipMemcpyDeviceToHost) {
        HIP_ASSERT(hipMemcpy(h_data.data(), d_data, vec_size, hipMemcpyDeviceToHost));
    }
    else throw std::invalid_argument("Unsupported memcpy kind for hipVectorCopy");
}


int main() {
    // Host allocated pointers
    std::vector<int> h_A_int(NUM);
       // Pointers for device global memory
    int* d_A_int;
    std::cout << "Initializing vectors.\n";

    initialize_array(h_A_int);
      std::cout << "Initialized vectors." << std::endl;
    std::cout << "Copying vectors to device.\n";

    hipVectorCopy(d_A_int, h_A_int, hipMemcpyHostToDevice);
      std::cout << "Copied vectors to device.\n";

    dim3 grid_dims(NUM / THREADS_PER_BLOCK + (NUM % THREADS_PER_BLOCK) ? 1 : 0);

    hipLaunchKernelGGL(k_vector_mult, 1,
                       1,
                       0, 0,
                       d_A_int);
    auto err = hipDeviceSynchronize( );
    if ( hipSuccess != err ) {
      std::cerr << "Kernel launch failed!. Error code: " << err << std::endl;
    }

    hipVectorCopy(d_A_int, h_A_int, hipMemcpyDeviceToHost);
   std::cout << h_A_int[0] << std::endl;
    HIP_ASSERT(hipFree(d_A_int));
    return 0;
}

2. Is it possible to perform HSA tasks (loading and freezing executables) inside the called host functions? Are there any implications I should be aware of?
   Tagging @b-sumner @arsenm @bwelton

Thanks in advance!

[b-sumner]

#1 Yes, after explicitly providing the prototype as you have. The address and arguments have to be supplied by the caller. Where it gets those is up to you. For example, you could pass the function address and some of the arguments into the kernel as kernel arguments. The host function can do whatever it needs with its 6 arguments including ignoring some of them. If any of them are GPU pointers then you need to ensure that they make sense to access on the host if the host function tries to access them.

#2 I'm not sure. Does it work to load and freeze executables for a device while a kernel is executing on the device? The wave executing the call to the host function will block for the return value, but all other waves will continue running.

[matinraayai]

@b-sumner thanks for the reply. How exactly can I pass the function pointer to the ockl call? I get the following "use of host function in kernel" error when I try to cast the host function's address to unsigned long:

vecadd.cpp:51:62: error: reference to __host__ function 'myFunc' in __global__ function
   51 |   __ockl_call_host_function(reinterpret_cast<unsigned long>(&myFunc), 0, 0, 0, 0, 0, 0, 0);

Sorry I missed the part where you said "have to be passed as kernel argument".

[b-sumner]

@matinraayai you can pass the function pointer in as a kernel argument or as a field in a kernel argument or hipMemcpy it into device memory somewhere, for example. You cannot form a reference to a host function in device code.

[matinraayai]

@b-sumner it works now that I pass the host function as a kernel argument:

#include <stdio.h>
#include <algorithm>
#include <stdlib.h>
#include <iostream>
#include "hip/hip_runtime.h"
#include <vector>
#include <random>

#define HIP_ASSERT(x) (assert((x)==hipSuccess))

#define NUM       1024

#define THREADS_PER_BLOCK 256

template<typename T>
void initialize_array(std::vector<T> vec) {
    std::random_device rd;
    std::mt19937 mt(rd());
    std::uniform_real_distribution<double> dist(0.0, 100.0);

    for (auto el: vec) {
        el = 0;
    }
}


extern "C" __device__ long2
__ockl_call_host_function(ulong fptr, ulong arg0, ulong arg1, ulong arg2,
                          ulong arg3, ulong arg4, ulong arg5, ulong arg6);

template<typename T>
int vector_verify(const std::vector<T> &a, const std::vector<T> &b, std::vector<T> &c) {
    int disp = 0;
    for (int i = 0; i < a.size(); i++) {
        if (c[i] != a[i] + b[i])
            disp++;
    }
    return disp;
}

void myFunc() {
  std::cout << "Called by device to run on host\n";
}


template<typename T>
__global__ void k_vector_mult(T* __restrict__ a, unsigned long myFuncPtr) {
//      int i = blockDim.x * blockIdx.x + threadIdx.x;
//      if ( i < vec_size)
        // printf("%lx\n", a);
  __ockl_call_host_function(myFuncPtr, 0, 0, 0, 0, 0, 0, 0);
        a[0] = a[0] + 42;
}

template<typename T>
void hipVectorCopy(T* &d_data, std::vector<T> &h_data, hipMemcpyKind memcpyKind) {
    size_t vec_size = sizeof(h_data.data());
    if (memcpyKind == hipMemcpyHostToDevice) {
        HIP_ASSERT(hipMalloc((void**)&d_data, vec_size));
        HIP_ASSERT(hipMemcpy(d_data, h_data.data(), vec_size, hipMemcpyHostToDevice));
    }
    else if (memcpyKind == hipMemcpyDeviceToHost) {
        HIP_ASSERT(hipMemcpy(h_data.data(), d_data, vec_size, hipMemcpyDeviceToHost));
    }
    else throw std::invalid_argument("Unsupported memcpy kind for hipVectorCopy");
}


int main() {
    // Host allocated pointers
    std::vector<int> h_A_int(NUM);
       // Pointers for device global memory
    int* d_A_int;
    std::cout << "Initializing vectors.\n";

    initialize_array(h_A_int);
      std::cout << "Initialized vectors." << std::endl;
    std::cout << "Copying vectors to device.\n";

    hipVectorCopy(d_A_int, h_A_int, hipMemcpyHostToDevice);
      std::cout << "Copied vectors to device.\n";

    dim3 grid_dims(NUM / THREADS_PER_BLOCK + (NUM % THREADS_PER_BLOCK) ? 1 : 0);

    hipLaunchKernelGGL(k_vector_mult<int>, 1,
                       64,
                       0, 0,
                       d_A_int, reinterpret_cast<unsigned long>(&myFunc));
    auto err = hipDeviceSynchronize( );
    if ( hipSuccess != err ) {
      std::cerr << "Kernel launch failed!. Error code: " << err << std::endl;
    }

    hipVectorCopy(d_A_int, h_A_int, hipMemcpyDeviceToHost);
   std::cout << h_A_int[0] << std::endl;
    HIP_ASSERT(hipFree(d_A_int));
    return 0;
}

Running it will end up printing Called by device to run on host\n 64 times.

It makes sense that you can't pass the host function pointer's symbol directly here since we don't know the exact load value of it on the host during compilation.

[matinraayai]

@b-sumner Regarding passing arguments, is it possible to pass arguments by value in the ockl host call? I'm trying to pass threadIdx.x as the first argument, but I'm getting an address value:

Called by thread 719e913f7300 device to run on host

Do I have to do a hipMemcpy before accessing these values?

[b-sumner]

@matinraayai the only argument passing here is by value. threadIdx.x is complicated and you will probably need to cast it to unsigned int or so before passing it.

[matinraayai]

@matinraayai the only argument passing here is by value. threadIdx.x is complicated and you will probably need to cast it to unsigned int or so before passing it.

Passing a constant also doesn't seem to work. It seems that I'm getting a host pointer in my host function. I modified the following part of my example:

unsigned long myFunc(unsigned long x) {
  hipPointerAttribute_t attrs;
  HIP_ASSERT(hipPointerGetAttributes(&attrs, reinterpret_cast<void*>(x)));
  std::cout << "Type: " << attrs.type << "\n";
  std::cout << "Dev ptr: " << attrs.devicePointer << " , host ptr: " << attrs.hostPointer << "\n";
  std::cout << *reinterpret_cast<unsigned long *>(x) << "\n";
  return 5;
}


template<typename T>
__global__ void k_vector_mult(T* __restrict__ a, unsigned long myFuncPtr) {
  __ockl_call_host_function(myFuncPtr, 5, 0, 0, 0, 0, 0, 0);
}

It ends up printing the following 64 times:

Type: 0
Dev ptr: 0 , host ptr: 0
Pointer: 7ff1903c3300, value: 0
Type: 0

Also it seems that I'm getting the same pointer address in all threads.

[b-sumner]

@matinraayai the host function prototype should look like this: void myFunc(unsigned long long *results, unsigned long long *args)

results points to up to 2 unsigned long long, and args points to up to 7 unsigned long long (or uint64_t if you prefer).

On the host, the args can be a void* or something else convenient, similarly the results. But the device needs to pack the arguments into those 7 longs and unpack the results.

[matinraayai]

@b-sumner Thanks I'm now able to extract the input and output arguments inside both the host and the kernel:

void myFunc(unsigned long long *results, unsigned long long *args) {
  for (int i = 0; i < 7; i++) {
    std::cout << args[i] << "\n";
  }
  results[0] = 100;
  results[1] = 200;
}


template<typename T>
__global__ void k_vector_mult(T* __restrict__ a, unsigned long myFuncPtr) {
  long2 out = __ockl_call_host_function(myFuncPtr, 1, 2, 3, 4, 5, 6, 7);
  a[0] = out.x + out.y;
}

This will print the following if the kernel is run with only a single thread:

1
2
3
4
5
6
7
300

With 300 being the result of 100 + 200.

threadIdx.x also seems to work. I think I have everything I need to test if I can load code onto the device inside myFunc.

@b-sumner Out of curiosity why doesn't this have a HIP binding like printf?

[b-sumner]

If Cuda had anything like this, or if this was a typical developer need, this or something like it might be considered for a HIP extension. But neither is true now. It's also undocumented and could potentially change without notice.

[matinraayai]

@b-sumner makes sense. I'm able to use hostcalls effectively on my end. I'm closing this issue. Thanks for your help.