1522 lines · c
1// REQUIRES: nvptx-registered-target2// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_70 -target-feature +ptx63 \3// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \4// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX63_SM70 -check-prefix=LP64 %s5// RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu sm_80 -target-feature +ptx70 \6// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \7// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX70_SM80 -check-prefix=LP32 %s8// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_80 -target-feature +ptx70 \9// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \10// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX70_SM80 -check-prefix=LP64 %s11// RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu sm_60 -target-feature +ptx62 \12// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \13// RUN: | FileCheck -check-prefix=CHECK -check-prefix=LP32 %s14// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_60 -target-feature +ptx62 \15// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \16// RUN: | FileCheck -check-prefix=CHECK -check-prefix=LP64 %s17// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_61 -target-feature +ptx62 \18// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \19// RUN: | FileCheck -check-prefix=CHECK -check-prefix=LP64 %s20// RUN: %clang_cc1 -triple nvptx-unknown-unknown -target-cpu sm_53 -target-feature +ptx62 \21// RUN: -DERROR_CHECK -fcuda-is-device -S -o /dev/null -x cuda -verify %s22// RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu sm_86 -target-feature +ptx72 \23// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \24// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX72_SM86 -check-prefix=LP32 %s25// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_86 -target-feature +ptx72 \26// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \27// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX72_SM86 -check-prefix=LP64 %s28// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_89 -target-feature +ptx81 -DPTX=81\29// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \30// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX81_SM89 %s31// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_80 -target-feature +ptx81 -DPTX=81 \32// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \33// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX81_SM80 %s34// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_90 -target-feature +ptx78 -DPTX=78 \35// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \36// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX78_SM90 %s37// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_100 -target-feature +ptx86 -DPTX=86 \38// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \39// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX86_SM100 %s40// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_100a -target-feature +ptx86 -DPTX=86 \41// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \42// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX86_SM100a %s43// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_101a -target-feature +ptx86 -DPTX=86 \44// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \45// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX86_SM101a %s46// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_120a -target-feature +ptx86 -DPTX=86 \47// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \48// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX86_SM120a %s49// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_103a -target-feature +ptx87 -DPTX=87 \50// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \51// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX87_SM103a %s52// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_100a -target-feature +ptx87 -DPTX=87 \53// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \54// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX87_SM100a %s55// ### The last run to check with the highest SM and PTX version available56// ### to make sure target builtins are still accepted.57// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_120a -target-feature +ptx87 -DPTX=87 \58// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \59// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX86_SM120a %s60 61#define __device__ __attribute__((device))62#define __global__ __attribute__((global))63#define __shared__ __attribute__((shared))64#define __constant__ __attribute__((constant))65 66__device__ int read_tid() {67 68// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.x()69// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.y()70// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.z()71// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.w()72 73 int x = __nvvm_read_ptx_sreg_tid_x();74 int y = __nvvm_read_ptx_sreg_tid_y();75 int z = __nvvm_read_ptx_sreg_tid_z();76 int w = __nvvm_read_ptx_sreg_tid_w();77 78 return x + y + z + w;79 80}81 82__device__ bool reflect() {83 84// CHECK: call i32 @llvm.nvvm.reflect(ptr {{.*}})85 86 unsigned x = __nvvm_reflect("__CUDA_ARCH");87 return x >= 700;88 89}90 91__device__ int read_ntid() {92 93// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()94// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.y()95// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.z()96// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.w()97 98 int x = __nvvm_read_ptx_sreg_ntid_x();99 int y = __nvvm_read_ptx_sreg_ntid_y();100 int z = __nvvm_read_ptx_sreg_ntid_z();101 int w = __nvvm_read_ptx_sreg_ntid_w();102 103 return x + y + z + w;104 105}106 107__device__ int read_ctaid() {108 109// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()110// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()111// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()112// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.w()113 114 int x = __nvvm_read_ptx_sreg_ctaid_x();115 int y = __nvvm_read_ptx_sreg_ctaid_y();116 int z = __nvvm_read_ptx_sreg_ctaid_z();117 int w = __nvvm_read_ptx_sreg_ctaid_w();118 119 return x + y + z + w;120 121}122 123__device__ int read_nctaid() {124 125// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()126// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()127// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()128// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.w()129 130 int x = __nvvm_read_ptx_sreg_nctaid_x();131 int y = __nvvm_read_ptx_sreg_nctaid_y();132 int z = __nvvm_read_ptx_sreg_nctaid_z();133 int w = __nvvm_read_ptx_sreg_nctaid_w();134 135 return x + y + z + w;136 137}138 139__device__ int read_ids() {140 141// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.laneid()142// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.warpid()143// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nwarpid()144// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.smid()145// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nsmid()146// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.gridid()147// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.warpsize()148 149 int a = __nvvm_read_ptx_sreg_laneid();150 int b = __nvvm_read_ptx_sreg_warpid();151 int c = __nvvm_read_ptx_sreg_nwarpid();152 int d = __nvvm_read_ptx_sreg_smid();153 int e = __nvvm_read_ptx_sreg_nsmid();154 int f = __nvvm_read_ptx_sreg_gridid();155 int g = __nvvm_read_ptx_sreg_warpsize();156 157 return a + b + c + d + e + f + g;158 159}160 161__device__ int read_lanemasks() {162 163// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.lanemask.eq()164// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.lanemask.le()165// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.lanemask.lt()166// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.lanemask.ge()167// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.lanemask.gt()168 169 int a = __nvvm_read_ptx_sreg_lanemask_eq();170 int b = __nvvm_read_ptx_sreg_lanemask_le();171 int c = __nvvm_read_ptx_sreg_lanemask_lt();172 int d = __nvvm_read_ptx_sreg_lanemask_ge();173 int e = __nvvm_read_ptx_sreg_lanemask_gt();174 175 return a + b + c + d + e;176 177}178 179__device__ long long read_clocks() {180 181// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.clock()182// CHECK: call i64 @llvm.nvvm.read.ptx.sreg.clock64()183// CHECK: call i64 @llvm.nvvm.read.ptx.sreg.globaltimer()184 185 int a = __nvvm_read_ptx_sreg_clock();186 long long b = __nvvm_read_ptx_sreg_clock64();187 long long c = __nvvm_read_ptx_sreg_globaltimer();188 189 return a + b + c;190}191 192__device__ int read_pms() {193 194// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.pm0()195// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.pm1()196// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.pm2()197// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.pm3()198 199 int a = __nvvm_read_ptx_sreg_pm0();200 int b = __nvvm_read_ptx_sreg_pm1();201 int c = __nvvm_read_ptx_sreg_pm2();202 int d = __nvvm_read_ptx_sreg_pm3();203 204 return a + b + c + d;205 206}207 208__device__ void sync() {209 210// CHECK: call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)211 212 __nvvm_bar_sync(0);213 214}215 216__device__ void activemask() {217 218// CHECK: call i32 @llvm.nvvm.activemask()219 220 __nvvm_activemask();221 222}223 224__device__ void exit() {225 226// CHECK: call void @llvm.nvvm.exit()227 228 __nvvm_exit();229 230}231 232// NVVM intrinsics233 234// The idea is not to test all intrinsics, just that Clang is recognizing the235// builtins defined in BuiltinsNVPTX.td236__device__ void nvvm_math(float f1, float f2, double d1, double d2) {237// CHECK: call float @llvm.nvvm.fmax.f238 float t1 = __nvvm_fmax_f(f1, f2);239// CHECK: call float @llvm.nvvm.fmin.f240 float t2 = __nvvm_fmin_f(f1, f2);241// CHECK: call float @llvm.nvvm.sqrt.rn.f242 float t3 = __nvvm_sqrt_rn_f(f1);243// CHECK: call float @llvm.nvvm.rcp.rn.f244 float t4 = __nvvm_rcp_rn_f(f2);245// CHECK: call float @llvm.nvvm.add.rn.f246 float t5 = __nvvm_add_rn_f(f1, f2);247 248// CHECK: call double @llvm.nvvm.fmax.d249 double td1 = __nvvm_fmax_d(d1, d2);250// CHECK: call double @llvm.nvvm.fmin.d251 double td2 = __nvvm_fmin_d(d1, d2);252// CHECK: call double @llvm.nvvm.sqrt.rn.d253 double td3 = __nvvm_sqrt_rn_d(d1);254// CHECK: call double @llvm.nvvm.rcp.rn.d255 double td4 = __nvvm_rcp_rn_d(d2);256 257// CHECK: call float @llvm.nvvm.fabs.f32258 float t6 = __nvvm_fabs_f(f1);259// CHECK: call float @llvm.nvvm.fabs.ftz.f32260 float t7 = __nvvm_fabs_ftz_f(f2);261 262// CHECK: call double @llvm.fabs.f64263 double td5 = __nvvm_fabs_d(d1);264 265// CHECK: call void @llvm.nvvm.membar.cta()266 __nvvm_membar_cta();267// CHECK: call void @llvm.nvvm.membar.gl()268 __nvvm_membar_gl();269// CHECK: call void @llvm.nvvm.membar.sys()270 __nvvm_membar_sys();271// CHECK: call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)272 __syncthreads();273}274 275__device__ int di;276__shared__ int si;277__device__ long dl;278__shared__ long sl;279__device__ long long dll;280__shared__ long long sll;281 282// Check for atomic intrinsics283// CHECK-LABEL: nvvm_atom284__device__ void nvvm_atom(float *fp, float f, double *dfp, double df,285 unsigned short *usp, unsigned short us, int *ip,286 int i, unsigned int *uip, unsigned ui, long *lp,287 long l, long long *llp, long long ll) {288 // CHECK: atomicrmw add ptr {{.*}} seq_cst, align 4289 __nvvm_atom_add_gen_i(ip, i);290 // CHECK: atomicrmw add ptr {{.*}} seq_cst, align {{4|8}}291 __nvvm_atom_add_gen_l(&dl, l);292 // CHECK: atomicrmw add ptr {{.*}} seq_cst, align 8293 __nvvm_atom_add_gen_ll(&sll, ll);294 295 // CHECK: atomicrmw sub ptr {{.*}} seq_cst, align 4296 __nvvm_atom_sub_gen_i(ip, i);297 // CHECK: atomicrmw sub ptr {{.*}} seq_cst, align {{4|8}}298 __nvvm_atom_sub_gen_l(&dl, l);299 // CHECK: atomicrmw sub ptr {{.*}} seq_cst, align 8300 __nvvm_atom_sub_gen_ll(&sll, ll);301 302 // CHECK: atomicrmw and ptr {{.*}} seq_cst, align 4303 __nvvm_atom_and_gen_i(ip, i);304 // CHECK: atomicrmw and ptr {{.*}} seq_cst, align {{4|8}}305 __nvvm_atom_and_gen_l(&dl, l);306 // CHECK: atomicrmw and ptr {{.*}} seq_cst, align 8307 __nvvm_atom_and_gen_ll(&sll, ll);308 309 // CHECK: atomicrmw or ptr {{.*}} seq_cst, align 4310 __nvvm_atom_or_gen_i(ip, i);311 // CHECK: atomicrmw or ptr {{.*}} seq_cst, align {{4|8}}312 __nvvm_atom_or_gen_l(&dl, l);313 // CHECK: atomicrmw or ptr {{.*}} seq_cst, align 8314 __nvvm_atom_or_gen_ll(&sll, ll);315 316 // CHECK: atomicrmw xor ptr {{.*}} seq_cst, align 4317 __nvvm_atom_xor_gen_i(ip, i);318 // CHECK: atomicrmw xor ptr {{.*}} seq_cst, align {{4|8}}319 __nvvm_atom_xor_gen_l(&dl, l);320 // CHECK: atomicrmw xor ptr {{.*}} seq_cst, align 8321 __nvvm_atom_xor_gen_ll(&sll, ll);322 323 // CHECK: atomicrmw xchg ptr {{.*}} seq_cst, align 4324 __nvvm_atom_xchg_gen_i(ip, i);325 // CHECK: atomicrmw xchg ptr {{.*}} seq_cst, align {{4|8}}326 __nvvm_atom_xchg_gen_l(&dl, l);327 // CHECK: atomicrmw xchg ptr {{.*}} seq_cst, align 8328 __nvvm_atom_xchg_gen_ll(&sll, ll);329 330 // CHECK: atomicrmw max ptr {{.*}} seq_cst, align 4331 __nvvm_atom_max_gen_i(ip, i);332 // CHECK: atomicrmw umax ptr {{.*}} seq_cst, align 4333 __nvvm_atom_max_gen_ui((unsigned int *)ip, i);334 // CHECK: atomicrmw max ptr {{.*}} seq_cst, align {{4|8}}335 __nvvm_atom_max_gen_l(&dl, l);336 // CHECK: atomicrmw umax ptr {{.*}} seq_cst, align {{4|8}}337 __nvvm_atom_max_gen_ul((unsigned long *)&dl, l);338 // CHECK: atomicrmw max ptr {{.*}} seq_cst, align 8339 __nvvm_atom_max_gen_ll(&sll, ll);340 // CHECK: atomicrmw umax ptr {{.*}} seq_cst, align 8341 __nvvm_atom_max_gen_ull((unsigned long long *)&sll, ll);342 343 // CHECK: atomicrmw min ptr {{.*}} seq_cst, align 4344 __nvvm_atom_min_gen_i(ip, i);345 // CHECK: atomicrmw umin ptr {{.*}} seq_cst, align 4346 __nvvm_atom_min_gen_ui((unsigned int *)ip, i);347 // CHECK: atomicrmw min ptr {{.*}} seq_cst, align {{4|8}}348 __nvvm_atom_min_gen_l(&dl, l);349 // CHECK: atomicrmw umin ptr {{.*}} seq_cst, align {{4|8}}350 __nvvm_atom_min_gen_ul((unsigned long *)&dl, l);351 // CHECK: atomicrmw min ptr {{.*}} seq_cst, align 8352 __nvvm_atom_min_gen_ll(&sll, ll);353 // CHECK: atomicrmw umin ptr {{.*}} seq_cst, align 8354 __nvvm_atom_min_gen_ull((unsigned long long *)&sll, ll);355 356 // CHECK: cmpxchg ptr {{.*}} seq_cst seq_cst, align 4357 // CHECK-NEXT: extractvalue { i32, i1 } {{%[0-9]+}}, 0358 __nvvm_atom_cas_gen_i(ip, 0, i);359 // CHECK: cmpxchg ptr {{.*}} seq_cst seq_cst, align {{4|8}}360 // CHECK-NEXT: extractvalue { {{i32|i64}}, i1 } {{%[0-9]+}}, 0361 __nvvm_atom_cas_gen_l(&dl, 0, l);362 // CHECK: cmpxchg ptr {{.*}} seq_cst seq_cst, align 8363 // CHECK-NEXT: extractvalue { i64, i1 } {{%[0-9]+}}, 0364 __nvvm_atom_cas_gen_ll(&sll, 0, ll);365 366 // CHECK: atomicrmw fadd ptr {{.*}} seq_cst, align 4367 __nvvm_atom_add_gen_f(fp, f);368 369 // CHECK: atomicrmw uinc_wrap ptr {{.*}} seq_cst, align 4370 __nvvm_atom_inc_gen_ui(uip, ui);371 372 // CHECK: atomicrmw udec_wrap ptr {{.*}} seq_cst, align 4373 __nvvm_atom_dec_gen_ui(uip, ui);374 375 376 //////////////////////////////////////////////////////////////////377 // Atomics with scope (only supported on sm_60+).378 379#if ERROR_CHECK || __CUDA_ARCH__ >= 600380 381 // CHECK: call i32 @llvm.nvvm.atomic.add.gen.i.cta.i32.p0382 // expected-error@+1 {{'__nvvm_atom_cta_add_gen_i' needs target feature sm_60}}383 __nvvm_atom_cta_add_gen_i(ip, i);384 // LP32: call i32 @llvm.nvvm.atomic.add.gen.i.cta.i32.p0385 // LP64: call i64 @llvm.nvvm.atomic.add.gen.i.cta.i64.p0386 // expected-error@+1 {{'__nvvm_atom_cta_add_gen_l' needs target feature sm_60}}387 __nvvm_atom_cta_add_gen_l(&dl, l);388 // CHECK: call i64 @llvm.nvvm.atomic.add.gen.i.cta.i64.p0389 // expected-error@+1 {{'__nvvm_atom_cta_add_gen_ll' needs target feature sm_60}}390 __nvvm_atom_cta_add_gen_ll(&sll, ll);391 // CHECK: call i32 @llvm.nvvm.atomic.add.gen.i.sys.i32.p0392 // expected-error@+1 {{'__nvvm_atom_sys_add_gen_i' needs target feature sm_60}}393 __nvvm_atom_sys_add_gen_i(ip, i);394 // LP32: call i32 @llvm.nvvm.atomic.add.gen.i.sys.i32.p0395 // LP64: call i64 @llvm.nvvm.atomic.add.gen.i.sys.i64.p0396 // expected-error@+1 {{'__nvvm_atom_sys_add_gen_l' needs target feature sm_60}}397 __nvvm_atom_sys_add_gen_l(&dl, l);398 // CHECK: call i64 @llvm.nvvm.atomic.add.gen.i.sys.i64.p0399 // expected-error@+1 {{'__nvvm_atom_sys_add_gen_ll' needs target feature sm_60}}400 __nvvm_atom_sys_add_gen_ll(&sll, ll);401 402 // CHECK: call float @llvm.nvvm.atomic.add.gen.f.cta.f32.p0403 // expected-error@+1 {{'__nvvm_atom_cta_add_gen_f' needs target feature sm_60}}404 __nvvm_atom_cta_add_gen_f(fp, f);405 // CHECK: call double @llvm.nvvm.atomic.add.gen.f.cta.f64.p0406 // expected-error@+1 {{'__nvvm_atom_cta_add_gen_d' needs target feature sm_60}}407 __nvvm_atom_cta_add_gen_d(dfp, df);408 // CHECK: call float @llvm.nvvm.atomic.add.gen.f.sys.f32.p0409 // expected-error@+1 {{'__nvvm_atom_sys_add_gen_f' needs target feature sm_60}}410 __nvvm_atom_sys_add_gen_f(fp, f);411 // CHECK: call double @llvm.nvvm.atomic.add.gen.f.sys.f64.p0412 // expected-error@+1 {{'__nvvm_atom_sys_add_gen_d' needs target feature sm_60}}413 __nvvm_atom_sys_add_gen_d(dfp, df);414 415 // CHECK: call i32 @llvm.nvvm.atomic.exch.gen.i.cta.i32.p0416 // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_i' needs target feature sm_60}}417 __nvvm_atom_cta_xchg_gen_i(ip, i);418 // LP32: call i32 @llvm.nvvm.atomic.exch.gen.i.cta.i32.p0419 // LP64: call i64 @llvm.nvvm.atomic.exch.gen.i.cta.i64.p0420 // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_l' needs target feature sm_60}}421 __nvvm_atom_cta_xchg_gen_l(&dl, l);422 // CHECK: call i64 @llvm.nvvm.atomic.exch.gen.i.cta.i64.p0423 // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_ll' needs target feature sm_60}}424 __nvvm_atom_cta_xchg_gen_ll(&sll, ll);425 426 // CHECK: call i32 @llvm.nvvm.atomic.exch.gen.i.sys.i32.p0427 // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_i' needs target feature sm_60}}428 __nvvm_atom_sys_xchg_gen_i(ip, i);429 // LP32: call i32 @llvm.nvvm.atomic.exch.gen.i.sys.i32.p0430 // LP64: call i64 @llvm.nvvm.atomic.exch.gen.i.sys.i64.p0431 // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_l' needs target feature sm_60}}432 __nvvm_atom_sys_xchg_gen_l(&dl, l);433 // CHECK: call i64 @llvm.nvvm.atomic.exch.gen.i.sys.i64.p0434 // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_ll' needs target feature sm_60}}435 __nvvm_atom_sys_xchg_gen_ll(&sll, ll);436 437 // CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0438 // expected-error@+1 {{'__nvvm_atom_cta_max_gen_i' needs target feature sm_60}}439 __nvvm_atom_cta_max_gen_i(ip, i);440 // CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0441 // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ui' needs target feature sm_60}}442 __nvvm_atom_cta_max_gen_ui((unsigned int *)ip, i);443 // LP32: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0444 // LP64: call i64 @llvm.nvvm.atomic.max.gen.i.cta.i64.p0445 // expected-error@+1 {{'__nvvm_atom_cta_max_gen_l' needs target feature sm_60}}446 __nvvm_atom_cta_max_gen_l(&dl, l);447 // LP32: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0448 // LP64: call i64 @llvm.nvvm.atomic.max.gen.i.cta.i64.p0449 // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ul' needs target feature sm_60}}450 __nvvm_atom_cta_max_gen_ul((unsigned long *)lp, l);451 // CHECK: call i64 @llvm.nvvm.atomic.max.gen.i.cta.i64.p0452 // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ll' needs target feature sm_60}}453 __nvvm_atom_cta_max_gen_ll(&sll, ll);454 // CHECK: call i64 @llvm.nvvm.atomic.max.gen.i.cta.i64.p0455 // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ull' needs target feature sm_60}}456 __nvvm_atom_cta_max_gen_ull((unsigned long long *)llp, ll);457 458 // CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.sys.i32.p0459 // expected-error@+1 {{'__nvvm_atom_sys_max_gen_i' needs target feature sm_60}}460 __nvvm_atom_sys_max_gen_i(ip, i);461 // CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.sys.i32.p0462 // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ui' needs target feature sm_60}}463 __nvvm_atom_sys_max_gen_ui((unsigned int *)ip, i);464 // LP32: call i32 @llvm.nvvm.atomic.max.gen.i.sys.i32.p0465 // LP64: call i64 @llvm.nvvm.atomic.max.gen.i.sys.i64.p0466 // expected-error@+1 {{'__nvvm_atom_sys_max_gen_l' needs target feature sm_60}}467 __nvvm_atom_sys_max_gen_l(&dl, l);468 // LP32: call i32 @llvm.nvvm.atomic.max.gen.i.sys.i32.p0469 // LP64: call i64 @llvm.nvvm.atomic.max.gen.i.sys.i64.p0470 // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ul' needs target feature sm_60}}471 __nvvm_atom_sys_max_gen_ul((unsigned long *)lp, l);472 // CHECK: call i64 @llvm.nvvm.atomic.max.gen.i.sys.i64.p0473 // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ll' needs target feature sm_60}}474 __nvvm_atom_sys_max_gen_ll(&sll, ll);475 // CHECK: call i64 @llvm.nvvm.atomic.max.gen.i.sys.i64.p0476 // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ull' needs target feature sm_60}}477 __nvvm_atom_sys_max_gen_ull((unsigned long long *)llp, ll);478 479 // CHECK: call i32 @llvm.nvvm.atomic.min.gen.i.cta.i32.p0480 // expected-error@+1 {{'__nvvm_atom_cta_min_gen_i' needs target feature sm_60}}481 __nvvm_atom_cta_min_gen_i(ip, i);482 // CHECK: call i32 @llvm.nvvm.atomic.min.gen.i.cta.i32.p0483 // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ui' needs target feature sm_60}}484 __nvvm_atom_cta_min_gen_ui((unsigned int *)ip, i);485 // LP32: call i32 @llvm.nvvm.atomic.min.gen.i.cta.i32.p0486 // LP64: call i64 @llvm.nvvm.atomic.min.gen.i.cta.i64.p0487 // expected-error@+1 {{'__nvvm_atom_cta_min_gen_l' needs target feature sm_60}}488 __nvvm_atom_cta_min_gen_l(&dl, l);489 // LP32: call i32 @llvm.nvvm.atomic.min.gen.i.cta.i32.p0490 // LP64: call i64 @llvm.nvvm.atomic.min.gen.i.cta.i64.p0491 // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ul' needs target feature sm_60}}492 __nvvm_atom_cta_min_gen_ul((unsigned long *)lp, l);493 // CHECK: call i64 @llvm.nvvm.atomic.min.gen.i.cta.i64.p0494 // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ll' needs target feature sm_60}}495 __nvvm_atom_cta_min_gen_ll(&sll, ll);496 // CHECK: call i64 @llvm.nvvm.atomic.min.gen.i.cta.i64.p0497 // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ull' needs target feature sm_60}}498 __nvvm_atom_cta_min_gen_ull((unsigned long long *)llp, ll);499 500 // CHECK: call i32 @llvm.nvvm.atomic.min.gen.i.sys.i32.p0501 // expected-error@+1 {{'__nvvm_atom_sys_min_gen_i' needs target feature sm_60}}502 __nvvm_atom_sys_min_gen_i(ip, i);503 // CHECK: call i32 @llvm.nvvm.atomic.min.gen.i.sys.i32.p0504 // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ui' needs target feature sm_60}}505 __nvvm_atom_sys_min_gen_ui((unsigned int *)ip, i);506 // LP32: call i32 @llvm.nvvm.atomic.min.gen.i.sys.i32.p0507 // LP64: call i64 @llvm.nvvm.atomic.min.gen.i.sys.i64.p0508 // expected-error@+1 {{'__nvvm_atom_sys_min_gen_l' needs target feature sm_60}}509 __nvvm_atom_sys_min_gen_l(&dl, l);510 // LP32: call i32 @llvm.nvvm.atomic.min.gen.i.sys.i32.p0511 // LP64: call i64 @llvm.nvvm.atomic.min.gen.i.sys.i64.p0512 // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ul' needs target feature sm_60}}513 __nvvm_atom_sys_min_gen_ul((unsigned long *)lp, l);514 // CHECK: call i64 @llvm.nvvm.atomic.min.gen.i.sys.i64.p0515 // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ll' needs target feature sm_60}}516 __nvvm_atom_sys_min_gen_ll(&sll, ll);517 // CHECK: call i64 @llvm.nvvm.atomic.min.gen.i.sys.i64.p0518 // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ull' needs target feature sm_60}}519 __nvvm_atom_sys_min_gen_ull((unsigned long long *)llp, ll);520 521 // CHECK: call i32 @llvm.nvvm.atomic.inc.gen.i.cta.i32.p0522 // expected-error@+1 {{'__nvvm_atom_cta_inc_gen_ui' needs target feature sm_60}}523 __nvvm_atom_cta_inc_gen_ui((unsigned int *)ip, i);524 // CHECK: call i32 @llvm.nvvm.atomic.inc.gen.i.sys.i32.p0525 // expected-error@+1 {{'__nvvm_atom_sys_inc_gen_ui' needs target feature sm_60}}526 __nvvm_atom_sys_inc_gen_ui((unsigned int *)ip, i);527 528 // CHECK: call i32 @llvm.nvvm.atomic.dec.gen.i.cta.i32.p0529 // expected-error@+1 {{'__nvvm_atom_cta_dec_gen_ui' needs target feature sm_60}}530 __nvvm_atom_cta_dec_gen_ui((unsigned int *)ip, i);531 // CHECK: call i32 @llvm.nvvm.atomic.dec.gen.i.sys.i32.p0532 // expected-error@+1 {{'__nvvm_atom_sys_dec_gen_ui' needs target feature sm_60}}533 __nvvm_atom_sys_dec_gen_ui((unsigned int *)ip, i);534 535 // CHECK: call i32 @llvm.nvvm.atomic.and.gen.i.cta.i32.p0536 // expected-error@+1 {{'__nvvm_atom_cta_and_gen_i' needs target feature sm_60}}537 __nvvm_atom_cta_and_gen_i(ip, i);538 // LP32: call i32 @llvm.nvvm.atomic.and.gen.i.cta.i32.p0539 // LP64: call i64 @llvm.nvvm.atomic.and.gen.i.cta.i64.p0540 // expected-error@+1 {{'__nvvm_atom_cta_and_gen_l' needs target feature sm_60}}541 __nvvm_atom_cta_and_gen_l(&dl, l);542 // CHECK: call i64 @llvm.nvvm.atomic.and.gen.i.cta.i64.p0543 // expected-error@+1 {{'__nvvm_atom_cta_and_gen_ll' needs target feature sm_60}}544 __nvvm_atom_cta_and_gen_ll(&sll, ll);545 546 // CHECK: call i32 @llvm.nvvm.atomic.and.gen.i.sys.i32.p0547 // expected-error@+1 {{'__nvvm_atom_sys_and_gen_i' needs target feature sm_60}}548 __nvvm_atom_sys_and_gen_i(ip, i);549 // LP32: call i32 @llvm.nvvm.atomic.and.gen.i.sys.i32.p0550 // LP64: call i64 @llvm.nvvm.atomic.and.gen.i.sys.i64.p0551 // expected-error@+1 {{'__nvvm_atom_sys_and_gen_l' needs target feature sm_60}}552 __nvvm_atom_sys_and_gen_l(&dl, l);553 // CHECK: call i64 @llvm.nvvm.atomic.and.gen.i.sys.i64.p0554 // expected-error@+1 {{'__nvvm_atom_sys_and_gen_ll' needs target feature sm_60}}555 __nvvm_atom_sys_and_gen_ll(&sll, ll);556 557 // CHECK: call i32 @llvm.nvvm.atomic.or.gen.i.cta.i32.p0558 // expected-error@+1 {{'__nvvm_atom_cta_or_gen_i' needs target feature sm_60}}559 __nvvm_atom_cta_or_gen_i(ip, i);560 // LP32: call i32 @llvm.nvvm.atomic.or.gen.i.cta.i32.p0561 // LP64: call i64 @llvm.nvvm.atomic.or.gen.i.cta.i64.p0562 // expected-error@+1 {{'__nvvm_atom_cta_or_gen_l' needs target feature sm_60}}563 __nvvm_atom_cta_or_gen_l(&dl, l);564 // CHECK: call i64 @llvm.nvvm.atomic.or.gen.i.cta.i64.p0565 // expected-error@+1 {{'__nvvm_atom_cta_or_gen_ll' needs target feature sm_60}}566 __nvvm_atom_cta_or_gen_ll(&sll, ll);567 568 // CHECK: call i32 @llvm.nvvm.atomic.or.gen.i.sys.i32.p0569 // expected-error@+1 {{'__nvvm_atom_sys_or_gen_i' needs target feature sm_60}}570 __nvvm_atom_sys_or_gen_i(ip, i);571 // LP32: call i32 @llvm.nvvm.atomic.or.gen.i.sys.i32.p0572 // LP64: call i64 @llvm.nvvm.atomic.or.gen.i.sys.i64.p0573 // expected-error@+1 {{'__nvvm_atom_sys_or_gen_l' needs target feature sm_60}}574 __nvvm_atom_sys_or_gen_l(&dl, l);575 // CHECK: call i64 @llvm.nvvm.atomic.or.gen.i.sys.i64.p0576 // expected-error@+1 {{'__nvvm_atom_sys_or_gen_ll' needs target feature sm_60}}577 __nvvm_atom_sys_or_gen_ll(&sll, ll);578 579 // CHECK: call i32 @llvm.nvvm.atomic.xor.gen.i.cta.i32.p0580 // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_i' needs target feature sm_60}}581 __nvvm_atom_cta_xor_gen_i(ip, i);582 // LP32: call i32 @llvm.nvvm.atomic.xor.gen.i.cta.i32.p0583 // LP64: call i64 @llvm.nvvm.atomic.xor.gen.i.cta.i64.p0584 // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_l' needs target feature sm_60}}585 __nvvm_atom_cta_xor_gen_l(&dl, l);586 // CHECK: call i64 @llvm.nvvm.atomic.xor.gen.i.cta.i64.p0587 // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_ll' needs target feature sm_60}}588 __nvvm_atom_cta_xor_gen_ll(&sll, ll);589 590 // CHECK: call i32 @llvm.nvvm.atomic.xor.gen.i.sys.i32.p0591 // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_i' needs target feature sm_60}}592 __nvvm_atom_sys_xor_gen_i(ip, i);593 // LP32: call i32 @llvm.nvvm.atomic.xor.gen.i.sys.i32.p0594 // LP64: call i64 @llvm.nvvm.atomic.xor.gen.i.sys.i64.p0595 // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_l' needs target feature sm_60}}596 __nvvm_atom_sys_xor_gen_l(&dl, l);597 // CHECK: call i64 @llvm.nvvm.atomic.xor.gen.i.sys.i64.p0598 // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_ll' needs target feature sm_60}}599 __nvvm_atom_sys_xor_gen_ll(&sll, ll);600 601 // CHECK: call i32 @llvm.nvvm.atomic.cas.gen.i.cta.i32.p0602 // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_i' needs target feature sm_60}}603 __nvvm_atom_cta_cas_gen_i(ip, i, 0);604 // LP32: call i32 @llvm.nvvm.atomic.cas.gen.i.cta.i32.p0605 // LP64: call i64 @llvm.nvvm.atomic.cas.gen.i.cta.i64.p0606 // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_l' needs target feature sm_60}}607 __nvvm_atom_cta_cas_gen_l(&dl, l, 0);608 // CHECK: call i64 @llvm.nvvm.atomic.cas.gen.i.cta.i64.p0609 // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_ll' needs target feature sm_60}}610 __nvvm_atom_cta_cas_gen_ll(&sll, ll, 0);611 612 // CHECK: call i32 @llvm.nvvm.atomic.cas.gen.i.sys.i32.p0613 // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_i' needs target feature sm_60}}614 __nvvm_atom_sys_cas_gen_i(ip, i, 0);615 // LP32: call i32 @llvm.nvvm.atomic.cas.gen.i.sys.i32.p0616 // LP64: call i64 @llvm.nvvm.atomic.cas.gen.i.sys.i64.p0617 // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_l' needs target feature sm_60}}618 __nvvm_atom_sys_cas_gen_l(&dl, l, 0);619 // CHECK: call i64 @llvm.nvvm.atomic.cas.gen.i.sys.i64.p0620 // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_ll' needs target feature sm_60}}621 __nvvm_atom_sys_cas_gen_ll(&sll, ll, 0);622#endif623 624#if __CUDA_ARCH__ >= 700625 // CHECK_PTX63_SM70: cmpxchg ptr {{.*}} seq_cst seq_cst, align 2626 // CHECK_PTX63_SM70-NEXT: extractvalue { i16, i1 } {{%[0-9]+}}, 0627 __nvvm_atom_cas_gen_us(usp, 0, us);628 // CHECK_PTX63_SM70: call i16 @llvm.nvvm.atomic.cas.gen.i.cta.i16.p0629 __nvvm_atom_cta_cas_gen_us(usp, 0, us);630 // CHECK_PTX63_SM70: call i16 @llvm.nvvm.atomic.cas.gen.i.sys.i16.p0631 __nvvm_atom_sys_cas_gen_us(usp, 0, us);632#endif633 634 // CHECK: ret635}636 637// CHECK-LABEL: nvvm_ldg638__device__ void nvvm_ldg(const void *p) {639 // CHECK: load i8, ptr addrspace(1) {{%[0-9]+}}, align 1, !invariant.load640 // CHECK: load i8, ptr addrspace(1) {{%[0-9]+}}, align 1, !invariant.load641 // CHECK: load i8, ptr addrspace(1) {{%[0-9]+}}, align 1, !invariant.load642 __nvvm_ldg_c((const char *)p);643 __nvvm_ldg_uc((const unsigned char *)p);644 __nvvm_ldg_sc((const signed char *)p);645 646 // CHECK: load i16, ptr addrspace(1) {{%[0-9]+}}, align 2, !invariant.load647 // CHECK: load i16, ptr addrspace(1) {{%[0-9]+}}, align 2, !invariant.load648 __nvvm_ldg_s((const short *)p);649 __nvvm_ldg_us((const unsigned short *)p);650 651 // CHECK: load i32, ptr addrspace(1) {{%[0-9]+}}, align 4, !invariant.load652 // CHECK: load i32, ptr addrspace(1) {{%[0-9]+}}, align 4, !invariant.load653 __nvvm_ldg_i((const int *)p);654 __nvvm_ldg_ui((const unsigned int *)p);655 656 // LP32: load i32, ptr addrspace(1) {{%[0-9]+}}, align 4, !invariant.load657 // LP32: load i32, ptr addrspace(1) {{%[0-9]+}}, align 4, !invariant.load658 // LP64: load i64, ptr addrspace(1) {{%[0-9]+}}, align 8, !invariant.load659 // LP64: load i64, ptr addrspace(1) {{%[0-9]+}}, align 8, !invariant.load660 __nvvm_ldg_l((const long *)p);661 __nvvm_ldg_ul((const unsigned long *)p);662 663 // CHECK: load float, ptr addrspace(1) {{%[0-9]+}}, align 4, !invariant.load664 __nvvm_ldg_f((const float *)p);665 // CHECK: load double, ptr addrspace(1) {{%[0-9]+}}, align 8, !invariant.load666 __nvvm_ldg_d((const double *)p);667 668 // In practice, the pointers we pass to __ldg will be aligned as appropriate669 // for the CUDA <type>N vector types (e.g. short4), which are not the same as670 // the LLVM vector types. However, each LLVM vector type has an alignment671 // less than or equal to its corresponding CUDA type, so we're OK.672 //673 // PTX Interoperability section 2.2: "For a vector with an even number of674 // elements, its alignment is set to number of elements times the alignment of675 // its member: n*alignof(t)."676 677 // CHECK: load <2 x i8>, ptr addrspace(1) {{%[0-9]+}}, align 2, !invariant.load678 // CHECK: load <2 x i8>, ptr addrspace(1) {{%[0-9]+}}, align 2, !invariant.load679 // CHECK: load <2 x i8>, ptr addrspace(1) {{%[0-9]+}}, align 2, !invariant.load680 typedef char char2 __attribute__((ext_vector_type(2)));681 typedef unsigned char uchar2 __attribute__((ext_vector_type(2)));682 typedef signed char schar2 __attribute__((ext_vector_type(2)));683 __nvvm_ldg_c2((const char2 *)p);684 __nvvm_ldg_uc2((const uchar2 *)p);685 __nvvm_ldg_sc2((const schar2 *)p);686 687 // CHECK: load <4 x i8>, ptr addrspace(1) {{%[0-9]+}}, align 4, !invariant.load688 // CHECK: load <4 x i8>, ptr addrspace(1) {{%[0-9]+}}, align 4, !invariant.load689 // CHECK: load <4 x i8>, ptr addrspace(1) {{%[0-9]+}}, align 4, !invariant.load690 typedef char char4 __attribute__((ext_vector_type(4)));691 typedef unsigned char uchar4 __attribute__((ext_vector_type(4)));692 typedef signed char schar4 __attribute__((ext_vector_type(4)));693 __nvvm_ldg_c4((const char4 *)p);694 __nvvm_ldg_uc4((const uchar4 *)p);695 __nvvm_ldg_sc4((const schar4 *)p);696 697 // CHECK: load <2 x i16>, ptr addrspace(1) {{%[0-9]+}}, align 4, !invariant.load698 // CHECK: load <2 x i16>, ptr addrspace(1) {{%[0-9]+}}, align 4, !invariant.load699 typedef short short2 __attribute__((ext_vector_type(2)));700 typedef unsigned short ushort2 __attribute__((ext_vector_type(2)));701 __nvvm_ldg_s2((const short2 *)p);702 __nvvm_ldg_us2((const ushort2 *)p);703 704 // CHECK: load <4 x i16>, ptr addrspace(1) {{%[0-9]+}}, align 8, !invariant.load705 // CHECK: load <4 x i16>, ptr addrspace(1) {{%[0-9]+}}, align 8, !invariant.load706 typedef short short4 __attribute__((ext_vector_type(4)));707 typedef unsigned short ushort4 __attribute__((ext_vector_type(4)));708 __nvvm_ldg_s4((const short4 *)p);709 __nvvm_ldg_us4((const ushort4 *)p);710 711 // CHECK: load <2 x i32>, ptr addrspace(1) {{%[0-9]+}}, align 8, !invariant.load712 // CHECK: load <2 x i32>, ptr addrspace(1) {{%[0-9]+}}, align 8, !invariant.load713 typedef int int2 __attribute__((ext_vector_type(2)));714 typedef unsigned int uint2 __attribute__((ext_vector_type(2)));715 __nvvm_ldg_i2((const int2 *)p);716 __nvvm_ldg_ui2((const uint2 *)p);717 718 // CHECK: load <4 x i32>, ptr addrspace(1) {{%[0-9]+}}, align 16, !invariant.load719 // CHECK: load <4 x i32>, ptr addrspace(1) {{%[0-9]+}}, align 16, !invariant.load720 typedef int int4 __attribute__((ext_vector_type(4)));721 typedef unsigned int uint4 __attribute__((ext_vector_type(4)));722 __nvvm_ldg_i4((const int4 *)p);723 __nvvm_ldg_ui4((const uint4 *)p);724 725 // LP32: load <2 x i32>, ptr addrspace(1) {{%[0-9]+}}, align 8, !invariant.load726 // LP32: load <2 x i32>, ptr addrspace(1) {{%[0-9]+}}, align 8, !invariant.load727 // LP64: load <2 x i64>, ptr addrspace(1) {{%[0-9]+}}, align 16, !invariant.load728 // LP64: load <2 x i64>, ptr addrspace(1) {{%[0-9]+}}, align 16, !invariant.load729 typedef long long2 __attribute__((ext_vector_type(2)));730 typedef unsigned long ulong2 __attribute__((ext_vector_type(2)));731 __nvvm_ldg_l2((const long2 *)p);732 __nvvm_ldg_ul2((const ulong2 *)p);733 734 // CHECK: load <2 x i64>, ptr addrspace(1) {{%[0-9]+}}, align 16, !invariant.load735 // CHECK: load <2 x i64>, ptr addrspace(1) {{%[0-9]+}}, align 16, !invariant.load736 typedef long long longlong2 __attribute__((ext_vector_type(2)));737 typedef unsigned long long ulonglong2 __attribute__((ext_vector_type(2)));738 __nvvm_ldg_ll2((const longlong2 *)p);739 __nvvm_ldg_ull2((const ulonglong2 *)p);740 741 // CHECK: load <2 x float>, ptr addrspace(1) {{%[0-9]+}}, align 8, !invariant.load742 typedef float float2 __attribute__((ext_vector_type(2)));743 __nvvm_ldg_f2((const float2 *)p);744 745 // CHECK: load <4 x float>, ptr addrspace(1) {{%[0-9]+}}, align 16, !invariant.load746 typedef float float4 __attribute__((ext_vector_type(4)));747 __nvvm_ldg_f4((const float4 *)p);748 749 // CHECK: load <2 x double>, ptr addrspace(1) {{%[0-9]+}}, align 16, !invariant.load750 typedef double double2 __attribute__((ext_vector_type(2)));751 __nvvm_ldg_d2((const double2 *)p);752}753 754// CHECK-LABEL: nvvm_ldu755__device__ void nvvm_ldu(const void *p) {756 // CHECK: call i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr {{%[0-9]+}}, i32 1)757 // CHECK: call i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr {{%[0-9]+}}, i32 1)758 // CHECK: call i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr {{%[0-9]+}}, i32 1)759 __nvvm_ldu_c((const char *)p);760 __nvvm_ldu_uc((const unsigned char *)p);761 __nvvm_ldu_sc((const signed char *)p);762 763 // CHECK: call i16 @llvm.nvvm.ldu.global.i.i16.p0(ptr {{%[0-9]+}}, i32 2)764 // CHECK: call i16 @llvm.nvvm.ldu.global.i.i16.p0(ptr {{%[0-9]+}}, i32 2)765 __nvvm_ldu_s((const short *)p);766 __nvvm_ldu_us((const unsigned short *)p);767 768 // CHECK: call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr {{%[0-9]+}}, i32 4)769 // CHECK: call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr {{%[0-9]+}}, i32 4)770 __nvvm_ldu_i((const int *)p);771 __nvvm_ldu_ui((const unsigned int *)p);772 773 // LP32: call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr {{%[0-9]+}}, i32 4)774 // LP32: call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr {{%[0-9]+}}, i32 4)775 // LP64: call i64 @llvm.nvvm.ldu.global.i.i64.p0(ptr {{%[0-9]+}}, i32 8)776 // LP64: call i64 @llvm.nvvm.ldu.global.i.i64.p0(ptr {{%[0-9]+}}, i32 8)777 __nvvm_ldu_l((const long *)p);778 __nvvm_ldu_ul((const unsigned long *)p);779 780 // CHECK: call float @llvm.nvvm.ldu.global.f.f32.p0(ptr {{%[0-9]+}}, i32 4)781 __nvvm_ldu_f((const float *)p);782 // CHECK: call double @llvm.nvvm.ldu.global.f.f64.p0(ptr {{%[0-9]+}}, i32 8)783 __nvvm_ldu_d((const double *)p);784 785 // CHECK: call <2 x i8> @llvm.nvvm.ldu.global.i.v2i8.p0(ptr {{%[0-9]+}}, i32 2)786 // CHECK: call <2 x i8> @llvm.nvvm.ldu.global.i.v2i8.p0(ptr {{%[0-9]+}}, i32 2)787 // CHECK: call <2 x i8> @llvm.nvvm.ldu.global.i.v2i8.p0(ptr {{%[0-9]+}}, i32 2)788 typedef char char2 __attribute__((ext_vector_type(2)));789 typedef unsigned char uchar2 __attribute__((ext_vector_type(2)));790 typedef signed char schar2 __attribute__((ext_vector_type(2)));791 __nvvm_ldu_c2((const char2 *)p);792 __nvvm_ldu_uc2((const uchar2 *)p);793 __nvvm_ldu_sc2((const schar2 *)p);794 795 // CHECK: call <4 x i8> @llvm.nvvm.ldu.global.i.v4i8.p0(ptr {{%[0-9]+}}, i32 4)796 // CHECK: call <4 x i8> @llvm.nvvm.ldu.global.i.v4i8.p0(ptr {{%[0-9]+}}, i32 4)797 // CHECK: call <4 x i8> @llvm.nvvm.ldu.global.i.v4i8.p0(ptr {{%[0-9]+}}, i32 4)798 typedef char char4 __attribute__((ext_vector_type(4)));799 typedef unsigned char uchar4 __attribute__((ext_vector_type(4)));800 typedef signed char schar4 __attribute__((ext_vector_type(4)));801 __nvvm_ldu_c4((const char4 *)p);802 __nvvm_ldu_uc4((const uchar4 *)p);803 __nvvm_ldu_sc4((const schar4 *)p);804 805 // CHECK: call <2 x i16> @llvm.nvvm.ldu.global.i.v2i16.p0(ptr {{%[0-9]+}}, i32 4)806 // CHECK: call <2 x i16> @llvm.nvvm.ldu.global.i.v2i16.p0(ptr {{%[0-9]+}}, i32 4)807 typedef short short2 __attribute__((ext_vector_type(2)));808 typedef unsigned short ushort2 __attribute__((ext_vector_type(2)));809 __nvvm_ldu_s2((const short2 *)p);810 __nvvm_ldu_us2((const ushort2 *)p);811 812 // CHECK: call <4 x i16> @llvm.nvvm.ldu.global.i.v4i16.p0(ptr {{%[0-9]+}}, i32 8)813 // CHECK: call <4 x i16> @llvm.nvvm.ldu.global.i.v4i16.p0(ptr {{%[0-9]+}}, i32 8)814 typedef short short4 __attribute__((ext_vector_type(4)));815 typedef unsigned short ushort4 __attribute__((ext_vector_type(4)));816 __nvvm_ldu_s4((const short4 *)p);817 __nvvm_ldu_us4((const ushort4 *)p);818 819 // CHECK: call <2 x i32> @llvm.nvvm.ldu.global.i.v2i32.p0(ptr {{%[0-9]+}}, i32 8)820 // CHECK: call <2 x i32> @llvm.nvvm.ldu.global.i.v2i32.p0(ptr {{%[0-9]+}}, i32 8)821 typedef int int2 __attribute__((ext_vector_type(2)));822 typedef unsigned int uint2 __attribute__((ext_vector_type(2)));823 __nvvm_ldu_i2((const int2 *)p);824 __nvvm_ldu_ui2((const uint2 *)p);825 826 // CHECK: call <4 x i32> @llvm.nvvm.ldu.global.i.v4i32.p0(ptr {{%[0-9]+}}, i32 16)827 // CHECK: call <4 x i32> @llvm.nvvm.ldu.global.i.v4i32.p0(ptr {{%[0-9]+}}, i32 16)828 typedef int int4 __attribute__((ext_vector_type(4)));829 typedef unsigned int uint4 __attribute__((ext_vector_type(4)));830 __nvvm_ldu_i4((const int4 *)p);831 __nvvm_ldu_ui4((const uint4 *)p);832 833 // LP32: call <2 x i32> @llvm.nvvm.ldu.global.i.v2i32.p0(ptr {{%[0-9]+}}, i32 8)834 // LP32: call <2 x i32> @llvm.nvvm.ldu.global.i.v2i32.p0(ptr {{%[0-9]+}}, i32 8)835 // LP64: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16)836 // LP64: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16)837 typedef long long2 __attribute__((ext_vector_type(2)));838 typedef unsigned long ulong2 __attribute__((ext_vector_type(2)));839 __nvvm_ldu_l2((const long2 *)p);840 __nvvm_ldu_ul2((const ulong2 *)p);841 842 // CHECK: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16)843 // CHECK: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16)844 typedef long long longlong2 __attribute__((ext_vector_type(2)));845 typedef unsigned long long ulonglong2 __attribute__((ext_vector_type(2)));846 __nvvm_ldu_ll2((const longlong2 *)p);847 __nvvm_ldu_ull2((const ulonglong2 *)p);848 849 // CHECK: call <2 x float> @llvm.nvvm.ldu.global.f.v2f32.p0(ptr {{%[0-9]+}}, i32 8)850 typedef float float2 __attribute__((ext_vector_type(2)));851 __nvvm_ldu_f2((const float2 *)p);852 853 // CHECK: call <4 x float> @llvm.nvvm.ldu.global.f.v4f32.p0(ptr {{%[0-9]+}}, i32 16)854 typedef float float4 __attribute__((ext_vector_type(4)));855 __nvvm_ldu_f4((const float4 *)p);856 857 // CHECK: call <2 x double> @llvm.nvvm.ldu.global.f.v2f64.p0(ptr {{%[0-9]+}}, i32 16)858 typedef double double2 __attribute__((ext_vector_type(2)));859 __nvvm_ldu_d2((const double2 *)p);860}861 862// CHECK-LABEL: nvvm_shfl863__device__ void nvvm_shfl(int i, float f, int a, int b) {864 // CHECK: call i32 @llvm.nvvm.shfl.down.i32(i32865 __nvvm_shfl_down_i32(i, a, b);866 // CHECK: call float @llvm.nvvm.shfl.down.f32(float867 __nvvm_shfl_down_f32(f, a, b);868 // CHECK: call i32 @llvm.nvvm.shfl.up.i32(i32869 __nvvm_shfl_up_i32(i, a, b);870 // CHECK: call float @llvm.nvvm.shfl.up.f32(float871 __nvvm_shfl_up_f32(f, a, b);872 // CHECK: call i32 @llvm.nvvm.shfl.bfly.i32(i32873 __nvvm_shfl_bfly_i32(i, a, b);874 // CHECK: call float @llvm.nvvm.shfl.bfly.f32(float875 __nvvm_shfl_bfly_f32(f, a, b);876 // CHECK: call i32 @llvm.nvvm.shfl.idx.i32(i32877 __nvvm_shfl_idx_i32(i, a, b);878 // CHECK: call float @llvm.nvvm.shfl.idx.f32(float879 __nvvm_shfl_idx_f32(f, a, b);880 // CHECK: ret void881}882 883__device__ void nvvm_vote(int pred) {884 // CHECK: call i1 @llvm.nvvm.vote.all(i1885 __nvvm_vote_all(pred);886 // CHECK: call i1 @llvm.nvvm.vote.any(i1887 __nvvm_vote_any(pred);888 // CHECK: call i1 @llvm.nvvm.vote.uni(i1889 __nvvm_vote_uni(pred);890 // CHECK: call i32 @llvm.nvvm.vote.ballot(i1891 __nvvm_vote_ballot(pred);892 // CHECK: ret void893}894 895// CHECK-LABEL: nvvm_pm_event_mask896__device__ void nvvm_pm_event_mask() {897 // CHECK: call void @llvm.nvvm.pm.event.mask(i16 255)898 __nvvm_pm_event_mask(255);899 // CHECK: ret void900}901 902// CHECK-LABEL: nvvm_nanosleep903__device__ void nvvm_nanosleep(int d) {904#if __CUDA_ARCH__ >= 700905 // CHECK_PTX70_SM80: call void @llvm.nvvm.nanosleep906 __nvvm_nanosleep(d);907 908 // CHECK_PTX70_SM80: call void @llvm.nvvm.nanosleep909 __nvvm_nanosleep(1);910#endif911}912 913// CHECK-LABEL: nvvm_mbarrier914__device__ void nvvm_mbarrier(long long* addr, __attribute__((address_space(3))) long long* sharedAddr, int count, long long state) {915 #if __CUDA_ARCH__ >= 800916 __nvvm_mbarrier_init(addr, count);917 // CHECK_PTX70_SM80: call void @llvm.nvvm.mbarrier.init918 __nvvm_mbarrier_init_shared(sharedAddr, count);919 // CHECK_PTX70_SM80: call void @llvm.nvvm.mbarrier.init.shared920 921 __nvvm_mbarrier_inval(addr);922 // CHECK_PTX70_SM80: call void @llvm.nvvm.mbarrier.inval923 __nvvm_mbarrier_inval_shared(sharedAddr);924 // CHECK_PTX70_SM80: call void @llvm.nvvm.mbarrier.inval.shared925 926 __nvvm_mbarrier_arrive(addr);927 // CHECK_PTX70_SM80: call i64 @llvm.nvvm.mbarrier.arrive928 __nvvm_mbarrier_arrive_shared(sharedAddr);929 // CHECK_PTX70_SM80: call i64 @llvm.nvvm.mbarrier.arrive.shared930 __nvvm_mbarrier_arrive_noComplete(addr, count);931 // CHECK_PTX70_SM80: call i64 @llvm.nvvm.mbarrier.arrive.noComplete932 __nvvm_mbarrier_arrive_noComplete_shared(sharedAddr, count);933 // CHECK_PTX70_SM80: call i64 @llvm.nvvm.mbarrier.arrive.noComplete.shared934 935 __nvvm_mbarrier_arrive_drop(addr);936 // CHECK_PTX70_SM80: call i64 @llvm.nvvm.mbarrier.arrive.drop937 __nvvm_mbarrier_arrive_drop_shared(sharedAddr);938 // CHECK_PTX70_SM80: call i64 @llvm.nvvm.mbarrier.arrive.drop.shared939 __nvvm_mbarrier_arrive_drop_noComplete(addr, count);940 // CHECK_PTX70_SM80: call i64 @llvm.nvvm.mbarrier.arrive.drop.noComplete941 __nvvm_mbarrier_arrive_drop_noComplete_shared(sharedAddr, count);942 // CHECK_PTX70_SM80: call i64 @llvm.nvvm.mbarrier.arrive.drop.noComplete.shared943 944 __nvvm_mbarrier_test_wait(addr, state);945 // CHECK_PTX70_SM80: call i1 @llvm.nvvm.mbarrier.test.wait946 __nvvm_mbarrier_test_wait_shared(sharedAddr, state);947 // CHECK_PTX70_SM80: call i1 @llvm.nvvm.mbarrier.test.wait.shared948 949 __nvvm_mbarrier_pending_count(state);950 // CHECK_PTX70_SM80: call i32 @llvm.nvvm.mbarrier.pending.count951 #endif952 // CHECK: ret void953}954 955// CHECK-LABEL: nvvm_async_copy956__device__ void nvvm_async_copy(__attribute__((address_space(3))) void* dst, __attribute__((address_space(1))) const void* src, long long* addr, __attribute__((address_space(3))) long long* sharedAddr) {957 #if __CUDA_ARCH__ >= 800958 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.mbarrier.arrive959 __nvvm_cp_async_mbarrier_arrive(addr);960 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.mbarrier.arrive.shared961 __nvvm_cp_async_mbarrier_arrive_shared(sharedAddr);962 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.mbarrier.arrive.noinc963 __nvvm_cp_async_mbarrier_arrive_noinc(addr);964 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.mbarrier.arrive.noinc.shared965 __nvvm_cp_async_mbarrier_arrive_noinc_shared(sharedAddr);966 967 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.4(968 __nvvm_cp_async_ca_shared_global_4(dst, src);969 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.8(970 __nvvm_cp_async_ca_shared_global_8(dst, src);971 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.16(972 __nvvm_cp_async_ca_shared_global_16(dst, src);973 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.cg.shared.global.16(974 __nvvm_cp_async_cg_shared_global_16(dst, src);975 976 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.4.s({{.*}}, i32 2)977 __nvvm_cp_async_ca_shared_global_4(dst, src, 2);978 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.8.s({{.*}}, i32 2)979 __nvvm_cp_async_ca_shared_global_8(dst, src, 2);980 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.16.s({{.*}}, i32 2)981 __nvvm_cp_async_ca_shared_global_16(dst, src, 2);982 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.cg.shared.global.16.s({{.*}}, i32 2)983 __nvvm_cp_async_cg_shared_global_16(dst, src, 2);984 985 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.commit.group986 __nvvm_cp_async_commit_group();987 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.wait.group(i32 0)988 __nvvm_cp_async_wait_group(0);989 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.wait.group(i32 8)990 __nvvm_cp_async_wait_group(8);991 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.wait.group(i32 16)992 __nvvm_cp_async_wait_group(16);993 // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.wait.all994 __nvvm_cp_async_wait_all();995 #endif996 // CHECK: ret void997}998 999// CHECK-LABEL: nvvm_cvt_sm801000__device__ void nvvm_cvt_sm80() {1001#if __CUDA_ARCH__ >= 8001002 // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn(float 1.000000e+00, float 1.000000e+00)1003 __nvvm_ff2bf16x2_rn(1, 1);1004 // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu(float 1.000000e+00, float 1.000000e+00)1005 __nvvm_ff2bf16x2_rn_relu(1, 1);1006 // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz(float 1.000000e+00, float 1.000000e+00)1007 __nvvm_ff2bf16x2_rz(1, 1);1008 // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu(float 1.000000e+00, float 1.000000e+00)1009 __nvvm_ff2bf16x2_rz_relu(1, 1);1010 #if PTX >= 811011 // CHECK_PTX81_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)1012 __nvvm_ff2bf16x2_rn_satfinite(1, 1);1013 // CHECK_PTX81_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)1014 __nvvm_ff2bf16x2_rn_relu_satfinite(1, 1);1015 // CHECK_PTX81_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.satfinite(float 1.000000e+00, float 1.000000e+00)1016 __nvvm_ff2bf16x2_rz_satfinite(1, 1);1017 // CHECK_PTX81_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu.satfinite(float 1.000000e+00, float 1.000000e+00)1018 __nvvm_ff2bf16x2_rz_relu_satfinite(1, 1);1019 #endif1020 1021 // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rn(float 1.000000e+00, float 1.000000e+00)1022 __nvvm_ff2f16x2_rn(1, 1);1023 // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rn.relu(float 1.000000e+00, float 1.000000e+00)1024 __nvvm_ff2f16x2_rn_relu(1, 1);1025 // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rz(float 1.000000e+00, float 1.000000e+00)1026 __nvvm_ff2f16x2_rz(1, 1);1027 // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rz.relu(float 1.000000e+00, float 1.000000e+00)1028 __nvvm_ff2f16x2_rz_relu(1, 1);1029 #if PTX >= 811030 // CHECK_PTX81_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)1031 __nvvm_ff2f16x2_rn_satfinite(1, 1);1032 // CHECK_PTX81_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)1033 __nvvm_ff2f16x2_rn_relu_satfinite(1, 1);1034 // CHECK_PTX81_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rz.satfinite(float 1.000000e+00, float 1.000000e+00)1035 __nvvm_ff2f16x2_rz_satfinite(1, 1);1036 // CHECK_PTX81_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rz.relu.satfinite(float 1.000000e+00, float 1.000000e+00)1037 __nvvm_ff2f16x2_rz_relu_satfinite(1, 1);1038 #endif1039 1040 // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.f2bf16.rn(float 1.000000e+00)1041 __nvvm_f2bf16_rn(1);1042 // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.f2bf16.rn.relu(float 1.000000e+00)1043 __nvvm_f2bf16_rn_relu(1);1044 // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.f2bf16.rz(float 1.000000e+00)1045 __nvvm_f2bf16_rz(1);1046 // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.f2bf16.rz.relu(float 1.000000e+00)1047 __nvvm_f2bf16_rz_relu(1);1048 #if PTX >= 811049 // CHECK_PTX81_SM80: call bfloat @llvm.nvvm.f2bf16.rn.satfinite(float 1.000000e+00)1050 __nvvm_f2bf16_rn_satfinite(1);1051 // CHECK_PTX81_SM80: call bfloat @llvm.nvvm.f2bf16.rn.relu.satfinite(float 1.000000e+00)1052 __nvvm_f2bf16_rn_relu_satfinite(1);1053 // CHECK_PTX81_SM80: call bfloat @llvm.nvvm.f2bf16.rz.satfinite(float 1.000000e+00)1054 __nvvm_f2bf16_rz_satfinite(1);1055 // CHECK_PTX81_SM80: call bfloat @llvm.nvvm.f2bf16.rz.relu.satfinite(float 1.000000e+00)1056 __nvvm_f2bf16_rz_relu_satfinite(1);1057 #endif1058 1059 // CHECK_PTX70_SM80: call half @llvm.nvvm.f2f16.rn(float 1.000000e+00)1060 __nvvm_f2f16_rn(1);1061 // CHECK_PTX70_SM80: call half @llvm.nvvm.f2f16.rn.relu(float 1.000000e+00)1062 __nvvm_f2f16_rn_relu(1);1063 // CHECK_PTX70_SM80: call half @llvm.nvvm.f2f16.rz(float 1.000000e+00)1064 __nvvm_f2f16_rz(1);1065 // CHECK_PTX70_SM80: call half @llvm.nvvm.f2f16.rz.relu(float 1.000000e+00)1066 __nvvm_f2f16_rz_relu(1);1067 #if PTX >= 811068 // CHECK_PTX81_SM80: call half @llvm.nvvm.f2f16.rn.satfinite(float 1.000000e+00)1069 __nvvm_f2f16_rn_satfinite(1);1070 // CHECK_PTX81_SM80: call half @llvm.nvvm.f2f16.rn.relu.satfinite(float 1.000000e+00)1071 __nvvm_f2f16_rn_relu_satfinite(1);1072 // CHECK_PTX81_SM80: call half @llvm.nvvm.f2f16.rz.satfinite(float 1.000000e+00)1073 __nvvm_f2f16_rz_satfinite(1);1074 // CHECK_PTX81_SM80: call half @llvm.nvvm.f2f16.rz.relu.satfinite(float 1.000000e+00)1075 __nvvm_f2f16_rz_relu_satfinite(1);1076 #endif1077 1078 // CHECK_PTX70_SM80: call i32 @llvm.nvvm.f2tf32.rna(float 1.000000e+00)1079 __nvvm_f2tf32_rna(1);1080 #if PTX >= 811081 // CHECK_PTX81_SM80: call i32 @llvm.nvvm.f2tf32.rna.satfinite(float 1.000000e+00)1082 __nvvm_f2tf32_rna_satfinite(1.0f);1083 #endif1084#endif1085 // CHECK: ret void1086}1087 1088// CHECK-LABEL: nvvm_cvt_sm891089__device__ void nvvm_cvt_sm89() {1090#if (PTX >= 81) && (__CUDA_ARCH__ >= 890)1091 // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e4m3x2.rn(float 1.000000e+00, float 1.000000e+00)1092 __nvvm_ff_to_e4m3x2_rn(1.0f, 1.0f);1093 // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e4m3x2.rn.relu(float 1.000000e+00, float 1.000000e+00)1094 __nvvm_ff_to_e4m3x2_rn_relu(1.0f, 1.0f);1095 // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e5m2x2.rn(float 1.000000e+00, float 1.000000e+00)1096 __nvvm_ff_to_e5m2x2_rn(1.0f, 1.0f);1097 // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e5m2x2.rn.relu(float 1.000000e+00, float 1.000000e+00)1098 __nvvm_ff_to_e5m2x2_rn_relu(1.0f, 1.0f);1099 1100 // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn(<2 x half> splat (half 0xH3C00))1101 __nvvm_f16x2_to_e4m3x2_rn({1.0f16, 1.0f16});1102 // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn.relu(<2 x half> splat (half 0xH3C00))1103 __nvvm_f16x2_to_e4m3x2_rn_relu({1.0f16, 1.0f16});1104 // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> splat (half 0xH3C00))1105 __nvvm_f16x2_to_e5m2x2_rn({1.0f16, 1.0f16});1106 // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn.relu(<2 x half> splat (half 0xH3C00))1107 __nvvm_f16x2_to_e5m2x2_rn_relu({1.0f16, 1.0f16});1108 1109 // CHECK_PTX81_SM89: call <2 x half> @llvm.nvvm.e4m3x2.to.f16x2.rn(i16 18504)1110 __nvvm_e4m3x2_to_f16x2_rn(0x4848);1111 // CHECK_PTX81_SM89: call <2 x half> @llvm.nvvm.e4m3x2.to.f16x2.rn.relu(i16 18504)1112 __nvvm_e4m3x2_to_f16x2_rn_relu(0x4848);1113 // CHECK_PTX81_SM89: call <2 x half> @llvm.nvvm.e5m2x2.to.f16x2.rn(i16 19532)1114 __nvvm_e5m2x2_to_f16x2_rn(0x4c4c);1115 // CHECK_PTX81_SM89: call <2 x half> @llvm.nvvm.e5m2x2.to.f16x2.rn.relu(i16 19532)1116 __nvvm_e5m2x2_to_f16x2_rn_relu(0x4c4c);1117#endif1118 // CHECK: ret void1119}1120 1121// CHECK-LABEL: nvvm_cvt_sm901122__device__ void nvvm_cvt_sm90() {1123#if (PTX >= 78) && (__CUDA_ARCH__ >= 900)1124 // CHECK_PTX78_SM90: call i32 @llvm.nvvm.f2tf32.rn(float 1.000000e+00)1125 __nvvm_f2tf32_rn(1.0f); 1126 // CHECK_PTX78_SM90: call i32 @llvm.nvvm.f2tf32.rn.relu(float 1.000000e+00)1127 __nvvm_f2tf32_rn_relu(1.0f); 1128 // CHECK_PTX78_SM90: call i32 @llvm.nvvm.f2tf32.rz(float 1.000000e+00)1129 __nvvm_f2tf32_rz(1.0f); 1130 // CHECK_PTX78_SM90: call i32 @llvm.nvvm.f2tf32.rz.relu(float 1.000000e+00)1131 __nvvm_f2tf32_rz_relu(1.0f); 1132#endif1133 // CHECK: ret void1134}1135 1136// CHECK-LABEL: nvvm_cvt_sm1001137__device__ void nvvm_cvt_sm100() {1138#if (PTX >= 86) && (__CUDA_ARCH__ >= 1000)1139 // CHECK_PTX86_SM100: call i32 @llvm.nvvm.f2tf32.rn.satfinite(float 1.000000e+00)1140 __nvvm_f2tf32_rn_satfinite(1.0f); 1141 // CHECK_PTX86_SM100: call i32 @llvm.nvvm.f2tf32.rn.relu.satfinite(float 1.000000e+00)1142 __nvvm_f2tf32_rn_relu_satfinite(1.0f); 1143 // CHECK_PTX86_SM100: call i32 @llvm.nvvm.f2tf32.rz.satfinite(float 1.000000e+00)1144 __nvvm_f2tf32_rz_satfinite(1.0f); 1145 // CHECK_PTX86_SM100: call i32 @llvm.nvvm.f2tf32.rz.relu.satfinite(float 1.000000e+00)1146 __nvvm_f2tf32_rz_relu_satfinite(1.0f); 1147#endif1148 // CHECK: ret void1149}1150 1151// CHECK-LABEL: nvvm_cvt_sm100a_sm101a_sm120a1152__device__ void nvvm_cvt_sm100a_sm101a_sm120a() {1153#if (PTX >= 86) && \1154 (__CUDA_ARCH_FEAT_SM100_ALL || __CUDA_ARCH_FEAT_SM101_ALL || \1155 __CUDA_ARCH_FEAT_SM120_ALL)1156 // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)1157 // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)1158 // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)1159 __nvvm_ff_to_e2m3x2_rn_satfinite(1.0f, 1.0f);1160 1161 // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)1162 // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)1163 // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)1164 __nvvm_ff_to_e2m3x2_rn_relu_satfinite(1.0f, 1.0f);1165 1166 // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)1167 // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)1168 // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)1169 __nvvm_ff_to_e3m2x2_rn_satfinite(1.0f, 1.0f);1170 1171 // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)1172 // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)1173 // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)1174 __nvvm_ff_to_e3m2x2_rn_relu_satfinite(1.0f, 1.0f);1175 1176 // CHECK_PTX86_SM100a: call <2 x half> @llvm.nvvm.e2m3x2.to.f16x2.rn(i16 19532)1177 // CHECK_PTX86_SM101a: call <2 x half> @llvm.nvvm.e2m3x2.to.f16x2.rn(i16 19532)1178 // CHECK_PTX86_SM120a: call <2 x half> @llvm.nvvm.e2m3x2.to.f16x2.rn(i16 19532)1179 __nvvm_e2m3x2_to_f16x2_rn(0x4C4C);1180 1181 // CHECK_PTX86_SM100a: call <2 x half> @llvm.nvvm.e2m3x2.to.f16x2.rn.relu(i16 18504)1182 // CHECK_PTX86_SM101a: call <2 x half> @llvm.nvvm.e2m3x2.to.f16x2.rn.relu(i16 18504)1183 // CHECK_PTX86_SM120a: call <2 x half> @llvm.nvvm.e2m3x2.to.f16x2.rn.relu(i16 18504)1184 __nvvm_e2m3x2_to_f16x2_rn_relu(0x4848);1185 1186 // CHECK_PTX86_SM100a: call <2 x half> @llvm.nvvm.e3m2x2.to.f16x2.rn(i16 18504)1187 // CHECK_PTX86_SM101a: call <2 x half> @llvm.nvvm.e3m2x2.to.f16x2.rn(i16 18504)1188 // CHECK_PTX86_SM120a: call <2 x half> @llvm.nvvm.e3m2x2.to.f16x2.rn(i16 18504)1189 __nvvm_e3m2x2_to_f16x2_rn(0x4848);1190 1191 // CHECK_PTX86_SM100a: call <2 x half> @llvm.nvvm.e3m2x2.to.f16x2.rn.relu(i16 19532)1192 // CHECK_PTX86_SM101a: call <2 x half> @llvm.nvvm.e3m2x2.to.f16x2.rn.relu(i16 19532)1193 // CHECK_PTX86_SM120a: call <2 x half> @llvm.nvvm.e3m2x2.to.f16x2.rn.relu(i16 19532)1194 __nvvm_e3m2x2_to_f16x2_rn_relu(0x4C4C);1195 1196 // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)1197 // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)1198 // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)1199 __nvvm_ff_to_e2m1x2_rn_satfinite(1.0f, 1.0f);1200 1201 // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)1202 // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)1203 // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)1204 __nvvm_ff_to_e2m1x2_rn_relu_satfinite(1.0f, 1.0f);1205 1206 // CHECK_PTX86_SM100a: call <2 x half> @llvm.nvvm.e2m1x2.to.f16x2.rn(i16 76)1207 // CHECK_PTX86_SM101a: call <2 x half> @llvm.nvvm.e2m1x2.to.f16x2.rn(i16 76)1208 // CHECK_PTX86_SM120a: call <2 x half> @llvm.nvvm.e2m1x2.to.f16x2.rn(i16 76)1209 __nvvm_e2m1x2_to_f16x2_rn(0x004C);1210 1211 // CHECK_PTX86_SM100a: call <2 x half> @llvm.nvvm.e2m1x2.to.f16x2.rn.relu(i16 76)1212 // CHECK_PTX86_SM101a: call <2 x half> @llvm.nvvm.e2m1x2.to.f16x2.rn.relu(i16 76)1213 // CHECK_PTX86_SM120a: call <2 x half> @llvm.nvvm.e2m1x2.to.f16x2.rn.relu(i16 76)1214 __nvvm_e2m1x2_to_f16x2_rn_relu(0x004C);1215 1216 // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.ue8m0x2.rz(float 1.000000e+00, float 1.000000e+00)1217 // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.ue8m0x2.rz(float 1.000000e+00, float 1.000000e+00)1218 // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.ue8m0x2.rz(float 1.000000e+00, float 1.000000e+00)1219 __nvvm_ff_to_ue8m0x2_rz(1.0f, 1.0f);1220 1221 // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.ue8m0x2.rz.satfinite(float 1.000000e+00, float 1.000000e+00)1222 // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.ue8m0x2.rz.satfinite(float 1.000000e+00, float 1.000000e+00)1223 // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.ue8m0x2.rz.satfinite(float 1.000000e+00, float 1.000000e+00)1224 __nvvm_ff_to_ue8m0x2_rz_satfinite(1.0f, 1.0f);1225 1226 // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.ue8m0x2.rp(float 1.000000e+00, float 1.000000e+00)1227 // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.ue8m0x2.rp(float 1.000000e+00, float 1.000000e+00)1228 // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.ue8m0x2.rp(float 1.000000e+00, float 1.000000e+00)1229 __nvvm_ff_to_ue8m0x2_rp(1.0f, 1.0f);1230 1231 // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.ue8m0x2.rp.satfinite(float 1.000000e+00, float 1.000000e+00)1232 // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.ue8m0x2.rp.satfinite(float 1.000000e+00, float 1.000000e+00)1233 // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.ue8m0x2.rp.satfinite(float 1.000000e+00, float 1.000000e+00)1234 __nvvm_ff_to_ue8m0x2_rp_satfinite(1.0f, 1.0f);1235 1236 // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rz(<2 x bfloat> splat (bfloat 0xR3DCD)1237 // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rz(<2 x bfloat> splat (bfloat 0xR3DCD)1238 // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rz(<2 x bfloat> splat (bfloat 0xR3DCD)1239 __nvvm_bf16x2_to_ue8m0x2_rz({(__bf16)0.1f, (__bf16)0.1f});1240 1241 // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rz.satfinite(<2 x bfloat> splat (bfloat 0xR3DCD)1242 // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rz.satfinite(<2 x bfloat> splat (bfloat 0xR3DCD)1243 // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rz.satfinite(<2 x bfloat> splat (bfloat 0xR3DCD)1244 __nvvm_bf16x2_to_ue8m0x2_rz_satfinite({(__bf16)0.1f, (__bf16)0.1f});1245 1246 // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rp(<2 x bfloat> splat (bfloat 0xR3DCD)1247 // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rp(<2 x bfloat> splat (bfloat 0xR3DCD)1248 // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rp(<2 x bfloat> splat (bfloat 0xR3DCD)1249 __nvvm_bf16x2_to_ue8m0x2_rp({(__bf16)0.1f, (__bf16)0.1f});1250 1251 // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rp.satfinite(<2 x bfloat> splat (bfloat 0xR3DCD)1252 // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rp.satfinite(<2 x bfloat> splat (bfloat 0xR3DCD)1253 // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rp.satfinite(<2 x bfloat> splat (bfloat 0xR3DCD)1254 __nvvm_bf16x2_to_ue8m0x2_rp_satfinite({(__bf16)0.1f, (__bf16)0.1f});1255 1256 // CHECK_PTX86_SM100a: call <2 x bfloat> @llvm.nvvm.ue8m0x2.to.bf16x2(i16 19532)1257 // CHECK_PTX86_SM101a: call <2 x bfloat> @llvm.nvvm.ue8m0x2.to.bf16x2(i16 19532)1258 // CHECK_PTX86_SM120a: call <2 x bfloat> @llvm.nvvm.ue8m0x2.to.bf16x2(i16 19532)1259 __nvvm_ue8m0x2_to_bf16x2(0x4C4C);1260 1261#endif1262 // CHECK: ret void1263}1264 1265__device__ void nvvm_cvt_sm100a_sm103a() {1266#if (PTX >= 87) && (__CUDA_ARCH_FEAT_SM100_ALL || __CUDA_ARCH_FEAT_SM103_ALL)1267 1268 typedef __fp16 f16x2 __attribute__((ext_vector_type(2)));1269 typedef __bf16 bf16x2 __attribute__((ext_vector_type(2)));1270 typedef char uint8x4 __attribute__((ext_vector_type(4)));1271 1272// CHECK_PTX87_SM100a: %[[R1:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0)1273// CHECK_PTX87_SM100a: store <2 x half> %[[R1]], ptr %r11274// CHECK_PTX87_SM103a: %[[R1:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0)1275// CHECK_PTX87_SM103a: store <2 x half> %[[R1]], ptr %r11276 f16x2 r1 = __nvvm_ff2f16x2_rs(1.0f, 1.0f, 0);1277 1278// CHECK_PTX87_SM100a: %[[R2:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0)1279// CHECK_PTX87_SM100a: store <2 x half> %[[R2]], ptr %r21280// CHECK_PTX87_SM103a: %[[R2:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0)1281// CHECK_PTX87_SM103a: store <2 x half> %[[R2]], ptr %r21282 f16x2 r2 = __nvvm_ff2f16x2_rs_relu(1.0f, 1.0f, 0);1283 1284// CHECK_PTX87_SM100a: %[[R3:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)1285// CHECK_PTX87_SM100a: store <2 x half> %[[R3]], ptr %r31286// CHECK_PTX87_SM103a: %[[R3:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)1287// CHECK_PTX87_SM103a: store <2 x half> %[[R3]], ptr %r31288 f16x2 r3 = __nvvm_ff2f16x2_rs_satfinite(1.0f, 1.0f, 0);1289 1290// CHECK_PTX87_SM100a: %[[R4:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)1291// CHECK_PTX87_SM100a: store <2 x half> %[[R4]], ptr %r41292// CHECK_PTX87_SM103a: %[[R4:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)1293// CHECK_PTX87_SM103a: store <2 x half> %[[R4]], ptr %r41294 f16x2 r4 = __nvvm_ff2f16x2_rs_relu_satfinite(1.0f, 1.0f, 0);1295 1296// CHECK_PTX87_SM100a: %[[R5:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0)1297// CHECK_PTX87_SM100a: store <2 x bfloat> %[[R5]], ptr %r51298// CHECK_PTX87_SM103a: %[[R5:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0)1299// CHECK_PTX87_SM103a: store <2 x bfloat> %[[R5]], ptr %r51300 bf16x2 r5 = __nvvm_ff2bf16x2_rs(1.0f, 1.0f, 0);1301 1302// CHECK_PTX87_SM100a: %[[R6:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0)1303// CHECK_PTX87_SM100a: store <2 x bfloat> %[[R6]], ptr %r61304// CHECK_PTX87_SM103a: %[[R6:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0)1305// CHECK_PTX87_SM103a: store <2 x bfloat> %[[R6]], ptr %r61306 bf16x2 r6 = __nvvm_ff2bf16x2_rs_relu(1.0f, 1.0f, 0);1307 1308// CHECK_PTX87_SM100a: %[[R7:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)1309// CHECK_PTX87_SM100a: store <2 x bfloat> %[[R7]], ptr %r71310// CHECK_PTX87_SM103a: %[[R7:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)1311// CHECK_PTX87_SM103a: store <2 x bfloat> %[[R7]], ptr %r71312 bf16x2 r7 = __nvvm_ff2bf16x2_rs_satfinite(1.0f, 1.0f, 0);1313 1314// CHECK_PTX87_SM100a: %[[R8:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)1315// CHECK_PTX87_SM100a: store <2 x bfloat> %[[R8]], ptr %r81316// CHECK_PTX87_SM103a: %[[R8:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)1317// CHECK_PTX87_SM103a: store <2 x bfloat> %[[R8]], ptr %r81318 bf16x2 r8 = __nvvm_ff2bf16x2_rs_relu_satfinite(1.0f, 1.0f, 0);1319 1320// CHECK_PTX87_SM100a: %[[R9:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e4m3x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1321// CHECK_PTX87_SM100a: store <4 x i8> %[[R9]], ptr %r91322// CHECK_PTX87_SM103a: %[[R9:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e4m3x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1323// CHECK_PTX87_SM103a: store <4 x i8> %[[R9]], ptr %r91324 uint8x4 r9 = __nvvm_f32x4_to_e4m3x4_rs_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);1325 1326// CHECK_PTX87_SM100a: %[[R10:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e4m3x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1327// CHECK_PTX87_SM100a: store <4 x i8> %[[R10]], ptr %r101328// CHECK_PTX87_SM103a: %[[R10:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e4m3x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1329// CHECK_PTX87_SM103a: store <4 x i8> %[[R10]], ptr %r101330 uint8x4 r10 = __nvvm_f32x4_to_e4m3x4_rs_relu_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);1331 1332// CHECK_PTX87_SM100a: %[[R11:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e5m2x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1333// CHECK_PTX87_SM100a: store <4 x i8> %[[R11]], ptr %r111334// CHECK_PTX87_SM103a: %[[R11:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e5m2x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1335// CHECK_PTX87_SM103a: store <4 x i8> %[[R11]], ptr %r111336 uint8x4 r11 = __nvvm_f32x4_to_e5m2x4_rs_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);1337 1338// CHECK_PTX87_SM100a: %[[R12:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e5m2x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1339// CHECK_PTX87_SM100a: store <4 x i8> %[[R12]], ptr %r121340// CHECK_PTX87_SM103a: %[[R12:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e5m2x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1341// CHECK_PTX87_SM103a: store <4 x i8> %[[R12]], ptr %r121342 uint8x4 r12 = __nvvm_f32x4_to_e5m2x4_rs_relu_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);1343 1344// CHECK_PTX87_SM100a: %[[R13:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e2m3x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1345// CHECK_PTX87_SM100a: store <4 x i8> %[[R13]], ptr %r131346// CHECK_PTX87_SM103a: %[[R13:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e2m3x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1347// CHECK_PTX87_SM103a: store <4 x i8> %[[R13]], ptr %r131348 uint8x4 r13 = __nvvm_f32x4_to_e2m3x4_rs_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0); 1349 1350// CHECK_PTX87_SM100a: %[[R14:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e2m3x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1351// CHECK_PTX87_SM100a: store <4 x i8> %[[R14]], ptr %r141352// CHECK_PTX87_SM103a: %[[R14:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e2m3x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1353// CHECK_PTX87_SM103a: store <4 x i8> %[[R14]], ptr %r141354 uint8x4 r14 = __nvvm_f32x4_to_e2m3x4_rs_relu_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);1355 1356// CHECK_PTX87_SM100a: %[[R15:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e3m2x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1357// CHECK_PTX87_SM100a: store <4 x i8> %[[R15]], ptr %r151358// CHECK_PTX87_SM103a: %[[R15:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e3m2x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1359// CHECK_PTX87_SM103a: store <4 x i8> %[[R15]], ptr %r151360 uint8x4 r15 = __nvvm_f32x4_to_e3m2x4_rs_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);1361 1362// CHECK_PTX87_SM100a: %[[R16:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e3m2x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1363// CHECK_PTX87_SM100a: store <4 x i8> %[[R16]], ptr %r161364// CHECK_PTX87_SM103a: %[[R16:.*]] = call <4 x i8> @llvm.nvvm.f32x4.to.e3m2x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1365// CHECK_PTX87_SM103a: store <4 x i8> %[[R16]], ptr %r161366 uint8x4 r16 = __nvvm_f32x4_to_e3m2x4_rs_relu_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);1367 1368// CHECK_PTX87_SM100a: %[[R17:.*]] = call i16 @llvm.nvvm.f32x4.to.e2m1x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1369// CHECK_PTX87_SM100a: store i16 %[[R17]], ptr %r171370// CHECK_PTX87_SM103a: %[[R17:.*]] = call i16 @llvm.nvvm.f32x4.to.e2m1x4.rs.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1371// CHECK_PTX87_SM103a: store i16 %[[R17]], ptr %r171372 short r17 = __nvvm_f32x4_to_e2m1x4_rs_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);1373 1374// CHECK_PTX87_SM100a: %[[R18:.*]] = call i16 @llvm.nvvm.f32x4.to.e2m1x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1375// CHECK_PTX87_SM100a: store i16 %[[R18]], ptr %r181376// CHECK_PTX87_SM103a: %[[R18:.*]] = call i16 @llvm.nvvm.f32x4.to.e2m1x4.rs.relu.satfinite(<4 x float> splat (float 1.000000e+00), i32 0)1377// CHECK_PTX87_SM103a: store i16 %[[R18]], ptr %r181378 short r18 = __nvvm_f32x4_to_e2m1x4_rs_relu_satfinite({1.0f, 1.0f, 1.0f, 1.0f}, 0);1379#endif1380}1381 1382#define NAN32 0x7FBFFFFF1383#define NAN16 (__bf16)0x7FBF1384#define BF16 (__bf16)0.1f1385#define BF16_2 (__bf16)0.2f1386#define NANBF16 (__bf16)0xFFC11387#define BF16X2 {(__bf16)0.1f, (__bf16)0.1f}1388#define BF16X2_2 {(__bf16)0.2f, (__bf16)0.2f}1389#define NANBF16X2 {NANBF16, NANBF16}1390 1391// CHECK-LABEL: nvvm_abs_neg_bf16_bf16x2_sm801392__device__ void nvvm_abs_neg_bf16_bf16x2_sm80() {1393#if __CUDA_ARCH__ >= 8001394 1395 // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.fabs.bf16(bfloat 0xR3DCD)1396 __nvvm_abs_bf16(BF16);1397 // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.fabs.v2bf16(<2 x bfloat> splat (bfloat 0xR3DCD))1398 __nvvm_abs_bf16x2(BF16X2);1399 1400 // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.neg.bf16(bfloat 0xR3DCD)1401 __nvvm_neg_bf16(BF16);1402 // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.neg.bf16x2(<2 x bfloat> splat (bfloat 0xR3DCD))1403 __nvvm_neg_bf16x2(BF16X2);1404#endif1405 // CHECK: ret void1406}1407 1408// CHECK-LABEL: nvvm_min_max_sm801409__device__ void nvvm_min_max_sm80() {1410#if __CUDA_ARCH__ >= 8001411 1412 // CHECK_PTX70_SM80: call float @llvm.nvvm.fmin.nan.f1413 __nvvm_fmin_nan_f(0.1f, (float)NAN32);1414 // CHECK_PTX70_SM80: call float @llvm.nvvm.fmin.ftz.nan.f1415 __nvvm_fmin_ftz_nan_f(0.1f, (float)NAN32);1416 1417 // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.fmin.bf161418 __nvvm_fmin_bf16(BF16, BF16_2);1419 // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.fmin.ftz.bf161420 __nvvm_fmin_ftz_bf16(BF16, BF16_2);1421 // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.fmin.nan.bf161422 __nvvm_fmin_nan_bf16(BF16, NANBF16);1423 // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.fmin.ftz.nan.bf161424 __nvvm_fmin_ftz_nan_bf16(BF16, NANBF16);1425 // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.fmin.bf16x21426 __nvvm_fmin_bf16x2(BF16X2, BF16X2_2);1427 // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.fmin.ftz.bf16x21428 __nvvm_fmin_ftz_bf16x2(BF16X2, BF16X2_2);1429 // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.fmin.nan.bf16x21430 __nvvm_fmin_nan_bf16x2(BF16X2, NANBF16X2);1431 // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.fmin.ftz.nan.bf16x21432 __nvvm_fmin_ftz_nan_bf16x2(BF16X2, NANBF16X2);1433 // CHECK_PTX70_SM80: call float @llvm.nvvm.fmax.nan.f1434 __nvvm_fmax_nan_f(0.1f, 0.11f);1435 // CHECK_PTX70_SM80: call float @llvm.nvvm.fmax.ftz.nan.f1436 __nvvm_fmax_ftz_nan_f(0.1f, (float)NAN32);1437 1438 // CHECK_PTX70_SM80: call float @llvm.nvvm.fmax.nan.f1439 __nvvm_fmax_nan_f(0.1f, (float)NAN32);1440 // CHECK_PTX70_SM80: call float @llvm.nvvm.fmax.ftz.nan.f1441 __nvvm_fmax_ftz_nan_f(0.1f, (float)NAN32);1442 // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.fmax.bf161443 __nvvm_fmax_bf16(BF16, BF16_2);1444 // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.fmax.ftz.bf161445 __nvvm_fmax_ftz_bf16(BF16, BF16_2);1446 // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.fmax.nan.bf161447 __nvvm_fmax_nan_bf16(BF16, NANBF16);1448 // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.fmax.ftz.nan.bf161449 __nvvm_fmax_ftz_nan_bf16(BF16, NANBF16);1450 // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.fmax.bf16x21451 __nvvm_fmax_bf16x2(BF16X2, BF16X2_2);1452 // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.fmax.ftz.bf16x21453 __nvvm_fmax_ftz_bf16x2(BF16X2, BF16X2_2);1454 // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.fmax.nan.bf16x21455 __nvvm_fmax_nan_bf16x2(NANBF16X2, BF16X2);1456 // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.fmax.ftz.nan.bf16x21457 __nvvm_fmax_ftz_nan_bf16x2(NANBF16X2, BF16X2);1458 // CHECK_PTX70_SM80: call float @llvm.nvvm.fmax.nan.f1459 __nvvm_fmax_nan_f(0.1f, (float)NAN32);1460 // CHECK_PTX70_SM80: call float @llvm.nvvm.fmax.ftz.nan.f1461 __nvvm_fmax_ftz_nan_f(0.1f, (float)NAN32);1462 1463#endif1464 // CHECK: ret void1465}1466 1467// CHECK-LABEL: nvvm_fma_bf16_bf16x2_sm801468__device__ void nvvm_fma_bf16_bf16x2_sm80() {1469#if __CUDA_ARCH__ >= 8001470 // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.fma.rn.bf161471 __nvvm_fma_rn_bf16(BF16, BF16_2, BF16_2);1472 // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.fma.rn.relu.bf161473 __nvvm_fma_rn_relu_bf16(BF16, BF16_2, BF16_2);1474 // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.fma.rn.bf16x21475 __nvvm_fma_rn_bf16x2(BF16X2, BF16X2_2, BF16X2_2);1476 // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.fma.rn.relu.bf16x21477 __nvvm_fma_rn_relu_bf16x2(BF16X2, BF16X2_2, BF16X2_2);1478#endif1479 // CHECK: ret void1480}1481 1482// CHECK-LABEL: nvvm_min_max_sm861483__device__ void nvvm_min_max_sm86() {1484#if __CUDA_ARCH__ >= 8601485 1486 // CHECK_PTX72_SM86: call bfloat @llvm.nvvm.fmin.xorsign.abs.bf161487 __nvvm_fmin_xorsign_abs_bf16(BF16, BF16_2);1488 // CHECK_PTX72_SM86: call bfloat @llvm.nvvm.fmin.nan.xorsign.abs.bf161489 __nvvm_fmin_nan_xorsign_abs_bf16(BF16, NANBF16);1490 // CHECK_PTX72_SM86: call <2 x bfloat> @llvm.nvvm.fmin.xorsign.abs.bf16x21491 __nvvm_fmin_xorsign_abs_bf16x2(BF16X2, BF16X2_2);1492 // CHECK_PTX72_SM86: call <2 x bfloat> @llvm.nvvm.fmin.nan.xorsign.abs.bf16x21493 __nvvm_fmin_nan_xorsign_abs_bf16x2(BF16X2, NANBF16X2);1494 // CHECK_PTX72_SM86: call float @llvm.nvvm.fmin.xorsign.abs.f1495 __nvvm_fmin_xorsign_abs_f(-0.1f, 0.1f);1496 // CHECK_PTX72_SM86: call float @llvm.nvvm.fmin.ftz.xorsign.abs.f1497 __nvvm_fmin_ftz_xorsign_abs_f(-0.1f, 0.1f);1498 // CHECK_PTX72_SM86: call float @llvm.nvvm.fmin.nan.xorsign.abs.f1499 __nvvm_fmin_nan_xorsign_abs_f(-0.1f, (float)NAN32);1500 // CHECK_PTX72_SM86: call float @llvm.nvvm.fmin.ftz.nan.xorsign.abs.f1501 __nvvm_fmin_ftz_nan_xorsign_abs_f(-0.1f, (float)NAN32);1502 1503 // CHECK_PTX72_SM86: call bfloat @llvm.nvvm.fmax.xorsign.abs.bf161504 __nvvm_fmax_xorsign_abs_bf16(BF16, BF16_2);1505 // CHECK_PTX72_SM86: call bfloat @llvm.nvvm.fmax.nan.xorsign.abs.bf161506 __nvvm_fmax_nan_xorsign_abs_bf16(BF16, NANBF16);1507 // CHECK_PTX72_SM86: call <2 x bfloat> @llvm.nvvm.fmax.xorsign.abs.bf16x21508 __nvvm_fmax_xorsign_abs_bf16x2(BF16X2, BF16X2_2);1509 // CHECK_PTX72_SM86: call <2 x bfloat> @llvm.nvvm.fmax.nan.xorsign.abs.bf16x21510 __nvvm_fmax_nan_xorsign_abs_bf16x2(BF16X2, NANBF16X2);1511 // CHECK_PTX72_SM86: call float @llvm.nvvm.fmax.xorsign.abs.f1512 __nvvm_fmax_xorsign_abs_f(-0.1f, 0.1f);1513 // CHECK_PTX72_SM86: call float @llvm.nvvm.fmax.ftz.xorsign.abs.f1514 __nvvm_fmax_ftz_xorsign_abs_f(-0.1f, 0.1f);1515 // CHECK_PTX72_SM86: call float @llvm.nvvm.fmax.nan.xorsign.abs.f1516 __nvvm_fmax_nan_xorsign_abs_f(-0.1f, (float)NAN32);1517 // CHECK_PTX72_SM86: call float @llvm.nvvm.fmax.ftz.nan.xorsign.abs.f1518 __nvvm_fmax_ftz_nan_xorsign_abs_f(-0.1f, (float)NAN32);1519#endif1520 // CHECK: ret void1521}1522