Skip to content

Commit 2d4f887

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web' (intel#6)
2 parents 5cd0db7 + d8f6ce5 commit 2d4f887

40 files changed

+356
-50
lines changed

‎clang/lib/Driver/Driver.cpp

Lines changed: 8 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -5859,8 +5859,12 @@ bool clang::driver::isOptimizationLevelFast(const ArgList &Args) {
58595859
}
58605860

58615861
bool clang::driver::isObjectFile(std::string FileName) {
5862-
return (llvm::sys::path::has_extension(FileName) &&
5863-
types::lookupTypeForExtension(
5864-
llvm::sys::path::extension(FileName).drop_front()) ==
5865-
types::TY_Object);
5862+
if (llvm::sys::path::has_extension(FileName)) {
5863+
std::string Ext(llvm::sys::path::extension(FileName).drop_front());
5864+
// We cannot rely on lookupTypeForExtension solely as that has 'lib'
5865+
// marked as an object.
5866+
return (Ext != "lib" &&
5867+
types::lookupTypeForExtension(Ext) == types::TY_Object);
5868+
}
5869+
return false;
58665870
}

‎clang/test/Driver/sycl-offload.c

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -583,5 +583,14 @@
583583
// FO-CHECK: clang{{.*}} "-include" "[[HEADER]]" {{.*}} "-o" "[[OUTPUT2:.+\.obj]]"
584584
// FO-CHECK: clang-offload-bundler{{.*}} "-outputs=somefile.obj" "-inputs=[[OUTPUT1]],[[OUTPUT2]]"
585585

586+
/// passing of a library should not trigger the unbundler
587+
// RUN: touch %t.a
588+
// RUN: touch %t.lib
589+
// RUN: %clang -ccc-print-phases -fsycl %t.a %s 2>&1 \
590+
// RUN: | FileCheck -check-prefix=LIB-UNBUNDLE-CHECK %s
591+
// RUN: %clang_cl -ccc-print-phases -fsycl %t.lib %s 2>&1 \
592+
// RUN: | FileCheck -check-prefix=LIB-UNBUNDLE-CHECK %s
593+
// LIB-UNBUNDLE-CHECK-NOT: clang-offload-unbundler
594+
586595
// TODO: SYCL specific fail - analyze and enable
587596
// XFAIL: windows-msvc

‎llvm-spirv/lib/SPIRV/SPIRVWriter.cpp

Lines changed: 11 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -223,8 +223,17 @@ static bool recursiveType(const StructType *ST, const Type *Ty) {
223223
StructTy->element_end();
224224
}
225225

226-
if (auto *PtrTy = dyn_cast<PointerType>(Ty))
227-
return Run(PtrTy->getPointerElementType());
226+
if (auto *PtrTy = dyn_cast<PointerType>(Ty)) {
227+
Type *ElTy = PtrTy->getPointerElementType();
228+
if (auto *FTy = dyn_cast<FunctionType>(ElTy)) {
229+
// If we have a function pointer, then argument types and return type of
230+
// the referenced function also need to be checked
231+
return Run(FTy->getReturnType()) ||
232+
any_of(FTy->param_begin(), FTy->param_end(), Run);
233+
}
234+
235+
return Run(ElTy);
236+
}
228237

229238
if (auto *ArrayTy = dyn_cast<ArrayType>(Ty))
230239
return Run(ArrayTy->getArrayElementType());

‎llvm-spirv/test/transcoding/SPV_INTEL_function_pointers/fp-in-recusive-type.ll

