Add a usecase discussion of how to deal with function pointers by crtrott · Pull Request #577 · kokkos/kokkos-core-wiki

crtrott · 2024-09-03T22:28:38Z

Here is the rendered version: https://github.com/crtrott/kokkos-core-wiki/blob/function-pointers/docs/source/usecases/FunctionPointers.md

JBludau

I think it is a good addition

JBludau · 2024-09-04T08:03:11Z

+I can use the function pointer in a host parallel_for!
+foo called from host
+Now I will crash if we compiled for CUDA/HIP
+cudaDeviceSynchronize() error( cudaErrorIllegalAddress): an illegal memory access was encountered /home/crtrott/Kokkos/kokkos/core/src/Cuda/Kokkos_Cuda_Instance.cpp:153


Suggested change

cudaDeviceSynchronize() error( cudaErrorIllegalAddress): an illegal memory access was encountered /home/crtrott/Kokkos/kokkos/core/src/Cuda/Kokkos_Cuda_Instance.cpp:153

cudaDeviceSynchronize() error( cudaErrorIllegalAddress): an illegal memory access was encountered /kokkos/core/src/Cuda/Kokkos_Cuda_Instance.cpp:153

JBludau · 2024-09-04T08:15:45Z

+};
+```
+
+This class is templated on the function pointer type, contains a pointer for both the host and the device version, and has a templated operator that forwards all the arguments, and calls the appropriate function pointer depending on call site.


Suggested change

This class is templated on the function pointer type, contains a pointer for both the host and the device version, and has a templated operator that forwards all the arguments, and calls the appropriate function pointer depending on call site.

This class is templated on the function pointer type, contains a pointer for both the host and the device version, and has a templated operator that copies all the arguments, and calls the appropriate function pointer depending on the call site.

I find it kind of important to make the distinction here as it adds another drawback.

JBludau · 2024-09-04T08:17:32Z

+};
+```
+
+However we still need to initialize the two function pointers, using the device initialization approach from before, and also initializing the host side. Note: the order of initialization and `deep_copy` matters, because `deep_copy` will overwrite both member poitners in our wrapper class:


Suggested change

However we still need to initialize the two function pointers, using the device initialization approach from before, and also initializing the host side. Note: the order of initialization and `deep_copy` matters, because `deep_copy` will overwrite both member poitners in our wrapper class:

However we still need to initialize the two function pointers using the device initialization approach from before, and also initializing the host side. Note: the order of initialization and `deep_copy` matters, because `deep_copy` will overwrite both member poitners in our wrapper class:

JBludau · 2024-09-04T08:19:45Z

+```
+
+While this generally works (and note this code will also work if you simply compile for host-only) there may be architectures in the future where this fails.
+Also we in principle strongly discourage the use of function pointers - the amount of possible problems you run into is fairly significant.


Suggested change

Also we in principle strongly discourage the use of function pointers - the amount of possible problems you run into is fairly significant.

Also we strongly discourage the use of function pointers - the amount of possible problems you run into is fairly significant.

I would just make that unconditional

masterleinad · 2024-09-04T11:39:40Z

+  KOKKOS_IF_ON_HOST(printf("foo called from host\n");)
+  KOKKOS_IF_ON_DEVICE(printf("foo called from device\n");)


Suggested change

KOKKOS_IF_ON_HOST(printf("foo called from host\n");)

KOKKOS_IF_ON_DEVICE(printf("foo called from device\n");)

KOKKOS_IF_ON_HOST(Kokkos::printf("foo called from host\n");)

KOKKOS_IF_ON_DEVICE(Kokkos::printf("foo called from device\n");)

masterleinad · 2024-09-04T11:40:05Z

+  KOKKOS_IF_ON_HOST(printf("foo called from host\n");)
+  KOKKOS_IF_ON_DEVICE(printf("foo called from device\n");)


Suggested change

KOKKOS_IF_ON_HOST(printf("foo called from host\n");)

KOKKOS_IF_ON_DEVICE(printf("foo called from device\n");)

KOKKOS_IF_ON_HOST(Kokkos::printf("foo called from host\n");)

KOKKOS_IF_ON_DEVICE(Kokkos::printf("foo called from device\n");)

masterleinad · 2024-09-04T11:43:43Z

+Aborted (core dumped)
+```
+
+*The function pointer we created was for a host function: you can not use it on device same as you can't dereference data pointers to host data!*


