736 lines · plain
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 52// REQUIRES: x86-registered-target3// REQUIRES: amdgpu-registered-target4 5// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck --check-prefixes=CHECK %s6// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck --check-prefixes=CHECK-SPIRV %s7// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -disable-O0-optnone -o - | opt -S -O2 | FileCheck %s --check-prefixes=OPT8// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -disable-O0-optnone -o - | opt -S -O2 | FileCheck %s --check-prefixes=OPT-SPIRV9// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -x hip %s -o - | FileCheck -check-prefix=HOST %s10 11#include "Inputs/cuda.h"12 13// Coerced struct from `struct S` without all generic pointers lowered into14// global ones.15 16// On the host-side compilation, generic pointer won't be coerced.17 18// CHECK-LABEL: define dso_local amdgpu_kernel void @_Z7kernel1Pi(19// CHECK-SAME: ptr addrspace(1) noundef [[X_COERCE:%.*]]) #[[ATTR0:[0-9]+]] {20// CHECK-NEXT: [[ENTRY:.*:]]21// CHECK-NEXT: [[X:%.*]] = alloca ptr, align 8, addrspace(5)22// CHECK-NEXT: [[X_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)23// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr24// CHECK-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr25// CHECK-NEXT: store ptr addrspace(1) [[X_COERCE]], ptr [[X_ASCAST]], align 826// CHECK-NEXT: [[X1:%.*]] = load ptr, ptr [[X_ASCAST]], align 827// CHECK-NEXT: store ptr [[X1]], ptr [[X_ADDR_ASCAST]], align 828// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[X_ADDR_ASCAST]], align 829// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 030// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX]], align 431// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 132// CHECK-NEXT: store i32 [[INC]], ptr [[ARRAYIDX]], align 433// CHECK-NEXT: ret void34//35// CHECK-SPIRV-LABEL: define spir_kernel void @_Z7kernel1Pi(36// CHECK-SPIRV-SAME: ptr addrspace(1) noundef [[X_COERCE:%.*]]) addrspace(4) #[[ATTR0:[0-9]+]] !max_work_group_size [[META5:![0-9]+]] {37// CHECK-SPIRV-NEXT: [[ENTRY:.*:]]38// CHECK-SPIRV-NEXT: [[X:%.*]] = alloca ptr addrspace(4), align 839// CHECK-SPIRV-NEXT: [[X_ADDR:%.*]] = alloca ptr addrspace(4), align 840// CHECK-SPIRV-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr [[X]] to ptr addrspace(4)41// CHECK-SPIRV-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr [[X_ADDR]] to ptr addrspace(4)42// CHECK-SPIRV-NEXT: store ptr addrspace(1) [[X_COERCE]], ptr addrspace(4) [[X_ASCAST]], align 843// CHECK-SPIRV-NEXT: [[X1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[X_ASCAST]], align 844// CHECK-SPIRV-NEXT: store ptr addrspace(4) [[X1]], ptr addrspace(4) [[X_ADDR_ASCAST]], align 845// CHECK-SPIRV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[X_ADDR_ASCAST]], align 846// CHECK-SPIRV-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr addrspace(4) [[TMP0]], i64 047// CHECK-SPIRV-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX]], align 448// CHECK-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 149// CHECK-SPIRV-NEXT: store i32 [[INC]], ptr addrspace(4) [[ARRAYIDX]], align 450// CHECK-SPIRV-NEXT: ret void51//52// OPT-LABEL: define dso_local amdgpu_kernel void @_Z7kernel1Pi(53// OPT-SAME: ptr addrspace(1) noundef captures(none) [[X_COERCE:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {54// OPT-NEXT: [[ENTRY:.*:]]55// OPT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(1) [[X_COERCE]], align 456// OPT-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 157// OPT-NEXT: store i32 [[INC]], ptr addrspace(1) [[X_COERCE]], align 458// OPT-NEXT: ret void59//60// OPT-SPIRV-LABEL: define spir_kernel void @_Z7kernel1Pi(61// OPT-SPIRV-SAME: ptr addrspace(1) noundef [[X_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0:[0-9]+]] !max_work_group_size [[META5:![0-9]+]] {62// OPT-SPIRV-NEXT: [[ENTRY:.*:]]63// OPT-SPIRV-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[X_COERCE]] to i6464// OPT-SPIRV-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4)65// OPT-SPIRV-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 466// OPT-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP2]], 167// OPT-SPIRV-NEXT: store i32 [[INC]], ptr addrspace(4) [[TMP1]], align 468// OPT-SPIRV-NEXT: ret void69//70// HOST-LABEL: define dso_local void @_Z22__device_stub__kernel1Pi(71// HOST-SAME: ptr noundef [[X:%.*]]) #[[ATTR0:[0-9]+]] {72// HOST-NEXT: [[ENTRY:.*:]]73// HOST-NEXT: [[X_ADDR:%.*]] = alloca ptr, align 874// HOST-NEXT: store ptr [[X]], ptr [[X_ADDR]], align 875// HOST-NEXT: [[TMP0:%.*]] = call i32 @hipSetupArgument(ptr [[X_ADDR]], i64 8, i64 0)76// HOST-NEXT: [[TMP1:%.*]] = icmp eq i32 [[TMP0]], 077// HOST-NEXT: br i1 [[TMP1]], label %[[SETUP_NEXT:.*]], label %[[SETUP_END:.*]]78// HOST: [[SETUP_NEXT]]:79// HOST-NEXT: [[TMP2:%.*]] = call i32 @hipLaunchByPtr(ptr @_Z7kernel1Pi)80// HOST-NEXT: br label %[[SETUP_END]]81// HOST: [[SETUP_END]]:82// HOST-NEXT: ret void83//84__global__ void kernel1(int *x) {85 x[0]++;86}87 88// CHECK-LABEL: define dso_local amdgpu_kernel void @_Z7kernel2Ri(89// CHECK-SAME: ptr addrspace(1) noundef nonnull align 4 dereferenceable(4) [[X_COERCE:%.*]]) #[[ATTR0]] {90// CHECK-NEXT: [[ENTRY:.*:]]91// CHECK-NEXT: [[X:%.*]] = alloca ptr, align 8, addrspace(5)92// CHECK-NEXT: [[X_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)93// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr94// CHECK-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr95// CHECK-NEXT: store ptr addrspace(1) [[X_COERCE]], ptr [[X_ASCAST]], align 896// CHECK-NEXT: [[X1:%.*]] = load ptr, ptr [[X_ASCAST]], align 897// CHECK-NEXT: store ptr [[X1]], ptr [[X_ADDR_ASCAST]], align 898// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[X_ADDR_ASCAST]], align 899// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4100// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1101// CHECK-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4102// CHECK-NEXT: ret void103//104// CHECK-SPIRV-LABEL: define spir_kernel void @_Z7kernel2Ri(105// CHECK-SPIRV-SAME: ptr addrspace(1) noundef align 4 dereferenceable(4) [[X_COERCE:%.*]]) addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] {106// CHECK-SPIRV-NEXT: [[ENTRY:.*:]]107// CHECK-SPIRV-NEXT: [[X:%.*]] = alloca ptr addrspace(4), align 8108// CHECK-SPIRV-NEXT: [[X_ADDR:%.*]] = alloca ptr addrspace(4), align 8109// CHECK-SPIRV-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr [[X]] to ptr addrspace(4)110// CHECK-SPIRV-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr [[X_ADDR]] to ptr addrspace(4)111// CHECK-SPIRV-NEXT: store ptr addrspace(1) [[X_COERCE]], ptr addrspace(4) [[X_ASCAST]], align 8112// CHECK-SPIRV-NEXT: [[X1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[X_ASCAST]], align 8113// CHECK-SPIRV-NEXT: store ptr addrspace(4) [[X1]], ptr addrspace(4) [[X_ADDR_ASCAST]], align 8114// CHECK-SPIRV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[X_ADDR_ASCAST]], align 8115// CHECK-SPIRV-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(4) [[TMP0]], align 4116// CHECK-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1117// CHECK-SPIRV-NEXT: store i32 [[INC]], ptr addrspace(4) [[TMP0]], align 4118// CHECK-SPIRV-NEXT: ret void119//120// OPT-LABEL: define dso_local amdgpu_kernel void @_Z7kernel2Ri(121// OPT-SAME: ptr addrspace(1) noundef nonnull align 4 captures(none) dereferenceable(4) [[X_COERCE:%.*]]) local_unnamed_addr #[[ATTR0]] {122// OPT-NEXT: [[ENTRY:.*:]]123// OPT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(1) [[X_COERCE]], align 4124// OPT-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1125// OPT-NEXT: store i32 [[INC]], ptr addrspace(1) [[X_COERCE]], align 4126// OPT-NEXT: ret void127//128// OPT-SPIRV-LABEL: define spir_kernel void @_Z7kernel2Ri(129// OPT-SPIRV-SAME: ptr addrspace(1) noundef align 4 dereferenceable(4) [[X_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] {130// OPT-SPIRV-NEXT: [[ENTRY:.*:]]131// OPT-SPIRV-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[X_COERCE]] to i64132// OPT-SPIRV-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4)133// OPT-SPIRV-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 4134// OPT-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP2]], 1135// OPT-SPIRV-NEXT: store i32 [[INC]], ptr addrspace(4) [[TMP1]], align 4136// OPT-SPIRV-NEXT: ret void137//138// HOST-LABEL: define dso_local void @_Z22__device_stub__kernel2Ri(139// HOST-SAME: ptr noundef nonnull align 4 dereferenceable(4) [[X:%.*]]) #[[ATTR0]] {140// HOST-NEXT: [[ENTRY:.*:]]141// HOST-NEXT: [[X_ADDR:%.*]] = alloca ptr, align 8142// HOST-NEXT: store ptr [[X]], ptr [[X_ADDR]], align 8143// HOST-NEXT: [[TMP0:%.*]] = call i32 @hipSetupArgument(ptr [[X_ADDR]], i64 8, i64 0)144// HOST-NEXT: [[TMP1:%.*]] = icmp eq i32 [[TMP0]], 0145// HOST-NEXT: br i1 [[TMP1]], label %[[SETUP_NEXT:.*]], label %[[SETUP_END:.*]]146// HOST: [[SETUP_NEXT]]:147// HOST-NEXT: [[TMP2:%.*]] = call i32 @hipLaunchByPtr(ptr @_Z7kernel2Ri)148// HOST-NEXT: br label %[[SETUP_END]]149// HOST: [[SETUP_END]]:150// HOST-NEXT: ret void151//152__global__ void kernel2(int &x) {153 x++;154}155 156// CHECK-LABEL: define dso_local amdgpu_kernel void @_Z7kernel3PU3AS2iPU3AS1i(157// CHECK-SAME: ptr addrspace(2) noundef [[X:%.*]], ptr addrspace(1) noundef [[Y:%.*]]) #[[ATTR0]] {158// CHECK-NEXT: [[ENTRY:.*:]]159// CHECK-NEXT: [[X_ADDR:%.*]] = alloca ptr addrspace(2), align 8, addrspace(5)160// CHECK-NEXT: [[Y_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)161// CHECK-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr162// CHECK-NEXT: [[Y_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[Y_ADDR]] to ptr163// CHECK-NEXT: store ptr addrspace(2) [[X]], ptr [[X_ADDR_ASCAST]], align 8164// CHECK-NEXT: store ptr addrspace(1) [[Y]], ptr [[Y_ADDR_ASCAST]], align 8165// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(2), ptr [[X_ADDR_ASCAST]], align 8166// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr addrspace(2) [[TMP0]], i64 0167// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(2) [[ARRAYIDX]], align 4168// CHECK-NEXT: [[TMP2:%.*]] = load ptr addrspace(1), ptr [[Y_ADDR_ASCAST]], align 8169// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[TMP2]], i64 0170// CHECK-NEXT: store i32 [[TMP1]], ptr addrspace(1) [[ARRAYIDX1]], align 4171// CHECK-NEXT: ret void172//173// CHECK-SPIRV-LABEL: define spir_kernel void @_Z7kernel3PU3AS2iPU3AS1i(174// CHECK-SPIRV-SAME: ptr addrspace(2) noundef [[X:%.*]], ptr addrspace(1) noundef [[Y:%.*]]) addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] {175// CHECK-SPIRV-NEXT: [[ENTRY:.*:]]176// CHECK-SPIRV-NEXT: [[X_ADDR:%.*]] = alloca ptr addrspace(2), align 8177// CHECK-SPIRV-NEXT: [[Y_ADDR:%.*]] = alloca ptr addrspace(1), align 8178// CHECK-SPIRV-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr [[X_ADDR]] to ptr addrspace(4)179// CHECK-SPIRV-NEXT: [[Y_ADDR_ASCAST:%.*]] = addrspacecast ptr [[Y_ADDR]] to ptr addrspace(4)180// CHECK-SPIRV-NEXT: store ptr addrspace(2) [[X]], ptr addrspace(4) [[X_ADDR_ASCAST]], align 8181// CHECK-SPIRV-NEXT: store ptr addrspace(1) [[Y]], ptr addrspace(4) [[Y_ADDR_ASCAST]], align 8182// CHECK-SPIRV-NEXT: [[TMP0:%.*]] = load ptr addrspace(2), ptr addrspace(4) [[X_ADDR_ASCAST]], align 8183// CHECK-SPIRV-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr addrspace(2) [[TMP0]], i64 0184// CHECK-SPIRV-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(2) [[ARRAYIDX]], align 4185// CHECK-SPIRV-NEXT: [[TMP2:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[Y_ADDR_ASCAST]], align 8186// CHECK-SPIRV-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[TMP2]], i64 0187// CHECK-SPIRV-NEXT: store i32 [[TMP1]], ptr addrspace(1) [[ARRAYIDX1]], align 4188// CHECK-SPIRV-NEXT: ret void189//190// OPT-LABEL: define dso_local amdgpu_kernel void @_Z7kernel3PU3AS2iPU3AS1i(191// OPT-SAME: ptr addrspace(2) noundef readonly captures(none) [[X:%.*]], ptr addrspace(1) noundef writeonly captures(none) initializes((0, 4)) [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] {192// OPT-NEXT: [[ENTRY:.*:]]193// OPT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(2) [[X]], align 4194// OPT-NEXT: store i32 [[TMP0]], ptr addrspace(1) [[Y]], align 4195// OPT-NEXT: ret void196//197// OPT-SPIRV-LABEL: define spir_kernel void @_Z7kernel3PU3AS2iPU3AS1i(198// OPT-SPIRV-SAME: ptr addrspace(2) noundef readonly captures(none) [[X:%.*]], ptr addrspace(1) noundef writeonly captures(none) initializes((0, 4)) [[Y:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR1:[0-9]+]] !max_work_group_size [[META5]] {199// OPT-SPIRV-NEXT: [[ENTRY:.*:]]200// OPT-SPIRV-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(2) [[X]], align 4201// OPT-SPIRV-NEXT: store i32 [[TMP0]], ptr addrspace(1) [[Y]], align 4202// OPT-SPIRV-NEXT: ret void203//204// HOST-LABEL: define dso_local void @_Z22__device_stub__kernel3PU3AS2iPU3AS1i(205// HOST-SAME: ptr addrspace(2) noundef [[X:%.*]], ptr addrspace(1) noundef [[Y:%.*]]) #[[ATTR0]] {206// HOST-NEXT: [[ENTRY:.*:]]207// HOST-NEXT: [[X_ADDR:%.*]] = alloca ptr addrspace(2), align 8208// HOST-NEXT: [[Y_ADDR:%.*]] = alloca ptr addrspace(1), align 8209// HOST-NEXT: store ptr addrspace(2) [[X]], ptr [[X_ADDR]], align 8210// HOST-NEXT: store ptr addrspace(1) [[Y]], ptr [[Y_ADDR]], align 8211// HOST-NEXT: [[TMP0:%.*]] = call i32 @hipSetupArgument(ptr [[X_ADDR]], i64 8, i64 0)212// HOST-NEXT: [[TMP1:%.*]] = icmp eq i32 [[TMP0]], 0213// HOST-NEXT: br i1 [[TMP1]], label %[[SETUP_NEXT:.*]], label %[[SETUP_END:.*]]214// HOST: [[SETUP_NEXT]]:215// HOST-NEXT: [[TMP2:%.*]] = call i32 @hipSetupArgument(ptr [[Y_ADDR]], i64 8, i64 8)216// HOST-NEXT: [[TMP3:%.*]] = icmp eq i32 [[TMP2]], 0217// HOST-NEXT: br i1 [[TMP3]], label %[[SETUP_NEXT1:.*]], label %[[SETUP_END]]218// HOST: [[SETUP_NEXT1]]:219// HOST-NEXT: [[TMP4:%.*]] = call i32 @hipLaunchByPtr(ptr @_Z7kernel3PU3AS2iPU3AS1i)220// HOST-NEXT: br label %[[SETUP_END]]221// HOST: [[SETUP_END]]:222// HOST-NEXT: ret void223//224__global__ void kernel3(__attribute__((address_space(2))) int *x,225 __attribute__((address_space(1))) int *y) {226 y[0] = x[0];227}228 229// CHECK-LABEL: define dso_local void @_Z4funcPi(230// CHECK-SAME: ptr noundef [[X:%.*]]) #[[ATTR1:[0-9]+]] {231// CHECK-NEXT: [[ENTRY:.*:]]232// CHECK-NEXT: [[X_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)233// CHECK-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr234// CHECK-NEXT: store ptr [[X]], ptr [[X_ADDR_ASCAST]], align 8235// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[X_ADDR_ASCAST]], align 8236// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 0237// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX]], align 4238// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1239// CHECK-NEXT: store i32 [[INC]], ptr [[ARRAYIDX]], align 4240// CHECK-NEXT: ret void241//242// CHECK-SPIRV-LABEL: define spir_func void @_Z4funcPi(243// CHECK-SPIRV-SAME: ptr addrspace(4) noundef [[X:%.*]]) addrspace(4) #[[ATTR1:[0-9]+]] {244// CHECK-SPIRV-NEXT: [[ENTRY:.*:]]245// CHECK-SPIRV-NEXT: [[X_ADDR:%.*]] = alloca ptr addrspace(4), align 8246// CHECK-SPIRV-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr [[X_ADDR]] to ptr addrspace(4)247// CHECK-SPIRV-NEXT: store ptr addrspace(4) [[X]], ptr addrspace(4) [[X_ADDR_ASCAST]], align 8248// CHECK-SPIRV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[X_ADDR_ASCAST]], align 8249// CHECK-SPIRV-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr addrspace(4) [[TMP0]], i64 0250// CHECK-SPIRV-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX]], align 4251// CHECK-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1252// CHECK-SPIRV-NEXT: store i32 [[INC]], ptr addrspace(4) [[ARRAYIDX]], align 4253// CHECK-SPIRV-NEXT: ret void254//255// OPT-LABEL: define dso_local void @_Z4funcPi(256// OPT-SAME: ptr noundef captures(none) [[X:%.*]]) local_unnamed_addr #[[ATTR1:[0-9]+]] {257// OPT-NEXT: [[ENTRY:.*:]]258// OPT-NEXT: [[TMP0:%.*]] = load i32, ptr [[X]], align 4259// OPT-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1260// OPT-NEXT: store i32 [[INC]], ptr [[X]], align 4261// OPT-NEXT: ret void262//263// OPT-SPIRV-LABEL: define spir_func void @_Z4funcPi(264// OPT-SPIRV-SAME: ptr addrspace(4) noundef captures(none) [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR2:[0-9]+]] {265// OPT-SPIRV-NEXT: [[ENTRY:.*:]]266// OPT-SPIRV-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(4) [[X]], align 4267// OPT-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1268// OPT-SPIRV-NEXT: store i32 [[INC]], ptr addrspace(4) [[X]], align 4269// OPT-SPIRV-NEXT: ret void270//271__device__ void func(int *x) {272 x[0]++;273}274 275struct S {276 int *x;277 float *y;278};279// `by-val` struct is passed by-indirect-alias (a mix of by-ref and indirect280// by-val). However, the enhanced address inferring pass should be able to281// assume they are global pointers.282// For SPIR-V, since byref is not supported at the moment, we pass it as direct.283 284// CHECK-LABEL: define dso_local amdgpu_kernel void @_Z7kernel41S(285// CHECK-SAME: ptr addrspace(4) noundef byref([[STRUCT_S:%.*]]) align 8 [[TMP0:%.*]]) #[[ATTR0]] {286// CHECK-NEXT: [[ENTRY:.*:]]287// CHECK-NEXT: [[COERCE:%.*]] = alloca [[STRUCT_S]], align 8, addrspace(5)288// CHECK-NEXT: [[S:%.*]] = addrspacecast ptr addrspace(5) [[COERCE]] to ptr289// CHECK-NEXT: call void @llvm.memcpy.p0.p4.i64(ptr align 8 [[S]], ptr addrspace(4) align 8 [[TMP0]], i64 16, i1 false)290// CHECK-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[S]], i32 0, i32 0291// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[X]], align 8292// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 0293// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[ARRAYIDX]], align 4294// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP2]], 1295// CHECK-NEXT: store i32 [[INC]], ptr [[ARRAYIDX]], align 4296// CHECK-NEXT: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[S]], i32 0, i32 1297// CHECK-NEXT: [[TMP3:%.*]] = load ptr, ptr [[Y]], align 8298// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds float, ptr [[TMP3]], i64 0299// CHECK-NEXT: [[TMP4:%.*]] = load float, ptr [[ARRAYIDX1]], align 4300// CHECK-NEXT: [[ADD:%.*]] = fadd contract float [[TMP4]], 1.000000e+00301// CHECK-NEXT: store float [[ADD]], ptr [[ARRAYIDX1]], align 4302// CHECK-NEXT: ret void303//304// CHECK-SPIRV-LABEL: define spir_kernel void @_Z7kernel41S(305// CHECK-SPIRV-SAME: [[STRUCT_S:%.*]] [[S_COERCE:%.*]]) addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] {306// CHECK-SPIRV-NEXT: [[ENTRY:.*:]]307// CHECK-SPIRV-NEXT: [[S:%.*]] = alloca [[STRUCT_S]], align 8308// CHECK-SPIRV-NEXT: [[S1:%.*]] = addrspacecast ptr [[S]] to ptr addrspace(4)309// CHECK-SPIRV-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr addrspace(4) [[S1]], i32 0, i32 0310// CHECK-SPIRV-NEXT: [[TMP1:%.*]] = extractvalue [[STRUCT_S]] [[S_COERCE]], 0311// CHECK-SPIRV-NEXT: store ptr addrspace(4) [[TMP1]], ptr addrspace(4) [[TMP0]], align 8312// CHECK-SPIRV-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr addrspace(4) [[S1]], i32 0, i32 1313// CHECK-SPIRV-NEXT: [[TMP3:%.*]] = extractvalue [[STRUCT_S]] [[S_COERCE]], 1314// CHECK-SPIRV-NEXT: store ptr addrspace(4) [[TMP3]], ptr addrspace(4) [[TMP2]], align 8315// CHECK-SPIRV-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr addrspace(4) [[S1]], i32 0, i32 0316// CHECK-SPIRV-NEXT: [[TMP4:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[X]], align 8317// CHECK-SPIRV-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr addrspace(4) [[TMP4]], i64 0318// CHECK-SPIRV-NEXT: [[TMP5:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX]], align 4319// CHECK-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP5]], 1320// CHECK-SPIRV-NEXT: store i32 [[INC]], ptr addrspace(4) [[ARRAYIDX]], align 4321// CHECK-SPIRV-NEXT: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr addrspace(4) [[S1]], i32 0, i32 1322// CHECK-SPIRV-NEXT: [[TMP6:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[Y]], align 8323// CHECK-SPIRV-NEXT: [[ARRAYIDX2:%.*]] = getelementptr inbounds float, ptr addrspace(4) [[TMP6]], i64 0324// CHECK-SPIRV-NEXT: [[TMP7:%.*]] = load float, ptr addrspace(4) [[ARRAYIDX2]], align 4325// CHECK-SPIRV-NEXT: [[ADD:%.*]] = fadd contract float [[TMP7]], 1.000000e+00326// CHECK-SPIRV-NEXT: store float [[ADD]], ptr addrspace(4) [[ARRAYIDX2]], align 4327// CHECK-SPIRV-NEXT: ret void328//329// OPT-LABEL: define dso_local amdgpu_kernel void @_Z7kernel41S(330// OPT-SAME: ptr addrspace(4) noundef readonly byref([[STRUCT_S:%.*]]) align 8 captures(none) [[TMP0:%.*]]) local_unnamed_addr #[[ATTR2:[0-9]+]] {331// OPT-NEXT: [[ENTRY:.*:]]332// OPT-NEXT: [[COERCE_SROA_0_0_COPYLOAD:%.*]] = load ptr, ptr addrspace(4) [[TMP0]], align 8, !amdgpu.noclobber [[META4:![0-9]+]]333// OPT-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[COERCE_SROA_0_0_COPYLOAD]] to ptr addrspace(1)334// OPT-NEXT: [[COERCE_SROA_2_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[TMP0]], i64 8335// OPT-NEXT: [[COERCE_SROA_2_0_COPYLOAD:%.*]] = load ptr, ptr addrspace(4) [[COERCE_SROA_2_0__SROA_IDX]], align 8, !amdgpu.noclobber [[META4]]336// OPT-NEXT: [[TMP2:%.*]] = addrspacecast ptr [[COERCE_SROA_2_0_COPYLOAD]] to ptr addrspace(1)337// OPT-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(1) [[TMP1]], align 4, !amdgpu.noclobber [[META4]]338// OPT-NEXT: [[INC:%.*]] = add nsw i32 [[TMP3]], 1339// OPT-NEXT: store i32 [[INC]], ptr addrspace(1) [[TMP1]], align 4340// OPT-NEXT: [[TMP4:%.*]] = load float, ptr addrspace(1) [[TMP2]], align 4341// OPT-NEXT: [[ADD:%.*]] = fadd contract float [[TMP4]], 1.000000e+00342// OPT-NEXT: store float [[ADD]], ptr addrspace(1) [[TMP2]], align 4343// OPT-NEXT: ret void344//345// OPT-SPIRV-LABEL: define spir_kernel void @_Z7kernel41S(346// OPT-SPIRV-SAME: [[STRUCT_S:%.*]] [[S_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] {347// OPT-SPIRV-NEXT: [[ENTRY:.*:]]348// OPT-SPIRV-NEXT: [[TMP0:%.*]] = extractvalue [[STRUCT_S]] [[S_COERCE]], 0349// OPT-SPIRV-NEXT: [[TMP1:%.*]] = extractvalue [[STRUCT_S]] [[S_COERCE]], 1350// OPT-SPIRV-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP0]], align 4351// OPT-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP2]], 1352// OPT-SPIRV-NEXT: store i32 [[INC]], ptr addrspace(4) [[TMP0]], align 4353// OPT-SPIRV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(4) [[TMP1]], align 4354// OPT-SPIRV-NEXT: [[ADD:%.*]] = fadd contract float [[TMP3]], 1.000000e+00355// OPT-SPIRV-NEXT: store float [[ADD]], ptr addrspace(4) [[TMP1]], align 4356// OPT-SPIRV-NEXT: ret void357//358// HOST-LABEL: define dso_local void @_Z22__device_stub__kernel41S(359// HOST-SAME: ptr [[S_COERCE0:%.*]], ptr [[S_COERCE1:%.*]]) #[[ATTR0]] {360// HOST-NEXT: [[ENTRY:.*:]]361// HOST-NEXT: [[S:%.*]] = alloca [[STRUCT_S:%.*]], align 8362// HOST-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw { ptr, ptr }, ptr [[S]], i32 0, i32 0363// HOST-NEXT: store ptr [[S_COERCE0]], ptr [[TMP0]], align 8364// HOST-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw { ptr, ptr }, ptr [[S]], i32 0, i32 1365// HOST-NEXT: store ptr [[S_COERCE1]], ptr [[TMP1]], align 8366// HOST-NEXT: [[TMP2:%.*]] = call i32 @hipSetupArgument(ptr [[S]], i64 16, i64 0)367// HOST-NEXT: [[TMP3:%.*]] = icmp eq i32 [[TMP2]], 0368// HOST-NEXT: br i1 [[TMP3]], label %[[SETUP_NEXT:.*]], label %[[SETUP_END:.*]]369// HOST: [[SETUP_NEXT]]:370// HOST-NEXT: [[TMP4:%.*]] = call i32 @hipLaunchByPtr(ptr @_Z7kernel41S)371// HOST-NEXT: br label %[[SETUP_END]]372// HOST: [[SETUP_END]]:373// HOST-NEXT: ret void374//375__global__ void kernel4(struct S s) {376 s.x[0]++;377 s.y[0] += 1.f;378}379 380// If a pointer to struct is passed, only the pointer itself is coerced into the global one.381 382// CHECK-LABEL: define dso_local amdgpu_kernel void @_Z7kernel5P1S(383// CHECK-SAME: ptr addrspace(1) noundef [[S_COERCE:%.*]]) #[[ATTR0]] {384// CHECK-NEXT: [[ENTRY:.*:]]385// CHECK-NEXT: [[S:%.*]] = alloca ptr, align 8, addrspace(5)386// CHECK-NEXT: [[S_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)387// CHECK-NEXT: [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr388// CHECK-NEXT: [[S_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S_ADDR]] to ptr389// CHECK-NEXT: store ptr addrspace(1) [[S_COERCE]], ptr [[S_ASCAST]], align 8390// CHECK-NEXT: [[S1:%.*]] = load ptr, ptr [[S_ASCAST]], align 8391// CHECK-NEXT: store ptr [[S1]], ptr [[S_ADDR_ASCAST]], align 8392// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[S_ADDR_ASCAST]], align 8393// CHECK-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[TMP0]], i32 0, i32 0394// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[X]], align 8395// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 0396// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[ARRAYIDX]], align 4397// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP2]], 1398// CHECK-NEXT: store i32 [[INC]], ptr [[ARRAYIDX]], align 4399// CHECK-NEXT: [[TMP3:%.*]] = load ptr, ptr [[S_ADDR_ASCAST]], align 8400// CHECK-NEXT: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[TMP3]], i32 0, i32 1401// CHECK-NEXT: [[TMP4:%.*]] = load ptr, ptr [[Y]], align 8402// CHECK-NEXT: [[ARRAYIDX2:%.*]] = getelementptr inbounds float, ptr [[TMP4]], i64 0403// CHECK-NEXT: [[TMP5:%.*]] = load float, ptr [[ARRAYIDX2]], align 4404// CHECK-NEXT: [[ADD:%.*]] = fadd contract float [[TMP5]], 1.000000e+00405// CHECK-NEXT: store float [[ADD]], ptr [[ARRAYIDX2]], align 4406// CHECK-NEXT: ret void407//408// CHECK-SPIRV-LABEL: define spir_kernel void @_Z7kernel5P1S(409// CHECK-SPIRV-SAME: ptr addrspace(1) noundef [[S_COERCE:%.*]]) addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] {410// CHECK-SPIRV-NEXT: [[ENTRY:.*:]]411// CHECK-SPIRV-NEXT: [[S:%.*]] = alloca ptr addrspace(4), align 8412// CHECK-SPIRV-NEXT: [[S_ADDR:%.*]] = alloca ptr addrspace(4), align 8413// CHECK-SPIRV-NEXT: [[S_ASCAST:%.*]] = addrspacecast ptr [[S]] to ptr addrspace(4)414// CHECK-SPIRV-NEXT: [[S_ADDR_ASCAST:%.*]] = addrspacecast ptr [[S_ADDR]] to ptr addrspace(4)415// CHECK-SPIRV-NEXT: store ptr addrspace(1) [[S_COERCE]], ptr addrspace(4) [[S_ASCAST]], align 8416// CHECK-SPIRV-NEXT: [[S1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[S_ASCAST]], align 8417// CHECK-SPIRV-NEXT: store ptr addrspace(4) [[S1]], ptr addrspace(4) [[S_ADDR_ASCAST]], align 8418// CHECK-SPIRV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[S_ADDR_ASCAST]], align 8419// CHECK-SPIRV-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr addrspace(4) [[TMP0]], i32 0, i32 0420// CHECK-SPIRV-NEXT: [[TMP1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[X]], align 8421// CHECK-SPIRV-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr addrspace(4) [[TMP1]], i64 0422// CHECK-SPIRV-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX]], align 4423// CHECK-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP2]], 1424// CHECK-SPIRV-NEXT: store i32 [[INC]], ptr addrspace(4) [[ARRAYIDX]], align 4425// CHECK-SPIRV-NEXT: [[TMP3:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[S_ADDR_ASCAST]], align 8426// CHECK-SPIRV-NEXT: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr addrspace(4) [[TMP3]], i32 0, i32 1427// CHECK-SPIRV-NEXT: [[TMP4:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[Y]], align 8428// CHECK-SPIRV-NEXT: [[ARRAYIDX2:%.*]] = getelementptr inbounds float, ptr addrspace(4) [[TMP4]], i64 0429// CHECK-SPIRV-NEXT: [[TMP5:%.*]] = load float, ptr addrspace(4) [[ARRAYIDX2]], align 4430// CHECK-SPIRV-NEXT: [[ADD:%.*]] = fadd contract float [[TMP5]], 1.000000e+00431// CHECK-SPIRV-NEXT: store float [[ADD]], ptr addrspace(4) [[ARRAYIDX2]], align 4432// CHECK-SPIRV-NEXT: ret void433//434// OPT-LABEL: define dso_local amdgpu_kernel void @_Z7kernel5P1S(435// OPT-SAME: ptr addrspace(1) noundef readonly captures(none) [[S_COERCE:%.*]]) local_unnamed_addr #[[ATTR2]] {436// OPT-NEXT: [[ENTRY:.*:]]437// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(1) [[S_COERCE]], align 8438// OPT-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4439// OPT-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1440// OPT-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4441// OPT-NEXT: [[Y:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(1) [[S_COERCE]], i64 8442// OPT-NEXT: [[TMP2:%.*]] = load ptr, ptr addrspace(1) [[Y]], align 8443// OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[TMP2]], align 4444// OPT-NEXT: [[ADD:%.*]] = fadd contract float [[TMP3]], 1.000000e+00445// OPT-NEXT: store float [[ADD]], ptr [[TMP2]], align 4446// OPT-NEXT: ret void447//448// OPT-SPIRV-LABEL: define spir_kernel void @_Z7kernel5P1S(449// OPT-SPIRV-SAME: ptr addrspace(1) noundef [[S_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] {450// OPT-SPIRV-NEXT: [[ENTRY:.*:]]451// OPT-SPIRV-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[S_COERCE]] to i64452// OPT-SPIRV-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4)453// OPT-SPIRV-NEXT: [[TMP2:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[TMP1]], align 8454// OPT-SPIRV-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(4) [[TMP2]], align 4455// OPT-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP3]], 1456// OPT-SPIRV-NEXT: store i32 [[INC]], ptr addrspace(4) [[TMP2]], align 4457// OPT-SPIRV-NEXT: [[Y:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[TMP1]], i64 8458// OPT-SPIRV-NEXT: [[TMP4:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[Y]], align 8459// OPT-SPIRV-NEXT: [[TMP5:%.*]] = load float, ptr addrspace(4) [[TMP4]], align 4460// OPT-SPIRV-NEXT: [[ADD:%.*]] = fadd contract float [[TMP5]], 1.000000e+00461// OPT-SPIRV-NEXT: store float [[ADD]], ptr addrspace(4) [[TMP4]], align 4462// OPT-SPIRV-NEXT: ret void463//464// HOST-LABEL: define dso_local void @_Z22__device_stub__kernel5P1S(465// HOST-SAME: ptr noundef [[S:%.*]]) #[[ATTR0]] {466// HOST-NEXT: [[ENTRY:.*:]]467// HOST-NEXT: [[S_ADDR:%.*]] = alloca ptr, align 8468// HOST-NEXT: store ptr [[S]], ptr [[S_ADDR]], align 8469// HOST-NEXT: [[TMP0:%.*]] = call i32 @hipSetupArgument(ptr [[S_ADDR]], i64 8, i64 0)470// HOST-NEXT: [[TMP1:%.*]] = icmp eq i32 [[TMP0]], 0471// HOST-NEXT: br i1 [[TMP1]], label %[[SETUP_NEXT:.*]], label %[[SETUP_END:.*]]472// HOST: [[SETUP_NEXT]]:473// HOST-NEXT: [[TMP2:%.*]] = call i32 @hipLaunchByPtr(ptr @_Z7kernel5P1S)474// HOST-NEXT: br label %[[SETUP_END]]475// HOST: [[SETUP_END]]:476// HOST-NEXT: ret void477//478__global__ void kernel5(struct S *s) {479 s->x[0]++;480 s->y[0] += 1.f;481}482 483struct T {484 float *x[2];485};486// `by-val` array is passed by-indirect-alias (a mix of by-ref and indirect487// by-val). However, the enhanced address inferring pass should be able to488// assume they are global pointers.489// For SPIR-V, since byref is not supported at the moment, we pass it as direct.490 491// CHECK-LABEL: define dso_local amdgpu_kernel void @_Z7kernel61T(492// CHECK-SAME: ptr addrspace(4) noundef byref([[STRUCT_T:%.*]]) align 8 [[TMP0:%.*]]) #[[ATTR0]] {493// CHECK-NEXT: [[ENTRY:.*:]]494// CHECK-NEXT: [[COERCE:%.*]] = alloca [[STRUCT_T]], align 8, addrspace(5)495// CHECK-NEXT: [[T:%.*]] = addrspacecast ptr addrspace(5) [[COERCE]] to ptr496// CHECK-NEXT: call void @llvm.memcpy.p0.p4.i64(ptr align 8 [[T]], ptr addrspace(4) align 8 [[TMP0]], i64 16, i1 false)497// CHECK-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_T]], ptr [[T]], i32 0, i32 0498// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x ptr], ptr [[X]], i64 0, i64 0499// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[ARRAYIDX]], align 8500// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds float, ptr [[TMP1]], i64 0501// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr [[ARRAYIDX1]], align 4502// CHECK-NEXT: [[ADD:%.*]] = fadd contract float [[TMP2]], 1.000000e+00503// CHECK-NEXT: store float [[ADD]], ptr [[ARRAYIDX1]], align 4504// CHECK-NEXT: [[X2:%.*]] = getelementptr inbounds nuw [[STRUCT_T]], ptr [[T]], i32 0, i32 0505// CHECK-NEXT: [[ARRAYIDX3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[X2]], i64 0, i64 1506// CHECK-NEXT: [[TMP3:%.*]] = load ptr, ptr [[ARRAYIDX3]], align 8507// CHECK-NEXT: [[ARRAYIDX4:%.*]] = getelementptr inbounds float, ptr [[TMP3]], i64 0508// CHECK-NEXT: [[TMP4:%.*]] = load float, ptr [[ARRAYIDX4]], align 4509// CHECK-NEXT: [[ADD5:%.*]] = fadd contract float [[TMP4]], 2.000000e+00510// CHECK-NEXT: store float [[ADD5]], ptr [[ARRAYIDX4]], align 4511// CHECK-NEXT: ret void512//513// CHECK-SPIRV-LABEL: define spir_kernel void @_Z7kernel61T(514// CHECK-SPIRV-SAME: [[STRUCT_T:%.*]] [[T_COERCE:%.*]]) addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] {515// CHECK-SPIRV-NEXT: [[ENTRY:.*:]]516// CHECK-SPIRV-NEXT: [[T:%.*]] = alloca [[STRUCT_T]], align 8517// CHECK-SPIRV-NEXT: [[T1:%.*]] = addrspacecast ptr [[T]] to ptr addrspace(4)518// CHECK-SPIRV-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT_T]], ptr addrspace(4) [[T1]], i32 0, i32 0519// CHECK-SPIRV-NEXT: [[TMP1:%.*]] = extractvalue [[STRUCT_T]] [[T_COERCE]], 0520// CHECK-SPIRV-NEXT: store [2 x ptr addrspace(4)] [[TMP1]], ptr addrspace(4) [[TMP0]], align 8521// CHECK-SPIRV-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_T]], ptr addrspace(4) [[T1]], i32 0, i32 0522// CHECK-SPIRV-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x ptr addrspace(4)], ptr addrspace(4) [[X]], i64 0, i64 0523// CHECK-SPIRV-NEXT: [[TMP2:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[ARRAYIDX]], align 8524// CHECK-SPIRV-NEXT: [[ARRAYIDX2:%.*]] = getelementptr inbounds float, ptr addrspace(4) [[TMP2]], i64 0525// CHECK-SPIRV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(4) [[ARRAYIDX2]], align 4526// CHECK-SPIRV-NEXT: [[ADD:%.*]] = fadd contract float [[TMP3]], 1.000000e+00527// CHECK-SPIRV-NEXT: store float [[ADD]], ptr addrspace(4) [[ARRAYIDX2]], align 4528// CHECK-SPIRV-NEXT: [[X3:%.*]] = getelementptr inbounds nuw [[STRUCT_T]], ptr addrspace(4) [[T1]], i32 0, i32 0529// CHECK-SPIRV-NEXT: [[ARRAYIDX4:%.*]] = getelementptr inbounds [2 x ptr addrspace(4)], ptr addrspace(4) [[X3]], i64 0, i64 1530// CHECK-SPIRV-NEXT: [[TMP4:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[ARRAYIDX4]], align 8531// CHECK-SPIRV-NEXT: [[ARRAYIDX5:%.*]] = getelementptr inbounds float, ptr addrspace(4) [[TMP4]], i64 0532// CHECK-SPIRV-NEXT: [[TMP5:%.*]] = load float, ptr addrspace(4) [[ARRAYIDX5]], align 4533// CHECK-SPIRV-NEXT: [[ADD6:%.*]] = fadd contract float [[TMP5]], 2.000000e+00534// CHECK-SPIRV-NEXT: store float [[ADD6]], ptr addrspace(4) [[ARRAYIDX5]], align 4535// CHECK-SPIRV-NEXT: ret void536//537// OPT-LABEL: define dso_local amdgpu_kernel void @_Z7kernel61T(538// OPT-SAME: ptr addrspace(4) noundef readonly byref([[STRUCT_T:%.*]]) align 8 captures(none) [[TMP0:%.*]]) local_unnamed_addr #[[ATTR2]] {539// OPT-NEXT: [[ENTRY:.*:]]540// OPT-NEXT: [[COERCE_SROA_0_0_COPYLOAD:%.*]] = load ptr, ptr addrspace(4) [[TMP0]], align 8, !amdgpu.noclobber [[META4]]541// OPT-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[COERCE_SROA_0_0_COPYLOAD]] to ptr addrspace(1)542// OPT-NEXT: [[COERCE_SROA_2_0__SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[TMP0]], i64 8543// OPT-NEXT: [[COERCE_SROA_2_0_COPYLOAD:%.*]] = load ptr, ptr addrspace(4) [[COERCE_SROA_2_0__SROA_IDX]], align 8, !amdgpu.noclobber [[META4]]544// OPT-NEXT: [[TMP2:%.*]] = addrspacecast ptr [[COERCE_SROA_2_0_COPYLOAD]] to ptr addrspace(1)545// OPT-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(1) [[TMP1]], align 4, !amdgpu.noclobber [[META4]]546// OPT-NEXT: [[ADD:%.*]] = fadd contract float [[TMP3]], 1.000000e+00547// OPT-NEXT: store float [[ADD]], ptr addrspace(1) [[TMP1]], align 4548// OPT-NEXT: [[TMP4:%.*]] = load float, ptr addrspace(1) [[TMP2]], align 4549// OPT-NEXT: [[ADD5:%.*]] = fadd contract float [[TMP4]], 2.000000e+00550// OPT-NEXT: store float [[ADD5]], ptr addrspace(1) [[TMP2]], align 4551// OPT-NEXT: ret void552//553// OPT-SPIRV-LABEL: define spir_kernel void @_Z7kernel61T(554// OPT-SPIRV-SAME: [[STRUCT_T:%.*]] [[T_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] {555// OPT-SPIRV-NEXT: [[ENTRY:.*:]]556// OPT-SPIRV-NEXT: [[TMP0:%.*]] = extractvalue [[STRUCT_T]] [[T_COERCE]], 0557// OPT-SPIRV-NEXT: [[DOTFCA_0_EXTRACT:%.*]] = extractvalue [2 x ptr addrspace(4)] [[TMP0]], 0558// OPT-SPIRV-NEXT: [[DOTFCA_1_EXTRACT:%.*]] = extractvalue [2 x ptr addrspace(4)] [[TMP0]], 1559// OPT-SPIRV-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(4) [[DOTFCA_0_EXTRACT]], align 4560// OPT-SPIRV-NEXT: [[ADD:%.*]] = fadd contract float [[TMP1]], 1.000000e+00561// OPT-SPIRV-NEXT: store float [[ADD]], ptr addrspace(4) [[DOTFCA_0_EXTRACT]], align 4562// OPT-SPIRV-NEXT: [[TMP2:%.*]] = load float, ptr addrspace(4) [[DOTFCA_1_EXTRACT]], align 4563// OPT-SPIRV-NEXT: [[ADD6:%.*]] = fadd contract float [[TMP2]], 2.000000e+00564// OPT-SPIRV-NEXT: store float [[ADD6]], ptr addrspace(4) [[DOTFCA_1_EXTRACT]], align 4565// OPT-SPIRV-NEXT: ret void566//567// HOST-LABEL: define dso_local void @_Z22__device_stub__kernel61T(568// HOST-SAME: ptr [[T_COERCE0:%.*]], ptr [[T_COERCE1:%.*]]) #[[ATTR0]] {569// HOST-NEXT: [[ENTRY:.*:]]570// HOST-NEXT: [[T:%.*]] = alloca [[STRUCT_T:%.*]], align 8571// HOST-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw { ptr, ptr }, ptr [[T]], i32 0, i32 0572// HOST-NEXT: store ptr [[T_COERCE0]], ptr [[TMP0]], align 8573// HOST-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw { ptr, ptr }, ptr [[T]], i32 0, i32 1574// HOST-NEXT: store ptr [[T_COERCE1]], ptr [[TMP1]], align 8575// HOST-NEXT: [[TMP2:%.*]] = call i32 @hipSetupArgument(ptr [[T]], i64 16, i64 0)576// HOST-NEXT: [[TMP3:%.*]] = icmp eq i32 [[TMP2]], 0577// HOST-NEXT: br i1 [[TMP3]], label %[[SETUP_NEXT:.*]], label %[[SETUP_END:.*]]578// HOST: [[SETUP_NEXT]]:579// HOST-NEXT: [[TMP4:%.*]] = call i32 @hipLaunchByPtr(ptr @_Z7kernel61T)580// HOST-NEXT: br label %[[SETUP_END]]581// HOST: [[SETUP_END]]:582// HOST-NEXT: ret void583//584__global__ void kernel6(struct T t) {585 t.x[0][0] += 1.f;586 t.x[1][0] += 2.f;587}588 589// Check that coerced pointers retain the noalias attribute when qualified with __restrict.590 591// CHECK-LABEL: define dso_local amdgpu_kernel void @_Z7kernel7Pi(592// CHECK-SAME: ptr addrspace(1) noalias noundef [[X_COERCE:%.*]]) #[[ATTR0]] {593// CHECK-NEXT: [[ENTRY:.*:]]594// CHECK-NEXT: [[X:%.*]] = alloca ptr, align 8, addrspace(5)595// CHECK-NEXT: [[X_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)596// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr597// CHECK-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr598// CHECK-NEXT: store ptr addrspace(1) [[X_COERCE]], ptr [[X_ASCAST]], align 8599// CHECK-NEXT: [[X1:%.*]] = load ptr, ptr [[X_ASCAST]], align 8600// CHECK-NEXT: store ptr [[X1]], ptr [[X_ADDR_ASCAST]], align 8601// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[X_ADDR_ASCAST]], align 8602// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 0603// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX]], align 4604// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1605// CHECK-NEXT: store i32 [[INC]], ptr [[ARRAYIDX]], align 4606// CHECK-NEXT: ret void607//608// CHECK-SPIRV-LABEL: define spir_kernel void @_Z7kernel7Pi(609// CHECK-SPIRV-SAME: ptr addrspace(1) noalias noundef [[X_COERCE:%.*]]) addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] {610// CHECK-SPIRV-NEXT: [[ENTRY:.*:]]611// CHECK-SPIRV-NEXT: [[X:%.*]] = alloca ptr addrspace(4), align 8612// CHECK-SPIRV-NEXT: [[X_ADDR:%.*]] = alloca ptr addrspace(4), align 8613// CHECK-SPIRV-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr [[X]] to ptr addrspace(4)614// CHECK-SPIRV-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr [[X_ADDR]] to ptr addrspace(4)615// CHECK-SPIRV-NEXT: store ptr addrspace(1) [[X_COERCE]], ptr addrspace(4) [[X_ASCAST]], align 8616// CHECK-SPIRV-NEXT: [[X1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[X_ASCAST]], align 8617// CHECK-SPIRV-NEXT: store ptr addrspace(4) [[X1]], ptr addrspace(4) [[X_ADDR_ASCAST]], align 8618// CHECK-SPIRV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[X_ADDR_ASCAST]], align 8619// CHECK-SPIRV-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr addrspace(4) [[TMP0]], i64 0620// CHECK-SPIRV-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX]], align 4621// CHECK-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1622// CHECK-SPIRV-NEXT: store i32 [[INC]], ptr addrspace(4) [[ARRAYIDX]], align 4623// CHECK-SPIRV-NEXT: ret void624//625// OPT-LABEL: define dso_local amdgpu_kernel void @_Z7kernel7Pi(626// OPT-SAME: ptr addrspace(1) noalias noundef captures(none) [[X_COERCE:%.*]]) local_unnamed_addr #[[ATTR0]] {627// OPT-NEXT: [[ENTRY:.*:]]628// OPT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(1) [[X_COERCE]], align 4629// OPT-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1630// OPT-NEXT: store i32 [[INC]], ptr addrspace(1) [[X_COERCE]], align 4631// OPT-NEXT: ret void632//633// OPT-SPIRV-LABEL: define spir_kernel void @_Z7kernel7Pi(634// OPT-SPIRV-SAME: ptr addrspace(1) noalias noundef [[X_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] {635// OPT-SPIRV-NEXT: [[ENTRY:.*:]]636// OPT-SPIRV-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[X_COERCE]] to i64637// OPT-SPIRV-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4)638// OPT-SPIRV-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 4639// OPT-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP2]], 1640// OPT-SPIRV-NEXT: store i32 [[INC]], ptr addrspace(4) [[TMP1]], align 4641// OPT-SPIRV-NEXT: ret void642//643// HOST-LABEL: define dso_local void @_Z22__device_stub__kernel7Pi(644// HOST-SAME: ptr noalias noundef [[X:%.*]]) #[[ATTR0]] {645// HOST-NEXT: [[ENTRY:.*:]]646// HOST-NEXT: [[X_ADDR:%.*]] = alloca ptr, align 8647// HOST-NEXT: store ptr [[X]], ptr [[X_ADDR]], align 8648// HOST-NEXT: [[TMP0:%.*]] = call i32 @hipSetupArgument(ptr [[X_ADDR]], i64 8, i64 0)649// HOST-NEXT: [[TMP1:%.*]] = icmp eq i32 [[TMP0]], 0650// HOST-NEXT: br i1 [[TMP1]], label %[[SETUP_NEXT:.*]], label %[[SETUP_END:.*]]651// HOST: [[SETUP_NEXT]]:652// HOST-NEXT: [[TMP2:%.*]] = call i32 @hipLaunchByPtr(ptr @_Z7kernel7Pi)653// HOST-NEXT: br label %[[SETUP_END]]654// HOST: [[SETUP_END]]:655// HOST-NEXT: ret void656//657__global__ void kernel7(int *__restrict x) {658 x[0]++;659}660 661// Single element struct.662struct SS {663 float *x;664};665// CHECK-LABEL: define dso_local amdgpu_kernel void @_Z7kernel82SS(666// CHECK-SAME: ptr addrspace(1) [[A_COERCE:%.*]]) #[[ATTR0]] {667// CHECK-NEXT: [[ENTRY:.*:]]668// CHECK-NEXT: [[A:%.*]] = alloca [[STRUCT_SS:%.*]], align 8, addrspace(5)669// CHECK-NEXT: [[A1:%.*]] = addrspacecast ptr addrspace(5) [[A]] to ptr670// CHECK-NEXT: [[COERCE_DIVE:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[A1]], i32 0, i32 0671// CHECK-NEXT: store ptr addrspace(1) [[A_COERCE]], ptr [[COERCE_DIVE]], align 8672// CHECK-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[A1]], i32 0, i32 0673// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[X]], align 8674// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[TMP0]], align 4675// CHECK-NEXT: [[ADD:%.*]] = fadd contract float [[TMP1]], 3.000000e+00676// CHECK-NEXT: store float [[ADD]], ptr [[TMP0]], align 4677// CHECK-NEXT: ret void678//679// CHECK-SPIRV-LABEL: define spir_kernel void @_Z7kernel82SS(680// CHECK-SPIRV-SAME: [[STRUCT_SS:%.*]] [[A_COERCE:%.*]]) addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] {681// CHECK-SPIRV-NEXT: [[ENTRY:.*:]]682// CHECK-SPIRV-NEXT: [[A:%.*]] = alloca [[STRUCT_SS]], align 8683// CHECK-SPIRV-NEXT: [[A1:%.*]] = addrspacecast ptr [[A]] to ptr addrspace(4)684// CHECK-SPIRV-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr addrspace(4) [[A1]], i32 0, i32 0685// CHECK-SPIRV-NEXT: [[TMP1:%.*]] = extractvalue [[STRUCT_SS]] [[A_COERCE]], 0686// CHECK-SPIRV-NEXT: store ptr addrspace(4) [[TMP1]], ptr addrspace(4) [[TMP0]], align 8687// CHECK-SPIRV-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr addrspace(4) [[A1]], i32 0, i32 0688// CHECK-SPIRV-NEXT: [[TMP2:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[X]], align 8689// CHECK-SPIRV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(4) [[TMP2]], align 4690// CHECK-SPIRV-NEXT: [[ADD:%.*]] = fadd contract float [[TMP3]], 3.000000e+00691// CHECK-SPIRV-NEXT: store float [[ADD]], ptr addrspace(4) [[TMP2]], align 4692// CHECK-SPIRV-NEXT: ret void693//694// OPT-LABEL: define dso_local amdgpu_kernel void @_Z7kernel82SS(695// OPT-SAME: ptr addrspace(1) captures(none) [[A_COERCE:%.*]]) local_unnamed_addr #[[ATTR0]] {696// OPT-NEXT: [[ENTRY:.*:]]697// OPT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(1) [[A_COERCE]], align 4698// OPT-NEXT: [[ADD:%.*]] = fadd contract float [[TMP0]], 3.000000e+00699// OPT-NEXT: store float [[ADD]], ptr addrspace(1) [[A_COERCE]], align 4700// OPT-NEXT: ret void701//702// OPT-SPIRV-LABEL: define spir_kernel void @_Z7kernel82SS(703// OPT-SPIRV-SAME: [[STRUCT_SS:%.*]] [[A_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] {704// OPT-SPIRV-NEXT: [[ENTRY:.*:]]705// OPT-SPIRV-NEXT: [[TMP0:%.*]] = extractvalue [[STRUCT_SS]] [[A_COERCE]], 0706// OPT-SPIRV-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(4) [[TMP0]], align 4707// OPT-SPIRV-NEXT: [[ADD:%.*]] = fadd contract float [[TMP1]], 3.000000e+00708// OPT-SPIRV-NEXT: store float [[ADD]], ptr addrspace(4) [[TMP0]], align 4709// OPT-SPIRV-NEXT: ret void710//711// HOST-LABEL: define dso_local void @_Z22__device_stub__kernel82SS(712// HOST-SAME: ptr [[A_COERCE:%.*]]) #[[ATTR0]] {713// HOST-NEXT: [[ENTRY:.*:]]714// HOST-NEXT: [[A:%.*]] = alloca [[STRUCT_SS:%.*]], align 8715// HOST-NEXT: [[COERCE_DIVE:%.*]] = getelementptr inbounds nuw [[STRUCT_SS]], ptr [[A]], i32 0, i32 0716// HOST-NEXT: store ptr [[A_COERCE]], ptr [[COERCE_DIVE]], align 8717// HOST-NEXT: [[TMP0:%.*]] = call i32 @hipSetupArgument(ptr [[A]], i64 8, i64 0)718// HOST-NEXT: [[TMP1:%.*]] = icmp eq i32 [[TMP0]], 0719// HOST-NEXT: br i1 [[TMP1]], label %[[SETUP_NEXT:.*]], label %[[SETUP_END:.*]]720// HOST: [[SETUP_NEXT]]:721// HOST-NEXT: [[TMP2:%.*]] = call i32 @hipLaunchByPtr(ptr @_Z7kernel82SS)722// HOST-NEXT: br label %[[SETUP_END]]723// HOST: [[SETUP_END]]:724// HOST-NEXT: ret void725//726__global__ void kernel8(struct SS a) {727 *a.x += 3.f;728}729//.730// CHECK-SPIRV: [[META5]] = !{i32 1024, i32 1, i32 1}731//.732// OPT: [[META4]] = !{}733//.734// OPT-SPIRV: [[META5]] = !{i32 1024, i32 1, i32 1}735//.736