[SYCLomatic] Support the atomicAdd for half type

Signed-off-by: Jiang, Zhiwei <[email protected]>
oneapi-src · Nov 27, 2024 · bf633b3 · bf633b3
1 parent 6ae8d31
commit bf633b3
Show file tree

Hide file tree

Showing 2 changed files with 3 additions and 8 deletions.
diff --git a/clang/lib/DPCT/RulesLang/RulesLangAtomic.cpp b/clang/lib/DPCT/RulesLang/RulesLangAtomic.cpp
@@ -50,12 +50,12 @@ void AtomicFunctionRule::registerMatcher(MatchFinder &MF) {
         new internal::HasNameMatcher(AtomicFuncNames));
   };
 
-  // Support all integer type, float, double and half2
-  // Type half is not supported
+  // Support all integer type, float, double half and half2.
   auto supportedTypes = [&]() {
     return anyOf(hasType(pointsTo(isInteger())),
                  hasType(pointsTo(asString("float"))),
                  hasType(pointsTo(asString("double"))),
+                 hasType(pointsTo(asString("__half"))),
                  hasType(pointsTo(asString("__half2"))));
   };
 

diff --git a/clang/test/dpct/math/half/half-cuda11.6-after.cu b/clang/test/dpct/math/half/half-cuda11.6-after.cu
@@ -23,13 +23,8 @@ __global__ void kernelFuncHalf(__half *deviceArrayHalf) {
   h_2 = __hmul_rn(h, h_1);
   // CHECK: h_2 = h - h_1;
   h_2 = __hsub_rn(h, h_1);
-#ifndef NO_BUILD_TEST
-  // CHECK: /*
-  // CHECK-NEXT: DPCT1007:{{[0-9]+}}: Migration of half version of atomicAdd is not supported.
-  // CHECK-NEXT: */
-  // CHECK-NEXT: atomicAdd(&h, h_1);
+  // CHECK: dpct::atomic_fetch_add<sycl::access::address_space::generic_space>(&h, h_1);
   atomicAdd(&h, h_1);
-#endif
 
   // Half2 Arithmetic Functions