maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx8.cl

Ignoring revisions in .git-blame-ignore-revs. Click here to bypass and see the normal blame view.

17 lines
683 B
Common Lisp
Raw Normal View History

[clang] Add clang builtins support for gfx90a Implement target builtins for gfx90a including fadd64, fadd32, add2h, max and min on various global, flat and ds address spaces for which intrinsics are implemented. Differential Revision: https://reviews.llvm.org/D106909
2021-08-05 16:08:06 +08:00
// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \
// RUN: %s -S -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \
// RUN: -S -o - %s | FileCheck -check-prefix=GFX8 %s
Mark tests as requiring AMDGPU target
2021-08-05 17:02:51 +08:00
// REQUIRES: amdgpu-registered-target
[clang] Add clang builtins support for gfx90a Implement target builtins for gfx90a including fadd64, fadd32, add2h, max and min on various global, flat and ds address spaces for which intrinsics are implemented. Differential Revision: https://reviews.llvm.org/D106909
2021-08-05 16:08:06 +08:00
// 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);
}