1//===-- spirvintrin.h - SPIR-V intrinsic functions ------------------------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8
9#ifndef __SPIRVINTRIN_H
10#define __SPIRVINTRIN_H
11
12#ifndef __SPIRV__
13#error "This file is intended for SPIR-V targets or offloading to SPIR-V"
14#endif
15
16#ifndef __GPUINTRIN_H
17#error "Never use <spirvintrin.h> directly; include <gpuintrin.h> instead"
18#endif
19
20_Pragma("omp begin declare target device_type(nohost)");
21_Pragma("omp begin declare variant match(device = {arch(spirv64)})");
22
23// Type aliases to the address spaces used by the SPIR-V backend.
24#define __gpu_private __attribute__((address_space(0)))
25#define __gpu_constant __attribute__((address_space(2)))
26#define __gpu_local __attribute__((address_space(3)))
27#define __gpu_global __attribute__((address_space(1)))
28#define __gpu_generic __attribute__((address_space(4)))
29
30// Attribute to declare a function as a kernel.
31#define __gpu_kernel __attribute__((device_kernel, visibility("protected")))
32
33// Returns the number of workgroups in the 'x' dimension of the grid.
34_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_blocks_x(void) {
35 return __builtin_spirv_num_workgroups(0);
36}
37
38// Returns the number of workgroups in the 'y' dimension of the grid.
39_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_blocks_y(void) {
40 return __builtin_spirv_num_workgroups(1);
41}
42
43// Returns the number of workgroups in the 'z' dimension of the grid.
44_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_blocks_z(void) {
45 return __builtin_spirv_num_workgroups(2);
46}
47
48// Returns the 'x' dimension of the current workgroup's id.
49_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_block_id_x(void) {
50 return __builtin_spirv_workgroup_id(0);
51}
52
53// Returns the 'y' dimension of the current workgroup's id.
54_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_block_id_y(void) {
55 return __builtin_spirv_workgroup_id(1);
56}
57
58// Returns the 'z' dimension of the current workgroup's id.
59_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_block_id_z(void) {
60 return __builtin_spirv_workgroup_id(2);
61}
62
63// Returns the number of workitems in the 'x' dimension.
64_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_threads_x(void) {
65 return __builtin_spirv_workgroup_size(0);
66}
67
68// Returns the number of workitems in the 'y' dimension.
69_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_threads_y(void) {
70 return __builtin_spirv_workgroup_size(1);
71}
72
73// Returns the number of workitems in the 'z' dimension.
74_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_threads_z(void) {
75 return __builtin_spirv_workgroup_size(2);
76}
77
78// Returns the 'x' dimension id of the workitem in the current workgroup.
79_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_thread_id_x(void) {
80 return __builtin_spirv_local_invocation_id(0);
81}
82
83// Returns the 'y' dimension id of the workitem in the current workgroup.
84_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_thread_id_y(void) {
85 return __builtin_spirv_local_invocation_id(1);
86}
87
88// Returns the 'z' dimension id of the workitem in the current workgroup.
89_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_thread_id_z(void) {
90 return __builtin_spirv_local_invocation_id(2);
91}
92
93// Returns the size of an wavefront, either 32 or 64 depending on hardware
94// and compilation options.
95_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_lanes(void) {
96 return __builtin_spirv_subgroup_size();
97}
98
99// Returns the id of the thread inside of an wavefront executing together.
100_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_lane_id(void) {
101 return __builtin_spirv_subgroup_local_invocation_id();
102}
103
104// Returns the bit-mask of active threads in the current wavefront. This
105// implementation is incorrect if the target uses more than 64 lanes.
106_DEFAULT_FN_ATTRS static __inline__ uint64_t __gpu_lane_mask(void) {
107 uint32_t [[clang::ext_vector_type(4)]] __mask =
108 __builtin_spirv_subgroup_ballot(1);
109 return __builtin_bit_cast(uint64_t,
110 __builtin_shufflevector(__mask, __mask, 0, 1));
111}
112
113// Copies the value from the first active thread in the wavefront to the rest.
114_DEFAULT_FN_ATTRS static __inline__ uint32_t
115__gpu_read_first_lane_u32(uint64_t __lane_mask, uint32_t __x) {
116 return __builtin_spirv_subgroup_shuffle(__x,
117 __builtin_ctzg(__gpu_lane_mask()));
118}
119
120// Returns a bitmask of threads in the current lane for which \p x is true. This
121// implementation is incorrect if the target uses more than 64 lanes.
122_DEFAULT_FN_ATTRS static __inline__ uint64_t __gpu_ballot(uint64_t __lane_mask,
123 bool __x) {
124 // The lane_mask & gives the nvptx semantics when lane_mask is a subset of
125 // the active threads.
126 uint32_t [[clang::ext_vector_type(4)]] __mask =
127 __builtin_spirv_subgroup_ballot(__x);
128 return __lane_mask & __builtin_bit_cast(uint64_t, __builtin_shufflevector(
129 __mask, __mask, 0, 1));
130}
131
132// Waits for all the threads in the block to converge and issues a fence.
133_DEFAULT_FN_ATTRS static __inline__ void __gpu_sync_threads(void) {
134 __builtin_spirv_group_barrier();
135}
136
137// Wait for all threads in the wavefront to converge, this is a noop on SPIR-V.
138_DEFAULT_FN_ATTRS static __inline__ void __gpu_sync_lane(uint64_t __lane_mask) {
139}
140
141// Shuffles the the lanes inside the wavefront according to the given index.
142_DEFAULT_FN_ATTRS static __inline__ uint32_t
143__gpu_shuffle_idx_u32(uint64_t __lane_mask, uint32_t __idx, uint32_t __x,
144 uint32_t __width) {
145 uint32_t __lane = __idx + (__gpu_lane_id() & ~(__width - 1));
146 return __builtin_spirv_subgroup_shuffle(__x, __lane);
147}
148
149// Returns a bitmask marking all lanes that have the same value of __x.
150_DEFAULT_FN_ATTRS static __inline__ uint64_t
151__gpu_match_any_u32(uint64_t __lane_mask, uint32_t __x) {
152 return __gpu_match_any_u32_impl(__lane_mask, __x);
153}
154
155// Returns a bitmask marking all lanes that have the same value of __x.
156_DEFAULT_FN_ATTRS static __inline__ uint64_t
157__gpu_match_any_u64(uint64_t __lane_mask, uint64_t __x) {
158 return __gpu_match_any_u64_impl(__lane_mask, __x);
159}
160
161// Returns the current lane mask if every lane contains __x.
162_DEFAULT_FN_ATTRS static __inline__ uint64_t
163__gpu_match_all_u32(uint64_t __lane_mask, uint32_t __x) {
164 return __gpu_match_all_u32_impl(__lane_mask, __x);
165}
166
167// Returns the current lane mask if every lane contains __x.
168_DEFAULT_FN_ATTRS static __inline__ uint64_t
169__gpu_match_all_u64(uint64_t __lane_mask, uint64_t __x) {
170 return __gpu_match_all_u64_impl(__lane_mask, __x);
171}
172
173// SPIR-V does not expose this, always return false.
174_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_local(void *ptr) {
175 return 0;
176}
177
178// SPIR-V does not expose this, always return false.
179_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_private(void *ptr) {
180 return 0;
181}
182
183// SPIR-V only supports 'OpTerminateInvocation' in fragment shaders.
184_DEFAULT_FN_ATTRS [[noreturn]] static __inline__ void __gpu_exit(void) {
185 __builtin_trap();
186}
187
188// This is a no-op as SPIR-V does not support it.
189_DEFAULT_FN_ATTRS static __inline__ void __gpu_thread_suspend(void) {}
190
191_Pragma("omp end declare variant");
192_Pragma("omp end declare target");
193
194#endif // __SPIRVINTRIN_H