Skip to content

Commit 5709b89

Browse files
badervladimirlaz
bader
authored and
vladimirlaz
committed
[SYCL][NFC] Format the code with clang-format
Signed-off-by: Alexey Bader <[email protected]>
1 parent 2ddf04e commit 5709b89

File tree

6 files changed

+40
-39
lines changed

6 files changed

+40
-39
lines changed

‎clang/lib/CodeGen/SYCLLowerIR/LowerWGScope.cpp

Lines changed: 21 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -442,13 +442,15 @@ static void copyBetweenPrivateAndShadow(Value *L, GlobalVariable *Shadow,
442442
//
443443
static void materializeLocalsInWIScopeBlocksImpl(
444444
const DenseMap<BasicBlock *, std::unique_ptr<LocalsSet>> &BB2MatLocals,
445-
const DenseMap<AllocaInst *, GlobalVariable *> &Local2Shadow, const Triple &TT) {
445+
const DenseMap<AllocaInst *, GlobalVariable *> &Local2Shadow,
446+
const Triple &TT) {
446447
for (auto &P : BB2MatLocals) {
447448
// generate LeaderBB and private<->shadow copies in proper BBs
448449
BasicBlock *LeaderBB = P.first;
449450
BasicBlock *BB = LeaderBB->splitBasicBlock(&LeaderBB->front(), "LeaderMat");
450451
// Add a barrier to the original block:
451-
Instruction *At = spirv::genWGBarrier(*BB->getFirstNonPHI(), TT)->getNextNode();
452+
Instruction *At =
453+
spirv::genWGBarrier(*BB->getFirstNonPHI(), TT)->getNextNode();
452454

453455
for (AllocaInst *L : *P.second.get()) {
454456
auto MapEntry = Local2Shadow.find(L);
@@ -533,10 +535,9 @@ static bool localMustBeMaterialized(const AllocaInst *L, const BasicBlock &BB) {
533535
// have any noticible effect, though, as reading from Shadow always goes to a
534536
// register file anyway.
535537
//
536-
void materializeLocalsInWIScopeBlocks(
537-
SmallPtrSetImpl<AllocaInst *> &Locals,
538-
SmallPtrSetImpl<BasicBlock *> &WIScopeBBs,
539-
const Triple &TT) {
538+
void materializeLocalsInWIScopeBlocks(SmallPtrSetImpl<AllocaInst *> &Locals,
539+
SmallPtrSetImpl<BasicBlock *> &WIScopeBBs,
540+
const Triple &TT) {
540541
// maps local variable to its "shadow" workgroup-shared global:
541542
DenseMap<AllocaInst *, GlobalVariable *> Local2Shadow;
542543
// records which locals must be materialized at the beginning of a block:
@@ -718,8 +719,7 @@ static void shareByValParams(Function &F, const Triple &TT) {
718719
spirv::genWGBarrier(MergeBB->front(), TT);
719720
}
720721

721-
PreservedAnalyses SYCLLowerWGScopePass::run(Function &F,
722-
const llvm::Triple &TT,
722+
PreservedAnalyses SYCLLowerWGScopePass::run(Function &F, const llvm::Triple &TT,
723723
FunctionAnalysisManager &FAM) {
724724
if (!F.getMetadata(WG_SCOPE_MD))
725725
return PreservedAnalyses::all();
@@ -876,16 +876,16 @@ Value *spirv::genLinearLocalID(Instruction &Before, const Triple &TT) {
876876
IRBuilder<> Bld(Ctx);
877877
Bld.SetInsertPoint(&Before);
878878

879-
#define CREATE_CALLEE(NAME, FN_NAME) \
880-
FunctionCallee FnCallee##NAME = M.getOrInsertFunction(FN_NAME, RetTy); \
881-
assert(FnCallee##NAME && "spirv intrinsic creation failed"); \
879+
#define CREATE_CALLEE(NAME, FN_NAME) \
880+
FunctionCallee FnCallee##NAME = M.getOrInsertFunction(FN_NAME, RetTy); \
881+
assert(FnCallee##NAME && "spirv intrinsic creation failed"); \
882882
auto NAME = Bld.CreateCall(FnCallee##NAME, {});
883883

884-
CREATE_CALLEE(LocalInvocationId_X, "_Z27__spirv_LocalInvocationId_xv");
885-
CREATE_CALLEE(LocalInvocationId_Y, "_Z27__spirv_LocalInvocationId_yv");
886-
CREATE_CALLEE(LocalInvocationId_Z, "_Z27__spirv_LocalInvocationId_zv");
887-
CREATE_CALLEE(WorkgroupSize_Y, "_Z23__spirv_WorkgroupSize_yv");
888-
CREATE_CALLEE(WorkgroupSize_Z, "_Z23__spirv_WorkgroupSize_zv");
884+
CREATE_CALLEE(LocalInvocationId_X, "_Z27__spirv_LocalInvocationId_xv");
885+
CREATE_CALLEE(LocalInvocationId_Y, "_Z27__spirv_LocalInvocationId_yv");
886+
CREATE_CALLEE(LocalInvocationId_Z, "_Z27__spirv_LocalInvocationId_zv");
887+
CREATE_CALLEE(WorkgroupSize_Y, "_Z23__spirv_WorkgroupSize_yv");
888+
CREATE_CALLEE(WorkgroupSize_Z, "_Z23__spirv_WorkgroupSize_zv");
889889

890890
#undef CREATE_CALLEE
891891

@@ -894,12 +894,11 @@ Value *spirv::genLinearLocalID(Instruction &Before, const Triple &TT) {
894894
// 3: + (__spirv_WorkgroupSize_z() * __spirv_LocalInvocationId_y())
895895
// 4: + (__spirv_LocalInvocationId_z())
896896
return Bld.CreateAdd(
897-
Bld.CreateAdd(
898-
Bld.CreateMul(
899-
Bld.CreateMul(WorkgroupSize_Y, WorkgroupSize_Z), // 1
900-
LocalInvocationId_X), // 2
901-
Bld.CreateMul(WorkgroupSize_Z, LocalInvocationId_Y)), // 3
902-
LocalInvocationId_Z); // 4
897+
Bld.CreateAdd(
898+
Bld.CreateMul(Bld.CreateMul(WorkgroupSize_Y, WorkgroupSize_Z), // 1
899+
LocalInvocationId_X), // 2
900+
Bld.CreateMul(WorkgroupSize_Z, LocalInvocationId_Y)), // 3
901+
LocalInvocationId_Z); // 4
903902
} else {
904903
StringRef Name = "__spirv_BuiltInLocalInvocationIndex";
905904
GlobalVariable *G = M.getGlobalVariable(Name);

‎sycl/include/CL/sycl/device.hpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -97,7 +97,8 @@ class device {
9797
/// exception must be thrown.
9898
///
9999
/// \param ComputeUnits is a desired count of compute units in each sub
100-
/// device. \return A vector class of sub devices partitioned from this SYCL
100+
/// device.
101+
/// \return A vector class of sub devices partitioned from this SYCL
101102
/// device equally based on the ComputeUnits parameter.
102103
template <info::partition_property prop>
103104
vector_class<device> create_sub_devices(size_t ComputeUnits) const;
@@ -126,7 +127,8 @@ class device {
126127
/// exception must be thrown.
127128
///
128129
/// \param AffinityDomain is one of the values described in Table 4.20 of SYCL
129-
/// Spec \return a vector class of sub devices partitioned from this SYCL
130+
/// Spec
131+
/// \return a vector class of sub devices partitioned from this SYCL
130132
/// device by affinity domain based on the AffinityDomain parameter
131133
template <info::partition_property prop>
132134
vector_class<device>

‎sycl/include/CL/sycl/queue.hpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -346,7 +346,7 @@ class queue {
346346
template <typename KernelName = detail::auto_name, typename KernelType,
347347
int Dims>
348348
event parallel_for(range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
349-
KernelType KernelFunc) {
349+
KernelType KernelFunc) {
350350
return submit([&](handler &CGH) {
351351
CGH.template parallel_for<KernelName, KernelType, Dims>(
352352
NumWorkItems, WorkItemOffset, KernelFunc);
@@ -415,8 +415,8 @@ class queue {
415415
/// \param KernelFunc is the Kernel functor or lambda
416416
template <typename KernelName = detail::auto_name, typename KernelType,
417417
int Dims>
418-
event parallel_for(nd_range<Dims> ExecutionRange,
419-
event DepEvent, KernelType KernelFunc) {
418+
event parallel_for(nd_range<Dims> ExecutionRange, event DepEvent,
419+
KernelType KernelFunc) {
420420
return submit([&](handler &CGH) {
421421
CGH.depends_on(DepEvent);
422422
CGH.template parallel_for<KernelName, KernelType, Dims>(ExecutionRange,

‎sycl/plugins/opencl/pi_opencl.cpp

Lines changed: 7 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -51,8 +51,6 @@ CONSTFIX char clGetMemAllocInfoName[] = "clGetMemAllocInfoINTEL";
5151

5252
#undef CONSTFIX
5353

54-
55-
5654
// USM helper function to get an extension function pointer
5755
template <const char *FuncName, typename T>
5856
static pi_result getExtFuncFromContext(pi_context context, T *fptr) {
@@ -88,7 +86,7 @@ static pi_result getExtFuncFromContext(pi_context context, T *fptr) {
8886
sizeof(cl_platform_id), &curPlatform, nullptr);
8987

9088
if (ret_err != CL_SUCCESS) {
91-
return PI_INVALID_CONTEXT;
89+
return PI_INVALID_CONTEXT;
9290
}
9391

9492
T FuncPtr =
@@ -123,7 +121,7 @@ static pi_result USMSetIndirectAccess(pi_kernel kernel) {
123121

124122
getExtFuncFromContext<clHostMemAllocName, clHostMemAllocINTEL_fn>(
125123
cast<pi_context>(CLContext), &HFunc);
126-
if (HFunc) {
124+
if (HFunc) {
127125
clSetKernelExecInfo(cast<cl_kernel>(kernel),
128126
CL_KERNEL_EXEC_INFO_INDIRECT_HOST_ACCESS_INTEL,
129127
sizeof(cl_bool), &TrueVal);
@@ -604,8 +602,8 @@ pi_result OCL(piextUSMHostAlloc)(void **result_ptr, pi_context context,
604602

605603
if (FuncPtr) {
606604
Ptr = FuncPtr(cast<cl_context>(context),
607-
cast<cl_mem_properties_intel *>(properties), size, alignment,
608-
cast<cl_int *>(&RetVal));
605+
cast<cl_mem_properties_intel *>(properties), size, alignment,
606+
cast<cl_int *>(&RetVal));
609607
}
610608

611609
*result_ptr = Ptr;
@@ -728,7 +726,8 @@ pi_result OCL(piextKernelSetArgPointer)(pi_kernel kernel, pi_uint32 arg_index,
728726
// This means we need to deref the arg to get the pointer value
729727
auto PtrToPtr = reinterpret_cast<const intptr_t *>(arg_value);
730728
auto DerefPtr = reinterpret_cast<void *>(*PtrToPtr);
731-
RetVal = cast<pi_result>(FuncPtr(cast<cl_kernel>(kernel), arg_index, DerefPtr));
729+
RetVal =
730+
cast<pi_result>(FuncPtr(cast<cl_kernel>(kernel), arg_index, DerefPtr));
732731
}
733732

734733
return RetVal;
@@ -1072,7 +1071,7 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
10721071
_PI_CL(piextUSMEnqueueMemAdvise, OCL(piextUSMEnqueueMemAdvise))
10731072
_PI_CL(piextUSMGetMemAllocInfo, OCL(piextUSMGetMemAllocInfo))
10741073

1075-
_PI_CL(piextKernelSetArgMemObj, OCL(piextKernelSetArgMemObj))
1074+
_PI_CL(piextKernelSetArgMemObj, OCL(piextKernelSetArgMemObj))
10761075

10771076
#undef _PI_CL
10781077

‎sycl/source/detail/platform_impl.hpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -75,8 +75,9 @@ class platform_impl {
7575

7676
bool is_cuda() const {
7777
const string_class CUDA_PLATFORM_STRING = "NVIDIA CUDA";
78-
const string_class PlatformName = get_platform_info<string_class,
79-
info::platform::name>::get(MPlatform, getPlugin());
78+
const string_class PlatformName =
79+
get_platform_info<string_class, info::platform::name>::get(MPlatform,
80+
getPlugin());
8081
return PlatformName == CUDA_PLATFORM_STRING;
8182
}
8283

‎sycl/source/detail/queue_impl.hpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -48,8 +48,8 @@ class queue_impl {
4848
queue_impl(DeviceImplPtr Device, async_handler AsyncHandler, QueueOrder Order,
4949
const property_list &PropList)
5050
: queue_impl(Device,
51-
detail::getSyclObjImpl(
52-
context(createSyclObjFromImpl<device>(Device), {}, true)),
51+
detail::getSyclObjImpl(context(
52+
createSyclObjFromImpl<device>(Device), {}, true)),
5353
AsyncHandler, Order, PropList){};
5454

5555
/// Constructs a SYCL queue with an async_handler and property_list provided

0 commit comments

Comments
 (0)