Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion GPU/Common/GPUCommonDefAPI.h
Original file line number Diff line number Diff line change
Expand Up @@ -295,5 +295,5 @@
#define get_group_id(dim) iBlock
#endif

// clang-format on
// clang-format on
#endif
2 changes: 1 addition & 1 deletion GPU/Common/GPUCommonDouble.h
Original file line number Diff line number Diff line change
Expand Up @@ -34,7 +34,7 @@ static_assert(alignof(double) == 8, "the emulated double must match the alignmen
template <>
GPUhdi() constexpr double GPUCommonMath::Abs<double>(double x)
{
return double::fromBits(x.bits() & ~GPUCA_B64_SIGN);
return double ::fromBits(x.bits() & ~GPUCA_B64_SIGN);
}

// metal::fabs is not constant-evaluable, so this also fails to compile if the
Expand Down
10 changes: 8 additions & 2 deletions GPU/Common/GPUCommonMath.h
Original file line number Diff line number Diff line change
Expand Up @@ -481,9 +481,15 @@ GPUhdi() constexpr int32_t GPUCommonMath::Abs<int32_t>(int32_t x)
// has the same size and alignment; the overloads keep the address space, which a
// generic pointer would not carry into atomic_*_explicit.
template <class T>
GPUdi() threadgroup metal::atomic<T>* GPUCommonMathMetalAtomic(threadgroup T* p) { return reinterpret_cast<threadgroup metal::atomic<T>*>(p); }
GPUdi() threadgroup metal::atomic<T>* GPUCommonMathMetalAtomic(threadgroup T* p)
{
return reinterpret_cast<threadgroup metal::atomic<T>*>(p);
}
template <class T>
GPUdi() device metal::atomic<T>* GPUCommonMathMetalAtomic(T* p) { return (device metal::atomic<T>*)p; }
GPUdi() device metal::atomic<T>* GPUCommonMathMetalAtomic(T* p)
{
return (device metal::atomic<T>*)p;
}
#endif

template <class S, class T>
Expand Down
2 changes: 1 addition & 1 deletion GPU/GPUTracking/Base/GPUReconstructionKernelMacros.h
Original file line number Diff line number Diff line change
Expand Up @@ -91,4 +91,4 @@
#define GPUCA_KRNL_LB(x_class, x_attributes, ...) GPUCA_KRNL(x_class, (REG, (GPUCA_M_CAT(GPUCA_LB_, GPUCA_M_KRNL_NAME(x_class))), GPUCA_M_STRIP(x_attributes)), __VA_ARGS__)

#endif // O2_GPU_GPURECONSTRUCTIONKERNELMACROS_H
// clang-format on
// clang-format on
3 changes: 2 additions & 1 deletion GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm
Original file line number Diff line number Diff line change
Expand Up @@ -399,7 +399,8 @@

// the section is part of the mapped image, so it outlives the dispatch_data_t
// and does not have to be copied
dispatch_data_t blob = dispatch_data_create(p, sz, nullptr, ^{});
dispatch_data_t blob = dispatch_data_create(p, sz, nullptr, ^{
});

NSError* error = nil;
mInternals->library = [mInternals->device newLibraryWithData:blob error:&error];
Expand Down
3 changes: 2 additions & 1 deletion GPU/GPUTracking/TPCClusterFinder/CfChargePos.h
Original file line number Diff line number Diff line change
Expand Up @@ -38,7 +38,8 @@ struct CfChargePos {
// INVALID_CHARGE_POS below lives in the constant address space, which a
// generic `this` does not reach in MSL.
constexpr GPUhdi() CfChargePos(tpccf::Row row, tpccf::Pad pad, tpccf::TPCFragmentTime t) constant
: gpad(tpcGlobalPadIdx(row, pad)), timePadded(t + GPUCF_PADDING_TIME)
: gpad(tpcGlobalPadIdx(row, pad)),
timePadded(t + GPUCF_PADDING_TIME)
{
}
#endif
Expand Down