986 lines · c
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs --version 52// RUN: %clang_cc1 -internal-isystem %S/Inputs/include \3// RUN: -internal-isystem %S/../../lib/Headers/ \4// RUN: -triple amdgcn-amd-amdhsa -emit-llvm %s -o - \5// RUN: | FileCheck %s --check-prefix=AMDGPU6//7// RUN: %clang_cc1 -internal-isystem %S/Inputs/include \8// RUN: -internal-isystem %S/../../lib/Headers/ \9// RUN: -target-feature +ptx62 \10// RUN: -triple nvptx64-nvidia-cuda -emit-llvm %s -o - \11// RUN: | FileCheck %s --check-prefix=NVPTX12 13#include <gpuintrin.h>14 15__gpu_kernel void foo() {16 __gpu_num_blocks_x();17 __gpu_num_blocks_y();18 __gpu_num_blocks_z();19 __gpu_num_blocks(0);20 __gpu_block_id_x();21 __gpu_block_id_y();22 __gpu_block_id_z();23 __gpu_block_id(0);24 __gpu_num_threads_x();25 __gpu_num_threads_y();26 __gpu_num_threads_z();27 __gpu_num_threads(0);28 __gpu_thread_id_x();29 __gpu_thread_id_y();30 __gpu_thread_id_z();31 __gpu_thread_id(0);32 __gpu_num_lanes();33 __gpu_lane_id();34 __gpu_lane_mask();35 __gpu_read_first_lane_u32(-1, -1);36 __gpu_read_first_lane_u64(-1, -1);37 __gpu_ballot(-1, 1);38 __gpu_sync_threads();39 __gpu_sync_lane(-1);40 __gpu_shuffle_idx_u32(-1, -1, -1, 0);41 __gpu_first_lane_id(-1);42 __gpu_is_first_in_lane(-1);43 __gpu_exit();44}45// AMDGPU-LABEL: define protected amdgpu_kernel void @foo(46// AMDGPU-SAME: ) #[[ATTR0:[0-9]+]] {47// AMDGPU-NEXT: [[ENTRY:.*:]]48// AMDGPU-NEXT: [[CALL:%.*]] = call i32 @__gpu_num_blocks_x() #[[ATTR7:[0-9]+]]49// AMDGPU-NEXT: [[CALL1:%.*]] = call i32 @__gpu_num_blocks_y() #[[ATTR7]]50// AMDGPU-NEXT: [[CALL2:%.*]] = call i32 @__gpu_num_blocks_z() #[[ATTR7]]51// AMDGPU-NEXT: [[CALL3:%.*]] = call i32 @__gpu_num_blocks(i32 noundef 0) #[[ATTR7]]52// AMDGPU-NEXT: [[CALL4:%.*]] = call i32 @__gpu_block_id_x() #[[ATTR7]]53// AMDGPU-NEXT: [[CALL5:%.*]] = call i32 @__gpu_block_id_y() #[[ATTR7]]54// AMDGPU-NEXT: [[CALL6:%.*]] = call i32 @__gpu_block_id_z() #[[ATTR7]]55// AMDGPU-NEXT: [[CALL7:%.*]] = call i32 @__gpu_block_id(i32 noundef 0) #[[ATTR7]]56// AMDGPU-NEXT: [[CALL8:%.*]] = call i32 @__gpu_num_threads_x() #[[ATTR7]]57// AMDGPU-NEXT: [[CALL9:%.*]] = call i32 @__gpu_num_threads_y() #[[ATTR7]]58// AMDGPU-NEXT: [[CALL10:%.*]] = call i32 @__gpu_num_threads_z() #[[ATTR7]]59// AMDGPU-NEXT: [[CALL11:%.*]] = call i32 @__gpu_num_threads(i32 noundef 0) #[[ATTR7]]60// AMDGPU-NEXT: [[CALL12:%.*]] = call i32 @__gpu_thread_id_x() #[[ATTR7]]61// AMDGPU-NEXT: [[CALL13:%.*]] = call i32 @__gpu_thread_id_y() #[[ATTR7]]62// AMDGPU-NEXT: [[CALL14:%.*]] = call i32 @__gpu_thread_id_z() #[[ATTR7]]63// AMDGPU-NEXT: [[CALL15:%.*]] = call i32 @__gpu_thread_id(i32 noundef 0) #[[ATTR7]]64// AMDGPU-NEXT: [[CALL16:%.*]] = call i32 @__gpu_num_lanes() #[[ATTR7]]65// AMDGPU-NEXT: [[CALL17:%.*]] = call i32 @__gpu_lane_id() #[[ATTR7]]66// AMDGPU-NEXT: [[CALL18:%.*]] = call i64 @__gpu_lane_mask() #[[ATTR7]]67// AMDGPU-NEXT: [[CALL19:%.*]] = call i32 @__gpu_read_first_lane_u32(i64 noundef -1, i32 noundef -1) #[[ATTR7]]68// AMDGPU-NEXT: [[CALL20:%.*]] = call i64 @__gpu_read_first_lane_u64(i64 noundef -1, i64 noundef -1) #[[ATTR7]]69// AMDGPU-NEXT: [[CALL21:%.*]] = call i64 @__gpu_ballot(i64 noundef -1, i1 noundef zeroext true) #[[ATTR7]]70// AMDGPU-NEXT: call void @__gpu_sync_threads() #[[ATTR7]]71// AMDGPU-NEXT: call void @__gpu_sync_lane(i64 noundef -1) #[[ATTR7]]72// AMDGPU-NEXT: [[CALL22:%.*]] = call i32 @__gpu_shuffle_idx_u32(i64 noundef -1, i32 noundef -1, i32 noundef -1, i32 noundef 0) #[[ATTR7]]73// AMDGPU-NEXT: [[CALL23:%.*]] = call i64 @__gpu_first_lane_id(i64 noundef -1) #[[ATTR7]]74// AMDGPU-NEXT: [[CALL24:%.*]] = call zeroext i1 @__gpu_is_first_in_lane(i64 noundef -1) #[[ATTR7]]75// AMDGPU-NEXT: call void @__gpu_exit() #[[ATTR8:[0-9]+]]76// AMDGPU-NEXT: unreachable77//78//79// AMDGPU-LABEL: define internal i32 @__gpu_num_blocks_x(80// AMDGPU-SAME: ) #[[ATTR0]] {81// AMDGPU-NEXT: [[ENTRY:.*:]]82// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)83// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr84// AMDGPU-NEXT: [[TMP0:%.*]] = call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()85// AMDGPU-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP0]], i32 1286// AMDGPU-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 4, !range [[RNG3:![0-9]+]], !invariant.load [[META4:![0-9]+]]87// AMDGPU-NEXT: [[TMP3:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()88// AMDGPU-NEXT: [[TMP4:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP3]], i32 1289// AMDGPU-NEXT: [[TMP5:%.*]] = load i16, ptr addrspace(4) [[TMP4]], align 2, !range [[RNG5:![0-9]+]], !invariant.load [[META4]], !noundef [[META4]]90// AMDGPU-NEXT: [[CONV:%.*]] = zext i16 [[TMP5]] to i3291// AMDGPU-NEXT: [[DIV:%.*]] = udiv i32 [[TMP2]], [[CONV]]92// AMDGPU-NEXT: ret i32 [[DIV]]93//94//95// AMDGPU-LABEL: define internal i32 @__gpu_num_blocks_y(96// AMDGPU-SAME: ) #[[ATTR0]] {97// AMDGPU-NEXT: [[ENTRY:.*:]]98// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)99// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr100// AMDGPU-NEXT: [[TMP0:%.*]] = call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()101// AMDGPU-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP0]], i32 16102// AMDGPU-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 4, !range [[RNG3]], !invariant.load [[META4]]103// AMDGPU-NEXT: [[TMP3:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()104// AMDGPU-NEXT: [[TMP4:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP3]], i32 14105// AMDGPU-NEXT: [[TMP5:%.*]] = load i16, ptr addrspace(4) [[TMP4]], align 2, !range [[RNG5]], !invariant.load [[META4]], !noundef [[META4]]106// AMDGPU-NEXT: [[CONV:%.*]] = zext i16 [[TMP5]] to i32107// AMDGPU-NEXT: [[DIV:%.*]] = udiv i32 [[TMP2]], [[CONV]]108// AMDGPU-NEXT: ret i32 [[DIV]]109//110//111// AMDGPU-LABEL: define internal i32 @__gpu_num_blocks_z(112// AMDGPU-SAME: ) #[[ATTR0]] {113// AMDGPU-NEXT: [[ENTRY:.*:]]114// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)115// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr116// AMDGPU-NEXT: [[TMP0:%.*]] = call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()117// AMDGPU-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP0]], i32 20118// AMDGPU-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 4, !range [[RNG3]], !invariant.load [[META4]]119// AMDGPU-NEXT: [[TMP3:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()120// AMDGPU-NEXT: [[TMP4:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP3]], i32 16121// AMDGPU-NEXT: [[TMP5:%.*]] = load i16, ptr addrspace(4) [[TMP4]], align 2, !range [[RNG5]], !invariant.load [[META4]], !noundef [[META4]]122// AMDGPU-NEXT: [[CONV:%.*]] = zext i16 [[TMP5]] to i32123// AMDGPU-NEXT: [[DIV:%.*]] = udiv i32 [[TMP2]], [[CONV]]124// AMDGPU-NEXT: ret i32 [[DIV]]125//126//127// AMDGPU-LABEL: define internal i32 @__gpu_num_blocks(128// AMDGPU-SAME: i32 noundef [[__DIM:%.*]]) #[[ATTR0]] {129// AMDGPU-NEXT: [[ENTRY:.*:]]130// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)131// AMDGPU-NEXT: [[__DIM_ADDR:%.*]] = alloca i32, align 4, addrspace(5)132// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr133// AMDGPU-NEXT: [[__DIM_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__DIM_ADDR]] to ptr134// AMDGPU-NEXT: store i32 [[__DIM]], ptr [[__DIM_ADDR_ASCAST]], align 4135// AMDGPU-NEXT: [[TMP0:%.*]] = load i32, ptr [[__DIM_ADDR_ASCAST]], align 4136// AMDGPU-NEXT: switch i32 [[TMP0]], label %[[SW_DEFAULT:.*]] [137// AMDGPU-NEXT: i32 0, label %[[SW_BB:.*]]138// AMDGPU-NEXT: i32 1, label %[[SW_BB1:.*]]139// AMDGPU-NEXT: i32 2, label %[[SW_BB3:.*]]140// AMDGPU-NEXT: ]141// AMDGPU: [[SW_BB]]:142// AMDGPU-NEXT: [[CALL:%.*]] = call i32 @__gpu_num_blocks_x() #[[ATTR7]]143// AMDGPU-NEXT: store i32 [[CALL]], ptr [[RETVAL_ASCAST]], align 4144// AMDGPU-NEXT: br label %[[RETURN:.*]]145// AMDGPU: [[SW_BB1]]:146// AMDGPU-NEXT: [[CALL2:%.*]] = call i32 @__gpu_num_blocks_y() #[[ATTR7]]147// AMDGPU-NEXT: store i32 [[CALL2]], ptr [[RETVAL_ASCAST]], align 4148// AMDGPU-NEXT: br label %[[RETURN]]149// AMDGPU: [[SW_BB3]]:150// AMDGPU-NEXT: [[CALL4:%.*]] = call i32 @__gpu_num_blocks_z() #[[ATTR7]]151// AMDGPU-NEXT: store i32 [[CALL4]], ptr [[RETVAL_ASCAST]], align 4152// AMDGPU-NEXT: br label %[[RETURN]]153// AMDGPU: [[SW_DEFAULT]]:154// AMDGPU-NEXT: unreachable155// AMDGPU: [[RETURN]]:156// AMDGPU-NEXT: [[TMP1:%.*]] = load i32, ptr [[RETVAL_ASCAST]], align 4157// AMDGPU-NEXT: ret i32 [[TMP1]]158//159//160// AMDGPU-LABEL: define internal i32 @__gpu_block_id_x(161// AMDGPU-SAME: ) #[[ATTR0]] {162// AMDGPU-NEXT: [[ENTRY:.*:]]163// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)164// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr165// AMDGPU-NEXT: [[TMP0:%.*]] = call i32 @llvm.amdgcn.workgroup.id.x()166// AMDGPU-NEXT: ret i32 [[TMP0]]167//168//169// AMDGPU-LABEL: define internal i32 @__gpu_block_id_y(170// AMDGPU-SAME: ) #[[ATTR0]] {171// AMDGPU-NEXT: [[ENTRY:.*:]]172// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)173// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr174// AMDGPU-NEXT: [[TMP0:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y()175// AMDGPU-NEXT: ret i32 [[TMP0]]176//177//178// AMDGPU-LABEL: define internal i32 @__gpu_block_id_z(179// AMDGPU-SAME: ) #[[ATTR0]] {180// AMDGPU-NEXT: [[ENTRY:.*:]]181// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)182// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr183// AMDGPU-NEXT: [[TMP0:%.*]] = call i32 @llvm.amdgcn.workgroup.id.z()184// AMDGPU-NEXT: ret i32 [[TMP0]]185//186//187// AMDGPU-LABEL: define internal i32 @__gpu_block_id(188// AMDGPU-SAME: i32 noundef [[__DIM:%.*]]) #[[ATTR0]] {189// AMDGPU-NEXT: [[ENTRY:.*:]]190// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)191// AMDGPU-NEXT: [[__DIM_ADDR:%.*]] = alloca i32, align 4, addrspace(5)192// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr193// AMDGPU-NEXT: [[__DIM_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__DIM_ADDR]] to ptr194// AMDGPU-NEXT: store i32 [[__DIM]], ptr [[__DIM_ADDR_ASCAST]], align 4195// AMDGPU-NEXT: [[TMP0:%.*]] = load i32, ptr [[__DIM_ADDR_ASCAST]], align 4196// AMDGPU-NEXT: switch i32 [[TMP0]], label %[[SW_DEFAULT:.*]] [197// AMDGPU-NEXT: i32 0, label %[[SW_BB:.*]]198// AMDGPU-NEXT: i32 1, label %[[SW_BB1:.*]]199// AMDGPU-NEXT: i32 2, label %[[SW_BB3:.*]]200// AMDGPU-NEXT: ]201// AMDGPU: [[SW_BB]]:202// AMDGPU-NEXT: [[CALL:%.*]] = call i32 @__gpu_block_id_x() #[[ATTR7]]203// AMDGPU-NEXT: store i32 [[CALL]], ptr [[RETVAL_ASCAST]], align 4204// AMDGPU-NEXT: br label %[[RETURN:.*]]205// AMDGPU: [[SW_BB1]]:206// AMDGPU-NEXT: [[CALL2:%.*]] = call i32 @__gpu_block_id_y() #[[ATTR7]]207// AMDGPU-NEXT: store i32 [[CALL2]], ptr [[RETVAL_ASCAST]], align 4208// AMDGPU-NEXT: br label %[[RETURN]]209// AMDGPU: [[SW_BB3]]:210// AMDGPU-NEXT: [[CALL4:%.*]] = call i32 @__gpu_block_id_z() #[[ATTR7]]211// AMDGPU-NEXT: store i32 [[CALL4]], ptr [[RETVAL_ASCAST]], align 4212// AMDGPU-NEXT: br label %[[RETURN]]213// AMDGPU: [[SW_DEFAULT]]:214// AMDGPU-NEXT: unreachable215// AMDGPU: [[RETURN]]:216// AMDGPU-NEXT: [[TMP1:%.*]] = load i32, ptr [[RETVAL_ASCAST]], align 4217// AMDGPU-NEXT: ret i32 [[TMP1]]218//219//220// AMDGPU-LABEL: define internal i32 @__gpu_num_threads_x(221// AMDGPU-SAME: ) #[[ATTR0]] {222// AMDGPU-NEXT: [[ENTRY:.*:]]223// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)224// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr225// AMDGPU-NEXT: [[TMP0:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()226// AMDGPU-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP0]], i32 12227// AMDGPU-NEXT: [[TMP2:%.*]] = load i16, ptr addrspace(4) [[TMP1]], align 2, !range [[RNG5]], !invariant.load [[META4]], !noundef [[META4]]228// AMDGPU-NEXT: [[CONV:%.*]] = zext i16 [[TMP2]] to i32229// AMDGPU-NEXT: ret i32 [[CONV]]230//231//232// AMDGPU-LABEL: define internal i32 @__gpu_num_threads_y(233// AMDGPU-SAME: ) #[[ATTR0]] {234// AMDGPU-NEXT: [[ENTRY:.*:]]235// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)236// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr237// AMDGPU-NEXT: [[TMP0:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()238// AMDGPU-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP0]], i32 14239// AMDGPU-NEXT: [[TMP2:%.*]] = load i16, ptr addrspace(4) [[TMP1]], align 2, !range [[RNG5]], !invariant.load [[META4]], !noundef [[META4]]240// AMDGPU-NEXT: [[CONV:%.*]] = zext i16 [[TMP2]] to i32241// AMDGPU-NEXT: ret i32 [[CONV]]242//243//244// AMDGPU-LABEL: define internal i32 @__gpu_num_threads_z(245// AMDGPU-SAME: ) #[[ATTR0]] {246// AMDGPU-NEXT: [[ENTRY:.*:]]247// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)248// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr249// AMDGPU-NEXT: [[TMP0:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()250// AMDGPU-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP0]], i32 16251// AMDGPU-NEXT: [[TMP2:%.*]] = load i16, ptr addrspace(4) [[TMP1]], align 2, !range [[RNG5]], !invariant.load [[META4]], !noundef [[META4]]252// AMDGPU-NEXT: [[CONV:%.*]] = zext i16 [[TMP2]] to i32253// AMDGPU-NEXT: ret i32 [[CONV]]254//255//256// AMDGPU-LABEL: define internal i32 @__gpu_num_threads(257// AMDGPU-SAME: i32 noundef [[__DIM:%.*]]) #[[ATTR0]] {258// AMDGPU-NEXT: [[ENTRY:.*:]]259// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)260// AMDGPU-NEXT: [[__DIM_ADDR:%.*]] = alloca i32, align 4, addrspace(5)261// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr262// AMDGPU-NEXT: [[__DIM_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__DIM_ADDR]] to ptr263// AMDGPU-NEXT: store i32 [[__DIM]], ptr [[__DIM_ADDR_ASCAST]], align 4264// AMDGPU-NEXT: [[TMP0:%.*]] = load i32, ptr [[__DIM_ADDR_ASCAST]], align 4265// AMDGPU-NEXT: switch i32 [[TMP0]], label %[[SW_DEFAULT:.*]] [266// AMDGPU-NEXT: i32 0, label %[[SW_BB:.*]]267// AMDGPU-NEXT: i32 1, label %[[SW_BB1:.*]]268// AMDGPU-NEXT: i32 2, label %[[SW_BB3:.*]]269// AMDGPU-NEXT: ]270// AMDGPU: [[SW_BB]]:271// AMDGPU-NEXT: [[CALL:%.*]] = call i32 @__gpu_num_threads_x() #[[ATTR7]]272// AMDGPU-NEXT: store i32 [[CALL]], ptr [[RETVAL_ASCAST]], align 4273// AMDGPU-NEXT: br label %[[RETURN:.*]]274// AMDGPU: [[SW_BB1]]:275// AMDGPU-NEXT: [[CALL2:%.*]] = call i32 @__gpu_num_threads_y() #[[ATTR7]]276// AMDGPU-NEXT: store i32 [[CALL2]], ptr [[RETVAL_ASCAST]], align 4277// AMDGPU-NEXT: br label %[[RETURN]]278// AMDGPU: [[SW_BB3]]:279// AMDGPU-NEXT: [[CALL4:%.*]] = call i32 @__gpu_num_threads_z() #[[ATTR7]]280// AMDGPU-NEXT: store i32 [[CALL4]], ptr [[RETVAL_ASCAST]], align 4281// AMDGPU-NEXT: br label %[[RETURN]]282// AMDGPU: [[SW_DEFAULT]]:283// AMDGPU-NEXT: unreachable284// AMDGPU: [[RETURN]]:285// AMDGPU-NEXT: [[TMP1:%.*]] = load i32, ptr [[RETVAL_ASCAST]], align 4286// AMDGPU-NEXT: ret i32 [[TMP1]]287//288//289// AMDGPU-LABEL: define internal i32 @__gpu_thread_id_x(290// AMDGPU-SAME: ) #[[ATTR0]] {291// AMDGPU-NEXT: [[ENTRY:.*:]]292// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)293// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr294// AMDGPU-NEXT: [[TMP0:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()295// AMDGPU-NEXT: ret i32 [[TMP0]]296//297//298// AMDGPU-LABEL: define internal i32 @__gpu_thread_id_y(299// AMDGPU-SAME: ) #[[ATTR0]] {300// AMDGPU-NEXT: [[ENTRY:.*:]]301// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)302// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr303// AMDGPU-NEXT: [[TMP0:%.*]] = call i32 @llvm.amdgcn.workitem.id.y()304// AMDGPU-NEXT: ret i32 [[TMP0]]305//306//307// AMDGPU-LABEL: define internal i32 @__gpu_thread_id_z(308// AMDGPU-SAME: ) #[[ATTR0]] {309// AMDGPU-NEXT: [[ENTRY:.*:]]310// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)311// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr312// AMDGPU-NEXT: [[TMP0:%.*]] = call i32 @llvm.amdgcn.workitem.id.z()313// AMDGPU-NEXT: ret i32 [[TMP0]]314//315//316// AMDGPU-LABEL: define internal i32 @__gpu_thread_id(317// AMDGPU-SAME: i32 noundef [[__DIM:%.*]]) #[[ATTR0]] {318// AMDGPU-NEXT: [[ENTRY:.*:]]319// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)320// AMDGPU-NEXT: [[__DIM_ADDR:%.*]] = alloca i32, align 4, addrspace(5)321// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr322// AMDGPU-NEXT: [[__DIM_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__DIM_ADDR]] to ptr323// AMDGPU-NEXT: store i32 [[__DIM]], ptr [[__DIM_ADDR_ASCAST]], align 4324// AMDGPU-NEXT: [[TMP0:%.*]] = load i32, ptr [[__DIM_ADDR_ASCAST]], align 4325// AMDGPU-NEXT: switch i32 [[TMP0]], label %[[SW_DEFAULT:.*]] [326// AMDGPU-NEXT: i32 0, label %[[SW_BB:.*]]327// AMDGPU-NEXT: i32 1, label %[[SW_BB1:.*]]328// AMDGPU-NEXT: i32 2, label %[[SW_BB3:.*]]329// AMDGPU-NEXT: ]330// AMDGPU: [[SW_BB]]:331// AMDGPU-NEXT: [[CALL:%.*]] = call i32 @__gpu_thread_id_x() #[[ATTR7]]332// AMDGPU-NEXT: store i32 [[CALL]], ptr [[RETVAL_ASCAST]], align 4333// AMDGPU-NEXT: br label %[[RETURN:.*]]334// AMDGPU: [[SW_BB1]]:335// AMDGPU-NEXT: [[CALL2:%.*]] = call i32 @__gpu_thread_id_y() #[[ATTR7]]336// AMDGPU-NEXT: store i32 [[CALL2]], ptr [[RETVAL_ASCAST]], align 4337// AMDGPU-NEXT: br label %[[RETURN]]338// AMDGPU: [[SW_BB3]]:339// AMDGPU-NEXT: [[CALL4:%.*]] = call i32 @__gpu_thread_id_z() #[[ATTR7]]340// AMDGPU-NEXT: store i32 [[CALL4]], ptr [[RETVAL_ASCAST]], align 4341// AMDGPU-NEXT: br label %[[RETURN]]342// AMDGPU: [[SW_DEFAULT]]:343// AMDGPU-NEXT: unreachable344// AMDGPU: [[RETURN]]:345// AMDGPU-NEXT: [[TMP1:%.*]] = load i32, ptr [[RETVAL_ASCAST]], align 4346// AMDGPU-NEXT: ret i32 [[TMP1]]347//348//349// AMDGPU-LABEL: define internal i32 @__gpu_num_lanes(350// AMDGPU-SAME: ) #[[ATTR0]] {351// AMDGPU-NEXT: [[ENTRY:.*:]]352// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)353// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr354// AMDGPU-NEXT: [[TMP0:%.*]] = call i32 @llvm.amdgcn.wavefrontsize()355// AMDGPU-NEXT: ret i32 [[TMP0]]356//357//358// AMDGPU-LABEL: define internal i32 @__gpu_lane_id(359// AMDGPU-SAME: ) #[[ATTR0]] {360// AMDGPU-NEXT: [[ENTRY:.*:]]361// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)362// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr363// AMDGPU-NEXT: [[TMP0:%.*]] = call i32 @llvm.amdgcn.mbcnt.lo(i32 -1, i32 0)364// AMDGPU-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.mbcnt.hi(i32 -1, i32 [[TMP0]])365// AMDGPU-NEXT: ret i32 [[TMP1]]366//367//368// AMDGPU-LABEL: define internal i64 @__gpu_lane_mask(369// AMDGPU-SAME: ) #[[ATTR0]] {370// AMDGPU-NEXT: [[ENTRY:.*:]]371// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i64, align 8, addrspace(5)372// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr373// AMDGPU-NEXT: [[TMP0:%.*]] = call i64 @llvm.amdgcn.ballot.i64(i1 true)374// AMDGPU-NEXT: ret i64 [[TMP0]]375//376//377// AMDGPU-LABEL: define internal i32 @__gpu_read_first_lane_u32(378// AMDGPU-SAME: i64 noundef [[__LANE_MASK:%.*]], i32 noundef [[__X:%.*]]) #[[ATTR0]] {379// AMDGPU-NEXT: [[ENTRY:.*:]]380// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)381// AMDGPU-NEXT: [[__LANE_MASK_ADDR:%.*]] = alloca i64, align 8, addrspace(5)382// AMDGPU-NEXT: [[__X_ADDR:%.*]] = alloca i32, align 4, addrspace(5)383// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr384// AMDGPU-NEXT: [[__LANE_MASK_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__LANE_MASK_ADDR]] to ptr385// AMDGPU-NEXT: [[__X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR]] to ptr386// AMDGPU-NEXT: store i64 [[__LANE_MASK]], ptr [[__LANE_MASK_ADDR_ASCAST]], align 8387// AMDGPU-NEXT: store i32 [[__X]], ptr [[__X_ADDR_ASCAST]], align 4388// AMDGPU-NEXT: [[TMP0:%.*]] = load i32, ptr [[__X_ADDR_ASCAST]], align 4389// AMDGPU-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.readfirstlane.i32(i32 [[TMP0]])390// AMDGPU-NEXT: ret i32 [[TMP1]]391//392//393// AMDGPU-LABEL: define internal i64 @__gpu_read_first_lane_u64(394// AMDGPU-SAME: i64 noundef [[__LANE_MASK:%.*]], i64 noundef [[__X:%.*]]) #[[ATTR0]] {395// AMDGPU-NEXT: [[ENTRY:.*:]]396// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i64, align 8, addrspace(5)397// AMDGPU-NEXT: [[__LANE_MASK_ADDR:%.*]] = alloca i64, align 8, addrspace(5)398// AMDGPU-NEXT: [[__X_ADDR:%.*]] = alloca i64, align 8, addrspace(5)399// AMDGPU-NEXT: [[__HI:%.*]] = alloca i32, align 4, addrspace(5)400// AMDGPU-NEXT: [[__LO:%.*]] = alloca i32, align 4, addrspace(5)401// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr402// AMDGPU-NEXT: [[__LANE_MASK_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__LANE_MASK_ADDR]] to ptr403// AMDGPU-NEXT: [[__X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR]] to ptr404// AMDGPU-NEXT: [[__HI_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__HI]] to ptr405// AMDGPU-NEXT: [[__LO_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__LO]] to ptr406// AMDGPU-NEXT: store i64 [[__LANE_MASK]], ptr [[__LANE_MASK_ADDR_ASCAST]], align 8407// AMDGPU-NEXT: store i64 [[__X]], ptr [[__X_ADDR_ASCAST]], align 8408// AMDGPU-NEXT: [[TMP0:%.*]] = load i64, ptr [[__X_ADDR_ASCAST]], align 8409// AMDGPU-NEXT: [[SHR:%.*]] = lshr i64 [[TMP0]], 32410// AMDGPU-NEXT: [[CONV:%.*]] = trunc i64 [[SHR]] to i32411// AMDGPU-NEXT: store i32 [[CONV]], ptr [[__HI_ASCAST]], align 4412// AMDGPU-NEXT: [[TMP1:%.*]] = load i64, ptr [[__X_ADDR_ASCAST]], align 8413// AMDGPU-NEXT: [[AND:%.*]] = and i64 [[TMP1]], 4294967295414// AMDGPU-NEXT: [[CONV1:%.*]] = trunc i64 [[AND]] to i32415// AMDGPU-NEXT: store i32 [[CONV1]], ptr [[__LO_ASCAST]], align 4416// AMDGPU-NEXT: [[TMP2:%.*]] = load i64, ptr [[__LANE_MASK_ADDR_ASCAST]], align 8417// AMDGPU-NEXT: [[TMP3:%.*]] = load i32, ptr [[__HI_ASCAST]], align 4418// AMDGPU-NEXT: [[CALL:%.*]] = call i32 @__gpu_read_first_lane_u32(i64 noundef [[TMP2]], i32 noundef [[TMP3]]) #[[ATTR7]]419// AMDGPU-NEXT: [[CONV2:%.*]] = zext i32 [[CALL]] to i64420// AMDGPU-NEXT: [[SHL:%.*]] = shl i64 [[CONV2]], 32421// AMDGPU-NEXT: [[TMP4:%.*]] = load i64, ptr [[__LANE_MASK_ADDR_ASCAST]], align 8422// AMDGPU-NEXT: [[TMP5:%.*]] = load i32, ptr [[__LO_ASCAST]], align 4423// AMDGPU-NEXT: [[CALL3:%.*]] = call i32 @__gpu_read_first_lane_u32(i64 noundef [[TMP4]], i32 noundef [[TMP5]]) #[[ATTR7]]424// AMDGPU-NEXT: [[CONV4:%.*]] = zext i32 [[CALL3]] to i64425// AMDGPU-NEXT: [[AND5:%.*]] = and i64 [[CONV4]], 4294967295426// AMDGPU-NEXT: [[OR:%.*]] = or i64 [[SHL]], [[AND5]]427// AMDGPU-NEXT: ret i64 [[OR]]428//429//430// AMDGPU-LABEL: define internal i64 @__gpu_ballot(431// AMDGPU-SAME: i64 noundef [[__LANE_MASK:%.*]], i1 noundef zeroext [[__X:%.*]]) #[[ATTR0]] {432// AMDGPU-NEXT: [[ENTRY:.*:]]433// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i64, align 8, addrspace(5)434// AMDGPU-NEXT: [[__LANE_MASK_ADDR:%.*]] = alloca i64, align 8, addrspace(5)435// AMDGPU-NEXT: [[__X_ADDR:%.*]] = alloca i8, align 1, addrspace(5)436// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr437// AMDGPU-NEXT: [[__LANE_MASK_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__LANE_MASK_ADDR]] to ptr438// AMDGPU-NEXT: [[__X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR]] to ptr439// AMDGPU-NEXT: store i64 [[__LANE_MASK]], ptr [[__LANE_MASK_ADDR_ASCAST]], align 8440// AMDGPU-NEXT: [[STOREDV:%.*]] = zext i1 [[__X]] to i8441// AMDGPU-NEXT: store i8 [[STOREDV]], ptr [[__X_ADDR_ASCAST]], align 1442// AMDGPU-NEXT: [[TMP0:%.*]] = load i64, ptr [[__LANE_MASK_ADDR_ASCAST]], align 8443// AMDGPU-NEXT: [[TMP1:%.*]] = load i8, ptr [[__X_ADDR_ASCAST]], align 1444// AMDGPU-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP1]] to i1445// AMDGPU-NEXT: [[TMP2:%.*]] = call i64 @llvm.amdgcn.ballot.i64(i1 [[LOADEDV]])446// AMDGPU-NEXT: [[AND:%.*]] = and i64 [[TMP0]], [[TMP2]]447// AMDGPU-NEXT: ret i64 [[AND]]448//449//450// AMDGPU-LABEL: define internal void @__gpu_sync_threads(451// AMDGPU-SAME: ) #[[ATTR0]] {452// AMDGPU-NEXT: [[ENTRY:.*:]]453// AMDGPU-NEXT: call void @llvm.amdgcn.s.barrier()454// AMDGPU-NEXT: fence syncscope("workgroup") seq_cst455// AMDGPU-NEXT: ret void456//457//458// AMDGPU-LABEL: define internal void @__gpu_sync_lane(459// AMDGPU-SAME: i64 noundef [[__LANE_MASK:%.*]]) #[[ATTR0]] {460// AMDGPU-NEXT: [[ENTRY:.*:]]461// AMDGPU-NEXT: [[__LANE_MASK_ADDR:%.*]] = alloca i64, align 8, addrspace(5)462// AMDGPU-NEXT: [[__LANE_MASK_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__LANE_MASK_ADDR]] to ptr463// AMDGPU-NEXT: store i64 [[__LANE_MASK]], ptr [[__LANE_MASK_ADDR_ASCAST]], align 8464// AMDGPU-NEXT: call void @llvm.amdgcn.wave.barrier()465// AMDGPU-NEXT: ret void466//467//468// AMDGPU-LABEL: define internal i32 @__gpu_shuffle_idx_u32(469// AMDGPU-SAME: i64 noundef [[__LANE_MASK:%.*]], i32 noundef [[__IDX:%.*]], i32 noundef [[__X:%.*]], i32 noundef [[__WIDTH:%.*]]) #[[ATTR0]] {470// AMDGPU-NEXT: [[ENTRY:.*:]]471// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)472// AMDGPU-NEXT: [[__LANE_MASK_ADDR:%.*]] = alloca i64, align 8, addrspace(5)473// AMDGPU-NEXT: [[__IDX_ADDR:%.*]] = alloca i32, align 4, addrspace(5)474// AMDGPU-NEXT: [[__X_ADDR:%.*]] = alloca i32, align 4, addrspace(5)475// AMDGPU-NEXT: [[__WIDTH_ADDR:%.*]] = alloca i32, align 4, addrspace(5)476// AMDGPU-NEXT: [[__LANE:%.*]] = alloca i32, align 4, addrspace(5)477// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr478// AMDGPU-NEXT: [[__LANE_MASK_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__LANE_MASK_ADDR]] to ptr479// AMDGPU-NEXT: [[__IDX_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__IDX_ADDR]] to ptr480// AMDGPU-NEXT: [[__X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR]] to ptr481// AMDGPU-NEXT: [[__WIDTH_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__WIDTH_ADDR]] to ptr482// AMDGPU-NEXT: [[__LANE_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__LANE]] to ptr483// AMDGPU-NEXT: store i64 [[__LANE_MASK]], ptr [[__LANE_MASK_ADDR_ASCAST]], align 8484// AMDGPU-NEXT: store i32 [[__IDX]], ptr [[__IDX_ADDR_ASCAST]], align 4485// AMDGPU-NEXT: store i32 [[__X]], ptr [[__X_ADDR_ASCAST]], align 4486// AMDGPU-NEXT: store i32 [[__WIDTH]], ptr [[__WIDTH_ADDR_ASCAST]], align 4487// AMDGPU-NEXT: [[TMP0:%.*]] = load i32, ptr [[__IDX_ADDR_ASCAST]], align 4488// AMDGPU-NEXT: [[CALL:%.*]] = call i32 @__gpu_lane_id() #[[ATTR7]]489// AMDGPU-NEXT: [[TMP1:%.*]] = load i32, ptr [[__WIDTH_ADDR_ASCAST]], align 4490// AMDGPU-NEXT: [[SUB:%.*]] = sub i32 [[TMP1]], 1491// AMDGPU-NEXT: [[NOT:%.*]] = xor i32 [[SUB]], -1492// AMDGPU-NEXT: [[AND:%.*]] = and i32 [[CALL]], [[NOT]]493// AMDGPU-NEXT: [[ADD:%.*]] = add i32 [[TMP0]], [[AND]]494// AMDGPU-NEXT: store i32 [[ADD]], ptr [[__LANE_ASCAST]], align 4495// AMDGPU-NEXT: [[TMP2:%.*]] = load i32, ptr [[__LANE_ASCAST]], align 4496// AMDGPU-NEXT: [[SHL:%.*]] = shl i32 [[TMP2]], 2497// AMDGPU-NEXT: [[TMP3:%.*]] = load i32, ptr [[__X_ADDR_ASCAST]], align 4498// AMDGPU-NEXT: [[TMP4:%.*]] = call i32 @llvm.amdgcn.ds.bpermute(i32 [[SHL]], i32 [[TMP3]])499// AMDGPU-NEXT: ret i32 [[TMP4]]500//501//502// AMDGPU-LABEL: define internal i64 @__gpu_first_lane_id(503// AMDGPU-SAME: i64 noundef [[__LANE_MASK:%.*]]) #[[ATTR0]] {504// AMDGPU-NEXT: [[ENTRY:.*:]]505// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i64, align 8, addrspace(5)506// AMDGPU-NEXT: [[__LANE_MASK_ADDR:%.*]] = alloca i64, align 8, addrspace(5)507// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr508// AMDGPU-NEXT: [[__LANE_MASK_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__LANE_MASK_ADDR]] to ptr509// AMDGPU-NEXT: store i64 [[__LANE_MASK]], ptr [[__LANE_MASK_ADDR_ASCAST]], align 8510// AMDGPU-NEXT: [[TMP0:%.*]] = load i64, ptr [[__LANE_MASK_ADDR_ASCAST]], align 8511// AMDGPU-NEXT: [[TMP1:%.*]] = call i64 @llvm.cttz.i64(i64 [[TMP0]], i1 true)512// AMDGPU-NEXT: [[TMP2:%.*]] = add i64 [[TMP1]], 1513// AMDGPU-NEXT: [[ISZERO:%.*]] = icmp eq i64 [[TMP0]], 0514// AMDGPU-NEXT: [[FFS:%.*]] = select i1 [[ISZERO]], i64 0, i64 [[TMP2]]515// AMDGPU-NEXT: [[CAST:%.*]] = trunc i64 [[FFS]] to i32516// AMDGPU-NEXT: [[SUB:%.*]] = sub nsw i32 [[CAST]], 1517// AMDGPU-NEXT: [[CONV:%.*]] = sext i32 [[SUB]] to i64518// AMDGPU-NEXT: ret i64 [[CONV]]519//520//521// AMDGPU-LABEL: define internal zeroext i1 @__gpu_is_first_in_lane(522// AMDGPU-SAME: i64 noundef [[__LANE_MASK:%.*]]) #[[ATTR0]] {523// AMDGPU-NEXT: [[ENTRY:.*:]]524// AMDGPU-NEXT: [[RETVAL:%.*]] = alloca i1, align 1, addrspace(5)525// AMDGPU-NEXT: [[__LANE_MASK_ADDR:%.*]] = alloca i64, align 8, addrspace(5)526// AMDGPU-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr527// AMDGPU-NEXT: [[__LANE_MASK_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[__LANE_MASK_ADDR]] to ptr528// AMDGPU-NEXT: store i64 [[__LANE_MASK]], ptr [[__LANE_MASK_ADDR_ASCAST]], align 8529// AMDGPU-NEXT: [[CALL:%.*]] = call i32 @__gpu_lane_id() #[[ATTR7]]530// AMDGPU-NEXT: [[CONV:%.*]] = zext i32 [[CALL]] to i64531// AMDGPU-NEXT: [[TMP0:%.*]] = load i64, ptr [[__LANE_MASK_ADDR_ASCAST]], align 8532// AMDGPU-NEXT: [[CALL1:%.*]] = call i64 @__gpu_first_lane_id(i64 noundef [[TMP0]]) #[[ATTR7]]533// AMDGPU-NEXT: [[CMP:%.*]] = icmp eq i64 [[CONV]], [[CALL1]]534// AMDGPU-NEXT: ret i1 [[CMP]]535//536//537// AMDGPU-LABEL: define internal void @__gpu_exit(538// AMDGPU-SAME: ) #[[ATTR1:[0-9]+]] {539// AMDGPU-NEXT: [[ENTRY:.*:]]540// AMDGPU-NEXT: call void @llvm.amdgcn.endpgm()541// AMDGPU-NEXT: ret void542//543//544// NVPTX-LABEL: define protected ptx_kernel void @foo(545// NVPTX-SAME: ) #[[ATTR0:[0-9]+]] {546// NVPTX-NEXT: [[ENTRY:.*:]]547// NVPTX-NEXT: [[CALL:%.*]] = call i32 @__gpu_num_blocks_x() #[[ATTR6:[0-9]+]]548// NVPTX-NEXT: [[CALL1:%.*]] = call i32 @__gpu_num_blocks_y() #[[ATTR6]]549// NVPTX-NEXT: [[CALL2:%.*]] = call i32 @__gpu_num_blocks_z() #[[ATTR6]]550// NVPTX-NEXT: [[CALL3:%.*]] = call i32 @__gpu_num_blocks(i32 noundef 0) #[[ATTR6]]551// NVPTX-NEXT: [[CALL4:%.*]] = call i32 @__gpu_block_id_x() #[[ATTR6]]552// NVPTX-NEXT: [[CALL5:%.*]] = call i32 @__gpu_block_id_y() #[[ATTR6]]553// NVPTX-NEXT: [[CALL6:%.*]] = call i32 @__gpu_block_id_z() #[[ATTR6]]554// NVPTX-NEXT: [[CALL7:%.*]] = call i32 @__gpu_block_id(i32 noundef 0) #[[ATTR6]]555// NVPTX-NEXT: [[CALL8:%.*]] = call i32 @__gpu_num_threads_x() #[[ATTR6]]556// NVPTX-NEXT: [[CALL9:%.*]] = call i32 @__gpu_num_threads_y() #[[ATTR6]]557// NVPTX-NEXT: [[CALL10:%.*]] = call i32 @__gpu_num_threads_z() #[[ATTR6]]558// NVPTX-NEXT: [[CALL11:%.*]] = call i32 @__gpu_num_threads(i32 noundef 0) #[[ATTR6]]559// NVPTX-NEXT: [[CALL12:%.*]] = call i32 @__gpu_thread_id_x() #[[ATTR6]]560// NVPTX-NEXT: [[CALL13:%.*]] = call i32 @__gpu_thread_id_y() #[[ATTR6]]561// NVPTX-NEXT: [[CALL14:%.*]] = call i32 @__gpu_thread_id_z() #[[ATTR6]]562// NVPTX-NEXT: [[CALL15:%.*]] = call i32 @__gpu_thread_id(i32 noundef 0) #[[ATTR6]]563// NVPTX-NEXT: [[CALL16:%.*]] = call i32 @__gpu_num_lanes() #[[ATTR6]]564// NVPTX-NEXT: [[CALL17:%.*]] = call i32 @__gpu_lane_id() #[[ATTR6]]565// NVPTX-NEXT: [[CALL18:%.*]] = call i64 @__gpu_lane_mask() #[[ATTR6]]566// NVPTX-NEXT: [[CALL19:%.*]] = call i32 @__gpu_read_first_lane_u32(i64 noundef -1, i32 noundef -1) #[[ATTR6]]567// NVPTX-NEXT: [[CALL20:%.*]] = call i64 @__gpu_read_first_lane_u64(i64 noundef -1, i64 noundef -1) #[[ATTR6]]568// NVPTX-NEXT: [[CALL21:%.*]] = call i64 @__gpu_ballot(i64 noundef -1, i1 noundef zeroext true) #[[ATTR6]]569// NVPTX-NEXT: call void @__gpu_sync_threads() #[[ATTR6]]570// NVPTX-NEXT: call void @__gpu_sync_lane(i64 noundef -1) #[[ATTR6]]571// NVPTX-NEXT: [[CALL22:%.*]] = call i32 @__gpu_shuffle_idx_u32(i64 noundef -1, i32 noundef -1, i32 noundef -1, i32 noundef 0) #[[ATTR6]]572// NVPTX-NEXT: [[CALL23:%.*]] = call i64 @__gpu_first_lane_id(i64 noundef -1) #[[ATTR6]]573// NVPTX-NEXT: [[CALL24:%.*]] = call zeroext i1 @__gpu_is_first_in_lane(i64 noundef -1) #[[ATTR6]]574// NVPTX-NEXT: call void @__gpu_exit() #[[ATTR7:[0-9]+]]575// NVPTX-NEXT: unreachable576//577//578// NVPTX-LABEL: define internal i32 @__gpu_num_blocks_x(579// NVPTX-SAME: ) #[[ATTR0]] {580// NVPTX-NEXT: [[ENTRY:.*:]]581// NVPTX-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()582// NVPTX-NEXT: ret i32 [[TMP0]]583//584//585// NVPTX-LABEL: define internal i32 @__gpu_num_blocks_y(586// NVPTX-SAME: ) #[[ATTR0]] {587// NVPTX-NEXT: [[ENTRY:.*:]]588// NVPTX-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()589// NVPTX-NEXT: ret i32 [[TMP0]]590//591//592// NVPTX-LABEL: define internal i32 @__gpu_num_blocks_z(593// NVPTX-SAME: ) #[[ATTR0]] {594// NVPTX-NEXT: [[ENTRY:.*:]]595// NVPTX-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()596// NVPTX-NEXT: ret i32 [[TMP0]]597//598//599// NVPTX-LABEL: define internal i32 @__gpu_num_blocks(600// NVPTX-SAME: i32 noundef [[__DIM:%.*]]) #[[ATTR0]] {601// NVPTX-NEXT: [[ENTRY:.*:]]602// NVPTX-NEXT: [[RETVAL:%.*]] = alloca i32, align 4603// NVPTX-NEXT: [[__DIM_ADDR:%.*]] = alloca i32, align 4604// NVPTX-NEXT: store i32 [[__DIM]], ptr [[__DIM_ADDR]], align 4605// NVPTX-NEXT: [[TMP0:%.*]] = load i32, ptr [[__DIM_ADDR]], align 4606// NVPTX-NEXT: switch i32 [[TMP0]], label %[[SW_DEFAULT:.*]] [607// NVPTX-NEXT: i32 0, label %[[SW_BB:.*]]608// NVPTX-NEXT: i32 1, label %[[SW_BB1:.*]]609// NVPTX-NEXT: i32 2, label %[[SW_BB3:.*]]610// NVPTX-NEXT: ]611// NVPTX: [[SW_BB]]:612// NVPTX-NEXT: [[CALL:%.*]] = call i32 @__gpu_num_blocks_x() #[[ATTR6]]613// NVPTX-NEXT: store i32 [[CALL]], ptr [[RETVAL]], align 4614// NVPTX-NEXT: br label %[[RETURN:.*]]615// NVPTX: [[SW_BB1]]:616// NVPTX-NEXT: [[CALL2:%.*]] = call i32 @__gpu_num_blocks_y() #[[ATTR6]]617// NVPTX-NEXT: store i32 [[CALL2]], ptr [[RETVAL]], align 4618// NVPTX-NEXT: br label %[[RETURN]]619// NVPTX: [[SW_BB3]]:620// NVPTX-NEXT: [[CALL4:%.*]] = call i32 @__gpu_num_blocks_z() #[[ATTR6]]621// NVPTX-NEXT: store i32 [[CALL4]], ptr [[RETVAL]], align 4622// NVPTX-NEXT: br label %[[RETURN]]623// NVPTX: [[SW_DEFAULT]]:624// NVPTX-NEXT: unreachable625// NVPTX: [[RETURN]]:626// NVPTX-NEXT: [[TMP1:%.*]] = load i32, ptr [[RETVAL]], align 4627// NVPTX-NEXT: ret i32 [[TMP1]]628//629//630// NVPTX-LABEL: define internal i32 @__gpu_block_id_x(631// NVPTX-SAME: ) #[[ATTR0]] {632// NVPTX-NEXT: [[ENTRY:.*:]]633// NVPTX-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()634// NVPTX-NEXT: ret i32 [[TMP0]]635//636//637// NVPTX-LABEL: define internal i32 @__gpu_block_id_y(638// NVPTX-SAME: ) #[[ATTR0]] {639// NVPTX-NEXT: [[ENTRY:.*:]]640// NVPTX-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()641// NVPTX-NEXT: ret i32 [[TMP0]]642//643//644// NVPTX-LABEL: define internal i32 @__gpu_block_id_z(645// NVPTX-SAME: ) #[[ATTR0]] {646// NVPTX-NEXT: [[ENTRY:.*:]]647// NVPTX-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()648// NVPTX-NEXT: ret i32 [[TMP0]]649//650//651// NVPTX-LABEL: define internal i32 @__gpu_block_id(652// NVPTX-SAME: i32 noundef [[__DIM:%.*]]) #[[ATTR0]] {653// NVPTX-NEXT: [[ENTRY:.*:]]654// NVPTX-NEXT: [[RETVAL:%.*]] = alloca i32, align 4655// NVPTX-NEXT: [[__DIM_ADDR:%.*]] = alloca i32, align 4656// NVPTX-NEXT: store i32 [[__DIM]], ptr [[__DIM_ADDR]], align 4657// NVPTX-NEXT: [[TMP0:%.*]] = load i32, ptr [[__DIM_ADDR]], align 4658// NVPTX-NEXT: switch i32 [[TMP0]], label %[[SW_DEFAULT:.*]] [659// NVPTX-NEXT: i32 0, label %[[SW_BB:.*]]660// NVPTX-NEXT: i32 1, label %[[SW_BB1:.*]]661// NVPTX-NEXT: i32 2, label %[[SW_BB3:.*]]662// NVPTX-NEXT: ]663// NVPTX: [[SW_BB]]:664// NVPTX-NEXT: [[CALL:%.*]] = call i32 @__gpu_block_id_x() #[[ATTR6]]665// NVPTX-NEXT: store i32 [[CALL]], ptr [[RETVAL]], align 4666// NVPTX-NEXT: br label %[[RETURN:.*]]667// NVPTX: [[SW_BB1]]:668// NVPTX-NEXT: [[CALL2:%.*]] = call i32 @__gpu_block_id_y() #[[ATTR6]]669// NVPTX-NEXT: store i32 [[CALL2]], ptr [[RETVAL]], align 4670// NVPTX-NEXT: br label %[[RETURN]]671// NVPTX: [[SW_BB3]]:672// NVPTX-NEXT: [[CALL4:%.*]] = call i32 @__gpu_block_id_z() #[[ATTR6]]673// NVPTX-NEXT: store i32 [[CALL4]], ptr [[RETVAL]], align 4674// NVPTX-NEXT: br label %[[RETURN]]675// NVPTX: [[SW_DEFAULT]]:676// NVPTX-NEXT: unreachable677// NVPTX: [[RETURN]]:678// NVPTX-NEXT: [[TMP1:%.*]] = load i32, ptr [[RETVAL]], align 4679// NVPTX-NEXT: ret i32 [[TMP1]]680//681//682// NVPTX-LABEL: define internal i32 @__gpu_num_threads_x(683// NVPTX-SAME: ) #[[ATTR0]] {684// NVPTX-NEXT: [[ENTRY:.*:]]685// NVPTX-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()686// NVPTX-NEXT: ret i32 [[TMP0]]687//688//689// NVPTX-LABEL: define internal i32 @__gpu_num_threads_y(690// NVPTX-SAME: ) #[[ATTR0]] {691// NVPTX-NEXT: [[ENTRY:.*:]]692// NVPTX-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.y()693// NVPTX-NEXT: ret i32 [[TMP0]]694//695//696// NVPTX-LABEL: define internal i32 @__gpu_num_threads_z(697// NVPTX-SAME: ) #[[ATTR0]] {698// NVPTX-NEXT: [[ENTRY:.*:]]699// NVPTX-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.z()700// NVPTX-NEXT: ret i32 [[TMP0]]701//702//703// NVPTX-LABEL: define internal i32 @__gpu_num_threads(704// NVPTX-SAME: i32 noundef [[__DIM:%.*]]) #[[ATTR0]] {705// NVPTX-NEXT: [[ENTRY:.*:]]706// NVPTX-NEXT: [[RETVAL:%.*]] = alloca i32, align 4707// NVPTX-NEXT: [[__DIM_ADDR:%.*]] = alloca i32, align 4708// NVPTX-NEXT: store i32 [[__DIM]], ptr [[__DIM_ADDR]], align 4709// NVPTX-NEXT: [[TMP0:%.*]] = load i32, ptr [[__DIM_ADDR]], align 4710// NVPTX-NEXT: switch i32 [[TMP0]], label %[[SW_DEFAULT:.*]] [711// NVPTX-NEXT: i32 0, label %[[SW_BB:.*]]712// NVPTX-NEXT: i32 1, label %[[SW_BB1:.*]]713// NVPTX-NEXT: i32 2, label %[[SW_BB3:.*]]714// NVPTX-NEXT: ]715// NVPTX: [[SW_BB]]:716// NVPTX-NEXT: [[CALL:%.*]] = call i32 @__gpu_num_threads_x() #[[ATTR6]]717// NVPTX-NEXT: store i32 [[CALL]], ptr [[RETVAL]], align 4718// NVPTX-NEXT: br label %[[RETURN:.*]]719// NVPTX: [[SW_BB1]]:720// NVPTX-NEXT: [[CALL2:%.*]] = call i32 @__gpu_num_threads_y() #[[ATTR6]]721// NVPTX-NEXT: store i32 [[CALL2]], ptr [[RETVAL]], align 4722// NVPTX-NEXT: br label %[[RETURN]]723// NVPTX: [[SW_BB3]]:724// NVPTX-NEXT: [[CALL4:%.*]] = call i32 @__gpu_num_threads_z() #[[ATTR6]]725// NVPTX-NEXT: store i32 [[CALL4]], ptr [[RETVAL]], align 4726// NVPTX-NEXT: br label %[[RETURN]]727// NVPTX: [[SW_DEFAULT]]:728// NVPTX-NEXT: unreachable729// NVPTX: [[RETURN]]:730// NVPTX-NEXT: [[TMP1:%.*]] = load i32, ptr [[RETVAL]], align 4731// NVPTX-NEXT: ret i32 [[TMP1]]732//733//734// NVPTX-LABEL: define internal i32 @__gpu_thread_id_x(735// NVPTX-SAME: ) #[[ATTR0]] {736// NVPTX-NEXT: [[ENTRY:.*:]]737// NVPTX-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()738// NVPTX-NEXT: ret i32 [[TMP0]]739//740//741// NVPTX-LABEL: define internal i32 @__gpu_thread_id_y(742// NVPTX-SAME: ) #[[ATTR0]] {743// NVPTX-NEXT: [[ENTRY:.*:]]744// NVPTX-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.y()745// NVPTX-NEXT: ret i32 [[TMP0]]746//747//748// NVPTX-LABEL: define internal i32 @__gpu_thread_id_z(749// NVPTX-SAME: ) #[[ATTR0]] {750// NVPTX-NEXT: [[ENTRY:.*:]]751// NVPTX-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.z()752// NVPTX-NEXT: ret i32 [[TMP0]]753//754//755// NVPTX-LABEL: define internal i32 @__gpu_thread_id(756// NVPTX-SAME: i32 noundef [[__DIM:%.*]]) #[[ATTR0]] {757// NVPTX-NEXT: [[ENTRY:.*:]]758// NVPTX-NEXT: [[RETVAL:%.*]] = alloca i32, align 4759// NVPTX-NEXT: [[__DIM_ADDR:%.*]] = alloca i32, align 4760// NVPTX-NEXT: store i32 [[__DIM]], ptr [[__DIM_ADDR]], align 4761// NVPTX-NEXT: [[TMP0:%.*]] = load i32, ptr [[__DIM_ADDR]], align 4762// NVPTX-NEXT: switch i32 [[TMP0]], label %[[SW_DEFAULT:.*]] [763// NVPTX-NEXT: i32 0, label %[[SW_BB:.*]]764// NVPTX-NEXT: i32 1, label %[[SW_BB1:.*]]765// NVPTX-NEXT: i32 2, label %[[SW_BB3:.*]]766// NVPTX-NEXT: ]767// NVPTX: [[SW_BB]]:768// NVPTX-NEXT: [[CALL:%.*]] = call i32 @__gpu_thread_id_x() #[[ATTR6]]769// NVPTX-NEXT: store i32 [[CALL]], ptr [[RETVAL]], align 4770// NVPTX-NEXT: br label %[[RETURN:.*]]771// NVPTX: [[SW_BB1]]:772// NVPTX-NEXT: [[CALL2:%.*]] = call i32 @__gpu_thread_id_y() #[[ATTR6]]773// NVPTX-NEXT: store i32 [[CALL2]], ptr [[RETVAL]], align 4774// NVPTX-NEXT: br label %[[RETURN]]775// NVPTX: [[SW_BB3]]:776// NVPTX-NEXT: [[CALL4:%.*]] = call i32 @__gpu_thread_id_z() #[[ATTR6]]777// NVPTX-NEXT: store i32 [[CALL4]], ptr [[RETVAL]], align 4778// NVPTX-NEXT: br label %[[RETURN]]779// NVPTX: [[SW_DEFAULT]]:780// NVPTX-NEXT: unreachable781// NVPTX: [[RETURN]]:782// NVPTX-NEXT: [[TMP1:%.*]] = load i32, ptr [[RETVAL]], align 4783// NVPTX-NEXT: ret i32 [[TMP1]]784//785//786// NVPTX-LABEL: define internal i32 @__gpu_num_lanes(787// NVPTX-SAME: ) #[[ATTR0]] {788// NVPTX-NEXT: [[ENTRY:.*:]]789// NVPTX-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()790// NVPTX-NEXT: ret i32 [[TMP0]]791//792//793// NVPTX-LABEL: define internal i32 @__gpu_lane_id(794// NVPTX-SAME: ) #[[ATTR0]] {795// NVPTX-NEXT: [[ENTRY:.*:]]796// NVPTX-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.laneid()797// NVPTX-NEXT: ret i32 [[TMP0]]798//799//800// NVPTX-LABEL: define internal i64 @__gpu_lane_mask(801// NVPTX-SAME: ) #[[ATTR0]] {802// NVPTX-NEXT: [[ENTRY:.*:]]803// NVPTX-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.activemask()804// NVPTX-NEXT: [[CONV:%.*]] = zext i32 [[TMP0]] to i64805// NVPTX-NEXT: ret i64 [[CONV]]806//807//808// NVPTX-LABEL: define internal i32 @__gpu_read_first_lane_u32(809// NVPTX-SAME: i64 noundef [[__LANE_MASK:%.*]], i32 noundef [[__X:%.*]]) #[[ATTR0]] {810// NVPTX-NEXT: [[ENTRY:.*:]]811// NVPTX-NEXT: [[__LANE_MASK_ADDR:%.*]] = alloca i64, align 8812// NVPTX-NEXT: [[__X_ADDR:%.*]] = alloca i32, align 4813// NVPTX-NEXT: [[__MASK:%.*]] = alloca i32, align 4814// NVPTX-NEXT: [[__ID:%.*]] = alloca i32, align 4815// NVPTX-NEXT: store i64 [[__LANE_MASK]], ptr [[__LANE_MASK_ADDR]], align 8816// NVPTX-NEXT: store i32 [[__X]], ptr [[__X_ADDR]], align 4817// NVPTX-NEXT: [[TMP0:%.*]] = load i64, ptr [[__LANE_MASK_ADDR]], align 8818// NVPTX-NEXT: [[CONV:%.*]] = trunc i64 [[TMP0]] to i32819// NVPTX-NEXT: store i32 [[CONV]], ptr [[__MASK]], align 4820// NVPTX-NEXT: [[TMP1:%.*]] = load i32, ptr [[__MASK]], align 4821// NVPTX-NEXT: [[TMP2:%.*]] = call i32 @llvm.cttz.i32(i32 [[TMP1]], i1 true)822// NVPTX-NEXT: [[TMP3:%.*]] = add i32 [[TMP2]], 1823// NVPTX-NEXT: [[ISZERO:%.*]] = icmp eq i32 [[TMP1]], 0824// NVPTX-NEXT: [[FFS:%.*]] = select i1 [[ISZERO]], i32 0, i32 [[TMP3]]825// NVPTX-NEXT: [[SUB:%.*]] = sub nsw i32 [[FFS]], 1826// NVPTX-NEXT: store i32 [[SUB]], ptr [[__ID]], align 4827// NVPTX-NEXT: [[TMP4:%.*]] = load i32, ptr [[__MASK]], align 4828// NVPTX-NEXT: [[TMP5:%.*]] = load i32, ptr [[__X_ADDR]], align 4829// NVPTX-NEXT: [[TMP6:%.*]] = load i32, ptr [[__ID]], align 4830// NVPTX-NEXT: [[CALL:%.*]] = call i32 @__gpu_num_lanes() #[[ATTR6]]831// NVPTX-NEXT: [[SUB1:%.*]] = sub i32 [[CALL]], 1832// NVPTX-NEXT: [[TMP7:%.*]] = call i32 @llvm.nvvm.shfl.sync.idx.i32(i32 [[TMP4]], i32 [[TMP5]], i32 [[TMP6]], i32 [[SUB1]])833// NVPTX-NEXT: ret i32 [[TMP7]]834//835//836// NVPTX-LABEL: define internal i64 @__gpu_read_first_lane_u64(837// NVPTX-SAME: i64 noundef [[__LANE_MASK:%.*]], i64 noundef [[__X:%.*]]) #[[ATTR0]] {838// NVPTX-NEXT: [[ENTRY:.*:]]839// NVPTX-NEXT: [[__LANE_MASK_ADDR:%.*]] = alloca i64, align 8840// NVPTX-NEXT: [[__X_ADDR:%.*]] = alloca i64, align 8841// NVPTX-NEXT: [[__HI:%.*]] = alloca i32, align 4842// NVPTX-NEXT: [[__LO:%.*]] = alloca i32, align 4843// NVPTX-NEXT: store i64 [[__LANE_MASK]], ptr [[__LANE_MASK_ADDR]], align 8844// NVPTX-NEXT: store i64 [[__X]], ptr [[__X_ADDR]], align 8845// NVPTX-NEXT: [[TMP0:%.*]] = load i64, ptr [[__X_ADDR]], align 8846// NVPTX-NEXT: [[SHR:%.*]] = lshr i64 [[TMP0]], 32847// NVPTX-NEXT: [[CONV:%.*]] = trunc i64 [[SHR]] to i32848// NVPTX-NEXT: store i32 [[CONV]], ptr [[__HI]], align 4849// NVPTX-NEXT: [[TMP1:%.*]] = load i64, ptr [[__X_ADDR]], align 8850// NVPTX-NEXT: [[AND:%.*]] = and i64 [[TMP1]], 4294967295851// NVPTX-NEXT: [[CONV1:%.*]] = trunc i64 [[AND]] to i32852// NVPTX-NEXT: store i32 [[CONV1]], ptr [[__LO]], align 4853// NVPTX-NEXT: [[TMP2:%.*]] = load i64, ptr [[__LANE_MASK_ADDR]], align 8854// NVPTX-NEXT: [[TMP3:%.*]] = load i32, ptr [[__HI]], align 4855// NVPTX-NEXT: [[CALL:%.*]] = call i32 @__gpu_read_first_lane_u32(i64 noundef [[TMP2]], i32 noundef [[TMP3]]) #[[ATTR6]]856// NVPTX-NEXT: [[CONV2:%.*]] = zext i32 [[CALL]] to i64857// NVPTX-NEXT: [[SHL:%.*]] = shl i64 [[CONV2]], 32858// NVPTX-NEXT: [[TMP4:%.*]] = load i64, ptr [[__LANE_MASK_ADDR]], align 8859// NVPTX-NEXT: [[TMP5:%.*]] = load i32, ptr [[__LO]], align 4860// NVPTX-NEXT: [[CALL3:%.*]] = call i32 @__gpu_read_first_lane_u32(i64 noundef [[TMP4]], i32 noundef [[TMP5]]) #[[ATTR6]]861// NVPTX-NEXT: [[CONV4:%.*]] = zext i32 [[CALL3]] to i64862// NVPTX-NEXT: [[AND5:%.*]] = and i64 [[CONV4]], 4294967295863// NVPTX-NEXT: [[OR:%.*]] = or i64 [[SHL]], [[AND5]]864// NVPTX-NEXT: ret i64 [[OR]]865//866//867// NVPTX-LABEL: define internal i64 @__gpu_ballot(868// NVPTX-SAME: i64 noundef [[__LANE_MASK:%.*]], i1 noundef zeroext [[__X:%.*]]) #[[ATTR0]] {869// NVPTX-NEXT: [[ENTRY:.*:]]870// NVPTX-NEXT: [[__LANE_MASK_ADDR:%.*]] = alloca i64, align 8871// NVPTX-NEXT: [[__X_ADDR:%.*]] = alloca i8, align 1872// NVPTX-NEXT: [[__MASK:%.*]] = alloca i32, align 4873// NVPTX-NEXT: store i64 [[__LANE_MASK]], ptr [[__LANE_MASK_ADDR]], align 8874// NVPTX-NEXT: [[STOREDV:%.*]] = zext i1 [[__X]] to i8875// NVPTX-NEXT: store i8 [[STOREDV]], ptr [[__X_ADDR]], align 1876// NVPTX-NEXT: [[TMP0:%.*]] = load i64, ptr [[__LANE_MASK_ADDR]], align 8877// NVPTX-NEXT: [[CONV:%.*]] = trunc i64 [[TMP0]] to i32878// NVPTX-NEXT: store i32 [[CONV]], ptr [[__MASK]], align 4879// NVPTX-NEXT: [[TMP1:%.*]] = load i32, ptr [[__MASK]], align 4880// NVPTX-NEXT: [[TMP2:%.*]] = load i8, ptr [[__X_ADDR]], align 1881// NVPTX-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP2]] to i1882// NVPTX-NEXT: [[TMP3:%.*]] = call i32 @llvm.nvvm.vote.ballot.sync(i32 [[TMP1]], i1 [[LOADEDV]])883// NVPTX-NEXT: [[CONV1:%.*]] = zext i32 [[TMP3]] to i64884// NVPTX-NEXT: ret i64 [[CONV1]]885//886//887// NVPTX-LABEL: define internal void @__gpu_sync_threads(888// NVPTX-SAME: ) #[[ATTR0]] {889// NVPTX-NEXT: [[ENTRY:.*:]]890// NVPTX-NEXT: call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)891// NVPTX-NEXT: ret void892//893//894// NVPTX-LABEL: define internal void @__gpu_sync_lane(895// NVPTX-SAME: i64 noundef [[__LANE_MASK:%.*]]) #[[ATTR0]] {896// NVPTX-NEXT: [[ENTRY:.*:]]897// NVPTX-NEXT: [[__LANE_MASK_ADDR:%.*]] = alloca i64, align 8898// NVPTX-NEXT: store i64 [[__LANE_MASK]], ptr [[__LANE_MASK_ADDR]], align 8899// NVPTX-NEXT: [[TMP0:%.*]] = load i64, ptr [[__LANE_MASK_ADDR]], align 8900// NVPTX-NEXT: [[CONV:%.*]] = trunc i64 [[TMP0]] to i32901// NVPTX-NEXT: call void @llvm.nvvm.bar.warp.sync(i32 [[CONV]])902// NVPTX-NEXT: ret void903//904//905// NVPTX-LABEL: define internal i32 @__gpu_shuffle_idx_u32(906// NVPTX-SAME: i64 noundef [[__LANE_MASK:%.*]], i32 noundef [[__IDX:%.*]], i32 noundef [[__X:%.*]], i32 noundef [[__WIDTH:%.*]]) #[[ATTR0]] {907// NVPTX-NEXT: [[ENTRY:.*:]]908// NVPTX-NEXT: [[__LANE_MASK_ADDR:%.*]] = alloca i64, align 8909// NVPTX-NEXT: [[__IDX_ADDR:%.*]] = alloca i32, align 4910// NVPTX-NEXT: [[__X_ADDR:%.*]] = alloca i32, align 4911// NVPTX-NEXT: [[__WIDTH_ADDR:%.*]] = alloca i32, align 4912// NVPTX-NEXT: [[__MASK:%.*]] = alloca i32, align 4913// NVPTX-NEXT: [[__BITMASK:%.*]] = alloca i8, align 1914// NVPTX-NEXT: store i64 [[__LANE_MASK]], ptr [[__LANE_MASK_ADDR]], align 8915// NVPTX-NEXT: store i32 [[__IDX]], ptr [[__IDX_ADDR]], align 4916// NVPTX-NEXT: store i32 [[__X]], ptr [[__X_ADDR]], align 4917// NVPTX-NEXT: store i32 [[__WIDTH]], ptr [[__WIDTH_ADDR]], align 4918// NVPTX-NEXT: [[TMP0:%.*]] = load i64, ptr [[__LANE_MASK_ADDR]], align 8919// NVPTX-NEXT: [[CONV:%.*]] = trunc i64 [[TMP0]] to i32920// NVPTX-NEXT: store i32 [[CONV]], ptr [[__MASK]], align 4921// NVPTX-NEXT: [[TMP1:%.*]] = load i32, ptr [[__IDX_ADDR]], align 4922// NVPTX-NEXT: [[SH_PROM:%.*]] = zext i32 [[TMP1]] to i64923// NVPTX-NEXT: [[SHL:%.*]] = shl i64 1, [[SH_PROM]]924// NVPTX-NEXT: [[TMP2:%.*]] = load i64, ptr [[__LANE_MASK_ADDR]], align 8925// NVPTX-NEXT: [[AND:%.*]] = and i64 [[SHL]], [[TMP2]]926// NVPTX-NEXT: [[TOBOOL:%.*]] = icmp ne i64 [[AND]], 0927// NVPTX-NEXT: [[STOREDV:%.*]] = zext i1 [[TOBOOL]] to i8928// NVPTX-NEXT: store i8 [[STOREDV]], ptr [[__BITMASK]], align 1929// NVPTX-NEXT: [[TMP3:%.*]] = load i8, ptr [[__BITMASK]], align 1930// NVPTX-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP3]] to i1931// NVPTX-NEXT: [[CONV1:%.*]] = zext i1 [[LOADEDV]] to i32932// NVPTX-NEXT: [[SUB:%.*]] = sub nsw i32 0, [[CONV1]]933// NVPTX-NEXT: [[TMP4:%.*]] = load i32, ptr [[__MASK]], align 4934// NVPTX-NEXT: [[TMP5:%.*]] = load i32, ptr [[__X_ADDR]], align 4935// NVPTX-NEXT: [[TMP6:%.*]] = load i32, ptr [[__IDX_ADDR]], align 4936// NVPTX-NEXT: [[CALL:%.*]] = call i32 @__gpu_num_lanes() #[[ATTR6]]937// NVPTX-NEXT: [[TMP7:%.*]] = load i32, ptr [[__WIDTH_ADDR]], align 4938// NVPTX-NEXT: [[SUB2:%.*]] = sub i32 [[CALL]], [[TMP7]]939// NVPTX-NEXT: [[SHL3:%.*]] = shl i32 [[SUB2]], 8940// NVPTX-NEXT: [[OR:%.*]] = or i32 [[SHL3]], 31941// NVPTX-NEXT: [[TMP8:%.*]] = call i32 @llvm.nvvm.shfl.sync.idx.i32(i32 [[TMP4]], i32 [[TMP5]], i32 [[TMP6]], i32 [[OR]])942// NVPTX-NEXT: [[AND4:%.*]] = and i32 [[SUB]], [[TMP8]]943// NVPTX-NEXT: ret i32 [[AND4]]944//945//946// NVPTX-LABEL: define internal i64 @__gpu_first_lane_id(947// NVPTX-SAME: i64 noundef [[__LANE_MASK:%.*]]) #[[ATTR0]] {948// NVPTX-NEXT: [[ENTRY:.*:]]949// NVPTX-NEXT: [[__LANE_MASK_ADDR:%.*]] = alloca i64, align 8950// NVPTX-NEXT: store i64 [[__LANE_MASK]], ptr [[__LANE_MASK_ADDR]], align 8951// NVPTX-NEXT: [[TMP0:%.*]] = load i64, ptr [[__LANE_MASK_ADDR]], align 8952// NVPTX-NEXT: [[TMP1:%.*]] = call i64 @llvm.cttz.i64(i64 [[TMP0]], i1 true)953// NVPTX-NEXT: [[TMP2:%.*]] = add i64 [[TMP1]], 1954// NVPTX-NEXT: [[ISZERO:%.*]] = icmp eq i64 [[TMP0]], 0955// NVPTX-NEXT: [[FFS:%.*]] = select i1 [[ISZERO]], i64 0, i64 [[TMP2]]956// NVPTX-NEXT: [[CAST:%.*]] = trunc i64 [[FFS]] to i32957// NVPTX-NEXT: [[SUB:%.*]] = sub nsw i32 [[CAST]], 1958// NVPTX-NEXT: [[CONV:%.*]] = sext i32 [[SUB]] to i64959// NVPTX-NEXT: ret i64 [[CONV]]960//961//962// NVPTX-LABEL: define internal zeroext i1 @__gpu_is_first_in_lane(963// NVPTX-SAME: i64 noundef [[__LANE_MASK:%.*]]) #[[ATTR0]] {964// NVPTX-NEXT: [[ENTRY:.*:]]965// NVPTX-NEXT: [[__LANE_MASK_ADDR:%.*]] = alloca i64, align 8966// NVPTX-NEXT: store i64 [[__LANE_MASK]], ptr [[__LANE_MASK_ADDR]], align 8967// NVPTX-NEXT: [[CALL:%.*]] = call i32 @__gpu_lane_id() #[[ATTR6]]968// NVPTX-NEXT: [[CONV:%.*]] = zext i32 [[CALL]] to i64969// NVPTX-NEXT: [[TMP0:%.*]] = load i64, ptr [[__LANE_MASK_ADDR]], align 8970// NVPTX-NEXT: [[CALL1:%.*]] = call i64 @__gpu_first_lane_id(i64 noundef [[TMP0]]) #[[ATTR6]]971// NVPTX-NEXT: [[CMP:%.*]] = icmp eq i64 [[CONV]], [[CALL1]]972// NVPTX-NEXT: ret i1 [[CMP]]973//974//975// NVPTX-LABEL: define internal void @__gpu_exit(976// NVPTX-SAME: ) #[[ATTR1:[0-9]+]] {977// NVPTX-NEXT: [[ENTRY:.*:]]978// NVPTX-NEXT: call void @llvm.nvvm.exit()979// NVPTX-NEXT: ret void980//981//.982// AMDGPU: [[RNG3]] = !{i32 1, i32 0}983// AMDGPU: [[META4]] = !{}984// AMDGPU: [[RNG5]] = !{i16 1, i16 1025}985//.986