Skip to content

Commit b8f35cf

Browse files
elizabethandrewselizabethandrews
authored
[SYCL] Fix param kind for annotated types (#7116)
Annotated types were generated with param kind `kind_sampler ` in integration header. This PR sets the param kind based on the type of the first parameter of the `__init()` function - either as `kind_std_pointer` or `kind_std_layout`. Signed-off-by: Elizabeth Andrews <[email protected]>
1 parent 7f1a6ef commit b8f35cf

File tree

10 files changed

+90
-47
lines changed

10 files changed

+90
-47
lines changed

‎clang/include/clang/Basic/Attr.td

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1302,11 +1302,13 @@ def SYCLType: InheritableAttr {
13021302
["accessor", "local_accessor", "spec_constant",
13031303
"specialization_id", "kernel_handler", "buffer_location",
13041304
"no_alias", "accessor_property_list", "group",
1305-
"private_memory", "aspect", "annotated_ptr", "annotated_arg"],
1305+
"private_memory", "aspect", "annotated_ptr", "annotated_arg",
1306+
"stream", "sampler"],
13061307
["accessor", "local_accessor", "spec_constant",
13071308
"specialization_id", "kernel_handler", "buffer_location",
13081309
"no_alias", "accessor_property_list", "group",
1309-
"private_memory", "aspect", "annotated_ptr", "annotated_arg"]>];
1310+
"private_memory", "aspect", "annotated_ptr", "annotated_arg",
1311+
"stream", "sampler"]>];
13101312
// Only used internally by SYCL implementation
13111313
let Documentation = [InternalOnly];
13121314
}

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 18 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -3616,17 +3616,25 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
36163616

36173617
Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,
36183618
CurOffset + offsetOf(FD, FieldTy));
3619+
} else if (isSyclType(FieldTy, SYCLTypeAttr::stream)) {
3620+
addParam(FD, FieldTy, SYCLIntegrationHeader::kind_stream);
3621+
} else if (isSyclType(FieldTy, SYCLTypeAttr::sampler) ||
3622+
isSyclType(FieldTy, SYCLTypeAttr::annotated_ptr) ||
3623+
isSyclType(FieldTy, SYCLTypeAttr::annotated_arg)) {
3624+
CXXMethodDecl *InitMethod = getMethodByName(ClassTy, InitMethodName);
3625+
assert(InitMethod && "type must have __init method");
3626+
const ParmVarDecl *InitArg = InitMethod->getParamDecl(0);
3627+
assert(InitArg && "Init method must have arguments");
3628+
QualType T = InitArg->getType();
3629+
SYCLIntegrationHeader::kernel_param_kind_t ParamKind =
3630+
isSyclType(FieldTy, SYCLTypeAttr::sampler)
3631+
? SYCLIntegrationHeader::kind_sampler
3632+
: (T->isPointerType() ? SYCLIntegrationHeader::kind_pointer
3633+
: SYCLIntegrationHeader::kind_std_layout);
3634+
addParam(T, ParamKind, offsetOf(FD, FieldTy));
36193635
} else {
3620-
if (getMethodByName(ClassTy, FinalizeMethodName))
3621-
addParam(FD, FieldTy, SYCLIntegrationHeader::kind_stream);
3622-
else {
3623-
CXXMethodDecl *InitMethod = getMethodByName(ClassTy, InitMethodName);
3624-
assert(InitMethod && "type must have __init method");
3625-
const ParmVarDecl *SamplerArg = InitMethod->getParamDecl(0);
3626-
assert(SamplerArg && "Init method must have arguments");
3627-
addParam(SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler,
3628-
offsetOf(FD, FieldTy));
3629-
}
3636+
llvm_unreachable(
3637+
"Unexpected SYCL special class when generating integration header");
36303638
}
36313639
return true;
36323640
}

‎clang/test/CodeGenSYCL/Inputs/sycl.hpp

Lines changed: 18 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,7 @@ struct sampler_impl {
1919
#endif
2020
};
2121

22-
class __attribute__((sycl_special_class)) sampler {
22+
class __attribute__((sycl_special_class)) __SYCL_TYPE(sampler) sampler {
2323
struct sampler_impl impl;
2424
#ifdef __SYCL_DEVICE_ONLY__
2525
void __init(__ocl_sampler_t Sampler) { impl.m_Sampler = Sampler; }
@@ -372,6 +372,22 @@ int printf(const __SYCL_CONSTANT_AS char *__format, Args... args) {
372372
#endif // defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__)
373373
}
374374

375+
template <typename T, typename... Props>
376+
class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) annotated_arg {
377+
T obj;
378+
#ifdef __SYCL_DEVICE_ONLY__
379+
void __init(T _obj) {}
380+
#endif
381+
};
382+
383+
template <typename T, typename... Props>
384+
class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_ptr) annotated_ptr {
385+
T* obj;
386+
#ifdef __SYCL_DEVICE_ONLY__
387+
void __init(T* _obj) {}
388+
#endif
389+
};
390+
375391
} // namespace experimental
376392
} // namespace oneapi
377393
} // namespace ext
@@ -494,7 +510,7 @@ class handler {
494510
}
495511
};
496512

