Skip to content

Commit aaea1a1

Browse files
daiyaan-ahmed6daiyaan-ahmed6
committed
test
Signed-off-by: Daiyaan Ahmed <[email protected]>
1 parent e6e72c2 commit aaea1a1

File tree

6 files changed

+69
-92
lines changed

6 files changed

+69
-92
lines changed

‎clang/lib/DPCT/RuleInfra/MapNames.cpp‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -640,6 +640,8 @@ void MapNames::setExplicitNamespaceMap(
640640
DpctGlobalInfo::useExtGraph()
641641
? getClNamespace() + "ext::oneapi::experimental::node_type"
642642
: "cudaGraphNodeType")},
643+
{"cudaGraphExecUpdateResultInfo", std::make_shared<TypeNameRule>("int")},
644+
{"cudaGraphExecUpdateResult", std::make_shared<TypeNameRule>("int")},
643645
{"CUmem_advise", std::make_shared<TypeNameRule>("int")},
644646
{"CUmemorytype",
645647
std::make_shared<TypeNameRule>(getClNamespace() + "usm::alloc")},

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

Lines changed: 0 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -104,25 +104,3 @@ ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
104104
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
105105
ARG("cudaGraphDestroy"),
106106
ARG("--use-experimental-features=graph"))))
107-
108-
ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
109-
UseExtGraph,
110-
CALL_FACTORY_ENTRY("cudaGraphKernelNodeGetParams",
111-
CALL(MapNames::getDpctNamespace() +
112-
"experimental::get_kernel_node_params",
113-
ARG(0), ARG(1))),
114-
UNSUPPORT_FACTORY_ENTRY("cudaGraphKernelNodeGetParams",
115-
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
116-
ARG("cudaGraphKernelNodeGetParams"),
117-
ARG("--use-experimental-features=graph"))))
118-
119-
ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
120-
UseExtGraph,
121-
CALL_FACTORY_ENTRY("cudaGraphKernelNodeSetParams",
122-
CALL(MapNames::getDpctNamespace() +
123-
"experimental::set_kernel_node_params",
124-
ARG(0), ARG(1))),
125-
UNSUPPORT_FACTORY_ENTRY("cudaGraphKernelNodeSetParams",
126-
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
127-
ARG("cudaGraphKernelNodeSetParams"),
128-
ARG("--use-experimental-features=graph"))))

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

