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..af82bd7adae1 100644 --- a/clang/test/dpct/clock.cu +++ b/clang/test/dpct/clock.cu @@ -1,42 +1,52 @@ // 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 +#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(); } -int main(int argc, char **argv) -{ - float *dinput = NULL; - float *doutput = NULL; - clock_t *dtimer = NULL; - timedReduction<<<64, 256, sizeof(float) * 2 * 256>>>(dinput, doutput, dtimer); - return 0; +// 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; +}