Skip to content

Fixes enabling hierarchical parallelism API. #239

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Jun 25, 2019
Merged

Fixes enabling hierarchical parallelism API. #239

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
7960931
[SYCL] Make item's and group's kernel initialization code reusable.
kbobrovs May 31, 2019
13a7ee5
[SYCL] Fix functor versions of cl::sycl::handler kernel invocation APIs.
kbobrovs May 31, 2019
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
45 changes: 45 additions & 0 deletions sycl/include/CL/__spirv/spirv_vars.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,13 +24,17 @@ extern "C" const __constant size_t_vec __spirv_BuiltInGlobalOffset;
template <> size_t get##POSTFIX<1>() { return __spirv_BuiltIn##POSTFIX.y; } \
template <> size_t get##POSTFIX<2>() { return __spirv_BuiltIn##POSTFIX.z; }

namespace __spirv {

DEFINE_INT_ID_TO_XYZ_CONVERTER(GlobalSize);
DEFINE_INT_ID_TO_XYZ_CONVERTER(GlobalInvocationId)
DEFINE_INT_ID_TO_XYZ_CONVERTER(WorkgroupSize)
DEFINE_INT_ID_TO_XYZ_CONVERTER(LocalInvocationId)
DEFINE_INT_ID_TO_XYZ_CONVERTER(WorkgroupId)
DEFINE_INT_ID_TO_XYZ_CONVERTER(GlobalOffset)

} // namespace __spirv

#undef DEFINE_INT_ID_TO_XYZ_CONVERTER

extern "C" const __constant uint32_t __spirv_BuiltInSubgroupSize;
Expand All @@ -40,4 +44,45 @@ extern "C" const __constant uint32_t __spirv_BuiltInNumEnqueuedSubgroups;
extern "C" const __constant uint32_t __spirv_BuiltInSubgroupId;
extern "C" const __constant uint32_t __spirv_BuiltInSubgroupLocalInvocationId;

#define DEFINE_INIT_SIZES(POSTFIX) \
\
template <int Dim, class DstT> struct InitSizesST##POSTFIX; \
\
template <class DstT> struct InitSizesST##POSTFIX<1, DstT> { \
static void initSize(DstT &Dst) { \
Dst[0] = get##POSTFIX<0>(); \
} \
}; \
\
template <class DstT> struct InitSizesST##POSTFIX<2, DstT> { \
static void initSize(DstT &Dst) { \
Dst[1] = get##POSTFIX<1>(); \
InitSizesST##POSTFIX<1, DstT>::initSize(Dst); \
} \
}; \
\
template <class DstT> struct InitSizesST##POSTFIX<3, DstT> { \
static void initSize(DstT &Dst) { \
Dst[2] = get##POSTFIX<2>(); \
InitSizesST##POSTFIX<2, DstT>::initSize(Dst); \
} \
}; \
\
template <int Dims, class DstT> static void init##POSTFIX(DstT &Dst) { \
InitSizesST##POSTFIX<Dims, DstT>::initSize(Dst); \
}

namespace __spirv {

DEFINE_INIT_SIZES(GlobalSize);
DEFINE_INIT_SIZES(GlobalInvocationId)
DEFINE_INIT_SIZES(WorkgroupSize)
DEFINE_INIT_SIZES(LocalInvocationId)
DEFINE_INIT_SIZES(WorkgroupId)
DEFINE_INIT_SIZES(GlobalOffset)

} // namespace __spirv

#undef DEFINE_INIT_SIZES

#endif // __SYCL_DEVICE_ONLY__
Loading