Skip to content

Commit e02e118

Browse files
author
Victor Lomuller
committed
[SYCL-PTX] Revert manual mangling of some SPIR-V builtins
Signed-off-by: Victor Lomuller <victor@codeplay.com>
1 parent a8482b1 commit e02e118

File tree

10 files changed

+43
-51
lines changed

10 files changed

+43
-51
lines changed

‎libclc/generic/include/func.h‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,5 +5,6 @@
55
#define _CLC_DECL
66
#define _CLC_DEF __attribute__((always_inline))
77
#define _CLC_INLINE __attribute__((always_inline)) inline
8+
#define _CLC_CONVERGENT __attribute__((convergent))
89

910
#endif // CLC_FUNC

‎libclc/generic/include/spirv/async/async_work_group_strided_copy.inc‎

Lines changed: 4 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -6,11 +6,7 @@
66
//
77
//===----------------------------------------------------------------------===//
88

9-
#define FN_NAME __CLC_XCONCAT(_Z22__spirv_GroupAsyncCopyI, __SPIRV_GENTYPE_MANGLED, E9ocl_eventN5__spv5ScopeEPU3, __SPIRV_DST_ADDR_SPACE_MANGLED, T_PU3, __SPIRV_SRC_ADDR_SPACE_MANGLED, S3_mmS0_)
10-
_CLC_DECL event_t FN_NAME (
11-
enum Scope scope,
12-
__SPIRV_DST_ADDR_SPACE __SPIRV_GENTYPE *dst,
13-
const __SPIRV_SRC_ADDR_SPACE __SPIRV_GENTYPE *src,
14-
size_t num_elements,
15-
size_t stride,
16-
event_t event);
9+
_CLC_OVERLOAD _CLC_DECL event_t __spirv_GroupAsyncCopy(
10+
unsigned int scope, __SPIRV_DST_ADDR_SPACE __SPIRV_GENTYPE *dst,
11+
const __SPIRV_SRC_ADDR_SPACE __SPIRV_GENTYPE *src, size_t num_elements,
12+
size_t stride, event_t event);

‎libclc/generic/include/spirv/async/wait_group_events.h‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,6 @@
66
//
77
//===----------------------------------------------------------------------===//
88

9-
// TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling.
10-
_CLC_DEF void _Z23__spirv_GroupWaitEventsN5__spv5ScopeEjP9ocl_event(
11-
enum Scope scope, int num_events, event_t *event_list);
9+
_CLC_OVERLOAD _CLC_DEF void __spirv_GroupWaitEvents(unsigned int scope,
10+
int num_events,
11+
event_t *event_list);

‎libclc/generic/include/spirv/synchronization/barrier.h‎

Lines changed: 5 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,8 @@
66
//
77
//===----------------------------------------------------------------------===//
88

9-
// TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling.
10-
_CLC_DECL void _Z22__spirv_ControlBarrierN5__spv5ScopeES0_j(enum Scope scope, enum Scope memory, unsigned int semantics);
11-
_CLC_DECL void _Z21__spirv_MemoryBarrierN5__spv5ScopeEj(enum Scope scope, unsigned int semantics);
9+
_CLC_OVERLOAD _CLC_DECL void __spirv_ControlBarrier(unsigned int scope,
10+
unsigned int memory,
11+
unsigned int semantics);
12+
_CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(unsigned int scope,
13+
unsigned int semantics);

‎libclc/generic/lib/async/async_work_group_strided_copy.inc‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,8 @@ _CLC_OVERLOAD _CLC_DEF event_t async_work_group_strided_copy(
88
size_t stride,
99
event_t event) {
1010

11-
return __CLC_XCONCAT(_Z22__spirv_GroupAsyncCopyI, __CLC_GENTYPE_MANGLED, E9ocl_eventN5__spv5ScopeEPU3AS1T_PU3AS3S3_mmS0_)(Workgroup, dst, src, num_gentypes, stride, event);
11+
return __spirv_GroupAsyncCopy(Workgroup, dst, src, num_gentypes, stride,
12+
event);
1213
}
1314

