Skip to content

Commit d543130

Browse files
author
Deepak Raj H R
authored
Merge branch 'oneapi-src:SYCLomatic' into 7284
2 parents a680f29 + da6345c commit d543130

File tree

15 files changed

+273
-64
lines changed

15 files changed

+273
-64
lines changed

‎clang/include/clang/Basic/DiagnosticDriverKinds.td‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -314,7 +314,9 @@ def err_arch_unsupported_isa
314314

315315
def err_drv_I_dash_not_supported : Error<
316316
"'%0' not supported, please use -iquote instead">;
317-
def err_drv_unknown_argument : Error<"unknown argument: '%0'">;
317+
// SYCLomatic_CUSTOMIZATION begin
318+
def err_drv_unknown_argument : Error<"dpct/c2s ignoring argument: '%0'">;
319+
// SYCLomatic_CUSTOMIZATION end
318320
def err_drv_unknown_argument_with_suggestion : Error<
319321
"unknown argument '%0'; did you mean '%1'?">;
320322
def warn_drv_unknown_argument_clang_cl : Warning<

‎clang/lib/DPCT/AnalysisInfo.cpp‎

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -5036,9 +5036,12 @@ void DeviceFunctionDecl::buildTextureObjectParamsInfo(
50365036
return;
50375037
for (unsigned Idx = 0; Idx < Parms.size(); ++Idx) {
50385038
auto Param = Parms[Idx];
5039-
if (DpctGlobalInfo::getUnqualifiedTypeName(Param->getType()) ==
5040-
"cudaTextureObject_t")
5039+
std::string ParamName =
5040+
DpctGlobalInfo::getUnqualifiedTypeName(Param->getType());
5041+
if (ParamName == "cudaTextureObject_t" ||
5042+
ParamName == "cudaSurfaceObject_t") {
50415043
TextureObjectList[Idx] = std::make_shared<TextureObjectInfo>(Param);
5044+
}
50425045
}
50435046
}
50445047
std::string DeviceFunctionDecl::getExtraParameters(LocInfo LI) {

‎clang/lib/DPCT/AnalysisInfo.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2054,7 +2054,7 @@ class TextureInfo {
20542054
bool isUseHelperFunc() { return true; }
20552055
};
20562056

2057-
// texture handle info
2057+
// texture object info can be used for CUDA texture and surface objects.
20582058
class TextureObjectInfo : public TextureInfo {
20592059
static const int ReplaceTypeLength;
20602060

‎clang/lib/DPCT/RulesLang/APINamesTexture.inc‎

Lines changed: 112 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -120,48 +120,101 @@ CONDITIONAL_FACTORY_ENTRY(
120120
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
121121
"cudaGetSurfaceObjectResourceDesc", DEREF(0),
122122
MEMBER_CALL(ARG(1), true, "get_data")))))
123-
124-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
125-
ENTRY_TEXTURE("tex1D", 0x01, 1))
123+
CONDITIONAL_FACTORY_ENTRY(
124+
UseExtBindlessImages,
125+
FEATURE_REQUEST_FACTORY(
126+
HelperFeatureEnum::device_ext,
127+
ENTRY_TEXTURE("tex1D", 0x01,
128+
MapNames::getClNamespace() +
129+
"ext::oneapi::experimental::sample_image",
130+
1)),
131+
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
132+
ENTRY_TEXTURE("tex1D", 0x01, "read", 1)))
126133

127134
CONDITIONAL_FACTORY_ENTRY(
128-
UseExtBindlessImages, ENTRY_TEXTURE("tex1DLod", 0x11, 1),
135+
UseExtBindlessImages,
136+
ENTRY_TEXTURE("tex1DLod", 0x11,
137+
MapNames::getClNamespace() +
138+
"ext::oneapi::experimental::sample_mipmap",
139+
1),
129140
UNSUPPORT_FACTORY_ENTRY("tex1DLod", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
130141
ARG("tex1DLod"),
131142
ARG("--use-experimental-features=bindless_images")))
132-
133-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
134-
ENTRY_TEXTURE("tex2D", 0x02, 1, 2))
143+
CONDITIONAL_FACTORY_ENTRY(
144+
UseExtBindlessImages,
145+
FEATURE_REQUEST_FACTORY(
146+
HelperFeatureEnum::device_ext,
147+
ENTRY_TEXTURE("tex2D", 0x02,
148+
MapNames::getClNamespace() +
149+
"ext::oneapi::experimental::sample_image",
150+
1, 2)),
151+
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
152+
ENTRY_TEXTURE("tex2D", 0x02, "read", 1, 2)))
135153

