Skip to content

Commit

Permalink
Merge remote-tracking branch 'origin/SYCLomatic' into support_more_ge…
Browse files Browse the repository at this point in the history 
…mm_batch
  • Loading branch information
@zhiweij1
zhiweij1 committed Jul 9, 2024
2 parents 258e107 + 398623b commit b8de47d
Show file tree
Hide file tree
Showing 70 changed files with 1,199 additions and 912 deletions.
2 changes: 2 additions & 0 deletions clang/examples/DPCT/Runtime/cudaStreamIsCapturing.cu
View file
Original file line number Diff line number Diff line change
@@ -1,3 +1,5 @@
// Option: --use-experimental-features=graph

void test(cudaStream_t s, enum cudaStreamCaptureStatus *ps) {
// Start
cudaStreamIsCapturing(s /*cudaStream_t*/,
Expand Down
3 changes: 2 additions & 1 deletion clang/lib/DPCT/APINames.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -102,7 +102,7 @@ ENTRY(cudaStreamGetCaptureInfo_v3, cudaStreamGetCaptureInfo_v3, false, NO_FLAG,
ENTRY(cudaStreamGetFlags, cudaStreamGetFlags, true, NO_FLAG, P4, "DPCT1014")
ENTRY(cudaStreamGetId, cudaStreamGetId, false, NO_FLAG, P4, "comment")
ENTRY(cudaStreamGetPriority, cudaStreamGetPriority, true, NO_FLAG, P4, "DPCT1014")
ENTRY(cudaStreamIsCapturing, cudaStreamIsCapturing, true, API_CALL_REMOVED, P4, "DPCT1026/DPCT1027")
ENTRY(cudaStreamIsCapturing, cudaStreamIsCapturing, true, NO_FLAG, P4, "Successful/DPCT1119")
ENTRY(cudaStreamQuery, cudaStreamQuery, true, API_CALL_REMOVED, P0, "DPCT1026/DPCT1027")
ENTRY(cudaStreamSetAttribute, cudaStreamSetAttribute, true, NO_FLAG, P7, "DPCT1026/DPCT1027")
ENTRY(cudaStreamSynchronize, cudaStreamSynchronize, true, NO_FLAG, P0, "Successful")
Expand Down Expand Up @@ -1429,6 +1429,7 @@ ENTRY(make_ushort4, make_ushort4, true, NO_FLAG, P0, "Successful")
ENTRY(__threadfence_block, __threadfence_block, true, NO_FLAG, P0, "Successful: DPCT1078")
ENTRY(__threadfence, __threadfence, true, NO_FLAG, P0, "Successful: DPCT1078")
ENTRY(__threadfence_system, __threadfence_system, true, NO_FLAG, P0, "Successful: DPCT1078")
ENTRY(__barrier_sync, __barrier_sync, true, NO_FLAG, P0, "Successful: DPCT1078")
ENTRY(__syncthreads, __syncthreads, true, NO_FLAG, P0, "Successful: DPCT1078")
ENTRY(__syncthreads_count, __syncthreads_count, true, NO_FLAG, P0, "Successful: DPCT1078")
ENTRY(__syncthreads_and, __syncthreads_and, true, NO_FLAG, P0, "Successful: DPCT1078")
Expand Down
10 changes: 6 additions & 4 deletions clang/lib/DPCT/APINamesCooperativeGroups.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -970,10 +970,12 @@ MEMBER_CALL_FACTORY_ENTRY("cooperative_groups::__v1::thread_group.num_threads",
MemberExprBase(), false, "get_local_linear_range")
MEMBER_CALL_FACTORY_ENTRY("cooperative_groups::__v1::thread_group.get_type",
MemberExprBase(), false, "get_type")
MEMBER_CALL_FACTORY_ENTRY("cooperative_groups::__v1::thread_block.group_index", MemberExprBase(),
false, "get_group_id")
MEMBER_CALL_FACTORY_ENTRY("cooperative_groups::__v1::thread_block.thread_index", MemberExprBase(),
false, "get_local_id")
CALL_FACTORY_ENTRY("cooperative_groups::__v1::thread_block.group_index",
CALL(MapNames::getDpctNamespace() + "dim3",
MEMBER_CALL(MemberExprBase(), false, "get_group_id")))
CALL_FACTORY_ENTRY("cooperative_groups::__v1::thread_block.thread_index",
CALL(MapNames::getDpctNamespace() + "dim3",
MEMBER_CALL(MemberExprBase(), false, "get_local_id")))

CONDITIONAL_FACTORY_ENTRY(
UseNonUniformGroups,
Expand Down
259 changes: 160 additions & 99 deletions clang/lib/DPCT/APINamesMemory.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -417,123 +417,184 @@ ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(

ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
HelperFeatureEnum::device_ext,
CALL_FACTORY_ENTRY(
"cuMemcpyAtoH_v2",
CALL(MapNames::getDpctNamespace() + "dpct_memcpy",
makeCallArgCreatorWithCall(0),
BO(BinaryOperatorKind::BO_Add,
makeCastExprCreator(
makeCharPtrCreator(),
MEMBER_CALL(MEMBER_CALL(makeCallArgCreatorWithCall(1), true,
"to_pitched_data"),
false, "get_data_ptr"),
true),
makeCallArgCreatorWithCall(2)),
makeCallArgCreatorWithCall(3)))))
CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
CALL_FACTORY_ENTRY(
"cuMemcpyAtoH_v2",
CALL(MapNames::getDpctNamespace() + "experimental::dpct_memcpy",
makeCallArgCreatorWithCall(0), makeCallArgCreatorWithCall(1),
makeCallArgCreatorWithCall(2), ARG("0"),
makeCallArgCreatorWithCall(3))),
CALL_FACTORY_ENTRY(
"cuMemcpyAtoH_v2",
CALL(MapNames::getDpctNamespace() + "dpct_memcpy",
makeCallArgCreatorWithCall(0),
BO(BinaryOperatorKind::BO_Add,
makeCastExprCreator(
makeCharPtrCreator(),
MEMBER_CALL(MEMBER_CALL(makeCallArgCreatorWithCall(1),
true, "to_pitched_data"),
false, "get_data_ptr"),
true),
makeCallArgCreatorWithCall(2)),
makeCallArgCreatorWithCall(3))))))

ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
HelperFeatureEnum::device_ext,
CALL_FACTORY_ENTRY(
"cuMemcpyHtoA_v2",
CALL(MapNames::getDpctNamespace() + "dpct_memcpy",
BO(BinaryOperatorKind::BO_Add,
makeCastExprCreator(
makeCharPtrCreator(),
MEMBER_CALL(MEMBER_CALL(makeCallArgCreatorWithCall(0), true,
"to_pitched_data"),
false, "get_data_ptr"),
true),
makeCallArgCreatorWithCall(1)),
makeCallArgCreatorWithCall(2), makeCallArgCreatorWithCall(3)))))
CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
CALL_FACTORY_ENTRY(
"cuMemcpyHtoA_v2",
CALL(MapNames::getDpctNamespace() + "experimental::dpct_memcpy",
makeCallArgCreatorWithCall(0), makeCallArgCreatorWithCall(1),
ARG("0"), makeCallArgCreatorWithCall(2),
makeCallArgCreatorWithCall(3))),
CALL_FACTORY_ENTRY(
"cuMemcpyHtoA_v2",
CALL(MapNames::getDpctNamespace() + "dpct_memcpy",
BO(BinaryOperatorKind::BO_Add,
makeCastExprCreator(
makeCharPtrCreator(),
MEMBER_CALL(MEMBER_CALL(makeCallArgCreatorWithCall(0),
true, "to_pitched_data"),
false, "get_data_ptr"),
true),
makeCallArgCreatorWithCall(1)),
makeCallArgCreatorWithCall(2),
makeCallArgCreatorWithCall(3))))))

ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
HelperFeatureEnum::device_ext,
CALL_FACTORY_ENTRY(
"cuMemcpyAtoHAsync_v2",
CALL(MapNames::getDpctNamespace() + "async_dpct_memcpy",
makeCallArgCreatorWithCall(0),
BO(BinaryOperatorKind::BO_Add,
makeCastExprCreator(
makeCharPtrCreator(),
MEMBER_CALL(MEMBER_CALL(makeCallArgCreatorWithCall(1), true,
"to_pitched_data"),
false, "get_data_ptr"),
true),
makeCallArgCreatorWithCall(2)),
makeCallArgCreatorWithCall(3),
ARG(MapNames::getDpctNamespace() + "automatic"),
DEREF(makeCallArgCreatorWithCall(4))))))
CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
CALL_FACTORY_ENTRY("cuMemcpyAtoHAsync_v2",
CALL(MapNames::getDpctNamespace() +
"experimental::async_dpct_memcpy",
makeCallArgCreatorWithCall(0),
makeCallArgCreatorWithCall(1),
makeCallArgCreatorWithCall(2), ARG("0"),
makeCallArgCreatorWithCall(3), STREAM(4))),
CALL_FACTORY_ENTRY(
"cuMemcpyAtoHAsync_v2",
CALL(MapNames::getDpctNamespace() + "async_dpct_memcpy",
makeCallArgCreatorWithCall(0),
BO(BinaryOperatorKind::BO_Add,
makeCastExprCreator(
makeCharPtrCreator(),
MEMBER_CALL(MEMBER_CALL(makeCallArgCreatorWithCall(1),
true, "to_pitched_data"),
false, "get_data_ptr"),
true),
makeCallArgCreatorWithCall(2)),
makeCallArgCreatorWithCall(3),
ARG(MapNames::getDpctNamespace() + "automatic"),
DEREF(makeCallArgCreatorWithCall(4)))))))

ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
HelperFeatureEnum::device_ext,
CALL_FACTORY_ENTRY(
"cuMemcpyHtoAAsync_v2",
CALL(MapNames::getDpctNamespace() + "async_dpct_memcpy",
BO(BinaryOperatorKind::BO_Add,
makeCastExprCreator(
makeCharPtrCreator(),
MEMBER_CALL(MEMBER_CALL(makeCallArgCreatorWithCall(0), true,
"to_pitched_data"),
false, "get_data_ptr"),
true),
makeCallArgCreatorWithCall(1)),
makeCallArgCreatorWithCall(2), makeCallArgCreatorWithCall(3),
ARG(MapNames::getDpctNamespace() + "automatic"),
DEREF(makeCallArgCreatorWithCall(4))))))
CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
CALL_FACTORY_ENTRY("cuMemcpyHtoAAsync_v2",
CALL(MapNames::getDpctNamespace() +
"experimental::async_dpct_memcpy",
makeCallArgCreatorWithCall(0),
makeCallArgCreatorWithCall(1), ARG("0"),
makeCallArgCreatorWithCall(2),
makeCallArgCreatorWithCall(3), STREAM(4))),
CALL_FACTORY_ENTRY(
"cuMemcpyHtoAAsync_v2",
CALL(MapNames::getDpctNamespace() + "async_dpct_memcpy",
BO(BinaryOperatorKind::BO_Add,
makeCastExprCreator(
makeCharPtrCreator(),
MEMBER_CALL(MEMBER_CALL(makeCallArgCreatorWithCall(0),
true, "to_pitched_data"),
false, "get_data_ptr"),
true),
makeCallArgCreatorWithCall(1)),
makeCallArgCreatorWithCall(2), makeCallArgCreatorWithCall(3),
ARG(MapNames::getDpctNamespace() + "automatic"),
DEREF(makeCallArgCreatorWithCall(4)))))))

ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
HelperFeatureEnum::device_ext,
CALL_FACTORY_ENTRY(
"cuMemcpyAtoD_v2",
CALL(MapNames::getDpctNamespace() + "dpct_memcpy",
makeCallArgCreatorWithCall(0),
BO(BinaryOperatorKind::BO_Add,
makeCastExprCreator(
makeCharPtrCreator(),
MEMBER_CALL(MEMBER_CALL(makeCallArgCreatorWithCall(1), true,
"to_pitched_data"),
false, "get_data_ptr"),
true),
makeCallArgCreatorWithCall(2)),
makeCallArgCreatorWithCall(3)))))
CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
CALL_FACTORY_ENTRY(
"cuMemcpyAtoD_v2",
CALL(MapNames::getDpctNamespace() + "experimental::dpct_memcpy",
makeCallArgCreatorWithCall(0), makeCallArgCreatorWithCall(1),
makeCallArgCreatorWithCall(2), ARG("0"),
makeCallArgCreatorWithCall(3))),
CALL_FACTORY_ENTRY(
"cuMemcpyAtoD_v2",
CALL(MapNames::getDpctNamespace() + "dpct_memcpy",
makeCallArgCreatorWithCall(0),
BO(BinaryOperatorKind::BO_Add,
makeCastExprCreator(
makeCharPtrCreator(),
MEMBER_CALL(MEMBER_CALL(makeCallArgCreatorWithCall(1),
true, "to_pitched_data"),
false, "get_data_ptr"),
true),
makeCallArgCreatorWithCall(2)),
makeCallArgCreatorWithCall(3))))))

ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
HelperFeatureEnum::device_ext,
CALL_FACTORY_ENTRY(
"cuMemcpyDtoA_v2",
CALL(MapNames::getDpctNamespace() + "dpct_memcpy",
BO(BinaryOperatorKind::BO_Add,
makeCastExprCreator(
makeCharPtrCreator(),
MEMBER_CALL(MEMBER_CALL(makeCallArgCreatorWithCall(0), true,
"to_pitched_data"),
false, "get_data_ptr"),
true),
makeCallArgCreatorWithCall(1)),
makeCallArgCreatorWithCall(2), makeCallArgCreatorWithCall(3)))))
CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
CALL_FACTORY_ENTRY(
"cuMemcpyDtoA_v2",
CALL(MapNames::getDpctNamespace() + "experimental::dpct_memcpy",
makeCallArgCreatorWithCall(0), makeCallArgCreatorWithCall(1),
ARG("0"), makeCallArgCreatorWithCall(2),
makeCallArgCreatorWithCall(3))),
CALL_FACTORY_ENTRY(
"cuMemcpyDtoA_v2",
CALL(MapNames::getDpctNamespace() + "dpct_memcpy",
BO(BinaryOperatorKind::BO_Add,
makeCastExprCreator(
makeCharPtrCreator(),
MEMBER_CALL(MEMBER_CALL(makeCallArgCreatorWithCall(0),
true, "to_pitched_data"),
false, "get_data_ptr"),
true),
makeCallArgCreatorWithCall(1)),
makeCallArgCreatorWithCall(2),
makeCallArgCreatorWithCall(3))))))

ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
HelperFeatureEnum::device_ext,
CALL_FACTORY_ENTRY(
"cuMemcpyAtoA_v2",
CALL(MapNames::getDpctNamespace() + "dpct_memcpy",
BO(BinaryOperatorKind::BO_Add,
makeCastExprCreator(
makeCharPtrCreator(),
MEMBER_CALL(MEMBER_CALL(makeCallArgCreatorWithCall(0), true,
"to_pitched_data"),
false, "get_data_ptr"),
true),
makeCallArgCreatorWithCall(1)),
BO(BinaryOperatorKind::BO_Add,
makeCastExprCreator(
makeCharPtrCreator(),
MEMBER_CALL(MEMBER_CALL(makeCallArgCreatorWithCall(2), true,
"to_pitched_data"),
false, "get_data_ptr"),
true),
makeCallArgCreatorWithCall(3)),
makeCallArgCreatorWithCall(4)))))
CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
CALL_FACTORY_ENTRY(
"cuMemcpyAtoA_v2",
CALL(MapNames::getDpctNamespace() + "experimental::dpct_memcpy",
makeCallArgCreatorWithCall(0), makeCallArgCreatorWithCall(1),
ARG("0"), makeCallArgCreatorWithCall(2),
makeCallArgCreatorWithCall(3), ARG("0"),
makeCallArgCreatorWithCall(4))),
CALL_FACTORY_ENTRY(
"cuMemcpyAtoA_v2",
CALL(MapNames::getDpctNamespace() + "dpct_memcpy",
BO(BinaryOperatorKind::BO_Add,
makeCastExprCreator(
makeCharPtrCreator(),
MEMBER_CALL(MEMBER_CALL(makeCallArgCreatorWithCall(0),
true, "to_pitched_data"),
false, "get_data_ptr"),
true),
makeCallArgCreatorWithCall(1)),
BO(BinaryOperatorKind::BO_Add,
makeCastExprCreator(
makeCharPtrCreator(),
MEMBER_CALL(MEMBER_CALL(makeCallArgCreatorWithCall(2),
true, "to_pitched_data"),
false, "get_data_ptr"),
true),
makeCallArgCreatorWithCall(3)),
makeCallArgCreatorWithCall(4))))))

ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY(
HelperFeatureEnum::device_ext,
Expand Down
10 changes: 10 additions & 0 deletions clang/lib/DPCT/APINamesStream.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -52,3 +52,13 @@ CONDITIONAL_FACTORY_ENTRY(
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
ARG("cudaStreamEndCapture"),
ARG("--use-experimental-features=graph")))