Lines changed: 36 additions & 35 deletions
Original file line numberDiff line numberDiff line change
@@ -289,35 +289,36 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
289289
"thrust::host_vector", "cublasHandle_t", "CUevent_st", "__half",
290290
"half", "__half2", "half2", "cudaMemoryAdvise", "cudaError_enum",
291291
"cudaDeviceProp", "cudaStreamCaptureStatus",
292-
"cudaGraphExecUpdateResult", "cudaPitchedPtr",
293-
"thrust::counting_iterator", "thrust::transform_iterator",
294-
"thrust::permutation_iterator", "thrust::iterator_difference",
295-
"cusolverDnHandle_t", "cusolverDnParams_t", "gesvdjInfo_t",
296-
"syevjInfo_t", "thrust::device_malloc_allocator",
297-
"thrust::divides", "thrust::tuple", "thrust::maximum",
298-
"thrust::multiplies", "thrust::plus", "cudaDataType_t",
299-
"cudaError_t", "CUresult", "CUdevice", "cudaEvent_t",
300-
"cublasStatus_t", "cuComplex", "cuFloatComplex",
301-
"cuDoubleComplex", "CUevent", "cublasFillMode_t",
302-
"cublasDiagType_t", "cublasSideMode_t", "cublasOperation_t",
303-
"cusolverStatus_t", "cusolverEigType_t", "cusolverEigMode_t",
304-
"curandStatus_t", "cudaStream_t", "cusparseStatus_t",
305-
"cusparseDiagType_t", "cusparseFillMode_t", "cusparseIndexBase_t",
306-
"cusparseMatrixType_t", "cusparseAlgMode_t",
307-
"cusparseOperation_t", "cusparseMatDescr_t", "cusparseHandle_t",
308-
"CUcontext", "cublasPointerMode_t", "cusparsePointerMode_t",
309-
"cublasGemmAlgo_t", "cusparseSolveAnalysisInfo_t", "cudaDataType",
310-
"cublasDataType_t", "curandState_t", "curandState",
311-
"curandStateXORWOW_t", "curandStateXORWOW",
312-
"curandStatePhilox4_32_10_t", "curandStatePhilox4_32_10",
313-
"curandStateMRG32k3a_t", "curandStateMRG32k3a", "thrust::minus",
314-
"thrust::negate", "thrust::logical_or", "thrust::equal_to",
315-
"thrust::less", "cudaSharedMemConfig", "curandGenerator_t",
316-
"curandRngType_t", "curandOrdering_t", "cufftHandle", "cufftReal",
317-
"cufftDoubleReal", "cufftComplex", "cufftDoubleComplex",
318-
"cufftResult_t", "cufftResult", "cufftType_t", "cufftType",
319-
"thrust::pair", "CUdeviceptr", "cudaDeviceAttr", "CUmodule",
320-
"CUjit_option", "CUfunction", "cudaMemcpyKind", "cudaComputeMode",
292+
"cudaGraphExecUpdateResult", "cudaGraphExecUpdateResultInfo",
293+
"cudaPitchedPtr", "thrust::counting_iterator",
294+
"thrust::transform_iterator", "thrust::permutation_iterator",
295+
"thrust::iterator_difference", "cusolverDnHandle_t",
296+
"cusolverDnParams_t", "gesvdjInfo_t", "syevjInfo_t",
297+
"thrust::device_malloc_allocator", "thrust::divides",
298+
"thrust::tuple", "thrust::maximum", "thrust::multiplies",
299+
"thrust::plus", "cudaDataType_t", "cudaError_t", "CUresult",
300+
"CUdevice", "cudaEvent_t", "cublasStatus_t", "cuComplex",
301+
"cuFloatComplex", "cuDoubleComplex", "CUevent",
302+
"cublasFillMode_t", "cublasDiagType_t", "cublasSideMode_t",
303+
"cublasOperation_t", "cusolverStatus_t", "cusolverEigType_t",
304+
"cusolverEigMode_t", "curandStatus_t", "cudaStream_t",
305+
"cusparseStatus_t", "cusparseDiagType_t", "cusparseFillMode_t",
306+
"cusparseIndexBase_t", "cusparseMatrixType_t",
307+
"cusparseAlgMode_t", "cusparseOperation_t", "cusparseMatDescr_t",
308+
"cusparseHandle_t", "CUcontext", "cublasPointerMode_t",
309+
"cusparsePointerMode_t", "cublasGemmAlgo_t",
310+
"cusparseSolveAnalysisInfo_t", "cudaDataType", "cublasDataType_t",
311+
"curandState_t", "curandState", "curandStateXORWOW_t",
312+
"curandStateXORWOW", "curandStatePhilox4_32_10_t",
313+
"curandStatePhilox4_32_10", "curandStateMRG32k3a_t",
314+
"curandStateMRG32k3a", "thrust::minus", "thrust::negate",
315+
"thrust::logical_or", "thrust::equal_to", "thrust::less",
316+
"cudaSharedMemConfig", "curandGenerator_t", "curandRngType_t",
317+
"curandOrdering_t", "cufftHandle", "cufftReal", "cufftDoubleReal",
318+
"cufftComplex", "cufftDoubleComplex", "cufftResult_t",
319+
"cufftResult", "cufftType_t", "cufftType", "thrust::pair",
320+
"CUdeviceptr", "cudaDeviceAttr", "CUmodule", "CUjit_option",
321+
"CUfunction", "cudaMemcpyKind", "cudaComputeMode",
321322
"__nv_bfloat16", "cooperative_groups::__v1::thread_group",
322323
"cooperative_groups::__v1::thread_block", "libraryPropertyType_t",
323324
"libraryPropertyType", "cudaDataType_t", "cudaDataType",
@@ -930,12 +931,6 @@ void TypeInDeclRule::runRule(const MatchFinder::MatchResult &Result) {
930931
}
931932
}
932933

