Mercurial > hg > CbC > CbC_llvm

diff clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx8.cl @ 236:c4bab56944e8 llvm-original

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
LLVM 16
author kono
date Wed, 09 Nov 2022 17:45:10 +0900
parents
children 1f2b6ac9f198
line wrap: on
line diff
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx8.cl	Wed Nov 09 17:45:10 2022 +0900
@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 -no-opaque-pointers -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \
+// RUN:   %s -S -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -no-opaque-pointers -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \
+// RUN:   -S -o - %s | FileCheck -check-prefix=GFX8 %s
+
+// REQUIRES: amdgpu-registered-target
+
+// CHECK-LABEL: test_fadd_local
+// CHECK: call float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %{{.*}}, float %{{.*}}, i32 0, i32 0, i1 false)
+// GFX8-LABEL: test_fadd_local$local:
+// GFX8: ds_add_rtn_f32 v2, v0, v1
+// GFX8: s_endpgm
+kernel void test_fadd_local(__local float *ptr, float val){
+    float *res;
+    *res = __builtin_amdgcn_ds_atomic_fadd_f32(ptr, val);
+}