Suggested change

*The function pointer we created was for a host function: you can not use it on device same as you can't dereference data pointers to host data!*

*The function pointer we created was for a host function: you can not use it on the device just like you can't dereference data pointers to host data!*

masterleinad · 2024-09-04T11:46:47Z

+    Kokkos::deep_copy(A, A_v);
+```
+
+We are leveraging here the fact that `deep_copy` allows you to copy from and to a host scalar value, from a `View` of rank-0.


Suggested change

We are leveraging here the fact that `deep_copy` allows you to copy from and to a host scalar value, from a `View` of rank-0.

We are leveraging here the fact that `deep_copy` allows you to copy in both direction between a host scalar value and a `View` of rank-0.

masterleinad · 2024-09-04T11:46:55Z

+
+We are leveraging here the fact that `deep_copy` allows you to copy from and to a host scalar value, from a `View` of rank-0.
+
+If we do this `A` will contain a function pointer to a device function. That means we can capture `A` into a parallel region and execute it on the device, but now it will crash on the host.


Suggested change

If we do this `A` will contain a function pointer to a device function. That means we can capture `A` into a parallel region and execute it on the device, but now it will crash on the host.

If we do this, `A` will contain a function pointer to a device function. That means we can capture `A` into a parallel region and execute it on the device, but now it will crash on the host.

masterleinad · 2024-09-04T11:47:59Z

+  KOKKOS_IF_ON_HOST(printf("foo called from host\n");)
+  KOKKOS_IF_ON_DEVICE(printf("foo called from device\n");)


Suggested change

KOKKOS_IF_ON_HOST(printf("foo called from host\n");)

KOKKOS_IF_ON_DEVICE(printf("foo called from device\n");)

KOKKOS_IF_ON_HOST(Kokkos::printf("foo called from host\n");)

KOKKOS_IF_ON_DEVICE(Kokkos::printf("foo called from device\n");)

masterleinad · 2024-09-04T11:49:05Z

+
+## Creating a dual function pointer object
+
+To do better than this we need something which contains both the device function pointer and the host function pointer.


Suggested change

To do better than this we need something which contains both the device function pointer and the host function pointer.

To do better than this, we need something which contains both the device function pointer and the host function pointer.

masterleinad · 2024-09-04T11:50:05Z

+};
+```
+
+However we still need to initialize the two function pointers, using the device initialization approach from before, and also initializing the host side. Note: the order of initialization and `deep_copy` matters, because `deep_copy` will overwrite both member poitners in our wrapper class:


Suggested change

However we still need to initialize the two function pointers, using the device initialization approach from before, and also initializing the host side. Note: the order of initialization and `deep_copy` matters, because `deep_copy` will overwrite both member poitners in our wrapper class:

However, we still need to initialize the two function pointers, using the device initialization approach from before, and also initializing the host side. Note: the order of initialization and `deep_copy` matters, because `deep_copy` will overwrite both member poitners in our wrapper class:

masterleinad · 2024-09-04T11:51:46Z

