Skip to content

Commit 810af7b

Browse files
rdeodharrdeodhar
committed
Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2
2 parents 0412db3 + e582224 commit 810af7b

File tree

196 files changed

+10364
-2044
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

196 files changed

+10364
-2044
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 22 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -308,11 +308,12 @@ static void reportConflictingAttrs(Sema &S, FunctionDecl *F, const Attr *A1,
308308
F->setInvalidDecl();
309309
}
310310

311-
// Returns the signed constant integer value represented by given expression.
312-
static int64_t getIntExprValue(Sema &S, const Expr *E) {
311+
/// Returns the signed constant integer value represented by given expression
312+
static int64_t getIntExprValue(const Expr *E, ASTContext &Ctx) {
313313
llvm::APSInt Val(32);
314-
bool IsValid = E->isIntegerConstantExpr(Val, S.getASTContext());
314+
bool IsValid = E->isIntegerConstantExpr(Val, Ctx);
315315
assert(IsValid && "expression must be constant integer");
316+
(void)IsValid;
316317
return Val.getSExtValue();
317318
}
318319

@@ -793,8 +794,8 @@ static void VisitField(CXXRecordDecl *Owner, RangeTy &&Item, QualType ItemTy,
793794
}
794795

795796
template <typename RangeTy, typename... Handlers>
796-
static void VisitScalarField(CXXRecordDecl *Owner, RangeTy &&Item, QualType ItemTy,
797-
Handlers &... handlers) {
797+
static void VisitScalarField(CXXRecordDecl *Owner, RangeTy &&Item,
798+
QualType ItemTy, Handlers &... handlers) {
798799
KF_FOR_EACH(handleScalarType, Item, ItemTy);
799800
}
800801

@@ -900,6 +901,9 @@ template <typename Derived> class SyclKernelFieldHandler {
900901
return true;
901902
}
902903
virtual bool handleSyclAccessorType(FieldDecl *, QualType) { return true; }
904+
virtual bool handleSyclSamplerType(const CXXBaseSpecifier &, QualType) {
905+
return true;
906+
}
903907
virtual bool handleSyclSamplerType(FieldDecl *, QualType) { return true; }
904908
virtual bool handleSyclSpecConstantType(FieldDecl *, QualType) {
905909
return true;
@@ -999,8 +1003,7 @@ class SyclKernelFieldChecker
9991003

10001004
public:
10011005
SyclKernelFieldChecker(Sema &S)
1002-
: SyclKernelFieldHandler(S), Diag(S.getASTContext().getDiagnostics()) {
1003-
}
1006+
: SyclKernelFieldHandler(S), Diag(S.getASTContext().getDiagnostics()) {}
10041007
bool isValid() { return !IsInvalid; }
10051008

10061009
bool handleReferenceType(FieldDecl *FD, QualType FieldTy) final {
@@ -1050,7 +1053,7 @@ class SyclKernelDeclCreator
10501053

10511054
void addParam(const FieldDecl *FD, QualType FieldTy) {
10521055
const ConstantArrayType *CAT =
1053-
SemaRef.getASTContext().getAsConstantArrayType(FieldTy);
1056+
SemaRef.getASTContext().getAsConstantArrayType(FieldTy);
10541057
if (CAT)
10551058
FieldTy = CAT->getElementType();
10561059
ParamDesc newParamDesc = makeParamDesc(FD, FieldTy);
@@ -1131,8 +1134,7 @@ class SyclKernelDeclCreator
11311134
: SyclKernelFieldHandler(S),
11321135
KernelDecl(createKernelDecl(S.getASTContext(), Name, Loc, IsInline,
11331136
IsSIMDKernel)),
1134-
ArgChecker(ArgChecker), FuncContext(SemaRef, KernelDecl) {
1135-
}
1137+
ArgChecker(ArgChecker), FuncContext(SemaRef, KernelDecl) {}
11361138

11371139
~SyclKernelDeclCreator() {
11381140
ASTContext &Ctx = SemaRef.getASTContext();
@@ -1195,7 +1197,7 @@ class SyclKernelDeclCreator
11951197
return true;
11961198
}
11971199

1198-
//FIXME Remove this function when structs are replaced by their fields
1200+
// FIXME Remove this function when structs are replaced by their fields
11991201
bool handleStructType(FieldDecl *FD, QualType FieldTy) final {
12001202
addParam(FD, FieldTy);
12011203
return true;
@@ -1222,6 +1224,7 @@ class SyclKernelDeclCreator
12221224
}
12231225

12241226
using SyclKernelFieldHandler::handleScalarType;
1227+
using SyclKernelFieldHandler::handleSyclSamplerType;
12251228
};
12261229

12271230
class SyclKernelBodyCreator
@@ -1491,7 +1494,7 @@ class SyclKernelBodyCreator
14911494
return true;
14921495
}
14931496

1494-
//FIXME Remove this function when structs are replaced by their fields
1497+
// FIXME Remove this function when structs are replaced by their fields
14951498
bool handleStructType(FieldDecl *FD, QualType FieldTy) final {
14961499
createExprForStructOrScalar(FD);
14971500
return true;
@@ -1543,6 +1546,7 @@ class SyclKernelBodyCreator
15431546
using SyclKernelFieldHandler::enterArray;
15441547
using SyclKernelFieldHandler::enterField;
15451548
using SyclKernelFieldHandler::handleScalarType;
1549+
using SyclKernelFieldHandler::handleSyclSamplerType;
15461550
using SyclKernelFieldHandler::leaveField;
15471551
};
15481552

@@ -1638,7 +1642,7 @@ class SyclKernelIntHeaderCreator
16381642
return true;
16391643
}
16401644

1641-
//FIXME Remove this function when structs are replaced by their fields
1645+
// FIXME Remove this function when structs are replaced by their fields
16421646
bool handleStructType(FieldDecl *FD, QualType FieldTy) final {
16431647
addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout);
16441648
return true;
@@ -1697,6 +1701,7 @@ class SyclKernelIntHeaderCreator
16971701
}
16981702

16991703
using SyclKernelFieldHandler::handleScalarType;
1704+
using SyclKernelFieldHandler::handleSyclSamplerType;
17001705
};
17011706
} // namespace
17021707

