Following-up Vlad's comments

- Relocating 'plugin_printers.hpp' back
- Changing cstdint inclusion path

Author: dongkyunahn-intel

sycl/include/CL/sycl/INTEL/esimd/detail/emu/esimd_emu_functions_v1.h

@@ -11,7 +11,9 @@
 /// \ingroup sycl_pi_esimd_cpu
 
 #pragma once
-#include <cstdint>
+
+// <cstdint> for 'uint32_t' type is included in upper-level device
+// interface file ('esimdcpu_device_interface.hpp')
 
 // This file defines function interfaces for ESIMD CPU Emulation
 // (ESIMD_CPU) to access LibCM CPU emulation functionalities from

sycl/include/CL/sycl/INTEL/esimd/detail/emu/esimdcpu_device_interface.hpp

@@ -20,6 +20,11 @@
 
 #include <CL/sycl/detail/pi.hpp>
 
+// cstdint-type fields such as 'uint32_t' are to be used in funtion
+// pointer table file ('esimd_emu_functions_v1.h') included in 'struct
+// ESIMDDeviceInterface' definition.
+#include <cstdint>
+
 __SYCL_INLINE_NAMESPACE(cl) {
   namespace sycl {
   namespace detail {

sycl/include/CL/sycl/INTEL/esimd/detail/esimd_memory_intrin.hpp

@@ -777,19 +777,22 @@ __esimd_dp4(sycl::INTEL::gpu::vector_type_t<Ty, N> v1,
 }
 
 inline void __esimd_slm_init(size_t size) {
-  ESIMDDeviceInterface *I = getESIMDDeviceInterface();
+  sycl::detail::ESIMDDeviceInterface *I =
+      sycl::detail::getESIMDDeviceInterface();
 
   I->cm_slm_init_ptr(size);
 }
 
 inline void __esimd_barrier() {
-  ESIMDDeviceInterface *I = getESIMDDeviceInterface();
+  sycl::detail::ESIMDDeviceInterface *I =
+      sycl::detail::getESIMDDeviceInterface();
 
   I->cm_barrier_ptr();
 }
 
 inline void __esimd_sbarrier(sycl::INTEL::gpu::EsimdSbarrierType flag) {
-  ESIMDDeviceInterface *I = getESIMDDeviceInterface();
+  sycl::detail::ESIMDDeviceInterface *I =
+      sycl::detail::getESIMDDeviceInterface();
 
   I->cm_sbarrier_ptr((uint)flag);
 }
@@ -801,7 +804,8 @@ inline sycl::INTEL::gpu::vector_type_t<Ty, N>
 __esimd_slm_read(sycl::INTEL::gpu::vector_type_t<uint32_t, N> addrs,
                  sycl::INTEL::gpu::vector_type_t<uint16_t, N> pred) {
   sycl::INTEL::gpu::vector_type_t<Ty, N> retv;
-  ESIMDDeviceInterface *I = getESIMDDeviceInterface();
+  sycl::detail::ESIMDDeviceInterface *I =
+      sycl::detail::getESIMDDeviceInterface();
 
   char *SlmBase = I->__cm_emu_get_slm_ptr();
 
@@ -816,7 +820,8 @@ inline void
 __esimd_slm_write(sycl::INTEL::gpu::vector_type_t<uint32_t, N> addrs,
                   sycl::INTEL::gpu::vector_type_t<Ty, N> vals,
                   sycl::INTEL::gpu::vector_type_t<uint16_t, N> pred) {
-  ESIMDDeviceInterface *I = getESIMDDeviceInterface();
+  sycl::detail::ESIMDDeviceInterface *I =
+      sycl::detail::getESIMDDeviceInterface();
 
   char *SlmBase = I->__cm_emu_get_slm_ptr();
 
@@ -828,7 +833,8 @@ template <typename Ty, int N>
 inline sycl::INTEL::gpu::vector_type_t<Ty, N>
 __esimd_slm_block_read(uint32_t addr) {
   sycl::INTEL::gpu::vector_type_t<Ty, N> retv;
-  ESIMDDeviceInterface *I = getESIMDDeviceInterface();
+  sycl::detail::ESIMDDeviceInterface *I =
+      sycl::detail::getESIMDDeviceInterface();
 
   char *SlmBase = I->__cm_emu_get_slm_ptr();
 
@@ -842,7 +848,8 @@ template <typename Ty, int N>
 inline void
 __esimd_slm_block_write(uint32_t addr,
                         sycl::INTEL::gpu::vector_type_t<Ty, N> vals) {
-  ESIMDDeviceInterface *I = getESIMDDeviceInterface();
+  sycl::detail::ESIMDDeviceInterface *I =
+      sycl::detail::getESIMDDeviceInterface();
 
   char *SlmBase = I->__cm_emu_get_slm_ptr();
 
@@ -855,7 +862,8 @@ inline sycl::INTEL::gpu::vector_type_t<Ty, N * NumChannels(Mask)>
 __esimd_slm_read4(sycl::INTEL::gpu::vector_type_t<uint32_t, N> addrs,
                   sycl::INTEL::gpu::vector_type_t<uint16_t, N> pred) {
   sycl::INTEL::gpu::vector_type_t<Ty, N * NumChannels(Mask)> retv;
-  ESIMDDeviceInterface *I = getESIMDDeviceInterface();
+  sycl::detail::ESIMDDeviceInterface *I =
+      sycl::detail::getESIMDDeviceInterface();
 
   char *ReadBase = I->__cm_emu_get_slm_ptr();
 
@@ -869,7 +877,8 @@ inline void __esimd_slm_write4(
     sycl::INTEL::gpu::vector_type_t<uint32_t, N> addrs,
     sycl::INTEL::gpu::vector_type_t<Ty, N * NumChannels(Mask)> vals,
     sycl::INTEL::gpu::vector_type_t<uint16_t, N> pred) {
-  ESIMDDeviceInterface *I = getESIMDDeviceInterface();
+  sycl::detail::ESIMDDeviceInterface *I =
+      sycl::detail::getESIMDDeviceInterface();
 
   char *WriteBase = I->__cm_emu_get_slm_ptr();
 

sycl/source/detail/pi.cpp

@@ -6,8 +6,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-/// \file pi.cpp/
-// Implementation of C++ wrappers for PI interface.
+/// \file pi.cpp
+/// Implementation of C++ wrappers for PI interface.
 ///
 /// \ingroup sycl_pi
 
@@ -60,7 +60,7 @@ void *getPluginOpaqueData(void *OpaqueDataParam) {
       cl::sycl::detail::pi::getPlugin<BE>();
 
   Plugin.call<cl::sycl::detail::PiApiKind::piextPluginGetOpaqueData>(
-      opaquedata_arg, &ReturnOpaqueData);
+      OpaqueDataParam, &ReturnOpaqueData);
 
   return ReturnOpaqueData;
 }

sycl/source/detail/plugin.hpp

@@ -10,8 +10,8 @@
 #include <CL/sycl/backend_types.hpp>
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/pi.hpp>
-#include <CL/sycl/detail/plugin_printers.hpp>
 #include <CL/sycl/stl.hpp>
+#include <detail/plugin_printers.hpp>
 #include <memory>
 #include <mutex>