Skip to content

Commit

Permalink
[SYCLomatic] Fix extra arg inserting location (in macro) (#2379)
Browse files Browse the repository at this point in the history
Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>
  • Loading branch information
zhiweij1 authored Oct 11, 2024
1 parent 52f7078 commit da8eaaa
Show file tree
Hide file tree
Showing 3 changed files with 18 additions and 16 deletions.
2 changes: 1 addition & 1 deletion clang/lib/DPCT/AnalysisInfo.h
Original file line number Diff line number Diff line change
Expand Up @@ -1484,7 +1484,7 @@ class DpctGlobalInfo {
return FD->getLocation();
}
static SourceLocation getLocation(const CallExpr *CE) {
return CE->getEndLoc();
return getDefinitionRange(CE->getBeginLoc(), CE->getEndLoc()).getEnd();
}
// The result will be also stored in KernelCallExpr.BeginLoc
static SourceLocation getLocation(const CUDAKernelCallExpr *CKC) {
Expand Down
18 changes: 3 additions & 15 deletions clang/test/dpct/macro_test.cu
Original file line number Diff line number Diff line change
Expand Up @@ -1385,25 +1385,13 @@ int foo39() {
return 0;
}

#define STRINGIFY_(...) #__VA_ARGS__
#define STRINGIFY(...) STRINGIFY_(__VA_ARGS__)

//CHECK: void foo40_dev_func(const char *file_name, const char *other,
//CHECK-NEXT: const sycl::stream &stream_ct1) {
//CHECK-NEXT: stream_ct1 << "sss\n";
//CHECK: void foo40_kernel(const sycl::stream &stream_ct1) {
//CHECK-NEXT: FOO40_MACRO;
//CHECK-NEXT: }
//CHECK-NEXT: #define FOO40_MACRO \
//CHECK-NEXT: foo40_dev_func(__FILE__, STRINGIFY(__OTHER_MACRO__), stream_ct1)
__device__ void foo40_dev_func(const char *file_name, const char *other) {
printf("sss\n");
}
#define FOO40_MACRO foo40_dev_func(__FILE__, STRINGIFY(__OTHER_MACRO__))

__global__ void foo40_kernel() {
FOO40_MACRO;
}
void foo40() {
foo40_kernel<<<1, 1>>>();
}

#endif
#endif
14 changes: 14 additions & 0 deletions clang/test/dpct/macro_test.h
Original file line number Diff line number Diff line change
@@ -1,2 +1,16 @@
//CHECK: #define THREAD_IDX_X item_ct1.get_local_id(2)
#define THREAD_IDX_X threadIdx.x

#define STRINGIFY_(...) #__VA_ARGS__
#define STRINGIFY(...) STRINGIFY_(__VA_ARGS__)

//CHECK: void foo40_dev_func(const char *file_name, const char *other,
//CHECK-NEXT: const sycl::stream &stream_ct1) {
//CHECK-NEXT: stream_ct1 << "sss\n";
//CHECK-NEXT: }
//CHECK-NEXT: #define FOO40_MACRO \
//CHECK-NEXT: foo40_dev_func(__FILE__, STRINGIFY(__OTHER_MACRO__), stream_ct1)
__device__ void foo40_dev_func(const char *file_name, const char *other) {
printf("sss\n");
}
#define FOO40_MACRO foo40_dev_func(__FILE__, STRINGIFY(__OTHER_MACRO__))

0 comments on commit da8eaaa

Please sign in to comment.