CONDITIONAL_FACTORY_ENTRY(
UseExtGraph,
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
"cudaStreamIsCapturing", DEREF(1),
MEMBER_CALL(ARG(0), true, "ext_oneapi_get_state"))),
UNSUPPORT_FACTORY_ENTRY("cudaStreamIsCapturing",
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
ARG("cudaStreamIsCapturing"),
ARG("--use-experimental-features=graph")))
6 changes: 5 additions & 1 deletion clang/lib/DPCT/APINamesTexture.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -184,7 +184,11 @@ ENTRY_UNSUPPORTED("cudaGetTextureObjectResourceViewDesc",
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
"cuArrayCreate_v2", DEREF(0),
NEW(MapNames::getDpctNamespace() + "image_matrix",
NEW(DpctGlobalInfo::useExtBindlessImages()
? MapNames::getDpctNamespace() +
"experimental::image_mem_wrapper"
: MapNames::getDpctNamespace() +
"image_matrix",
ARG(1)))))
ASSIGNABLE_FACTORY(DELETER_FACTORY_ENTRY("cuArrayDestroy", ARG(0)))
ENTRY_UNSUPPORTED("cuTexObjectGetResourceViewDesc", Diagnostics::API_NOT_MIGRATED)
Expand Down
Loading

0 comments on commit b8de47d

Please sign in to comment.