From 5f8f211a5466eeb884230f04aeaf67a258a2e430 Mon Sep 17 00:00:00 2001 From: "Ahmed, Daiyaan" Date: Mon, 10 Jul 2023 22:07:18 +0800 Subject: [PATCH 1/3] Remove warning for clock function generated in host Signed-off-by: Ahmed, Daiyaan --- clang/lib/DPCT/ASTTraversal.cpp | 13 ++++++++----- clang/test/dpct/clock.cu | 14 +++++++++++++- 2 files changed, 21 insertions(+), 6 deletions(-) diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index c906e13fc000..a4422db759ae 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -6894,11 +6894,14 @@ void FunctionCallRule::runRule(const MatchFinder::MatchResult &Result) { emplaceTransformation(EA.getReplacement()); EA.applyAllSubExprRepl(); } else if (FuncName == "clock" || FuncName == "clock64") { - report(CE->getBeginLoc(), Diagnostics::API_NOT_MIGRATED_SYCL_UNDEF, false, - FuncName); - // Add '#include ' directive to the file only once - auto Loc = CE->getBeginLoc(); - DpctGlobalInfo::getInstance().insertHeader(Loc, HT_Time); + if (CE->getCalleeDecl()->hasAttr() || + CE->getCalleeDecl()->hasAttr()) { + report(CE->getBeginLoc(), Diagnostics::API_NOT_MIGRATED_SYCL_UNDEF, false, + FuncName); + // Add '#include ' directive to the file only once + auto Loc = CE->getBeginLoc(); + DpctGlobalInfo::getInstance().insertHeader(Loc, HT_Time); + } } else if (FuncName == "cudaDeviceSetLimit" || FuncName == "cudaThreadSetLimit" || FuncName == "cudaDeviceSetCacheConfig" || diff --git a/clang/test/dpct/clock.cu b/clang/test/dpct/clock.cu index 308597e6fb53..261d77131088 100644 --- a/clang/test/dpct/clock.cu +++ b/clang/test/dpct/clock.cu @@ -31,12 +31,24 @@ __global__ static void timedReduction(const float *input, float *output, clock_t clock64(); } +// CHECK: void timedAddition(const float *input, float *output, clock_t *timer){ +// CHECK-NEXT: *timer = clock(); +// CHECK-NEXT: clock(); +// CHECK-NEXT:} +void timedAddition(const float *input, float *output, clock_t *timer){ + *timer = clock(); + clock(); +} + int main(int argc, char **argv) { float *dinput = NULL; + float *hinput = NULL; float *doutput = NULL; + float *houtput = NULL; clock_t *dtimer = NULL; + clock_t *htimer = NULL; timedReduction<<<64, 256, sizeof(float) * 2 * 256>>>(dinput, doutput, dtimer); + timedAddition(hinput, houtput, htimer); return 0; } - From a4acb06ee85d34f05c9dae911d85f7d40dbc4c22 Mon Sep 17 00:00:00 2001 From: "Ahmed, Daiyaan" Date: Tue, 1 Aug 2023 13:11:00 +0800 Subject: [PATCH 2/3] Apply clang-format and rebase Signed-off-by: Ahmed, Daiyaan --- clang/test/dpct/clock.cu | 64 +++++++++++++++++++--------------------- 1 file changed, 31 insertions(+), 33 deletions(-) diff --git a/clang/test/dpct/clock.cu b/clang/test/dpct/clock.cu index 261d77131088..4ddc965e22a2 100644 --- a/clang/test/dpct/clock.cu +++ b/clang/test/dpct/clock.cu @@ -4,51 +4,49 @@ // CHECK: #include // CHECK-NEXT: #include // CHECK-NEXT: #include -#include #include +#include #include -__global__ static void timedReduction(const float *input, float *output, clock_t *timer) -{ - // CHECK: /* - // CHECK-NEXT: DPCT1008:{{[0-9]+}}: clock function is not defined in SYCL. This is a hardware-specific feature. Consult with your hardware vendor to find a replacement. - // CHECK-NEXT: */ - *timer = clock(); +__global__ static void timedReduction(const float *input, float *output, clock_t *timer) { + // CHECK: /* + // CHECK-NEXT: DPCT1008:{{[0-9]+}}: clock function is not defined in SYCL. This is a hardware-specific feature. Consult with your hardware vendor to find a replacement. + // CHECK-NEXT: */ + *timer = clock(); - // CHECK: /* - // CHECK-NEXT: DPCT1008:{{[0-9]+}}: clock function is not defined in SYCL. This is a hardware-specific feature. Consult with your hardware vendor to find a replacement. - // CHECK-NEXT: */ - clock(); + // CHECK: /* + // CHECK-NEXT: DPCT1008:{{[0-9]+}}: clock function is not defined in SYCL. This is a hardware-specific feature. Consult with your hardware vendor to find a replacement. + // CHECK-NEXT: */ + clock(); - // CHECK: /* - // CHECK-NEXT: DPCT1008:{{[0-9]+}}: clock64 function is not defined in SYCL. This is a hardware-specific feature. Consult with your hardware vendor to find a replacement. - // CHECK-NEXT: */ - *timer = clock64(); + // CHECK: /* + // CHECK-NEXT: DPCT1008:{{[0-9]+}}: clock64 function is not defined in SYCL. This is a hardware-specific feature. Consult with your hardware vendor to find a replacement. + // CHECK-NEXT: */ + *timer = clock64(); - // CHECK: /* - // CHECK-NEXT: DPCT1008:{{[0-9]+}}: clock64 function is not defined in SYCL. This is a hardware-specific feature. Consult with your hardware vendor to find a replacement. - // CHECK-NEXT: */ - clock64(); + // CHECK: /* + // CHECK-NEXT: DPCT1008:{{[0-9]+}}: clock64 function is not defined in SYCL. This is a hardware-specific feature. Consult with your hardware vendor to find a replacement. + // CHECK-NEXT: */ + clock64(); } // CHECK: void timedAddition(const float *input, float *output, clock_t *timer){ // CHECK-NEXT: *timer = clock(); // CHECK-NEXT: clock(); // CHECK-NEXT:} -void timedAddition(const float *input, float *output, clock_t *timer){ - *timer = clock(); - clock(); +void timedAddition(const float *input, float *output, clock_t *timer) { + *timer = clock(); + clock(); } -int main(int argc, char **argv) -{ - float *dinput = NULL; - float *hinput = NULL; - float *doutput = NULL; - float *houtput = NULL; - clock_t *dtimer = NULL; - clock_t *htimer = NULL; - timedReduction<<<64, 256, sizeof(float) * 2 * 256>>>(dinput, doutput, dtimer); - timedAddition(hinput, houtput, htimer); - return 0; +int main(int argc, char **argv) { + float *dinput = NULL; + float *hinput = NULL; + float *doutput = NULL; + float *houtput = NULL; + clock_t *dtimer = NULL; + clock_t *htimer = NULL; + timedReduction<<<64, 256, sizeof(float) * 2 * 256>>>(dinput, doutput, dtimer); + timedAddition(hinput, houtput, htimer); + return 0; } From 7fa38274dc6404ed4b081c4915951f11cf18ffd1 Mon Sep 17 00:00:00 2001 From: "Ahmed, Daiyaan" Date: Wed, 2 Aug 2023 13:04:58 +0800 Subject: [PATCH 3/3] Fix clock LIT test Signed-off-by: Ahmed, Daiyaan --- clang/test/dpct/clock.cu | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/test/dpct/clock.cu b/clang/test/dpct/clock.cu index 4ddc965e22a2..af82bd7adae1 100644 --- a/clang/test/dpct/clock.cu +++ b/clang/test/dpct/clock.cu @@ -1,8 +1,8 @@ // RUN: dpct --format-range=none -out-root %T/clock %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only // RUN: FileCheck %s --match-full-lines --input-file %T/clock/clock.dp.cpp -// CHECK: #include -// CHECK-NEXT: #include +// CHECK: #include +// CHECK-NEXT: #include // CHECK-NEXT: #include #include #include @@ -30,7 +30,7 @@ __global__ static void timedReduction(const float *input, float *output, clock_t clock64(); } -// CHECK: void timedAddition(const float *input, float *output, clock_t *timer){ +// CHECK: void timedAddition(const float *input, float *output, clock_t *timer) { // CHECK-NEXT: *timer = clock(); // CHECK-NEXT: clock(); // CHECK-NEXT:}