136154
CONDITIONAL_FACTORY_ENTRY(
137-
UseExtBindlessImages, ENTRY_TEXTURE("tex2DLod", 0x12, 1, 2),
155+
UseExtBindlessImages,
156+
ENTRY_TEXTURE("tex2DLod", 0x12,
157+
MapNames::getClNamespace() +
158+
"ext::oneapi::experimental::sample_mipmap",
159+
1, 2),
138160
UNSUPPORT_FACTORY_ENTRY("tex2DLod", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
139161
ARG("tex2DLod"),
140162
ARG("--use-experimental-features=bindless_images")))
141163

142-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
143-
ENTRY_TEXTURE("tex3D", 0x03, 1, 2, 3))
164+
FEATURE_REQUEST_FACTORY(
165+
HelperFeatureEnum::device_ext,
166+
ENTRY_TEXTURE("tex3D", 0x03,
167+
MapNames::getClNamespace() +
168+
"ext::oneapi::experimental::sample_image",
169+
1, 2, 3))
144170

145171
CONDITIONAL_FACTORY_ENTRY(
146-
UseExtBindlessImages, ENTRY_TEXTURE("tex3DLod", 0x13, 1, 2, 3),
172+
UseExtBindlessImages,
173+
ENTRY_TEXTURE("tex3DLod", 0x13,
174+
MapNames::getClNamespace() +
175+
"ext::oneapi::experimental::sample_mipmap",
176+
1, 2, 3),
147177
UNSUPPORT_FACTORY_ENTRY("tex3DLod", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
148178
ARG("tex3DLod"),
149179
ARG("--use-experimental-features=bindless_images")))
150180

151-
FEATURE_REQUEST_FACTORY(
152-
HelperFeatureEnum::device_ext,
153-
WARNING_FACTORY_ENTRY("tex1Dfetch", ENTRY_TEXTURE("tex1Dfetch", 0x01, 1),
154-
Diagnostics::TEX_FETCH))
155-
CONDITIONAL_FACTORY_ENTRY(UseExtBindlessImages,
156-
ENTRY_TEXTURE("tex1DLayered", 0xF1, 1),
157-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
158-
ENTRY_TEXTURE("tex1DLayered",
159-
0xF1, 2, 1)))
160-
CONDITIONAL_FACTORY_ENTRY(UseExtBindlessImages,
161-
ENTRY_TEXTURE("tex2DLayered", 0xF2, 1, 2),
162-
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
163-
ENTRY_TEXTURE("tex2DLayered",
164-
0xF2, 3, 1, 2)))
181+
CONDITIONAL_FACTORY_ENTRY(
182+
UseExtBindlessImages,
183+
FEATURE_REQUEST_FACTORY(
184+
HelperFeatureEnum::device_ext,
185+
WARNING_FACTORY_ENTRY(
186+
"tex1Dfetch",
187+
ENTRY_TEXTURE("tex1Dfetch", 0x01,
188+
MapNames::getClNamespace() +
189+
"ext::oneapi::experimental::sample_image",
190+
1),
191+
Diagnostics::TEX_FETCH)),
192+
FEATURE_REQUEST_FACTORY(
193+
HelperFeatureEnum::device_ext,
194+
WARNING_FACTORY_ENTRY("tex1Dfetch",
195+
ENTRY_TEXTURE("tex1Dfetch", 0x01, "read", 1),
196+
Diagnostics::TEX_FETCH)))
197+
198+
CONDITIONAL_FACTORY_ENTRY(
199+
UseExtBindlessImages,
200+
FEATURE_REQUEST_FACTORY(
201+
HelperFeatureEnum::device_ext,
202+
ENTRY_TEXTURE("tex1DLayered", 0xF1,
203+
MapNames::getClNamespace() +
204+
"ext::oneapi::experimental::sample_image_array",
205+
1)),
206+
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
207+
ENTRY_TEXTURE("tex1DLayered", 0xF1, "read", 2, 1)))
208+
209+
CONDITIONAL_FACTORY_ENTRY(
210+
UseExtBindlessImages,
211+
ENTRY_TEXTURE("tex2DLayered", 0xF2,
212+
MapNames::getClNamespace() +
213+
"ext::oneapi::experimental::sample_image_array",
214+
1, 2),
215+
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
216+
ENTRY_TEXTURE("tex2DLayered", 0xF2, "read", 3, 1,
217+
2)))
165218