@@ -1787,15 +1792,16 @@ void Sema::MarkDevice(void) {
17871792
KernelBody ? KernelBody->getAttr<SYCLSimdAttr>() : nullptr;
17881793
if (auto *Existing =
17891794
SYCLKernel->getAttr<IntelReqdSubGroupSizeAttr>()) {
1790-
if (Existing->getSubGroupSize() != Attr->getSubGroupSize()) {
1795+
if (getIntExprValue(Existing->getSubGroupSize(), getASTContext()) !=
1796+
getIntExprValue(Attr->getSubGroupSize(), getASTContext())) {
17911797
Diag(SYCLKernel->getLocation(),
17921798
diag::err_conflicting_sycl_kernel_attributes);
17931799
Diag(Existing->getLocation(), diag::note_conflicting_attribute);
17941800
Diag(Attr->getLocation(), diag::note_conflicting_attribute);
17951801
SYCLKernel->setInvalidDecl();
17961802
}
1797-
} else if (KBSimdAttr &&
1798-
(getIntExprValue(*this, Attr->getSubGroupSize()) != 1)) {
1803+
} else if (KBSimdAttr && (getIntExprValue(Attr->getSubGroupSize(),
1804+
getASTContext()) != 1)) {
17991805
reportConflictingAttrs(*this, KernelBody, KBSimdAttr, Attr);
18001806
} else {
18011807
SYCLKernel->addAttr(A);

‎clang/test/CodeGenSYCL/sampler.cpp

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -11,8 +11,20 @@
1111
// CHECK-NEXT: [[GEPCAST:%[0-9]+]] = addrspacecast %"class{{.*}}.cl::sycl::sampler"* [[GEP]] to %"class{{.*}}.cl::sycl::sampler" addrspace(4)*
1212
// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* [[GEPCAST]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]])
1313
//
14+
15+
// CHECK: define spir_kernel void @{{[a-zA-Z0-9_]+}}(%struct{{.*}}sampler_wrapper{{.*}} %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED:%[a-zA-Z0-9_]+]])
16+
// CHECK: [[SAMPLER_ARG_WRAPPED]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8
17+
// CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED]], %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG_WRAPPED]].addr, align 8
18+
// CHECK: [[LOAD_SAMPLER_ARG_WRAPPED:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG_WRAPPED]].addr, align 8
19+
// CHECK: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* {{.*}}, %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG_WRAPPED]])
20+
//
1421
#include "sycl.hpp"
1522

23+
struct sampler_wrapper {
24+
cl::sycl::sampler smpl;
25+
int a;
26+
};
27+
1628
template <typename KernelName, typename KernelType>
1729
__attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) {
1830
kernelFunc();
@@ -24,5 +36,10 @@ int main() {
2436
smplr.use();
2537
});
2638

39+
sampler_wrapper wrappedSampler = {smplr, 1};
40+
kernel_single_task<class second_kernel>([=]() {
41+
wrappedSampler.smpl.use();
42+
});
43+
2744
return 0;
2845
}

‎clang/test/SemaSYCL/reqd-sub-group-size-device.cpp

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -45,7 +45,16 @@ void bar() {
4545
baz();
4646
});
4747
#endif
48+
4849
kernel<class kernel_name5>([]() [[cl::intel_reqd_sub_group_size(2)]] { });
50+
kernel<class kernel_name6>([]() [[cl::intel_reqd_sub_group_size(4)]] { foo(); });
51+
}
52+
53+
[[cl::intel_reqd_sub_group_size(16)]] SYCL_EXTERNAL void B();
54+
[[cl::intel_reqd_sub_group_size(16)]] void A() {
55+
}
56+
[[cl::intel_reqd_sub_group_size(16)]] SYCL_EXTERNAL void B() {
57+
A();
4958
}
5059

