llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll

   1 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
   2 ; RUN: opt < %s -S -nvptx-lower-args --mtriple nvptx64-nvidia-cuda -mcpu=sm_70 -mattr=+ptx77 | FileCheck %s --check-prefixes OPT
   3 ; RUN: llc < %s --mtriple nvptx64-nvidia-cuda -mcpu=sm_70 -mattr=+ptx77 | FileCheck %s --check-prefixes PTX
   4
   5 %struct.uint4 = type { i32, i32, i32, i32 }
   6
   7 @gi = dso_local addrspace(1) externally_initialized global %struct.uint4 { i32 50462976, i32 117835012, i32 185207048, i32 252579084 }, align 16
   8
   9 ; Function Attrs: mustprogress nofree noinline norecurse nosync nounwind willreturn memory(read, inaccessiblemem: none)
  10 ; Regular functions mus still make a copy. `cvta.param` does not always work there.
  11 define dso_local noundef i32 @non_kernel_function(ptr nocapture noundef readonly byval(%struct.uint4) align 16 %a, i1 noundef zeroext %b, i32 noundef %c) local_unnamed_addr #0 {
  12 ; OPT-LABEL: define dso_local noundef i32 @non_kernel_function(
  13 ; OPT-SAME: ptr nocapture noundef readonly byval([[STRUCT_UINT4:%.*]]) align 16 [[A:%.*]], i1 noundef zeroext [[B:%.*]], i32 noundef [[C:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
  14 ; OPT-NEXT:  [[ENTRY:.*:]]
  15 ; OPT-NEXT:    [[A1:%.*]] = alloca [[STRUCT_UINT4]], align 16
  16 ; OPT-NEXT:    [[A2:%.*]] = addrspacecast ptr [[A]] to ptr addrspace(101)
  17 ; OPT-NEXT:    call void @llvm.memcpy.p0.p101.i64(ptr align 16 [[A1]], ptr addrspace(101) align 16 [[A2]], i64 16, i1 false)
  18 ; OPT-NEXT:    [[A_:%.*]] = select i1 [[B]], ptr [[A1]], ptr addrspacecast (ptr addrspace(1) @gi to ptr)
  19 ; OPT-NEXT:    [[IDX_EXT:%.*]] = sext i32 [[C]] to i64
  20 ; OPT-NEXT:    [[ADD_PTR:%.*]] = getelementptr inbounds i8, ptr [[A_]], i64 [[IDX_EXT]]
  21 ; OPT-NEXT:    [[TMP0:%.*]] = load i32, ptr [[ADD_PTR]], align 1
  22 ; OPT-NEXT:    ret i32 [[TMP0]]
  23 ;
  24 ; PTX-LABEL: non_kernel_function(
  25 ; PTX:       {
  26 ; PTX-NEXT:    .local .align 16 .b8 __local_depot0[16];
  27 ; PTX-NEXT:    .reg .b64 %SP;
  28 ; PTX-NEXT:    .reg .b64 %SPL;
  29 ; PTX-NEXT:    .reg .pred %p<2>;
  30 ; PTX-NEXT:    .reg .b16 %rs<3>;
  31 ; PTX-NEXT:    .reg .b32 %r<11>;
  32 ; PTX-NEXT:    .reg .b64 %rd<10>;
  33 ; PTX-EMPTY:
  34 ; PTX-NEXT:  // %bb.0: // %entry
  35 ; PTX-NEXT:    mov.u64 %SPL, __local_depot0;
  36 ; PTX-NEXT:    cvta.local.u64 %SP, %SPL;
  37 ; PTX-NEXT:    ld.param.u8 %rs1, [non_kernel_function_param_1];
  38 ; PTX-NEXT:    and.b16 %rs2, %rs1, 1;
  39 ; PTX-NEXT:    setp.eq.b16 %p1, %rs2, 1;
  40 ; PTX-NEXT:    ld.param.s32 %rd1, [non_kernel_function_param_2];
  41 ; PTX-NEXT:    add.u64 %rd2, %SP, 0;
  42 ; PTX-NEXT:    or.b64 %rd3, %rd2, 8;
  43 ; PTX-NEXT:    ld.param.u64 %rd4, [non_kernel_function_param_0+8];
  44 ; PTX-NEXT:    st.u64 [%rd3], %rd4;
  45 ; PTX-NEXT:    ld.param.u64 %rd5, [non_kernel_function_param_0];
  46 ; PTX-NEXT:    st.u64 [%SP], %rd5;
  47 ; PTX-NEXT:    mov.u64 %rd6, gi;
  48 ; PTX-NEXT:    cvta.global.u64 %rd7, %rd6;
  49 ; PTX-NEXT:    selp.b64 %rd8, %rd2, %rd7, %p1;
  50 ; PTX-NEXT:    add.s64 %rd9, %rd8, %rd1;
  51 ; PTX-NEXT:    ld.u8 %r1, [%rd9];
  52 ; PTX-NEXT:    ld.u8 %r2, [%rd9+1];
  53 ; PTX-NEXT:    shl.b32 %r3, %r2, 8;
  54 ; PTX-NEXT:    or.b32 %r4, %r3, %r1;
  55 ; PTX-NEXT:    ld.u8 %r5, [%rd9+2];
  56 ; PTX-NEXT:    shl.b32 %r6, %r5, 16;
  57 ; PTX-NEXT:    ld.u8 %r7, [%rd9+3];
  58 ; PTX-NEXT:    shl.b32 %r8, %r7, 24;
  59 ; PTX-NEXT:    or.b32 %r9, %r8, %r6;
  60 ; PTX-NEXT:    or.b32 %r10, %r9, %r4;
  61 ; PTX-NEXT:    st.param.b32 [func_retval0], %r10;
  62 ; PTX-NEXT:    ret;
  63 entry:
  64   %a. = select i1 %b, ptr %a, ptr addrspacecast (ptr addrspace(1) @gi to ptr), !dbg !17
  65   %idx.ext = sext i32 %c to i64, !dbg !18
  66   %add.ptr = getelementptr inbounds i8, ptr %a., i64 %idx.ext, !dbg !18
  67   %0 = load i32, ptr %add.ptr, align 1, !dbg !19
  68   ret i32 %0, !dbg !23
  69 }
  70
  71 define void @grid_const_int(ptr byval(i32) align 4 %input1, i32 %input2, ptr %out, i32 %n) {
  72 ; PTX-LABEL: grid_const_int(
  73 ; PTX:       {
  74 ; PTX-NEXT:    .reg .b32 %r<4>;
  75 ; PTX-NEXT:    .reg .b64 %rd<3>;
  76 ; PTX-EMPTY:
  77 ; PTX-NEXT:  // %bb.0:
  78 ; PTX-NEXT:    ld.param.u64 %rd1, [grid_const_int_param_2];
  79 ; PTX-NEXT:    cvta.to.global.u64 %rd2, %rd1;
  80 ; PTX-NEXT:    ld.param.u32 %r1, [grid_const_int_param_1];
  81 ; PTX-NEXT:    ld.param.u32 %r2, [grid_const_int_param_0];
  82 ; PTX-NEXT:    add.s32 %r3, %r2, %r1;
  83 ; PTX-NEXT:    st.global.u32 [%rd2], %r3;
  84 ; PTX-NEXT:    ret;
  85 ; OPT-LABEL: define void @grid_const_int(
  86 ; OPT-SAME: ptr byval(i32) align 4 [[INPUT1:%.*]], i32 [[INPUT2:%.*]], ptr [[OUT:%.*]], i32 [[N:%.*]]) #[[ATTR0]] {
  87 ; OPT-NEXT:    [[OUT2:%.*]] = addrspacecast ptr [[OUT]] to ptr addrspace(1)
  88 ; OPT-NEXT:    [[OUT3:%.*]] = addrspacecast ptr addrspace(1) [[OUT2]] to ptr
  89 ; OPT-NEXT:    [[INPUT11:%.*]] = addrspacecast ptr [[INPUT1]] to ptr addrspace(101)
  90 ; OPT-NEXT:    [[TMP:%.*]] = load i32, ptr addrspace(101) [[INPUT11]], align 4
  91 ; OPT-NEXT:    [[ADD:%.*]] = add i32 [[TMP]], [[INPUT2]]
  92 ; OPT-NEXT:    store i32 [[ADD]], ptr [[OUT3]], align 4
  93 ; OPT-NEXT:    ret void
  94   %tmp = load i32, ptr %input1, align 4
  95   %add = add i32 %tmp, %input2
  96   store i32 %add, ptr %out
  97   ret void
  98 }
  99
 100 %struct.s = type { i32, i32 }
 101
 102 define void @grid_const_struct(ptr byval(%struct.s) align 4 %input, ptr %out){
 103 ; PTX-LABEL: grid_const_struct(
 104 ; PTX:       {
 105 ; PTX-NEXT:    .reg .b32 %r<4>;
 106 ; PTX-NEXT:    .reg .b64 %rd<3>;
 107 ; PTX-EMPTY:
 108 ; PTX-NEXT:  // %bb.0:
 109 ; PTX-NEXT:    ld.param.u64 %rd1, [grid_const_struct_param_1];
 110 ; PTX-NEXT:    cvta.to.global.u64 %rd2, %rd1;
 111 ; PTX-NEXT:    ld.param.u32 %r1, [grid_const_struct_param_0];
 112 ; PTX-NEXT:    ld.param.u32 %r2, [grid_const_struct_param_0+4];
 113 ; PTX-NEXT:    add.s32 %r3, %r1, %r2;
 114 ; PTX-NEXT:    st.global.u32 [%rd2], %r3;
 115 ; PTX-NEXT:    ret;
 116 ; OPT-LABEL: define void @grid_const_struct(
 117 ; OPT-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 [[INPUT:%.*]], ptr [[OUT:%.*]]) #[[ATTR0]] {
 118 ; OPT-NEXT:    [[OUT4:%.*]] = addrspacecast ptr [[OUT]] to ptr addrspace(1)
 119 ; OPT-NEXT:    [[OUT5:%.*]] = addrspacecast ptr addrspace(1) [[OUT4]] to ptr
 120 ; OPT-NEXT:    [[INPUT1:%.*]] = addrspacecast ptr [[INPUT]] to ptr addrspace(101)
 121 ; OPT-NEXT:    [[GEP13:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr addrspace(101) [[INPUT1]], i32 0, i32 0
 122 ; OPT-NEXT:    [[GEP22:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr addrspace(101) [[INPUT1]], i32 0, i32 1
 123 ; OPT-NEXT:    [[TMP1:%.*]] = load i32, ptr addrspace(101) [[GEP13]], align 4
 124 ; OPT-NEXT:    [[TMP2:%.*]] = load i32, ptr addrspace(101) [[GEP22]], align 4
 125 ; OPT-NEXT:    [[ADD:%.*]] = add i32 [[TMP1]], [[TMP2]]
 126 ; OPT-NEXT:    store i32 [[ADD]], ptr [[OUT5]], align 4
 127 ; OPT-NEXT:    ret void
 128   %gep1 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 0
 129   %gep2 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 1
 130   %int1 = load i32, ptr %gep1
 131   %int2 = load i32, ptr %gep2
 132   %add = add i32 %int1, %int2
 133   store i32 %add, ptr %out
 134   ret void
 135 }
 136
 137 define void @grid_const_escape(ptr byval(%struct.s) align 4 %input) {
 138 ; PTX-LABEL: grid_const_escape(
 139 ; PTX:       {
 140 ; PTX-NEXT:    .reg .b32 %r<3>;
 141 ; PTX-NEXT:    .reg .b64 %rd<5>;
 142 ; PTX-EMPTY:
 143 ; PTX-NEXT:  // %bb.0:
 144 ; PTX-NEXT:    mov.b64 %rd2, grid_const_escape_param_0;
 145 ; PTX-NEXT:    mov.u64 %rd3, %rd2;
 146 ; PTX-NEXT:    cvta.param.u64 %rd4, %rd3;
 147 ; PTX-NEXT:    mov.u64 %rd1, escape;
 148 ; PTX-NEXT:    { // callseq 0, 0
 149 ; PTX-NEXT:    .param .b64 param0;
 150 ; PTX-NEXT:    st.param.b64 [param0], %rd4;
 151 ; PTX-NEXT:    .param .b32 retval0;
 152 ; PTX-NEXT:    prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _);
 153 ; PTX-NEXT:    call (retval0),
 154 ; PTX-NEXT:    %rd1,
 155 ; PTX-NEXT:    (
 156 ; PTX-NEXT:    param0
 157 ; PTX-NEXT:    )
 158 ; PTX-NEXT:    , prototype_0;
 159 ; PTX-NEXT:    ld.param.b32 %r1, [retval0];
 160 ; PTX-NEXT:    } // callseq 0
 161 ; PTX-NEXT:    ret;
 162 ; OPT-LABEL: define void @grid_const_escape(
 163 ; OPT-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 [[INPUT:%.*]]) #[[ATTR0]] {
 164 ; OPT-NEXT:    [[INPUT_PARAM:%.*]] = addrspacecast ptr [[INPUT]] to ptr addrspace(101)
 165 ; OPT-NEXT:    [[INPUT_PARAM_GEN:%.*]] = call ptr @llvm.nvvm.ptr.param.to.gen.p0.p101(ptr addrspace(101) [[INPUT_PARAM]])
 166 ; OPT-NEXT:    [[CALL:%.*]] = call i32 @escape(ptr [[INPUT_PARAM_GEN]])
 167 ; OPT-NEXT:    ret void
 168   %call = call i32 @escape(ptr %input)
 169   ret void
 170 }
 171
 172 define void @multiple_grid_const_escape(ptr byval(%struct.s) align 4 %input, i32 %a, ptr byval(i32) align 4 %b) {
 173 ; PTX-LABEL: multiple_grid_const_escape(
 174 ; PTX:       {
 175 ; PTX-NEXT:    .local .align 4 .b8 __local_depot4[4];
 176 ; PTX-NEXT:    .reg .b64 %SP;
 177 ; PTX-NEXT:    .reg .b64 %SPL;
 178 ; PTX-NEXT:    .reg .b32 %r<4>;
 179 ; PTX-NEXT:    .reg .b64 %rd<10>;
 180 ; PTX-EMPTY:
 181 ; PTX-NEXT:  // %bb.0:
 182 ; PTX-NEXT:    mov.u64 %SPL, __local_depot4;
 183 ; PTX-NEXT:    cvta.local.u64 %SP, %SPL;
 184 ; PTX-NEXT:    mov.b64 %rd2, multiple_grid_const_escape_param_0;
 185 ; PTX-NEXT:    mov.b64 %rd3, multiple_grid_const_escape_param_2;
 186 ; PTX-NEXT:    mov.u64 %rd4, %rd3;
 187 ; PTX-NEXT:    ld.param.u32 %r1, [multiple_grid_const_escape_param_1];
 188 ; PTX-NEXT:    cvta.param.u64 %rd5, %rd4;
 189 ; PTX-NEXT:    mov.u64 %rd6, %rd2;
 190 ; PTX-NEXT:    cvta.param.u64 %rd7, %rd6;
 191 ; PTX-NEXT:    add.u64 %rd8, %SP, 0;
 192 ; PTX-NEXT:    add.u64 %rd9, %SPL, 0;
 193 ; PTX-NEXT:    st.local.u32 [%rd9], %r1;
 194 ; PTX-NEXT:    mov.u64 %rd1, escape3;
 195 ; PTX-NEXT:    { // callseq 1, 0
 196 ; PTX-NEXT:    .param .b64 param0;
 197 ; PTX-NEXT:    st.param.b64 [param0], %rd7;
 198 ; PTX-NEXT:    .param .b64 param1;
 199 ; PTX-NEXT:    st.param.b64 [param1], %rd8;
 200 ; PTX-NEXT:    .param .b64 param2;
 201 ; PTX-NEXT:    st.param.b64 [param2], %rd5;
 202 ; PTX-NEXT:    .param .b32 retval0;
 203 ; PTX-NEXT:    prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
 204 ; PTX-NEXT:    call (retval0),
 205 ; PTX-NEXT:    %rd1,
 206 ; PTX-NEXT:    (
 207 ; PTX-NEXT:    param0,
 208 ; PTX-NEXT:    param1,
 209 ; PTX-NEXT:    param2
 210 ; PTX-NEXT:    )
 211 ; PTX-NEXT:    , prototype_1;
 212 ; PTX-NEXT:    ld.param.b32 %r2, [retval0];
 213 ; PTX-NEXT:    } // callseq 1
 214 ; PTX-NEXT:    ret;
 215 ; OPT-LABEL: define void @multiple_grid_const_escape(
 216 ; OPT-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 [[INPUT:%.*]], i32 [[A:%.*]], ptr byval(i32) align 4 [[B:%.*]]) #[[ATTR0]] {
 217 ; OPT-NEXT:    [[B_PARAM:%.*]] = addrspacecast ptr [[B]] to ptr addrspace(101)
 218 ; OPT-NEXT:    [[B_PARAM_GEN:%.*]] = call ptr @llvm.nvvm.ptr.param.to.gen.p0.p101(ptr addrspace(101) [[B_PARAM]])
 219 ; OPT-NEXT:    [[INPUT_PARAM:%.*]] = addrspacecast ptr [[INPUT]] to ptr addrspace(101)
 220 ; OPT-NEXT:    [[INPUT_PARAM_GEN:%.*]] = call ptr @llvm.nvvm.ptr.param.to.gen.p0.p101(ptr addrspace(101) [[INPUT_PARAM]])
 221 ; OPT-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4
 222 ; OPT-NEXT:    store i32 [[A]], ptr [[A_ADDR]], align 4
 223 ; OPT-NEXT:    [[CALL:%.*]] = call i32 @escape3(ptr [[INPUT_PARAM_GEN]], ptr [[A_ADDR]], ptr [[B_PARAM_GEN]])
 224 ; OPT-NEXT:    ret void
 225   %a.addr = alloca i32, align 4
 226   store i32 %a, ptr %a.addr, align 4
 227   %call = call i32 @escape3(ptr %input, ptr %a.addr, ptr %b)
 228   ret void
 229 }
 230
 231 define void @grid_const_memory_escape(ptr byval(%struct.s) align 4 %input, ptr %addr) {
 232 ; PTX-LABEL: grid_const_memory_escape(
 233 ; PTX:       {
 234 ; PTX-NEXT:    .reg .b64 %rd<6>;
 235 ; PTX-EMPTY:
 236 ; PTX-NEXT:  // %bb.0:
 237 ; PTX-NEXT:    mov.b64 %rd1, grid_const_memory_escape_param_0;
 238 ; PTX-NEXT:    ld.param.u64 %rd2, [grid_const_memory_escape_param_1];
 239 ; PTX-NEXT:    cvta.to.global.u64 %rd3, %rd2;
 240 ; PTX-NEXT:    mov.u64 %rd4, %rd1;
 241 ; PTX-NEXT:    cvta.param.u64 %rd5, %rd4;
 242 ; PTX-NEXT:    st.global.u64 [%rd3], %rd5;
 243 ; PTX-NEXT:    ret;
 244 ; OPT-LABEL: define void @grid_const_memory_escape(
 245 ; OPT-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 [[INPUT:%.*]], ptr [[ADDR:%.*]]) #[[ATTR0]] {
 246 ; OPT-NEXT:    [[ADDR4:%.*]] = addrspacecast ptr [[ADDR]] to ptr addrspace(1)
 247 ; OPT-NEXT:    [[ADDR5:%.*]] = addrspacecast ptr addrspace(1) [[ADDR4]] to ptr
 248 ; OPT-NEXT:    [[INPUT_PARAM:%.*]] = addrspacecast ptr [[INPUT]] to ptr addrspace(101)
 249 ; OPT-NEXT:    [[INPUT1:%.*]] = call ptr @llvm.nvvm.ptr.param.to.gen.p0.p101(ptr addrspace(101) [[INPUT_PARAM]])
 250 ; OPT-NEXT:    store ptr [[INPUT1]], ptr [[ADDR5]], align 8
 251 ; OPT-NEXT:    ret void
 252   store ptr %input, ptr %addr, align 8
 253   ret void
 254 }
 255
 256 define void @grid_const_inlineasm_escape(ptr byval(%struct.s) align 4 %input, ptr %result) {
 257 ; PTX-LABEL: grid_const_inlineasm_escape(
 258 ; PTX:       {
 259 ; PTX-NEXT:    .reg .b64 %rd<8>;
 260 ; PTX-EMPTY:
 261 ; PTX-NEXT:  // %bb.0:
 262 ; PTX-NEXT:    mov.b64 %rd4, grid_const_inlineasm_escape_param_0;
 263 ; PTX-NEXT:    ld.param.u64 %rd5, [grid_const_inlineasm_escape_param_1];
 264 ; PTX-NEXT:    cvta.to.global.u64 %rd6, %rd5;
 265 ; PTX-NEXT:    mov.u64 %rd7, %rd4;
 266 ; PTX-NEXT:    cvta.param.u64 %rd2, %rd7;
 267 ; PTX-NEXT:    add.s64 %rd3, %rd2, 4;
 268 ; PTX-NEXT:    // begin inline asm
 269 ; PTX-NEXT:    add.s64 %rd1, %rd2, %rd3;
 270 ; PTX-NEXT:    // end inline asm
 271 ; PTX-NEXT:    st.global.u64 [%rd6], %rd1;
 272 ; PTX-NEXT:    ret;
 273 ; PTX-NOT      .local
 274 ; OPT-LABEL: define void @grid_const_inlineasm_escape(
 275 ; OPT-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 [[INPUT:%.*]], ptr [[RESULT:%.*]]) #[[ATTR0]] {
 276 ; OPT-NEXT:    [[RESULT4:%.*]] = addrspacecast ptr [[RESULT]] to ptr addrspace(1)
 277 ; OPT-NEXT:    [[RESULT5:%.*]] = addrspacecast ptr addrspace(1) [[RESULT4]] to ptr
 278 ; OPT-NEXT:    [[INPUT_PARAM:%.*]] = addrspacecast ptr [[INPUT]] to ptr addrspace(101)
 279 ; OPT-NEXT:    [[INPUT1:%.*]] = call ptr @llvm.nvvm.ptr.param.to.gen.p0.p101(ptr addrspace(101) [[INPUT_PARAM]])
 280 ; OPT-NEXT:    [[TMPPTR1:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT1]], i32 0, i32 0
 281 ; OPT-NEXT:    [[TMPPTR2:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT1]], i32 0, i32 1
 282 ; OPT-NEXT:    [[TMP2:%.*]] = call i64 asm "add.s64 $0, $1, $2
 283 ; OPT-NEXT:    store i64 [[TMP2]], ptr [[RESULT5]], align 8
 284 ; OPT-NEXT:    ret void
 285   %tmpptr1 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 0
 286   %tmpptr2 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 1
 287   %1 = call i64 asm "add.s64 $0, $1, $2;", "=l,l,l"(ptr %tmpptr1, ptr %tmpptr2) #1
 288   store i64 %1, ptr %result, align 8
 289   ret void
 290 }
 291
 292 define void @grid_const_partial_escape(ptr byval(i32) %input, ptr %output) {
 293 ; PTX-LABEL: grid_const_partial_escape(
 294 ; PTX:       {
 295 ; PTX-NEXT:    .reg .b32 %r<5>;
 296 ; PTX-NEXT:    .reg .b64 %rd<7>;
 297 ; PTX-EMPTY:
 298 ; PTX-NEXT:  // %bb.0:
 299 ; PTX-NEXT:    mov.b64 %rd2, grid_const_partial_escape_param_0;
 300 ; PTX-NEXT:    ld.param.u64 %rd3, [grid_const_partial_escape_param_1];
 301 ; PTX-NEXT:    cvta.to.global.u64 %rd4, %rd3;
 302 ; PTX-NEXT:    mov.u64 %rd5, %rd2;
 303 ; PTX-NEXT:    cvta.param.u64 %rd6, %rd5;
 304 ; PTX-NEXT:    ld.u32 %r1, [%rd6];
 305 ; PTX-NEXT:    add.s32 %r2, %r1, %r1;
 306 ; PTX-NEXT:    st.global.u32 [%rd4], %r2;
 307 ; PTX-NEXT:    mov.u64 %rd1, escape;
 308 ; PTX-NEXT:    { // callseq 2, 0
 309 ; PTX-NEXT:    .param .b64 param0;
 310 ; PTX-NEXT:    st.param.b64 [param0], %rd6;
 311 ; PTX-NEXT:    .param .b32 retval0;
 312 ; PTX-NEXT:    prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _);
 313 ; PTX-NEXT:    call (retval0),
 314 ; PTX-NEXT:    %rd1,
 315 ; PTX-NEXT:    (
 316 ; PTX-NEXT:    param0
 317 ; PTX-NEXT:    )
 318 ; PTX-NEXT:    , prototype_2;
 319 ; PTX-NEXT:    ld.param.b32 %r3, [retval0];
 320 ; PTX-NEXT:    } // callseq 2
 321 ; PTX-NEXT:    ret;
 322 ; OPT-LABEL: define void @grid_const_partial_escape(
 323 ; OPT-SAME: ptr byval(i32) [[INPUT:%.*]], ptr [[OUTPUT:%.*]]) #[[ATTR0]] {
 324 ; OPT-NEXT:    [[OUTPUT4:%.*]] = addrspacecast ptr [[OUTPUT]] to ptr addrspace(1)
 325 ; OPT-NEXT:    [[OUTPUT5:%.*]] = addrspacecast ptr addrspace(1) [[OUTPUT4]] to ptr
 326 ; OPT-NEXT:    [[INPUT1:%.*]] = addrspacecast ptr [[INPUT]] to ptr addrspace(101)
 327 ; OPT-NEXT:    [[INPUT1_GEN:%.*]] = call ptr @llvm.nvvm.ptr.param.to.gen.p0.p101(ptr addrspace(101) [[INPUT1]])
 328 ; OPT-NEXT:    [[VAL1:%.*]] = load i32, ptr [[INPUT1_GEN]], align 4
 329 ; OPT-NEXT:    [[TWICE:%.*]] = add i32 [[VAL1]], [[VAL1]]
 330 ; OPT-NEXT:    store i32 [[TWICE]], ptr [[OUTPUT5]], align 4
 331 ; OPT-NEXT:    [[CALL:%.*]] = call i32 @escape(ptr [[INPUT1_GEN]])
 332 ; OPT-NEXT:    ret void
 333   %val = load i32, ptr %input
 334   %twice = add i32 %val, %val
 335   store i32 %twice, ptr %output
 336   %call = call i32 @escape(ptr %input)
 337   ret void
 338 }
 339
 340 define i32 @grid_const_partial_escapemem(ptr byval(%struct.s) %input, ptr %output) {
 341 ; PTX-LABEL: grid_const_partial_escapemem(
 342 ; PTX:       {
 343 ; PTX-NEXT:    .reg .b32 %r<6>;
 344 ; PTX-NEXT:    .reg .b64 %rd<7>;
 345 ; PTX-EMPTY:
 346 ; PTX-NEXT:  // %bb.0:
 347 ; PTX-NEXT:    mov.b64 %rd2, grid_const_partial_escapemem_param_0;
 348 ; PTX-NEXT:    ld.param.u64 %rd3, [grid_const_partial_escapemem_param_1];
 349 ; PTX-NEXT:    cvta.to.global.u64 %rd4, %rd3;
 350 ; PTX-NEXT:    mov.u64 %rd5, %rd2;
 351 ; PTX-NEXT:    cvta.param.u64 %rd6, %rd5;
 352 ; PTX-NEXT:    ld.u32 %r1, [%rd6];
 353 ; PTX-NEXT:    ld.u32 %r2, [%rd6+4];
 354 ; PTX-NEXT:    st.global.u64 [%rd4], %rd6;
 355 ; PTX-NEXT:    add.s32 %r3, %r1, %r2;
 356 ; PTX-NEXT:    mov.u64 %rd1, escape;
 357 ; PTX-NEXT:    { // callseq 3, 0
 358 ; PTX-NEXT:    .param .b64 param0;
 359 ; PTX-NEXT:    st.param.b64 [param0], %rd6;
 360 ; PTX-NEXT:    .param .b32 retval0;
 361 ; PTX-NEXT:    prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _);
 362 ; PTX-NEXT:    call (retval0),
 363 ; PTX-NEXT:    %rd1,
 364 ; PTX-NEXT:    (
 365 ; PTX-NEXT:    param0
 366 ; PTX-NEXT:    )
 367 ; PTX-NEXT:    , prototype_3;
 368 ; PTX-NEXT:    ld.param.b32 %r4, [retval0];
 369 ; PTX-NEXT:    } // callseq 3
 370 ; PTX-NEXT:    st.param.b32 [func_retval0], %r3;
 371 ; PTX-NEXT:    ret;
 372 ; OPT-LABEL: define i32 @grid_const_partial_escapemem(
 373 ; OPT-SAME: ptr byval([[STRUCT_S:%.*]]) [[INPUT:%.*]], ptr [[OUTPUT:%.*]]) #[[ATTR0]] {
 374 ; OPT-NEXT:    [[OUTPUT4:%.*]] = addrspacecast ptr [[OUTPUT]] to ptr addrspace(1)
 375 ; OPT-NEXT:    [[OUTPUT5:%.*]] = addrspacecast ptr addrspace(1) [[OUTPUT4]] to ptr
 376 ; OPT-NEXT:    [[INPUT2:%.*]] = addrspacecast ptr [[INPUT]] to ptr addrspace(101)
 377 ; OPT-NEXT:    [[INPUT1:%.*]] = call ptr @llvm.nvvm.ptr.param.to.gen.p0.p101(ptr addrspace(101) [[INPUT2]])
 378 ; OPT-NEXT:    [[PTR1:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT1]], i32 0, i32 0
 379 ; OPT-NEXT:    [[VAL1:%.*]] = load i32, ptr [[PTR1]], align 4
 380 ; OPT-NEXT:    [[PTR2:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT1]], i32 0, i32 1
 381 ; OPT-NEXT:    [[VAL2:%.*]] = load i32, ptr [[PTR2]], align 4
 382 ; OPT-NEXT:    store ptr [[INPUT1]], ptr [[OUTPUT5]], align 8
 383 ; OPT-NEXT:    [[ADD:%.*]] = add i32 [[VAL1]], [[VAL2]]
 384 ; OPT-NEXT:    [[CALL2:%.*]] = call i32 @escape(ptr [[PTR1]])
 385 ; OPT-NEXT:    ret i32 [[ADD]]
 386   %ptr1 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 0
 387   %val1 = load i32, ptr %ptr1
 388   %ptr2 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 1
 389   %val2 = load i32, ptr %ptr2
 390   store ptr %input, ptr %output
 391   %add = add i32 %val1, %val2
 392   %call2 = call i32 @escape(ptr %ptr1)
 393   ret i32 %add
 394 }
 395
 396 define void @grid_const_phi(ptr byval(%struct.s) align 4 %input1, ptr %inout) {
 397 ; PTX-LABEL: grid_const_phi(
 398 ; PTX:       {
 399 ; PTX-NEXT:    .reg .pred %p<2>;
 400 ; PTX-NEXT:    .reg .b32 %r<3>;
 401 ; PTX-NEXT:    .reg .b64 %rd<9>;
 402 ; PTX-EMPTY:
 403 ; PTX-NEXT:  // %bb.0:
 404 ; PTX-NEXT:    mov.b64 %rd5, grid_const_phi_param_0;
 405 ; PTX-NEXT:    ld.param.u64 %rd6, [grid_const_phi_param_1];
 406 ; PTX-NEXT:    cvta.to.global.u64 %rd1, %rd6;
 407 ; PTX-NEXT:    mov.u64 %rd7, %rd5;
 408 ; PTX-NEXT:    cvta.param.u64 %rd8, %rd7;
 409 ; PTX-NEXT:    ld.global.u32 %r1, [%rd1];
 410 ; PTX-NEXT:    setp.lt.s32 %p1, %r1, 0;
 411 ; PTX-NEXT:    @%p1 bra $L__BB9_2;
 412 ; PTX-NEXT:  // %bb.1: // %second
 413 ; PTX-NEXT:    add.s64 %rd8, %rd8, 4;
 414 ; PTX-NEXT:  $L__BB9_2: // %merge
 415 ; PTX-NEXT:    ld.u32 %r2, [%rd8];
 416 ; PTX-NEXT:    st.global.u32 [%rd1], %r2;
 417 ; PTX-NEXT:    ret;
 418 ; OPT-LABEL: define void @grid_const_phi(
 419 ; OPT-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 [[INPUT1:%.*]], ptr [[INOUT:%.*]]) #[[ATTR0]] {
 420 ; OPT-NEXT:    [[INOUT1:%.*]] = addrspacecast ptr [[INOUT]] to ptr addrspace(1)
 421 ; OPT-NEXT:    [[INOUT2:%.*]] = addrspacecast ptr addrspace(1) [[INOUT1]] to ptr
 422 ; OPT-NEXT:    [[INPUT1_PARAM:%.*]] = addrspacecast ptr [[INPUT1]] to ptr addrspace(101)
 423 ; OPT-NEXT:    [[INPUT1_PARAM_GEN:%.*]] = call ptr @llvm.nvvm.ptr.param.to.gen.p0.p101(ptr addrspace(101) [[INPUT1_PARAM]])
 424 ; OPT-NEXT:    [[VAL:%.*]] = load i32, ptr [[INOUT2]], align 4
 425 ; OPT-NEXT:    [[LESS:%.*]] = icmp slt i32 [[VAL]], 0
 426 ; OPT-NEXT:    br i1 [[LESS]], label %[[FIRST:.*]], label %[[SECOND:.*]]
 427 ; OPT:       [[FIRST]]:
 428 ; OPT-NEXT:    [[PTR1:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT1_PARAM_GEN]], i32 0, i32 0
 429 ; OPT-NEXT:    br label %[[MERGE:.*]]
 430 ; OPT:       [[SECOND]]:
 431 ; OPT-NEXT:    [[PTR2:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT1_PARAM_GEN]], i32 0, i32 1
 432 ; OPT-NEXT:    br label %[[MERGE]]
 433 ; OPT:       [[MERGE]]:
 434 ; OPT-NEXT:    [[PTRNEW:%.*]] = phi ptr [ [[PTR1]], %[[FIRST]] ], [ [[PTR2]], %[[SECOND]] ]
 435 ; OPT-NEXT:    [[VALLOADED:%.*]] = load i32, ptr [[PTRNEW]], align 4
 436 ; OPT-NEXT:    store i32 [[VALLOADED]], ptr [[INOUT2]], align 4
 437 ; OPT-NEXT:    ret void
 438
 439   %val = load i32, ptr %inout
 440   %less = icmp slt i32 %val, 0
 441   br i1 %less, label %first, label %second
 442 first:
 443   %ptr1 = getelementptr inbounds %struct.s, ptr %input1, i32 0, i32 0
 444   br label %merge
 445 second:
 446   %ptr2 = getelementptr inbounds %struct.s, ptr %input1, i32 0, i32 1
 447   br label %merge
 448 merge:
 449   %ptrnew = phi ptr [%ptr1, %first], [%ptr2, %second]
 450   %valloaded = load i32, ptr %ptrnew
 451   store i32 %valloaded, ptr %inout
 452   ret void
 453 }
 454
 455 ; NOTE: %input2 is *not* grid_constant
 456 define void @grid_const_phi_ngc(ptr byval(%struct.s) align 4 %input1, ptr byval(%struct.s) %input2, ptr %inout) {
 457 ; PTX-LABEL: grid_const_phi_ngc(
 458 ; PTX:       {
 459 ; PTX-NEXT:    .reg .pred %p<2>;
 460 ; PTX-NEXT:    .reg .b32 %r<3>;
 461 ; PTX-NEXT:    .reg .b64 %rd<12>;
 462 ; PTX-EMPTY:
 463 ; PTX-NEXT:  // %bb.0:
 464 ; PTX-NEXT:    mov.b64 %rd6, grid_const_phi_ngc_param_0;
 465 ; PTX-NEXT:    ld.param.u64 %rd7, [grid_const_phi_ngc_param_2];
 466 ; PTX-NEXT:    cvta.to.global.u64 %rd1, %rd7;
 467 ; PTX-NEXT:    mov.u64 %rd10, %rd6;
 468 ; PTX-NEXT:    cvta.param.u64 %rd11, %rd10;
 469 ; PTX-NEXT:    ld.global.u32 %r1, [%rd1];
 470 ; PTX-NEXT:    setp.lt.s32 %p1, %r1, 0;
 471 ; PTX-NEXT:    @%p1 bra $L__BB10_2;
 472 ; PTX-NEXT:  // %bb.1: // %second
 473 ; PTX-NEXT:    mov.b64 %rd8, grid_const_phi_ngc_param_1;
 474 ; PTX-NEXT:    mov.u64 %rd9, %rd8;
 475 ; PTX-NEXT:    cvta.param.u64 %rd2, %rd9;
 476 ; PTX-NEXT:    add.s64 %rd11, %rd2, 4;
 477 ; PTX-NEXT:  $L__BB10_2: // %merge
 478 ; PTX-NEXT:    ld.u32 %r2, [%rd11];
 479 ; PTX-NEXT:    st.global.u32 [%rd1], %r2;
 480 ; PTX-NEXT:    ret;
 481 ; OPT-LABEL: define void @grid_const_phi_ngc(
 482 ; OPT-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 [[INPUT1:%.*]], ptr byval([[STRUCT_S]]) [[INPUT2:%.*]], ptr [[INOUT:%.*]]) #[[ATTR0]] {
 483 ; OPT-NEXT:    [[INOUT1:%.*]] = addrspacecast ptr [[INOUT]] to ptr addrspace(1)
 484 ; OPT-NEXT:    [[INOUT2:%.*]] = addrspacecast ptr addrspace(1) [[INOUT1]] to ptr
 485 ; OPT-NEXT:    [[INPUT2_PARAM:%.*]] = addrspacecast ptr [[INPUT2]] to ptr addrspace(101)
 486 ; OPT-NEXT:    [[INPUT2_PARAM_GEN:%.*]] = call ptr @llvm.nvvm.ptr.param.to.gen.p0.p101(ptr addrspace(101) [[INPUT2_PARAM]])
 487 ; OPT-NEXT:    [[INPUT1_PARAM:%.*]] = addrspacecast ptr [[INPUT1]] to ptr addrspace(101)
 488 ; OPT-NEXT:    [[INPUT1_PARAM_GEN:%.*]] = call ptr @llvm.nvvm.ptr.param.to.gen.p0.p101(ptr addrspace(101) [[INPUT1_PARAM]])
 489 ; OPT-NEXT:    [[VAL:%.*]] = load i32, ptr [[INOUT2]], align 4
 490 ; OPT-NEXT:    [[LESS:%.*]] = icmp slt i32 [[VAL]], 0
 491 ; OPT-NEXT:    br i1 [[LESS]], label %[[FIRST:.*]], label %[[SECOND:.*]]
 492 ; OPT:       [[FIRST]]:
 493 ; OPT-NEXT:    [[PTR1:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT1_PARAM_GEN]], i32 0, i32 0
 494 ; OPT-NEXT:    br label %[[MERGE:.*]]
 495 ; OPT:       [[SECOND]]:
 496 ; OPT-NEXT:    [[PTR2:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT2_PARAM_GEN]], i32 0, i32 1
 497 ; OPT-NEXT:    br label %[[MERGE]]
 498 ; OPT:       [[MERGE]]:
 499 ; OPT-NEXT:    [[PTRNEW:%.*]] = phi ptr [ [[PTR1]], %[[FIRST]] ], [ [[PTR2]], %[[SECOND]] ]
 500 ; OPT-NEXT:    [[VALLOADED:%.*]] = load i32, ptr [[PTRNEW]], align 4
 501 ; OPT-NEXT:    store i32 [[VALLOADED]], ptr [[INOUT2]], align 4
 502 ; OPT-NEXT:    ret void
 503   %val = load i32, ptr %inout
 504   %less = icmp slt i32 %val, 0
 505   br i1 %less, label %first, label %second
 506 first:
 507   %ptr1 = getelementptr inbounds %struct.s, ptr %input1, i32 0, i32 0
 508   br label %merge
 509 second:
 510   %ptr2 = getelementptr inbounds %struct.s, ptr %input2, i32 0, i32 1
 511   br label %merge
 512 merge:
 513   %ptrnew = phi ptr [%ptr1, %first], [%ptr2, %second]
 514   %valloaded = load i32, ptr %ptrnew
 515   store i32 %valloaded, ptr %inout
 516   ret void
 517 }
 518
 519 ; NOTE: %input2 is *not* grid_constant
 520 define void @grid_const_select(ptr byval(i32) align 4 %input1, ptr byval(i32) %input2, ptr %inout) {
 521 ; PTX-LABEL: grid_const_select(
 522 ; PTX:       {
 523 ; PTX-NEXT:    .reg .pred %p<2>;
 524 ; PTX-NEXT:    .reg .b32 %r<3>;
 525 ; PTX-NEXT:    .reg .b64 %rd<10>;
 526 ; PTX-EMPTY:
 527 ; PTX-NEXT:  // %bb.0:
 528 ; PTX-NEXT:    mov.b64 %rd1, grid_const_select_param_0;
 529 ; PTX-NEXT:    ld.param.u64 %rd2, [grid_const_select_param_2];
 530 ; PTX-NEXT:    cvta.to.global.u64 %rd3, %rd2;
 531 ; PTX-NEXT:    mov.b64 %rd4, grid_const_select_param_1;
 532 ; PTX-NEXT:    mov.u64 %rd5, %rd4;
 533 ; PTX-NEXT:    cvta.param.u64 %rd6, %rd5;
 534 ; PTX-NEXT:    mov.u64 %rd7, %rd1;
 535 ; PTX-NEXT:    cvta.param.u64 %rd8, %rd7;
 536 ; PTX-NEXT:    ld.global.u32 %r1, [%rd3];
 537 ; PTX-NEXT:    setp.lt.s32 %p1, %r1, 0;
 538 ; PTX-NEXT:    selp.b64 %rd9, %rd8, %rd6, %p1;
 539 ; PTX-NEXT:    ld.u32 %r2, [%rd9];
 540 ; PTX-NEXT:    st.global.u32 [%rd3], %r2;
 541 ; PTX-NEXT:    ret;
 542 ; OPT-LABEL: define void @grid_const_select(
 543 ; OPT-SAME: ptr byval(i32) align 4 [[INPUT1:%.*]], ptr byval(i32) [[INPUT2:%.*]], ptr [[INOUT:%.*]]) #[[ATTR0]] {
 544 ; OPT-NEXT:    [[INOUT1:%.*]] = addrspacecast ptr [[INOUT]] to ptr addrspace(1)
 545 ; OPT-NEXT:    [[INOUT2:%.*]] = addrspacecast ptr addrspace(1) [[INOUT1]] to ptr
 546 ; OPT-NEXT:    [[INPUT2_PARAM:%.*]] = addrspacecast ptr [[INPUT2]] to ptr addrspace(101)
 547 ; OPT-NEXT:    [[INPUT2_PARAM_GEN:%.*]] = call ptr @llvm.nvvm.ptr.param.to.gen.p0.p101(ptr addrspace(101) [[INPUT2_PARAM]])
 548 ; OPT-NEXT:    [[INPUT1_PARAM:%.*]] = addrspacecast ptr [[INPUT1]] to ptr addrspace(101)
 549 ; OPT-NEXT:    [[INPUT1_PARAM_GEN:%.*]] = call ptr @llvm.nvvm.ptr.param.to.gen.p0.p101(ptr addrspace(101) [[INPUT1_PARAM]])
 550 ; OPT-NEXT:    [[VAL:%.*]] = load i32, ptr [[INOUT2]], align 4
 551 ; OPT-NEXT:    [[LESS:%.*]] = icmp slt i32 [[VAL]], 0
 552 ; OPT-NEXT:    [[PTRNEW:%.*]] = select i1 [[LESS]], ptr [[INPUT1_PARAM_GEN]], ptr [[INPUT2_PARAM_GEN]]
 553 ; OPT-NEXT:    [[VALLOADED:%.*]] = load i32, ptr [[PTRNEW]], align 4
 554 ; OPT-NEXT:    store i32 [[VALLOADED]], ptr [[INOUT2]], align 4
 555 ; OPT-NEXT:    ret void
 556   %val = load i32, ptr %inout
 557   %less = icmp slt i32 %val, 0
 558   %ptrnew = select i1 %less, ptr %input1, ptr %input2
 559   %valloaded = load i32, ptr %ptrnew
 560   store i32 %valloaded, ptr %inout
 561   ret void
 562 }
 563
 564 define i32 @grid_const_ptrtoint(ptr byval(i32) %input) {
 565 ; PTX-LABEL: grid_const_ptrtoint(
 566 ; PTX:       {
 567 ; PTX-NEXT:    .reg .b32 %r<4>;
 568 ; PTX-NEXT:    .reg .b64 %rd<4>;
 569 ; PTX-EMPTY:
 570 ; PTX-NEXT:  // %bb.0:
 571 ; PTX-NEXT:    mov.b64 %rd1, grid_const_ptrtoint_param_0;
 572 ; PTX-NEXT:    mov.u64 %rd2, %rd1;
 573 ; PTX-NEXT:    ld.param.u32 %r1, [grid_const_ptrtoint_param_0];
 574 ; PTX-NEXT:    cvta.param.u64 %rd3, %rd2;
 575 ; PTX-NEXT:    cvt.u32.u64 %r2, %rd3;
 576 ; PTX-NEXT:    add.s32 %r3, %r1, %r2;
 577 ; PTX-NEXT:    st.param.b32 [func_retval0], %r3;
 578 ; PTX-NEXT:    ret;
 579 ; OPT-LABEL: define i32 @grid_const_ptrtoint(
 580 ; OPT-SAME: ptr byval(i32) align 4 [[INPUT:%.*]]) #[[ATTR0]] {
 581 ; OPT-NEXT:    [[INPUT2:%.*]] = addrspacecast ptr [[INPUT]] to ptr addrspace(101)
 582 ; OPT-NEXT:    [[INPUT3:%.*]] = load i32, ptr addrspace(101) [[INPUT2]], align 4
 583 ; OPT-NEXT:    [[INPUT1:%.*]] = call ptr @llvm.nvvm.ptr.param.to.gen.p0.p101(ptr addrspace(101) [[INPUT2]])
 584 ; OPT-NEXT:    [[PTRVAL:%.*]] = ptrtoint ptr [[INPUT1]] to i32
 585 ; OPT-NEXT:    [[KEEPALIVE:%.*]] = add i32 [[INPUT3]], [[PTRVAL]]
 586 ; OPT-NEXT:    ret i32 [[KEEPALIVE]]
 587   %val = load i32, ptr %input
 588   %ptrval = ptrtoint ptr %input to i32
 589   %keepalive = add i32 %val, %ptrval
 590   ret i32 %keepalive
 591 }
 592
 593
 594
 595 declare dso_local void @dummy() local_unnamed_addr
 596 declare dso_local ptr @escape(ptr) local_unnamed_addr
 597 declare dso_local ptr @escape3(ptr, ptr, ptr) local_unnamed_addr
 598
 599 !nvvm.annotations = !{!0, !1, !2, !3, !4, !5, !6, !7, !8, !9, !10, !11, !12, !13, !14, !15, !16, !17, !18, !19, !20, !21, !22, !23}
 600
 601 !0 = !{ptr @grid_const_int, !"kernel", i32 1, !"grid_constant", !1}
 602 !1 = !{i32 1}
 603
 604 !2 = !{ptr @grid_const_struct, !"kernel", i32 1, !"grid_constant", !3}
 605 !3 = !{i32 1}
 606
 607 !4 = !{ptr @grid_const_escape, !"kernel", i32 1, !"grid_constant", !5}
 608 !5 = !{i32 1}
 609
 610 !6 = !{ptr @multiple_grid_const_escape, !"kernel", i32 1, !"grid_constant", !7}
 611 !7 = !{i32 1, i32 3}
 612
 613 !8 = !{ptr @grid_const_memory_escape, !"kernel", i32 1, !"grid_constant", !9}
 614 !9 = !{i32 1}
 615
 616 !10 = !{ptr @grid_const_inlineasm_escape, !"kernel", i32 1, !"grid_constant", !11}
 617 !11 = !{i32 1}
 618
 619 !12 = !{ptr @grid_const_partial_escape, !"kernel", i32 1, !"grid_constant", !13}
 620 !13 = !{i32 1}
 621
 622 !14 = !{ptr @grid_const_partial_escapemem, !"kernel", i32 1, !"grid_constant", !15}
 623 !15 = !{i32 1}
 624
 625 !16 = !{ptr @grid_const_phi, !"kernel", i32 1, !"grid_constant", !17}
 626 !17 = !{i32 1}
 627
 628 !18 = !{ptr @grid_const_phi_ngc, !"kernel", i32 1, !"grid_constant", !19}
 629 !19 = !{i32 1}
 630
 631 !20 = !{ptr @grid_const_select, !"kernel", i32 1, !"grid_constant", !21}
 632 !21 = !{i32 1}
 633
 634 !22 = !{ptr @grid_const_ptrtoint, !"kernel", i32 1, !"grid_constant", !23}
 635 !23 = !{i32 1}
 636
 637