+}
+```
+
+While this generally works (and note this code will also work if you simply compile for host-only) there may be architectures in the future where this fails.


Suggested change

While this generally works (and note this code will also work if you simply compile for host-only) there may be architectures in the future where this fails.

While this generally works (and note this code will also work if you simply compile for host-only), there may be architectures in the future where this fails. SYCL (and therefore the SYCL backend) doesn't allow using function pointers on the device for example.

Add a usecase discussion of how to deal with function pointers

dbbe563

crtrott force-pushed the function-pointers branch from 38b379e to dbbe563 Compare September 3, 2024 22:56

JBludau reviewed Sep 4, 2024

View reviewed changes

masterleinad reviewed Sep 4, 2024

View reviewed changes

	cudaDeviceSynchronize() error( cudaErrorIllegalAddress): an illegal memory access was encountered /home/crtrott/Kokkos/kokkos/core/src/Cuda/Kokkos_Cuda_Instance.cpp:153
	cudaDeviceSynchronize() error( cudaErrorIllegalAddress): an illegal memory access was encountered /kokkos/core/src/Cuda/Kokkos_Cuda_Instance.cpp:153

	This class is templated on the function pointer type, contains a pointer for both the host and the device version, and has a templated operator that forwards all the arguments, and calls the appropriate function pointer depending on call site.
	This class is templated on the function pointer type, contains a pointer for both the host and the device version, and has a templated operator that copies all the arguments, and calls the appropriate function pointer depending on the call site.

	However we still need to initialize the two function pointers, using the device initialization approach from before, and also initializing the host side. Note: the order of initialization and `deep_copy` matters, because `deep_copy` will overwrite both member poitners in our wrapper class:
	However we still need to initialize the two function pointers using the device initialization approach from before, and also initializing the host side. Note: the order of initialization and `deep_copy` matters, because `deep_copy` will overwrite both member poitners in our wrapper class:

	Also we in principle strongly discourage the use of function pointers - the amount of possible problems you run into is fairly significant.
	Also we strongly discourage the use of function pointers - the amount of possible problems you run into is fairly significant.

		KOKKOS_IF_ON_HOST(printf("foo called from host\n");)
		KOKKOS_IF_ON_DEVICE(printf("foo called from device\n");)

	The function pointer we created was for a host function: you can not use it on device same as you can't dereference data pointers to host data!
	The function pointer we created was for a host function: you can not use it on the device just like you can't dereference data pointers to host data!

	We are leveraging here the fact that `deep_copy` allows you to copy from and to a host scalar value, from a `View` of rank-0.
	We are leveraging here the fact that `deep_copy` allows you to copy in both direction between a host scalar value and a `View` of rank-0.


		We are leveraging here the fact that `deep_copy` allows you to copy from and to a host scalar value, from a `View` of rank-0.

		If we do this `A` will contain a function pointer to a device function. That means we can capture `A` into a parallel region and execute it on the device, but now it will crash on the host.

	If we do this `A` will contain a function pointer to a device function. That means we can capture `A` into a parallel region and execute it on the device, but now it will crash on the host.
	If we do this, `A` will contain a function pointer to a device function. That means we can capture `A` into a parallel region and execute it on the device, but now it will crash on the host.


		## Creating a dual function pointer object

		To do better than this we need something which contains both the device function pointer and the host function pointer.

	To do better than this we need something which contains both the device function pointer and the host function pointer.
	To do better than this, we need something which contains both the device function pointer and the host function pointer.

	However we still need to initialize the two function pointers, using the device initialization approach from before, and also initializing the host side. Note: the order of initialization and `deep_copy` matters, because `deep_copy` will overwrite both member poitners in our wrapper class:
	However, we still need to initialize the two function pointers, using the device initialization approach from before, and also initializing the host side. Note: the order of initialization and `deep_copy` matters, because `deep_copy` will overwrite both member poitners in our wrapper class:

	While this generally works (and note this code will also work if you simply compile for host-only) there may be architectures in the future where this fails.
	While this generally works (and note this code will also work if you simply compile for host-only), there may be architectures in the future where this fails. SYCL (and therefore the SYCL backend) doesn't allow using function pointers on the device for example.

Conversation

crtrott commented Sep 3, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

JBludau left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

crtrott commented Sep 3, 2024 •

edited

Loading