497-
class __attribute__((sycl_special_class)) stream {
513+
class __attribute__((sycl_special_class)) __SYCL_TYPE(stream) stream {
498514
public:
499515
stream(unsigned long BufferSize, unsigned long MaxStatementSize,
500516
handler &CGH) {}

‎clang/test/CodeGenSYCL/add_ir_attributes_kernel_parameter.cpp

Lines changed: 14 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,7 @@
1111

1212
// One __init parameter with add_ir_attributes_kernel_parameter attribute.
1313

14-
template <typename... Properties> class __attribute__((sycl_special_class)) g {
14+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) g {
1515
public:
1616
int *x;
1717

@@ -27,7 +27,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) g {
2727
#endif
2828
};
2929

30-
class __attribute__((sycl_special_class)) h {
30+
class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) h {
3131
public:
3232
int *x;
3333

@@ -44,7 +44,7 @@ class __attribute__((sycl_special_class)) h {
4444
#endif
4545
};
4646

47-
template <typename... Properties> class __attribute__((sycl_special_class)) gh {
47+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) gh {
4848
public:
4949
int *x;
5050

@@ -61,7 +61,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) gh {
6161
#endif
6262
};
6363

64-
template <typename... Properties> class __attribute__((sycl_special_class)) hg {
64+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) hg {
6565
public:
6666
int *x;
6767

@@ -80,7 +80,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) hg {
8080

8181
// Two __init parameters, one with add_ir_attributes_kernel_parameter attribute.
8282

83-
template <typename... Properties> class __attribute__((sycl_special_class)) k {
83+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) k {
8484
public:
8585
int *x;
8686
float *y;
@@ -99,7 +99,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) k {
9999
#endif
100100
};
101101

102-
class __attribute__((sycl_special_class)) l {
102+
class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) l {
103103
public:
104104
int *x;
105105
float *y;
@@ -119,7 +119,7 @@ class __attribute__((sycl_special_class)) l {
119119
#endif
120120
};
121121

122-
template <typename... Properties> class __attribute__((sycl_special_class)) kl {
122+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) kl {
123123
public:
124124
int *x;
125125
float *y;
@@ -139,7 +139,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) kl {
139139
#endif
140140
};
141141

142-
template <typename... Properties> class __attribute__((sycl_special_class)) lk {
142+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) lk {
143143
public:
144144
int *x;
145145
float *y;
@@ -161,7 +161,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) lk {
161161

162162
// Two __init parameters, both with add_ir_attributes_kernel_parameter attribute.
163163

164-
template <typename... Properties> class __attribute__((sycl_special_class)) m {
164+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) m {
165165
public:
166166
int *x;
167167
float *y;
@@ -181,7 +181,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) m {
181181
#endif
182182
};
183183

184-
class __attribute__((sycl_special_class)) n {
184+
class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) n {
185185
public:
186186
int *x;
187187
float *y;
@@ -203,7 +203,7 @@ class __attribute__((sycl_special_class)) n {
203203
#endif
204204
};
205205

206-
template <typename... Properties> class __attribute__((sycl_special_class)) mn {
206+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) mn {
207207
public:
208208
int *x;
209209
float *y;
@@ -225,7 +225,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) mn {
225225
#endif
226226
};
227227

228-
template <typename... Properties> class __attribute__((sycl_special_class)) nm {
228+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) nm {
229229
public:
230230
int *x;
231231
float *y;
@@ -249,7 +249,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) nm {
249249

250250
// Empty attribute names.
251251

252-
class __attribute__((sycl_special_class)) np {
252+
class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) np {
253253
public:
254254
int *x;
255255

@@ -266,7 +266,7 @@ class __attribute__((sycl_special_class)) np {
266266
#endif
267267
};
268268

269-
class __attribute__((sycl_special_class)) mp {
269+
class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) mp {
270270
public:
271271
int *x;
272272

‎clang/test/CodeGenSYCL/add_ir_attributes_kernel_parameter_filter.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@
99
#include "mock_properties.hpp"
1010
#include "sycl.hpp"
1111

12-
template <typename... Properties> class __attribute__((sycl_special_class)) g {
12+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) g {
1313
public:
1414
int *x;
1515

‎clang/test/CodeGenSYCL/integration_header.cpp

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -18,6 +18,7 @@
1818
// CHECK-NEXT: "_ZTSN16second_namespace13second_kernelIcEE",
1919
// CHECK-NEXT: "_ZTS13fourth_kernelIJN15template_arg_ns14namespaced_argILi1EEEEE"
2020
// CHECK-NEXT: "_ZTSZ4mainE16accessor_in_base"
21+
// CHECK-NEXT: "_ZTSZ4mainE15annotated_types"
2122
// CHECK-NEXT: };
2223
//
2324
// CHECK: static constexpr
@@ -48,6 +49,11 @@
4849
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 40 },
4950
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 52 },
5051
// CHECK-EMPTY:
52+
// CHECK-NEXT: //--- _ZTSZ4mainE15annotated_types
53+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
54+
// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 8 },
55+
// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 16 },
56+
// CHECK-EMPTY:
5157
// CHECK-NEXT: { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
5258
// CHECK-NEXT: };
5359
//
@@ -102,6 +108,8 @@ struct captured : base, base2 {
102108

103109
}; // namespace accessor_in_base
104110

111+
struct MockProperty {};
112+
105113
int main() {
106114

107115
sycl::accessor<char, 1, sycl::access::mode::read> acc1;
@@ -143,5 +151,14 @@ int main() {
143151
c.use();
144152
});
145153

154+
sycl::ext::oneapi::experimental::annotated_arg<int, MockProperty> AA1;
155+
sycl::ext::oneapi::experimental::annotated_ptr<int, MockProperty> AP2;
156+
sycl::ext::oneapi::experimental::annotated_arg<float*, MockProperty> AP3;
157+
kernel_single_task<class annotated_types>([=]() {
158+
(void)AA1;
159+
(void)AP2;
160+
(void)AP3;
161+
});
162+
146163
return 0;
147164
}

‎clang/test/CodeGenSYCL/no_opaque_add_ir_attributes_kernel_parameter.cpp

Lines changed: 14 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,7 @@
1111

1212
// One __init parameter with add_ir_attributes_kernel_parameter attribute.
1313

14-
template <typename... Properties> class __attribute__((sycl_special_class)) g {
14+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) g {
1515
public:
1616
int *x;
1717

@@ -27,7 +27,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) g {
2727
#endif
2828
};
2929

30-
class __attribute__((sycl_special_class)) h {
30+
class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) h {
3131
public:
3232
int *x;
3333

@@ -44,7 +44,7 @@ class __attribute__((sycl_special_class)) h {
4444
#endif
4545
};
4646

47-
template <typename... Properties> class __attribute__((sycl_special_class)) gh {
47+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) gh {
4848
public:
4949
int *x;
5050

@@ -61,7 +61,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) gh {
6161
#endif
6262
};
6363

64-
template <typename... Properties> class __attribute__((sycl_special_class)) hg {
64+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) hg {
6565
public:
6666
int *x;
6767

@@ -80,7 +80,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) hg {
8080

8181
// Two __init parameters, one with add_ir_attributes_kernel_parameter attribute.
8282

83-
template <typename... Properties> class __attribute__((sycl_special_class)) k {
83+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) k {
8484
public:
8585
int *x;
8686
float *y;
@@ -99,7 +99,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) k {
9999
#endif
100100
};
101101

102-
class __attribute__((sycl_special_class)) l {
102+
class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) l {
103103
public:
104104
int *x;
105105
float *y;
@@ -119,7 +119,7 @@ class __attribute__((sycl_special_class)) l {
119119
#endif
120120
};
121121

122-
template <typename... Properties> class __attribute__((sycl_special_class)) kl {
122+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) kl {
123123
public:
124124
int *x;
125125
float *y;
@@ -139,7 +139,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) kl {
139139
#endif
140140
};
141141

142-
template <typename... Properties> class __attribute__((sycl_special_class)) lk {
142+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) lk {
143143
public:
144144
int *x;
145145
float *y;
@@ -161,7 +161,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) lk {
161161

162162
// Two __init parameters, both with add_ir_attributes_kernel_parameter attribute.
163163

164-
template <typename... Properties> class __attribute__((sycl_special_class)) m {
164+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) m {
165165
public:
166166
int *x;
167167
float *y;
@@ -181,7 +181,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) m {
181181
#endif
182182
};
183183

184-
class __attribute__((sycl_special_class)) n {
184+
class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) n {
185185
public:
186186
int *x;
187187
float *y;
@@ -203,7 +203,7 @@ class __attribute__((sycl_special_class)) n {
203203
#endif
204204
};
205205

206-
template <typename... Properties> class __attribute__((sycl_special_class)) mn {
206+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) mn {
207207
public:
208208
int *x;
209209
float *y;
@@ -225,7 +225,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) mn {
225225
#endif
226226
};
227227

228-
template <typename... Properties> class __attribute__((sycl_special_class)) nm {
228+
template <typename... Properties> class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) nm {
229229
public:
230230
int *x;
231231
float *y;
@@ -249,7 +249,7 @@ template <typename... Properties> class __attribute__((sycl_special_class)) nm {
249249

250250
// Empty attribute names.
251251

252-
class __attribute__((sycl_special_class)) np {
252+
class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) np {
253253
public:
254254
int *x;
255255

@@ -266,7 +266,7 @@ class __attribute__((sycl_special_class)) np {
266266
#endif
267267
};
268268

269-
class __attribute__((sycl_special_class)) mp {
269+
class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) mp {
270270
public:
271271
int *x;
272272

0 commit comments

Comments
 (0)