1415
_CLC_OVERLOAD _CLC_DEF event_t async_work_group_strided_copy(
@@ -17,7 +18,8 @@ _CLC_OVERLOAD _CLC_DEF event_t async_work_group_strided_copy(
1718
size_t num_gentypes,
1819
size_t stride,
1920
event_t event) {
20-
return __CLC_XCONCAT(_Z22__spirv_GroupAsyncCopyI, __CLC_GENTYPE_MANGLED, E9ocl_eventN5__spv5ScopeEPU3AS3T_PU3AS1S3_mmS0_)(Workgroup, dst, src, num_gentypes, stride, event);
21+
return __spirv_GroupAsyncCopy(Workgroup, dst, src, num_gentypes, stride,
22+
event);
2123
}
2224

2325
#undef __CLC_XCONCAT

‎libclc/generic/lib/async/wait_group_events.cl‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,5 +2,5 @@
22
#include <clc/clc.h>
33

44
_CLC_DEF void wait_group_events(int num_events, event_t *event_list) {
5-
_Z23__spirv_GroupWaitEventsN5__spv5ScopeEjP9ocl_event(Workgroup, num_events, event_list);
5+
__spirv_GroupWaitEvents(Workgroup, num_events, event_list);
66
}

‎libclc/generic/lib/synchronization/barrier.cl‎

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -10,8 +10,9 @@
1010
#include <clc/clc.h>
1111

1212
_CLC_DEF void barrier(cl_mem_fence_flags flags) {
13-
unsigned int mem_semantic = (flag & CLK_GLOBAL_MEM_FENCE ? 0x200 : 0) |
14-
(flag & CLK_LOCAL_MEM_FENCE ? 0x100 : 0)
15-
// TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling.
16-
_Z22__spirv_ControlBarrierN5__spv5ScopeES0_j(Workgroup, Workgroup, mem_semantic);
13+
unsigned int mem_semantic =
14+
SequentiallyConsistent |
15+
(flag & CLK_GLOBAL_MEM_FENCE ? CrossWorkgroupMemory : 0) |
16+
(flag & CLK_LOCAL_MEM_FENCE ? WorkgroupMemory : 0)
17+
__spirv_ControlBarrier(Workgroup, Workgroup, SequentiallyConsistent);
1718
}

‎libclc/generic/libspirv/async/async_work_group_strided_copy.inc‎

Lines changed: 8 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -21,30 +21,18 @@
2121
dst[i * DST_STRIDE] = src[i * SRC_STRIDE]; \
2222
}
2323

24-
#define __CLC_CONCAT(a, b, c) a ## b ## c
25-
#define __CLC_XCONCAT(a, b, c) __CLC_CONCAT(a, b, c)
26-
27-
_CLC_DEF event_t __CLC_XCONCAT(_Z22__spirv_GroupAsyncCopyI, __CLC_GENTYPE_MANGLED, E9ocl_eventN5__spv5ScopeEPU3AS1T_PU3AS3S3_mmS0_) (
28-
enum Scope scope,
29-
global __CLC_GENTYPE *dst,
30-
const local __CLC_GENTYPE *src,
31-
size_t num_gentypes,
32-
size_t stride,
33-
event_t event) {
24+
_CLC_OVERLOAD _CLC_DEF event_t
25+
__spirv_GroupAsyncCopy(unsigned int scope, global __CLC_GENTYPE *dst,
26+
const local __CLC_GENTYPE *src, size_t num_gentypes,
27+
size_t stride, event_t event) {
3428
STRIDED_COPY(global, local, stride, 1);
3529
return event;
3630
}
3731

38-
_CLC_DEF event_t __CLC_XCONCAT(_Z22__spirv_GroupAsyncCopyI, __CLC_GENTYPE_MANGLED, E9ocl_eventN5__spv5ScopeEPU3AS3T_PU3AS1S3_mmS0_) (
39-
enum Scope scope,
40-
local __CLC_GENTYPE *dst,
41-
const global __CLC_GENTYPE *src,
42-
size_t num_gentypes,
43-
size_t stride,
44-
event_t event) {
32+
_CLC_OVERLOAD _CLC_DEF event_t
33+
__spirv_GroupAsyncCopy(unsigned int scope, local __CLC_GENTYPE *dst,
34+
const global __CLC_GENTYPE *src, size_t num_gentypes,
35+
size_t stride, event_t event) {
4536
STRIDED_COPY(local, global, 1, stride);
4637
return event;
4738
}
48-
49-
#undef __CLC_XCONCAT
50-
#undef __CLC_CONCAT

‎libclc/generic/libspirv/async/wait_group_events.cl‎

Lines changed: 4 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -8,9 +8,8 @@
88

99
#include <spirv/spirv.h>
1010

11-
// TODO: Stop manually mangling these names. Need C++ namespaces to get the
12-
// exact mangling.
13-
_CLC_DEF void _Z23__spirv_GroupWaitEventsN5__spv5ScopeEjP9ocl_event(
14-
enum Scope scope, int num_events, event_t *event_list) {
15-
_Z22__spirv_ControlBarrierN5__spv5ScopeES0_j(scope, Workgroup, 0x200 | 0x100);
11+
_CLC_OVERLOAD _CLC_DEF void __spirv_GroupWaitEvents(unsigned int scope,
12+
int num_events,
13+
event_t *event_list) {
14+
__spirv_ControlBarrier(scope, Workgroup, SequentiallyConsistent);
1615
}

‎libclc/ptx-nvidiacl/libspirv/synchronization/barrier.cl‎

Lines changed: 8 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -8,11 +8,14 @@
88

99
#include <spirv/spirv.h>
1010

11-
// TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling.
12-
_CLC_DEF void _Z22__spirv_ControlBarrierN5__spv5ScopeES0_j(enum Scope scope, enum Scope memory, unsigned int semantics) {
13-
__syncthreads();
11+
_CLC_OVERLOAD _CLC_DEF void __spirv_MemoryBarrier(unsigned int memory,
12+
unsigned int semantics) {
13+
__nvvm_membar_cta();
1414
}
1515

16-
// TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling.
17-
_CLC_DEF void _Z21__spirv_MemoryBarrierN5__spv5ScopeEj(enum Scope scope, unsigned int semantics) {
16+
_CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT void
17+
__spirv_ControlBarrier(unsigned int scope, unsigned int memory,
18+
unsigned int semantics) {
19+
__syncthreads();
20+
__spirv_MemoryBarrier(memory, semantics);
1821
}

0 commit comments

Comments
 (0)