Skip to content

Commit 0822fbe

Browse files
committed
Add migration of red.relaxed.gpu.global.min.u32
Signed-off-by: chenwei.sun <[email protected]>
1 parent 76b902c commit 0822fbe

File tree

2 files changed

+12
-0
lines changed

2 files changed

+12
-0
lines changed

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2775,6 +2775,11 @@ class SYCLGen : public SYCLGenBase {
27752775
<< ")";
27762776
endstmt();
27772777
return SYCLGenSuccess();
2778+
} else if (Inst->hasAttr(InstAttr::min)) {
2779+
OS() << " = " << MapNames::getClNamespace() + "min(" << a << ", " << b
2780+
<< ")";
2781+
endstmt();
2782+
return SYCLGenSuccess();
27782783
} else
27792784
return SYCLGenError();
27802785

clang/test/dpct/asm/red.cu

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -57,4 +57,11 @@ __global__ void atomicMaxKernel(uint32_t* lock, uint32_t val) {
5757
::"l"(lock),"r"(val):"memory");
5858
}
5959

60+
// CHECK: void atomicMinKernel(uint32_t* lock, uint32_t val) {
61+
// CHECK-NEXT: *lock = sycl::min(*lock, val);
62+
// CHECK-NEXT: }
63+
__global__ void atomicMinKernel(uint32_t* lock, uint32_t val) {
64+
asm volatile("red.relaxed.gpu.global.min.u32 [%0], %1;\n"
65+
::"l"(lock),"r"(val):"memory");
66+
}
6067
// clang-format on

0 commit comments

Comments
 (0)