166219
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
167220
ASSIGNABLE_FACTORY(ENTRY_BIND("cudaBindTexture",
@@ -287,3 +340,37 @@ FEATURE_REQUEST_FACTORY(
287340
"cuTexRefSetFormat", true, false, true, true,
288341
MEMBER_CALL(ARG(0), true, "set_channel_type", ARG(1)),
289342
MEMBER_CALL(ARG(0), true, "set_channel_num", ARG(2))))
343+
344+
CONDITIONAL_FACTORY_ENTRY(
345+
UseExtBindlessImages,
346+
FEATURE_REQUEST_FACTORY(
347+
HelperFeatureEnum::device_ext,
348+
ENTRY_TEXTURE("surf1Dread", 0x01,
349+
MapNames::getLibraryHelperNamespace() +
350+
"experimental::sample_image_by_byte",
351+
1)),
352+
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
353+
ENTRY_TEXTURE("surf1Dread", 0x01, "read_byte", 1)))
354+
355+
CONDITIONAL_FACTORY_ENTRY(
356+
UseExtBindlessImages,
357+
FEATURE_REQUEST_FACTORY(
358+
HelperFeatureEnum::device_ext,
359+
ENTRY_TEXTURE("surf2Dread", 0x02,
360+
MapNames::getLibraryHelperNamespace() +
361+
"experimental::sample_image_by_byte",
362+
1, 2)),
363+
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
364+
ENTRY_TEXTURE("surf2Dread", 0x02, "read_byte", 1,
365+
2)))
366+
CONDITIONAL_FACTORY_ENTRY(
367+
UseExtBindlessImages,
368+
FEATURE_REQUEST_FACTORY(
369+
HelperFeatureEnum::device_ext,
370+
ENTRY_TEXTURE("surf3Dread", 0x03,
371+
MapNames::getLibraryHelperNamespace() +
372+
"experimental::sample_image_by_byte",
373+
1, 2, 3)),
374+
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
375+
ENTRY_TEXTURE("surf3Dread", 0x03, "read_byte", 1, 2,
376+
3)))

‎clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp‎

Lines changed: 26 additions & 29 deletions
Original file line numberDiff line numberDiff line change
@@ -16,25 +16,25 @@ template <size_t... Idx>
1616
class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
1717
std::string Source;
1818
int TexType;
19+
std::string TargetName;
1920

2021
inline int getDim() const { return TexType & 0x0f; }
2122

2223
template <class BaseT>
2324
std::shared_ptr<CallExprRewriter>
2425
createRewriter(const CallExpr *C, bool RetAssign, BaseT Base) const {
25-
const static std::string MemberName = "read";
2626
using ReaderPrinter = decltype(makeMemberCallCreator<false>(
2727
std::declval<std::function<BaseT(const CallExpr *)>>(), false,
28-
MemberName, makeCallArgCreatorWithCall(Idx)...)(C));
28+
TargetName, makeCallArgCreatorWithCall(Idx)...)(C));
2929
if (RetAssign) {
3030
return std::make_shared<PrinterRewriter<
3131
BinaryOperatorPrinter<BO_Assign, DerefExpr, ReaderPrinter>>>(
3232
C, Source, DerefExpr(C->getArg(0), C),
33-
ReaderPrinter(std::move(Base), false, MemberName,
33+
ReaderPrinter(std::move(Base), false, TargetName,
3434
std::make_pair(C, C->getArg(Idx + 1))...));
3535
}
3636
return std::make_shared<PrinterRewriter<ReaderPrinter>>(
37-
C, Source, Base, false, MemberName,
37+
C, Source, Base, false, TargetName,
3838
std::make_pair(C, C->getArg(Idx))...);
3939
}
4040