5160
#ifdef TRIGGER_ERROR

‎libclc/CMakeLists.txt

Lines changed: 13 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -227,8 +227,8 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
227227
LIB_DIR lib
228228
DIRS ${dirs} ${DARCH} ${DARCH}-${OS} ${DARCH}-${VENDOR}-${OS}
229229
DEPS convert-clc.cl )
230-
set( libspirv_files )
231-
libclc_configure_lib_source(libspirv_files
230+
set( libspirv_files_base )
231+
libclc_configure_lib_source(libspirv_files_base
232232
LIB_DIR libspirv
233233
DIRS ${dirs} ${DARCH} ${DARCH}-${OS} ${DARCH}-${VENDOR}-${OS}
234234
DEPS convert-spirv.cl convert-core.cl)
@@ -244,6 +244,17 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
244244
endif()
245245
message( " DEVICE: ${d} ( ${${d}_aliases} )" )
246246

247+
# FIXME: this is a hack, remove once we can use sycldevice in the triple
248+
# without changing the language
249+
# see issue: https://github.com/intel/llvm/issues/1814
250+
set(libspirv_files ${libspirv_files_base})
251+
if( ${ARCH} STREQUAL nvptx OR ${ARCH} STREQUAL nvptx64 )
252+
add_libclc_sycl_binding(libspirv_files
253+
TRIPLE ${t}
254+
COMPILE_OPT ${mcpu}
255+
FILES generic/libspirv/sycldevice-binding.cpp)
256+
endif()
257+
247258
add_libclc_builtin_set(libspirv-${arch_suffix}
248259
TRIPLE ${t}
249260
TARGET_ENV libspirv

‎libclc/cmake/modules/AddLibclc.cmake

Lines changed: 52 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -177,3 +177,55 @@ function(libclc_configure_lib_source OUT_LIST)
177177
set( ${OUT_LIST} ${rel_files} PARENT_SCOPE )
178178

179179
endfunction(libclc_configure_lib_source OUT_LIST)
180+
181+
# add_libclc_sycl_binding(arch_suffix
182+
# TRIPLE string
183+
# Triple used to compile
184+
# FILES string ...
185+
# List of file that should be built for this library
186+
# COMPILE_OPT
187+
# Compilation options
188+
# )
189+
#
190+
# Build the sycl binding file for SYCLDEVICE.
191+
# The path to the generated object file are appended in OUT_LIST.
192+
#
193+
# The mangling for sycl device is not yet fully
194+
# compatible with standard mangling.
195+
# For various reason, we need a mangling specific
196+
# for the Default address space (mapping to generic in SYCL).
197+
# The Default address space is not accessible in CL mode,
198+
# so we build this file in sycl mode for mangling purposes.
199+
#
200+
# FIXME: all the files should be compiled with the sycldevice triple
201+
# but this is not possible at the moment as this will trigger
202+
# the SYCL mode which we don't want.
203+
#
204+
function(add_libclc_sycl_binding OUT_LIST)
205+
cmake_parse_arguments(ARG
206+
""
207+
"TRIPLE"
208+
"FILES;COMPILE_OPT"
209+
${ARGN})
210+
211+
foreach( file ${ARG_FILES} )
212+
file( TO_CMAKE_PATH ${LIBCLC_ROOT_DIR}/${file} SYCLDEVICE_BINDING )
213+
if( EXISTS ${SYCLDEVICE_BINDING} )
214+
set( SYCLDEVICE_BINDING_OUT ${CMAKE_CURRENT_BINARY_DIR}/sycldevice-binding-${ARG_TRIPLE}/sycldevice-binding.bc )
215+
add_custom_command( OUTPUT ${SYCLDEVICE_BINDING_OUT}
216+
COMMAND ${LLVM_CLANG}
217+
-target ${ARG_TRIPLE}-sycldevice
218+
-fsycl
219+
-fsycl-device-only
220+
-Dcl_khr_fp64
221+
-I${LIBCLC_ROOT_DIR}/generic/include
222+
${ARG_COMPILE_OPT}
223+
${SYCLDEVICE_BINDING}
224+
-o ${SYCLDEVICE_BINDING_OUT}
225+
MAIN_DEPENDENCY ${SYCLDEVICE_BINDING}
226+
DEPENDS ${SYCLDEVICE_BINDING} ${LLVM_CLANG}
227+
VERBATIM )
228+
set( ${OUT_LIST} "${${OUT_LIST}};${SYCLDEVICE_BINDING_OUT}" PARENT_SCOPE )
229+
endif()
230+
endforeach()
231+
endfunction(add_libclc_sycl_binding OUT_LIST)

‎libclc/generic/include/clc/relational/floatn.inc

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,10 @@
1+
//===----------------------------------------------------------------------===//
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+
//===----------------------------------------------------------------------===//
18

29
#define __CLC_FLOATN float
310
#define __CLC_INTN int

‎libclc/generic/include/clcmacro.h

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,14 @@
1+
//===----------------------------------------------------------------------===//
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 __CLC_MACRO_H
10+
#define __CLC_MACRO_H
11+
112
#define _CLC_UNARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE) \
213
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x) { \
314
return (RET_TYPE##2)(FUNCTION(x.x), FUNCTION(x.y)); \
@@ -177,3 +188,5 @@
177188
#define _CLC_DEFINE_UNARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE) \
178189
_CLC_DEF _CLC_OVERLOAD RET_TYPE FUNCTION(ARG1_TYPE x) { return BUILTIN(x); } \
179190
_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE)
191+
192+
#endif // !__CLC_MACRO_H

‎libclc/generic/include/lp64_types.h

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -108,6 +108,10 @@ typedef half __clc_vec16_fp16_t __attribute__((ext_vector_type(16)));
108108

109109
typedef __clc_int64_t __clc_size_t;
110110

111+
#ifdef __SYCL_DEVICE_ONLY__
112+
typedef __ocl_event_t __clc_event_t;
113+
#else
111114
typedef event_t __clc_event_t;
115+
#endif
112116

113117
#endif // CLC_LP64_TYPES

‎libclc/ptx-nvidiacl/libspirv/math/binary_builtin.inc renamed to ‎libclc/generic/include/math/binary_builtin.inc

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

9-
#include "../../../generic/lib/clcmacro.h"
109
#include "utils.h"
10+
#include <clcmacro.h>
11+
12+
#ifndef __CLC_BUILTIN
13+
#define __CLC_BUILTIN __CLC_XCONCAT(__clc_, __CLC_FUNCTION)
14+
#endif
15+
16+
#ifndef __CLC_BUILTIN_D
17+
#define __CLC_BUILTIN_D __CLC_BUILTIN
18+
#endif
1119

1220
#ifndef __CLC_BUILTIN_F
13-
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, f)
21+
#define __CLC_BUILTIN_F __CLC_BUILTIN
22+
#endif
23+
24+
#ifndef __CLC_BUILTIN_H
25+
#define __CLC_BUILTIN_H __CLC_BUILTIN_F
1426
#endif
1527

1628
_CLC_DEFINE_BINARY_BUILTIN(float, __CLC_FUNCTION, __CLC_BUILTIN_F, float, float)
@@ -21,7 +33,7 @@ _CLC_DEFINE_BINARY_BUILTIN(float, __CLC_FUNCTION, __CLC_BUILTIN_F, float, float)
2133

2234
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
2335

24-
_CLC_DEFINE_BINARY_BUILTIN(double, __CLC_FUNCTION, __CLC_BUILTIN, double,
36+
_CLC_DEFINE_BINARY_BUILTIN(double, __CLC_FUNCTION, __CLC_BUILTIN_D, double,
2537
double)
2638

2739
#endif
@@ -30,12 +42,8 @@ _CLC_DEFINE_BINARY_BUILTIN(double, __CLC_FUNCTION, __CLC_BUILTIN, double,
3042

3143
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
3244

33-
_CLC_DEFINE_BINARY_BUILTIN(half, __CLC_FUNCTION, __CLC_BUILTIN, half, half)
45+
_CLC_DEFINE_BINARY_BUILTIN(half, __CLC_FUNCTION, __CLC_BUILTIN_H, half, half)
3446

3547
#endif
3648

3749
#endif
38-
39-
#undef __CLC_BUILTIN
40-
#undef __CLC_BUILTIN_F
41-
#undef __CLC_FUNCTION

0 commit comments

Comments
 (0)