kernel argument setup

[pvelesko]

I would like to clarification on how argument setup works.

Consider the following kernel:

static __global__ void simpleKernel(float* out_d, float* in_d) {
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx < N) out_d[idx] = CONSTANT * in_d[idx];
}

and I will be passing the following pointers:

  float *in_d, *out_d;
  HIP_CHECK(hipMalloc(&in_d, N * sizeof(float)));
  HIP_CHECK(hipMalloc(&out_d, N * sizeof(float)));

We setup arguments using void **Args pointer which is generated by Clang, correct?
It should be of size 2 and each value should correspond to &in_d and &out_d ?
I should be able to dereference these pointers to get addresses in the GPU space. I can do that when I allocate, but I am having issues doing that using **Args ptrs. What am I missing here?

@pjaaskel @franz

[linehill]

We setup arguments using void **Args pointer which is generated by Clang, correct?

Yes, the clang creates an array of pointers and each of them points to the value of the corresponding argument. Args[0] = a pointer to the value of the first argument, Arg[1] = pointer to the value of next argument, etc.

It should be of size 2 and each value should correspond to &in_d and &out_d?

The size is two, but the clang created array does not hold &in_d and &out_d pointers in the current state. Instead, the array has pointers to local copies of the arguments whose lifetime extends to the duration of the kernel launch. For example, for simpleKernel<<<...>>>(out_d, in_d) expression, clang generates a call to a kernel launch stub (also generated by the clang) which looks roughly like the following:

void __device_stub__simpleKernel(float *arg0, float *arg1) {
  float *local_arg0 = arg0;
  float *local_arg1 = arg1;
  void *Args[2] = {&local_arg0, &local_arg1};
  hipLaunchKernel(..., Args, ...);
}

... I can do that when I allocate, but I am having issues doing that using **Args ptrs.

Where and when are you trying to access the argument values? Could the pointers of the Args array be dereferenced after the __device_stub__simpleKernel call (in which point the pointers to the local_* variables are invalid)?

[pvelesko]

  float *local_arg0 = arg0;
  float *local_arg1 = arg1;

1. What is the purpose of making these local copies instead of using arguments directly?
2. Even with these local pointers, I should be able to de-reference them to get the original addresses that were passed to the kernel invocation, no?

Where and when are you trying to access the argument values?

Much later. This is for graphs so first we accumulate all the launches and then merge them into a single command list.

Could the pointers of the Args array be dereferenced after the _device_stub__simpleKernel call (in which point the pointers to the local* variables are invalid)?

That is exactly the case. I found a solution of setting up CHIPExecItem right away (instead of doing so during the launch) and it works but I was curious as to why I wasn't able to get the device pointers from **Args

[linehill]

What is the purpose of making these local copies instead of using arguments directly?

I got the kernel launch stub a bit wrong. More accurately it looks like this:

void __device_stub__simpleKernel(float *arg0, float *arg1) {
  void *Args[2] = {&arg0, &arg1};
  hipLaunchKernel(..., Args, ...);
}

Even with these local pointers, I should be able to de-reference them to get the original addresses that were passed to the kernel invocation, no?

Yes, they can be dereferenced, but until to the point the caller returns from the stub, and you do get to access the arguments of the kernel invocation in the state they are passed to the invocation.

I’m starting to see that I misread the original “It should be of size 2 and each value should correspond to &in_d and &out_d?” question (I was thrown off by the &in_d and &out_d). I thought you referred to the addresses of the in_d and out_d pointer variables.