@@ -43,8 +43,6 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
4343
createbindlessRewriterNormal(const CallExpr *C, bool RetAssign,
4444
const TemplateArgumentInfo &TAI,
4545
const std::string &VecTypeName) const {
46-
const static std::string FuncName =
47-
MapNames::getClNamespace() + "ext::oneapi::experimental::sample_image";
4846
using FuncNamePrinter =
4947
TemplatedNamePrinter<StringRef, std::vector<TemplateArgumentInfo>>;
5048
using ReaderPrinter =
@@ -55,11 +53,12 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
5553
BinaryOperatorPrinter<BO_Assign, DerefExpr, ReaderPrinter>>>(
5654
C, Source, DerefExpr(C->getArg(0), C),
5755
ReaderPrinter(
58-
FuncNamePrinter(FuncName, {TAI}), std::make_pair(C, C->getArg(1)),
56+
FuncNamePrinter(TargetName, {TAI}),
57+
std::make_pair(C, C->getArg(1)),
5958
VecType(VecTypeName, std::make_pair(C, C->getArg(Idx + 1))...)));
6059
}
6160
return std::make_shared<PrinterRewriter<ReaderPrinter>>(
62-
C, Source, FuncNamePrinter(FuncName, {TAI}),
61+
C, Source, FuncNamePrinter(TargetName, {TAI}),
6362
std::make_pair(C, C->getArg(0)),
6463
VecType(VecTypeName, std::make_pair(C, C->getArg(Idx))...));
6564
}
@@ -69,8 +68,6 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
6968
createbindlessRewriterLod(const CallExpr *C, bool RetAssign,
7069
const TemplateArgumentInfo &TAI,
7170
const std::string &VecTypeName) const {
72-
const static std::string FuncName =
73-
MapNames::getClNamespace() + "ext::oneapi::experimental::sample_mipmap";
7471
using FuncNamePrinter =
7572
TemplatedNamePrinter<StringRef, std::vector<TemplateArgumentInfo>>;
7673
using ReaderPrinter =
@@ -82,12 +79,13 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
8279
BinaryOperatorPrinter<BO_Assign, DerefExpr, ReaderPrinter>>>(
8380
C, Source, DerefExpr(C->getArg(0), C),
8481
ReaderPrinter(
85-
FuncNamePrinter(FuncName, {TAI}), std::make_pair(C, C->getArg(1)),
82+
FuncNamePrinter(TargetName, {TAI}),
83+
std::make_pair(C, C->getArg(1)),
8684
VecType(VecTypeName, std::make_pair(C, C->getArg(Idx + 1))...),
8785
std::make_pair(C, C->getArg(C->getNumArgs() - 1))));
8886
}
8987
return std::make_shared<PrinterRewriter<ReaderPrinter>>(
90-
C, Source, FuncNamePrinter(FuncName, {TAI}),
88+
C, Source, FuncNamePrinter(TargetName, {TAI}),
9189
std::make_pair(C, C->getArg(0)),
9290
VecType(VecTypeName, std::make_pair(C, C->getArg(Idx))...),
9391
std::make_pair(C, C->getArg(C->getNumArgs() - 1)));
@@ -98,9 +96,6 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
9896
createbindlessRewriterLayered(const CallExpr *C, bool RetAssign,
9997
const TemplateArgumentInfo &TAI,
10098
const std::string &VecTypeName) const {
101-
const static std::string FuncName =
102-
MapNames::getClNamespace() +
103-
"ext::oneapi::experimental::sample_image_array";
10499
using FuncNamePrinter =
105100
TemplatedNamePrinter<StringRef, std::vector<TemplateArgumentInfo>>;
106101
using ReaderPrinter =
@@ -112,17 +107,17 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
112107
BinaryOperatorPrinter<BO_Assign, DerefExpr, ReaderPrinter>>>(
113108
C, Source, DerefExpr(C->getArg(0), C),
114109
ReaderPrinter(
115-
FuncNamePrinter(FuncName, {TAI}), std::make_pair(C, C->getArg(1)),
110+
FuncNamePrinter(TargetName, {TAI}),
111+
std::make_pair(C, C->getArg(1)),
116112
VecType(VecTypeName, std::make_pair(C, C->getArg(Idx + 1))...),
117113
std::make_pair(C, C->getArg(C->getNumArgs() - 1))));
118114
}
119115
return std::make_shared<PrinterRewriter<ReaderPrinter>>(
120-
C, Source, FuncNamePrinter(FuncName, {TAI}),
116+
C, Source, FuncNamePrinter(TargetName, {TAI}),
121117
std::make_pair(C, C->getArg(0)),
122118
VecType(VecTypeName, std::make_pair(C, C->getArg(Idx))...),
123119
std::make_pair(C, C->getArg(C->getNumArgs() - 1)));
124120
}
125-
126121
std::shared_ptr<CallExprRewriter>
127122
createbindlessRewriter(const CallExpr *C, bool RetAssign,
128123
QualType TargetType) const {
@@ -160,8 +155,8 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
160155
}
161156

