132 lines · plain
1// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \2// RUN: %s -emit-llvm -o - | FileCheck %s -check-prefix=CHECK3 4// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \5// RUN: -S -o - %s | FileCheck -check-prefix=GFX90A %s6 7// REQUIRES: amdgpu-registered-target8 9typedef half __attribute__((ext_vector_type(2))) half2;10 11// CHECK-LABEL: test_global_add_f6412// CHECK: = atomicrmw fadd ptr addrspace(1) %{{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}13// GFX90A-LABEL: test_global_add_f64$local:14// GFX90A: global_atomic_add_f6415void test_global_add_f64(__global double *addr, double x) {16 double *rtn;17 *rtn = __builtin_amdgcn_global_atomic_fadd_f64(addr, x);18}19 20// CHECK-LABEL: test_global_add_half221// CHECK: = atomicrmw fadd ptr addrspace(1) %{{.+}}, <2 x half> %{{.+}} syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}22// GFX90A-LABEL: test_global_add_half223// GFX90A: global_atomic_pk_add_f16 v2, v[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, off glc24void test_global_add_half2(__global half2 *addr, half2 x) {25 half2 *rtn;26 *rtn = __builtin_amdgcn_global_atomic_fadd_v2f16(addr, x);27}28 29// CHECK-LABEL: test_global_global_min_f6430// CHECK: = atomicrmw fmin ptr addrspace(1) {{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}31 32// GFX90A-LABEL: test_global_global_min_f64$local33// GFX90A: global_atomic_min_f6434void test_global_global_min_f64(__global double *addr, double x){35 double *rtn;36 *rtn = __builtin_amdgcn_global_atomic_fmin_f64(addr, x);37}38 39// CHECK-LABEL: test_global_max_f6440// CHECK: = atomicrmw fmax ptr addrspace(1) {{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}41 42// GFX90A-LABEL: test_global_max_f64$local43// GFX90A: global_atomic_max_f6444void test_global_max_f64(__global double *addr, double x){45 double *rtn;46 *rtn = __builtin_amdgcn_global_atomic_fmax_f64(addr, x);47}48 49// CHECK-LABEL: test_flat_add_local_f6450// CHECK: = atomicrmw fadd ptr addrspace(3) %{{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8{{$}}51 52// GFX90A-LABEL: test_flat_add_local_f64$local53// GFX90A: ds_add_rtn_f6454void test_flat_add_local_f64(__local double *addr, double x){55 double *rtn;56 *rtn = __builtin_amdgcn_flat_atomic_fadd_f64(addr, x);57}58 59// CHECK-LABEL: test_flat_global_add_f6460// CHECK: = atomicrmw fadd ptr addrspace(1) {{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}61 62// GFX90A-LABEL: test_flat_global_add_f64$local63// GFX90A: global_atomic_add_f6464void test_flat_global_add_f64(__global double *addr, double x){65 double *rtn;66 *rtn = __builtin_amdgcn_flat_atomic_fadd_f64(addr, x);67}68 69// CHECK-LABEL: test_flat_min_flat_f6470// CHECK: = atomicrmw fmin ptr {{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}71 72// GFX90A-LABEL: test_flat_min_flat_f64$local73// GFX90A: flat_atomic_min_f6474void test_flat_min_flat_f64(__generic double *addr, double x){75 double *rtn;76 *rtn = __builtin_amdgcn_flat_atomic_fmin_f64(addr, x);77}78 79// CHECK-LABEL: test_flat_global_min_f6480// CHECK: = atomicrmw fmin ptr addrspace(1) {{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}81 82// GFX90A: test_flat_global_min_f64$local83// GFX90A: global_atomic_min_f6484void test_flat_global_min_f64(__global double *addr, double x){85 double *rtn;86 *rtn = __builtin_amdgcn_flat_atomic_fmin_f64(addr, x);87}88 89// CHECK-LABEL: test_flat_max_flat_f6490// CHECK: = atomicrmw fmax ptr {{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}91 92// GFX90A-LABEL: test_flat_max_flat_f64$local93// GFX90A: flat_atomic_max_f6494void test_flat_max_flat_f64(__generic double *addr, double x){95 double *rtn;96 *rtn = __builtin_amdgcn_flat_atomic_fmax_f64(addr, x);97}98 99// CHECK-LABEL: test_flat_global_max_f64100// CHECK: = atomicrmw fmax ptr addrspace(1) {{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}101 102// GFX90A-LABEL: test_flat_global_max_f64$local103// GFX90A: global_atomic_max_f64104void test_flat_global_max_f64(__global double *addr, double x){105 double *rtn;106 *rtn = __builtin_amdgcn_flat_atomic_fmax_f64(addr, x);107}108 109// CHECK-LABEL: test_ds_add_local_f64110// CHECK: = atomicrmw fadd ptr addrspace(3) %{{.+}}, double %{{.+}} monotonic, align 8111// GFX90A: test_ds_add_local_f64$local112// GFX90A: ds_add_rtn_f64113void test_ds_add_local_f64(__local double *addr, double x){114 double *rtn;115 *rtn = __builtin_amdgcn_ds_atomic_fadd_f64(addr, x);116}117 118// CHECK-LABEL: test_ds_addf_local_f32119// CHECK: = atomicrmw fadd ptr addrspace(3) %{{.+}}, float %{{.+}} monotonic, align 4120// GFX90A-LABEL: test_ds_addf_local_f32$local121// GFX90A: ds_add_rtn_f32122void test_ds_addf_local_f32(__local float *addr, float x){123 float *rtn;124 *rtn = __builtin_amdgcn_ds_atomic_fadd_f32(addr, x);125}126 127// CHECK-LABEL: @test_global_add_f32128// CHECK: = atomicrmw fadd ptr addrspace(1) %{{.+}}, float %{{.+}} syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+}}, !amdgpu.ignore.denormal.mode !{{[0-9]+$}}129void test_global_add_f32(float *rtn, global float *addr, float x) {130 *rtn = __builtin_amdgcn_global_atomic_fadd_f32(addr, x);131}132