Mercurial > hg > CbC > CbC_llvm

annotate 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
Ignore whitespace changes - Everywhere: Within whitespace: At end of lines:
rev   line source
236
c4bab56944e8 LLVM 16
kono
parents:
diff changeset
1 // RUN: %clang_cc1 -no-opaque-pointers -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \
c4bab56944e8 LLVM 16
kono
parents:
diff changeset
2 // RUN: %s -S -emit-llvm -o - | FileCheck %s
c4bab56944e8 LLVM 16
kono
parents:
diff changeset
3 // RUN: %clang_cc1 -no-opaque-pointers -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \
c4bab56944e8 LLVM 16
kono
parents:
diff changeset
4 // RUN: -S -o - %s | FileCheck -check-prefix=GFX8 %s
c4bab56944e8 LLVM 16
kono
parents:
diff changeset
5
c4bab56944e8 LLVM 16
kono
parents:
diff changeset
6 // REQUIRES: amdgpu-registered-target
c4bab56944e8 LLVM 16
kono
parents:
diff changeset
7
c4bab56944e8 LLVM 16
kono
parents:
diff changeset
8 // CHECK-LABEL: test_fadd_local
c4bab56944e8 LLVM 16
kono
parents:
diff changeset
9 // CHECK: call float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %{{.*}}, float %{{.*}}, i32 0, i32 0, i1 false)
c4bab56944e8 LLVM 16
kono
parents:
diff changeset
10 // GFX8-LABEL: test_fadd_local$local:
c4bab56944e8 LLVM 16
kono
parents:
diff changeset
11 // GFX8: ds_add_rtn_f32 v2, v0, v1
c4bab56944e8 LLVM 16
kono
parents:
diff changeset
12 // GFX8: s_endpgm
c4bab56944e8 LLVM 16
kono
parents:
diff changeset
13 kernel void test_fadd_local(__local float *ptr, float val){
c4bab56944e8 LLVM 16
kono
parents:
diff changeset
14 float *res;
c4bab56944e8 LLVM 16
kono
parents:
diff changeset
15 *res = __builtin_amdgcn_ds_atomic_fadd_f32(ptr, val);
c4bab56944e8 LLVM 16
kono
parents:
diff changeset
16 }