Lines changed: 85 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,85 @@
1+
; RUN: llvm-as %s -o %t.bc
2+
; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv
3+
; RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
4+
;
5+
; This tests that translator is able to convert the following recursive data
6+
; type which contains function pointers:
7+
;
8+
; typedef void(*fp_t)(struct Desc *);
9+
; typedef Desc *(*fp2_t)();
10+
;
11+
; struct Type;
12+
;
13+
; struct Desc {
14+
; int B;
15+
; Type *Ptr;
16+
; };
17+
;
18+
; struct Type {
19+
; int A;
20+
; fp_t FP;
21+
; };
22+
;
23+
; __kernel void foo() {
24+
; Type T;
25+
; fp2_t ptr;
26+
; }
27+
;
28+
; Without proper recursive types detection, the translator crashes while trying
29+
; to translate this LLVM IR example with the following error:
30+
;
31+
; virtual SPIRV::SPIRVEntry* SPIRV::SPIRVModuleImpl::getEntry(SPIRV::SPIRVId) const: Assertion `Loc != IdEntryMap.end() && "Id is not in map"' failed.
32+
;
33+
; CHECK-SPIRV: TypeInt [[INTTY:[0-9]+]] 32 0
34+
; CHECK-SPIRV: TypeVoid [[VOIDTY:[0-9]+]]
35+
; CHECK-SPIRV: TypeStruct [[TYPETY:[0-9]+]] [[INTTY]] [[FPTRTY:[0-9]+]]
36+
; CHECK-SPIRV: TypeStruct [[DESCTY:[0-9]+]] [[INTTY]] [[TYPEPTRTY:[0-9]+]]
37+
; CHECK-SPIRV: TypePointer [[TYPEPTRTY]] {{[0-9]+}} [[TYPETY]]
38+
; CHECK-SPIRV: TypePointer [[DESCPTRTY:[0-9]+]] {{[0-9]+}} [[DESCTY]]
39+
; CHECK-SPIRV: TypeFunction [[FTY:[0-9]+]] [[VOIDTY]] [[DESCPTRTY]]
40+
; CHECK-SPIRV: TypePointer [[FPTRTY]] {{[0-9]+}} [[FTY]]
41+
; CHECK-SPIRV: TypePointer [[DESCPTR1TY:[0-9]+]] {{[0-9]+}} [[TYPETY]]
42+
; CHECK-SPIRV: TypeFunction [[F2TY:[0-9]+]] [[DESCPTRTY]]
43+
; CHECK-SPIRV: TypePointer [[FPTR2TY:[0-9]+]] {{[0-9]+}} [[F2TY]]
44+
; CHECK-SPIRV: TypePointer [[FPTR2ALLOCATY:[0-9]+]] {{[0-9]+}} [[FPTR2TY]]
45+
; CHECK-SPIRV: Variable [[DESCPTR1TY]]
46+
; CHECK-SPIRV: Variable [[FPTR2ALLOCATY]]
47+
48+
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
49+
target triple = "spir64-unknown-unknown"
50+
51+
%struct.Type = type { i32, void (%struct.Desc addrspace(4)*) addrspace(4)* }
52+
%struct.Desc = type { i32, %struct.Type addrspace(4)* }
53+
54+
define spir_kernel void @foo() #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !4 !kernel_arg_type !4 !kernel_arg_base_type !4 !kernel_arg_type_qual !4 {
55+
entry:
56+
%t = alloca %struct.Type, align 8
57+
%ptr = alloca %struct.Desc addrspace(4)* () *, align 8
58+
ret void
59+
}
60+
61+
; Function Attrs: argmemonly nounwind willreturn
62+
declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1
63+
64+
; Function Attrs: argmemonly nounwind willreturn
65+
declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1
66+
67+
attributes #0 = { "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
68+
attributes #1 = { argmemonly nounwind willreturn }
69+
attributes #2 = { inlinehint nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
70+
attributes #3 = { nounwind }
71+
72+
!llvm.module.flags = !{!0}
73+
!opencl.spir.version = !{!1}
74+
!spirv.Source = !{!2}
75+
!llvm.ident = !{!3}
76+
77+
!0 = !{i32 1, !"wchar_size", i32 4}
78+
!1 = !{i32 1, i32 2}
79+
!2 = !{i32 4, i32 100000}
80+
!3 = !{!"clang version 10.0.0 "}
81+
!4 = !{}
82+
!5 = !{!6, !6, i64 0}
83+
!6 = !{!"any pointer", !7, i64 0}
84+
!7 = !{!"omnipotent char", !8, i64 0}
85+
!8 = !{!"Simple C++ TBAA"}

‎sycl/include/CL/sycl/detail/aligned_allocator.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -14,8 +14,8 @@
1414
#include <CL/sycl/range.hpp>
1515

1616
#include <algorithm>
17-
#include <cstring>
1817
#include <cstdlib>
18+
#include <cstring>
1919
#include <memory>
2020
#include <vector>
2121

‎sycl/include/CL/sycl/detail/array.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -7,8 +7,8 @@
77
//===----------------------------------------------------------------------===//
88

99
#pragma once
10-
#include <CL/sycl/exception.hpp>
1110
#include <CL/sycl/detail/type_traits.hpp>
11+
#include <CL/sycl/exception.hpp>
1212
#include <functional>
1313
#include <stdexcept>
1414

‎sycl/include/CL/sycl/detail/buffer_impl.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -13,9 +13,9 @@
1313
#include <CL/sycl/context.hpp>
1414
#include <CL/sycl/detail/aligned_allocator.hpp>
1515
#include <CL/sycl/detail/common.hpp>
16-
#include <CL/sycl/detail/pi.hpp>
1716
#include <CL/sycl/detail/helpers.hpp>
1817
#include <CL/sycl/detail/memory_manager.hpp>
18+
#include <CL/sycl/detail/pi.hpp>
1919
#include <CL/sycl/detail/scheduler/scheduler.hpp>
2020
#include <CL/sycl/detail/sycl_mem_obj_t.hpp>
2121
#include <CL/sycl/handler.hpp>

‎sycl/include/CL/sycl/detail/cg.hpp

Lines changed: 22 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -340,7 +340,8 @@ class CG {
340340
UPDATE_HOST,
341341
RUN_ON_HOST_INTEL,
342342
COPY_USM,
343-
FILL_USM
343+
FILL_USM,
344+
PREFETCH_USM
344345
};
345346

346347
CG(CGTYPE Type, std::vector<std::vector<char>> ArgsStorage,
@@ -529,6 +530,26 @@ class CGFillUSM : public CG {
529530
int getFill() { return MPattern[0]; }
530531
};
531532

533+
// The class which represents "prefetch" command group for USM pointers.
534+
class CGPrefetchUSM : public CG {
535+
void *MDst;
536+
size_t MLength;
537+
538+
public:
539+
CGPrefetchUSM(void *DstPtr, size_t Length,
540+
std::vector<std::vector<char>> ArgsStorage,
541+
std::vector<detail::AccessorImplPtr> AccStorage,
542+
std::vector<std::shared_ptr<const void>> SharedPtrStorage,
543+
std::vector<Requirement *> Requirements,
544+
std::vector<detail::EventImplPtr> Events)
545+
: CG(PREFETCH_USM, std::move(ArgsStorage), std::move(AccStorage),
546+
std::move(SharedPtrStorage), std::move(Requirements),
547+
std::move(Events)),
548+
MDst(DstPtr), MLength(Length) {}
549+
void *getDst() { return MDst; }
550+
size_t getLength() { return MLength; }
551+
};
552+
532553
} // namespace detail
533554
} // namespace sycl
534555
} // namespace cl

‎sycl/include/CL/sycl/detail/clusm.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -10,8 +10,8 @@
1010

1111
#include <CL/cl.h>
1212
#include <CL/cl_usm_ext.h>
13-
#include <CL/sycl/detail/pi.hpp>
1413
#include <CL/sycl/detail/os_util.hpp>
14+
#include <CL/sycl/detail/pi.hpp>
1515

1616
#include <map>
1717
#include <mutex>

‎sycl/include/CL/sycl/detail/device_impl.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,8 +8,8 @@
88

99
#pragma once
1010

11-
#include <CL/sycl/detail/pi.hpp>
1211
#include <CL/sycl/detail/device_info.hpp>
12+
#include <CL/sycl/detail/pi.hpp>
1313
#include <CL/sycl/stl.hpp>
1414
#include <algorithm>
1515
#include <memory>

‎sycl/include/CL/sycl/detail/event_impl.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,8 +9,8 @@
99
#pragma once
1010

1111
#include <CL/sycl/detail/common.hpp>
12-
#include <CL/sycl/detail/pi.hpp>
1312
#include <CL/sycl/detail/event_info.hpp>
13+
#include <CL/sycl/detail/pi.hpp>
1414
#include <CL/sycl/stl.hpp>
1515

1616
#include <cassert>

‎sycl/include/CL/sycl/detail/kernel_impl.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -10,8 +10,8 @@
1010

1111
#include <CL/sycl/context.hpp>
1212
#include <CL/sycl/detail/common.hpp>
13-
#include <CL/sycl/detail/pi.hpp>
1413
#include <CL/sycl/detail/kernel_info.hpp>
14+
#include <CL/sycl/detail/pi.hpp>
1515
#include <CL/sycl/device.hpp>
1616
#include <CL/sycl/info/info_desc.hpp>
1717

‎sycl/include/CL/sycl/detail/kernel_info.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,8 +9,8 @@
99
#pragma once
1010

1111
#include <CL/sycl/detail/common.hpp>
12-
#include <CL/sycl/detail/pi.hpp>
1312
#include <CL/sycl/detail/common_info.hpp>
13+
#include <CL/sycl/detail/pi.hpp>
1414
#include <CL/sycl/info/info_desc.hpp>
1515

1616
namespace cl {

‎sycl/include/CL/sycl/detail/memory_manager.hpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -129,6 +129,10 @@ class MemoryManager {
129129
int Pattern, std::vector<RT::PiEvent> DepEvents,
130130
RT::PiEvent &OutEvent);
131131

132+
static void prefetch_usm(void *Ptr, QueueImplPtr Queue, size_t Len,
133+
std::vector<RT::PiEvent> DepEvents,
134+
RT::PiEvent &OutEvent);
135+
132136
};
133137
} // namespace detail
134138
} // namespace sycl

‎sycl/include/CL/sycl/detail/os_util.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -10,8 +10,8 @@
1010

1111
#pragma once
1212

13-
#include <stdlib.h>
1413
#include <cstdint>
14+
#include <stdlib.h>
1515

1616
#ifdef _WIN32
1717
#define SYCL_RT_OS_WINDOWS

‎sycl/include/CL/sycl/detail/pi.hpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -10,9 +10,9 @@
1010
//
1111
#pragma once
1212

13+
#include <CL/sycl/detail/common.hpp>
1314
#include <CL/sycl/detail/os_util.hpp>
1415
#include <CL/sycl/detail/pi.h>
15-
#include <CL/sycl/detail/common.hpp>
1616

1717
namespace cl {
1818
namespace sycl {
@@ -67,8 +67,8 @@ namespace pi {
6767
To cast(From value);
6868

6969
// Forward declarations of the PI dispatch entries.
70-
#define _PI_API(api) __SYCL_EXPORTED extern decltype(::api) * api;
71-
#include <CL/sycl/detail/pi.def>
70+
#define _PI_API(api) __SYCL_EXPORTED extern decltype(::api) *(api);
71+
#include <CL/sycl/detail/pi.def>
7272

7373
// Performs PI one-time initialization.
7474
void initialize();

‎sycl/include/CL/sycl/detail/platform_impl.hpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -8,10 +8,10 @@
88

99
#pragma once
1010
#include <CL/sycl/detail/common.hpp>
11-
#include <CL/sycl/detail/pi.hpp>
1211
#include <CL/sycl/detail/force_device.hpp>
13-
#include <CL/sycl/info/info_desc.hpp>
12+
#include <CL/sycl/detail/pi.hpp>
1413
#include <CL/sycl/detail/platform_info.hpp>
14+
#include <CL/sycl/info/info_desc.hpp>
1515
#include <CL/sycl/stl.hpp>
1616

1717
// 4.6.2 Platform class

‎sycl/include/CL/sycl/detail/program_manager/program_manager.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,9 +8,9 @@
88

99
#pragma once
1010

11-
#include <CL/sycl/detail/pi.hpp>
1211
#include <CL/sycl/detail/common.hpp>
1312
#include <CL/sycl/detail/os_util.hpp>
13+
#include <CL/sycl/detail/pi.hpp>
1414
#include <CL/sycl/stl.hpp>
1515

1616
#include <map>

‎sycl/include/CL/sycl/detail/stream_impl.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -12,8 +12,8 @@
1212
#include <CL/sycl/builtins.hpp>
1313
#include <CL/sycl/detail/array.hpp>
1414
#include <CL/sycl/device_selector.hpp>
15-
#include <CL/sycl/queue.hpp>
1615
#include <CL/sycl/ordered_queue.hpp>
16+
#include <CL/sycl/queue.hpp>
1717

1818
namespace cl {
1919
namespace sycl {

‎sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -101,6 +101,11 @@ template <typename AllocatorT> class SYCLMemObjT : public SYCLMemObjI {
101101
}
102102

103103
size_t get_size() const { return MSizeInBytes; }
104+
size_t get_count() const {
105+
auto constexpr AllocatorValueSize =
106+
sizeof(allocator_value_type_t<AllocatorT>);
107+
return (get_size() + AllocatorValueSize - 1) / AllocatorValueSize;
108+
}
104109

105110
template <typename propertyT> bool has_property() const {
106111
return MProps.has_property<propertyT>();
@@ -112,16 +117,11 @@ template <typename AllocatorT> class SYCLMemObjT : public SYCLMemObjI {
112117

113118
AllocatorT get_allocator() const { return MAllocator; }
114119

115-
void *allocateHostMem() override {
116-
size_t AllocatorValueSize = sizeof(allocator_value_type_t<AllocatorT>);
117-
size_t AllocationSize = get_size() / AllocatorValueSize;
118-
AllocationSize += (get_size() % AllocatorValueSize) ? 1 : 0;
119-
return MAllocator.allocate(AllocationSize);
120-
}
120+
void *allocateHostMem() override { return MAllocator.allocate(get_count()); }
121121

122122
void releaseHostMem(void *Ptr) override {
123123
if (Ptr)
124-
MAllocator.deallocate(allocator_pointer_t<AllocatorT>(Ptr), get_size());
124+
MAllocator.deallocate(allocator_pointer_t<AllocatorT>(Ptr), get_count());
125125
}
126126

127127
void releaseMem(ContextImplPtr Context, void *MemAllocation) override {

‎sycl/include/CL/sycl/detail/usm_dispatch.hpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -49,6 +49,9 @@ class USMDispatcher {
4949
void *ParamValue, size_t *ParamValueSizeRet);
5050
void memAdvise(pi_queue Queue, const void *Ptr, size_t Length, int Advice,
5151
pi_event *Event);
52+
pi_result enqueuePrefetch(pi_queue Queue, void *Ptr, size_t Size,
53+
pi_uint32 NumEventsInWaitList,
54+
const pi_event *EventWaitList, pi_event *Event);
5255

5356
private:
5457
bool mEmulated = false;

‎sycl/include/CL/sycl/half_type.hpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -127,9 +127,9 @@ template <> struct hash<half> {
127127

128128
#define SYCL_HLF_MAX_EXP 16
129129

130-
#define SYCL_HLF_MIN_10_EXP -4
130+
#define SYCL_HLF_MIN_10_EXP (-4)
131131

132-
#define SYCL_HLF_MIN_EXP -13
132+
#define SYCL_HLF_MIN_EXP (-13)
133133

134134
#define SYCL_HLF_MANT_DIG 11
135135

0 commit comments

Comments
 (0)