162157
public:
163-
TextureReadRewriterFactory(std::string Name, int Tex)
164-
: Source(std::move(Name)), TexType(Tex) {}
158+
TextureReadRewriterFactory(std::string Name, int Tex, std::string TargetName)
159+
: Source(std::move(Name)), TexType(Tex), TargetName(TargetName) {}
165160
std::shared_ptr<CallExprRewriter>
166161
create(const CallExpr *Call) const override {
167162
const Expr *SourceExpr = Call->getArg(0);
@@ -227,14 +222,15 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
227222
/// original call expr.
228223
template <size_t... Idx>
229224
inline std::shared_ptr<CallExprRewriterFactoryBase>
230-
createTextureReaderRewriterFactory(const std::string &Source, int TextureType) {
231-
return std::make_shared<TextureReadRewriterFactory<Idx...>>(Source,
232-
TextureType);
225+
createTextureReaderRewriterFactory(const std::string &Source, int TextureType,
226+
const std::string &TargetName) {
227+
return std::make_shared<TextureReadRewriterFactory<Idx...>>(
228+
Source, TextureType, TargetName);
233229
}
234230

235-
#define TEX_FUNCTION_FACTORY_ENTRY(FuncName, TexType, ...) \
236-
{FuncName, \
237-
createTextureReaderRewriterFactory<__VA_ARGS__>(FuncName, TexType)},
231+
#define TEX_FUNCTION_FACTORY_ENTRY(FuncName, TexType, TargetName, ...) \
232+
{FuncName, createTextureReaderRewriterFactory<__VA_ARGS__>( \
233+
FuncName, TexType, TargetName)},
238234
#define BIND_TEXTURE_FACTORY_ENTRY(FuncName, ...) \
239235
{FuncName, createBindTextureRewriterFactory<__VA_ARGS__>(FuncName)},
240236

@@ -244,16 +240,17 @@ createTextureReaderRewriterFactory(const std::string &Source, int TextureType) {
244240
REWRITER_FACTORY_ENTRY(FuncName, FuncCallExprRewriterFactory, RewriterName)
245241
#define UNSUPPORTED_FACTORY_ENTRY(FuncName, MsgID) \
246242
REWRITER_FACTORY_ENTRY(FuncName, \
247-
UnsupportFunctionRewriterFactory<std::string>, MsgID, FuncName)
243+
UnsupportFunctionRewriterFactory<std::string>, MsgID, \
244+
FuncName)
248245

249246
void CallExprRewriterFactoryBase::initRewriterMapTexture() {
250247
RewriterMap->merge(
251248
std::unordered_map<std::string,
252249
std::shared_ptr<CallExprRewriterFactoryBase>>({
253250
#define ENTRY_RENAMED(SOURCEAPINAME, TARGETAPINAME) \
254251
FUNC_NAME_FACTORY_ENTRY(SOURCEAPINAME, TARGETAPINAME)
255-
#define ENTRY_TEXTURE(SOURCEAPINAME, TEXTYPE, ...) \
256-
TEX_FUNCTION_FACTORY_ENTRY(SOURCEAPINAME, TEXTYPE, __VA_ARGS__)
252+
#define ENTRY_TEXTURE(SOURCEAPINAME, TEXTYPE, TARGETAPINAME, ...) \
253+
TEX_FUNCTION_FACTORY_ENTRY(SOURCEAPINAME, TEXTYPE, TARGETAPINAME, __VA_ARGS__)
257254
#define ENTRY_UNSUPPORTED(SOURCEAPINAME, MSGID) \
258255
UNSUPPORTED_FACTORY_ENTRY(SOURCEAPINAME, MSGID)
259256
#define ENTRY_BIND(SOURCEAPINAME, ...) \

‎clang/lib/DPCT/RulesLang/RulesLangTexture.cpp‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -560,6 +560,9 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
560560
"tex1Dfetch",
561561
"tex1DLayered",
562562
"tex2DLayered",
563+
"surf1Dread",
564+
"surf2Dread",
565+
"surf3Dread",
563566
"cudaCreateTextureObject",
564567
"cudaDestroyTextureObject",
565568
"cudaGetTextureObjectResourceDesc",

‎clang/lib/DPCT/SrcAPI/APINames.inc‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1484,11 +1484,11 @@ ENTRY(texCubemapLod, texCubemapLod, false, NO_FLAG, P4, "comment")
14841484
ENTRY(texCubemapLayered, texCubemapLayered, false, NO_FLAG, P4, "comment")
14851485
ENTRY(texCubemapLayeredLod, texCubemapLayeredLod, false, NO_FLAG, P4, "comment")
14861486
ENTRY(tex2Dgather, tex2Dgather, false, NO_FLAG, P0, "comment")
1487-
ENTRY(surf1Dread, surf1Dread, false, NO_FLAG, P4, "comment")
1487+
ENTRY(surf1Dread, surf1Dread, true, NO_FLAG, P4, "Successful")
14881488
ENTRY(surf1Dwrite, surf1Dwrite, false, NO_FLAG, P0, "comment")
1489-
ENTRY(surf2Dread, surf2Dread, false, NO_FLAG, P4, "comment")
1489+
ENTRY(surf2Dread, surf2Dread, true, NO_FLAG, P4, "Successful")
14901490
ENTRY(surf2Dwrite, surf2Dwrite, false, NO_FLAG, P0, "comment")
1491-
ENTRY(surf3Dread, surf3Dread, false, NO_FLAG, P4, "comment")
1491+
ENTRY(surf3Dread, surf3Dread, true, NO_FLAG, P4, "Successful")
14921492
ENTRY(surf3Dwrite, surf3Dwrite, false, NO_FLAG, P0, "comment")
14931493
ENTRY(surf1DLayeredread, surf1DLayeredread, false, NO_FLAG, P4, "comment")
14941494
ENTRY(surf1DLayeredwrite, surf1DLayeredwrite, false, NO_FLAG, P4, "comment")

‎clang/lib/Headers/__clang_cuda_runtime_wrapper.h‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -425,6 +425,9 @@ __host__ __device__ void __nv_tex_surf_handler(const char *name, T *ptr,
425425
#endif // __cplusplus >= 201103L && CUDA_VERSION >= 9000
426426
#include "texture_fetch_functions.h"
427427
#include "texture_indirect_functions.h"
428+
#ifdef SYCLomatic_CUSTOMIZATION
429+
#include "surface_indirect_functions.h"
430+
#endif // SYCLomatic_CUSTOMIZATION
428431

429432
// Restore state of __CUDA_ARCH__ and __THROW we had on entry.
430433
#pragma pop_macro("__CUDA_ARCH__")

0 commit comments

Comments
 (0)