933-
if (CanonicalTypeStr == "cudaGraphExecUpdateResult") {
934-
report(TL->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false,
935-
CanonicalTypeStr);
936-
return;
937-
}
938-
939934
if (CanonicalTypeStr == "cudaGraphicsRegisterFlags" ||
940935
CanonicalTypeStr == "cudaGraphicsMapFlags") {
941936
if (!DpctGlobalInfo::useExtBindlessImages()) {
@@ -4577,6 +4572,9 @@ void KernelCallRefRule::insertWrapperPostfix(const T *Node,
45774572
bool isInsertWrapperRegister) {
45784573
auto NLoc = DpctGlobalInfo::getSourceManager().getSpellingLoc(
45794574
Node->getNameInfo().getBeginLoc());
4575+
4576+
std::cout <<"WRAPPER APPENDED: " << "\n";
4577+
45804578
emplaceTransformation(new InsertText(
45814579
NLoc.getLocWithOffset(Node->getNameInfo().getAsString().length()),
45824580
"_wrapper"));
@@ -7182,11 +7180,14 @@ ReplaceMemberAssignAsSetMethod(const Expr *E, const MemberExpr *ME,
71827180
StringRef ExtraArg, StringRef ExtraFeild) {
71837181
if (ReplacedArg.empty()) {
71847182
if (auto RHS = getRhs(E)) {
7183+
StringRef c = ExprAnalysis::ref(RHS);
7184+
std::cout <<"Replaced String: "<< c.str() <<"\n";
71857185
return ReplaceMemberAssignAsSetMethod(
71867186
getStmtExpansionSourceRange(E).getEnd(), ME, MethodName,
71877187
ExprAnalysis::ref(RHS), ExtraArg, ExtraFeild);
71887188
}
71897189
}
7190+
std::cout << "Coming her!!!!!!!!!e\n";
71907191
return ReplaceMemberAssignAsSetMethod(getStmtExpansionSourceRange(E).getEnd(),
71917192
ME, MethodName, ReplacedArg, ExtraArg);
71927193
}

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

Lines changed: 24 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -34,8 +34,7 @@ void GraphRule::registerMatcher(MatchFinder &MF) {
3434
"cudaGraphExecDestroy", "cudaGraphAddEmptyNode",
3535
"cudaGraphAddDependencies", "cudaGraphExecUpdate",
3636
"cudaGraphNodeGetType", "cudaGraphGetNodes",
37-
"cudaGraphGetRootNodes", "cudaGraphDestroy", "cudaGraphKernelNodeGetParams",
38-
"cudaGraphKernelNodeSetParams");
37+
"cudaGraphGetRootNodes", "cudaGraphDestroy");
3938
};
4039
MF.addMatcher(
4140
callExpr(callee(functionDecl(functionName()))).bind("FunctionCall"),
@@ -56,15 +55,36 @@ void GraphRule::runRule(const MatchFinder::MatchResult &Result) {
5655
*Result.Context);
5756
auto MemberName = ME->getMemberNameInfo().getAsString();
5857
if (BaseTy == "cudaKernelNodeParams") {
59-
6058
auto FieldName = KernelNodeParamNames[MemberName];
6159
if (FieldName.empty()) {
6260
report(ME->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false,
6361
DpctGlobalInfo::getOriginalTypeName(ME->getBase()->getType()) +
6462
"::" + ME->getMemberDecl()->getName().str());
6563
return;
6664
}
67-
requestFeature(HelperFeatureEnum::device_ext);
65+
// if(FieldName == "func"){
66+
// if(auto BO = dyn_cast<BinaryOperator>(getParentAsAssignedBO(ME, *Result.Context))){
67+
// const Expr *RHS = BO->getRHS();
68+
// const Expr *StrippedRHS = RHS->IgnoreParenCasts();
69+
// std::string RHSStr;
70+
// llvm::raw_string_ostream OS(RHSStr);
71+
// std::cout <<"RHSSTR: " <<RHSStr << "\n";
72+
// StrippedRHS->printPretty(OS, nullptr, Result.Context->getPrintingPolicy());
73+
74+
75+
// // Create the replacement string using dpct::wrapper_register
76+
// auto ReplacementStr = "set_func.dpct::wrapper_register(&" + RHSStr + "_wrapper).get()";
77+
// std::cout<< "ReplacementSTR:" << ReplacementStr << "\n";
78+
79+
// // Replace the assignment with the set_func method call
80+
// // emplaceTransformation(ReplaceMemberAssignAsSetMethod(
81+
// // BO, ME, FieldName, ReplacementStr));
82+
// emplaceTransformation(new ReplaceText(getStmtExpansionSourceRange(RHS).getBegin(),
83+
// ReplacementStr.length(),
84+
// std::move(ReplacementStr)));
85+
// return;
86+
// }
87+
// }
6888
if (auto BO = getParentAsAssignedBO(ME, *Result.Context)) {
6989
StringRef ReplacedArg = "";
7090
emplaceTransformation(

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -456,12 +456,12 @@ ENTRY(cudaGraphInstantiateWithFlags, cudaGraphInstantiateWithFlags, false, NO_FL
456456
ENTRY(cudaGraphInstantiateWithParams, cudaGraphInstantiateWithParams, false, NO_FLAG, P4, "comment")
457457
ENTRY(cudaGraphKernelNodeCopyAttributes, cudaGraphKernelNodeCopyAttributes, false, NO_FLAG, P4, "comment")
458458
ENTRY(cudaGraphKernelNodeGetAttribute, cudaGraphKernelNodeGetAttribute, false, NO_FLAG, P4, "comment")
459-
ENTRY(cudaGraphKernelNodeGetParams, cudaGraphKernelNodeGetParams, true, NO_FLAG, P4, "Successful/DPCT1119")
459+
ENTRY(cudaGraphKernelNodeGetParams, cudaGraphKernelNodeGetParams, false, NO_FLAG, P4, "comment")
460460
ENTRY(cudaGraphKernelNodeSetAttribute, cudaGraphKernelNodeSetAttribute, false, NO_FLAG, P4, "comment")
461461
ENTRY(cudaGraphKernelNodeSetEnabled, cudaGraphKernelNodeSetEnabled, false, NO_FLAG, P4, "comment")
462462
ENTRY(cudaGraphKernelNodeSetGridDim, cudaGraphKernelNodeSetGridDim, false, NO_FLAG, P4, "comment")
463463
ENTRY(cudaGraphKernelNodeSetParam, cudaGraphKernelNodeSetParam, false, NO_FLAG, P4, "comment")
464-
ENTRY(cudaGraphKernelNodeSetParams, cudaGraphKernelNodeSetParams, true, NO_FLAG, P4, "Successful/DPCT1119")
464+
ENTRY(cudaGraphKernelNodeSetParams, cudaGraphKernelNodeSetParams, false, NO_FLAG, P4, "comment")
465465
ENTRY(cudaGraphKernelNodeUpdatesApply, cudaGraphKernelNodeUpdatesApply, false, NO_FLAG, P4, "comment")
466466
ENTRY(cudaGraphLaunch, cudaGraphLaunch, true, NO_FLAG, P4, "Successful/DPCT1119")
467467
ENTRY(cudaGraphMemAllocNodeGetParams, cudaGraphMemAllocNodeGetParams, false, NO_FLAG, P4, "comment")

‎clang/runtime/dpct-rt/include/dpct/graph.hpp‎

Lines changed: 5 additions & 29 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,7 @@
1010

1111
#include "dpct/util.hpp"
1212
#include "sycl/handler.hpp"
13+
#include <cstddef>
1314
#include <sycl/ext/oneapi/experimental/graph.hpp>
1415
#include <sycl/sycl.hpp>
1516
#include <unordered_map>
@@ -31,21 +32,21 @@ typedef sycl::ext::oneapi::experimental::node *node_ptr;
3132
struct kernel_node_params {
3233
dpct::dim3 block_dim;
3334
dpct::dim3 grid_dim;
34-
void *kernel_params;
35+
void **kernel_params;
3536
void* func;
3637
unsigned int shared_mem_bytes;
3738

3839
public:
3940
void set_block_dim(dpct::dim3 block_dim) { block_dim = block_dim; }
4041
void set_grid_dim(dpct::dim3 grid_dim) { grid_dim = grid_dim; }
41-
void set_kernel_params(void *kernel_params) { kernel_params = kernel_params; }
42+
void set_kernel_params(void **kernel_params) { kernel_params = kernel_params; }
4243
void set_func(void *func) { func = func; }
4344
void set_shared_mem_bytes(unsigned int shared_mem_bytes) {
4445
shared_mem_bytes = shared_mem_bytes;
4546
}
4647
dpct::dim3 get_block_dim() { return block_dim; }
4748
dpct::dim3 get_grid_dim() { return grid_dim; }
48-
void *get_kernel_params() { return kernel_params; }
49+
void **get_kernel_params() { return kernel_params; }
4950
void *get_func() { return func; }
5051
unsigned int get_shared_mem_bytes() { return shared_mem_bytes; }
5152
};
@@ -123,18 +124,6 @@ class graph_mgr {
123124
}
124125
}
125126

126-
void kernel_node_set_params(
127-
dpct::experimental::node_ptr node,
128-
dpct::experimental::kernel_node_params *kernel_node_params) {
129-
kernel_node_params_map[node] = kernel_node_params;
130-
}
131-
132-
void get_kernel_node_get_params(
133-
dpct::experimental::node_ptr node,
134-
dpct::experimental::kernel_node_params *kernel_node_params) {
135-
kernel_node_params = kernel_node_params_map[node];
136-
}
137-
138127
private:
139128
std::unordered_map<sycl::queue *, command_graph_ptr> queue_graph_map;
140129
std::unordered_map<dpct::experimental::command_graph_ptr,
@@ -232,21 +221,8 @@ detail::graph_mgr::instance().get_root_nodes(graph, nodesArray,
232221
numberOfNodes);
233222
}
234223

235-
static void
236-
kernel_node_set_params(dpct::experimental::node_ptr node,
237-
dpct::experimental::kernel_node_params *params) {
238-
detail::graph_mgr::instance().kernel_node_set_params(node, params);
239-
}
240-
241-
static void
242-
kernel_node_get_params(dpct::experimental::node_ptr node,
243-
dpct::experimental::kernel_node_params *params) {
244-
detail::graph_mgr::instance().kernel_node_set_params(node, params);
245-
}
246-
224+
static void add_kernel_node(dpct::experimental::node_ptr* node, dpct::experimental::node_ptr* dependencies, std::size_t &numberOfDependencies, dpct::experimental::kernel_node_params* params){
247225

248-
static void add_kernel_node(dpct::experimental::node_ptr node, dpct::experimental::command_graph_ptr graph, dpct::experimental::node_ptr *dependencies, std::size_t numberOfDependencies, dpct::experimental::kernel_node_params &kernelNodeParams){
249-
250226
}
251227

252228
} // namespace experimental

0 commit comments

Comments
 (0)