llvm/test/CodeGen/NVPTX/ldu-ldg.ll

   1 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_32 | FileCheck %s
   2 ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_32 | %ptxas-verify %}
   3
   4
   5 declare i8 @llvm.nvvm.ldu.global.i.i8.p1(ptr addrspace(1) %ptr, i32 %align)
   6 declare i16 @llvm.nvvm.ldu.global.i.i16.p1(ptr addrspace(1) %ptr, i32 %align)
   7 declare i32 @llvm.nvvm.ldu.global.i.i32.p1(ptr addrspace(1) %ptr, i32 %align)
   8 declare i64 @llvm.nvvm.ldu.global.i.i64.p1(ptr addrspace(1) %ptr, i32 %align)
   9 declare ptr @llvm.nvvm.ldu.global.p.p1i8(ptr addrspace(1) %ptr, i32 %align)
  10 declare float @llvm.nvvm.ldu.global.f.f32.p1(ptr addrspace(1) %ptr, i32 %align)
  11 declare double @llvm.nvvm.ldu.global.f.f64.p1(ptr addrspace(1) %ptr, i32 %align)
  12 declare half @llvm.nvvm.ldu.global.f.f16.p1(ptr addrspace(1) %ptr, i32 %align)
  13 declare <2 x half> @llvm.nvvm.ldu.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 %align)
  14
  15 declare i8 @llvm.nvvm.ldg.global.i.i8.p1(ptr addrspace(1) %ptr, i32 %align)
  16 declare i16 @llvm.nvvm.ldg.global.i.i16.p1(ptr addrspace(1) %ptr, i32 %align)
  17 declare i32 @llvm.nvvm.ldg.global.i.i32.p1(ptr addrspace(1) %ptr, i32 %align)
  18 declare i64 @llvm.nvvm.ldg.global.i.i64.p1(ptr addrspace(1) %ptr, i32 %align)
  19 declare ptr @llvm.nvvm.ldg.global.p.p1i8(ptr addrspace(1) %ptr, i32 %align)
  20 declare float @llvm.nvvm.ldg.global.f.f32.p1(ptr addrspace(1) %ptr, i32 %align)
  21 declare double @llvm.nvvm.ldg.global.f.f64.p1(ptr addrspace(1) %ptr, i32 %align)
  22 declare half @llvm.nvvm.ldg.global.f.f16.p1(ptr addrspace(1) %ptr, i32 %align)
  23 declare <2 x half> @llvm.nvvm.ldg.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 %align)
  24
  25 ; CHECK-LABEL: test_ldu_i8
  26 define i8 @test_ldu_i8(ptr addrspace(1) %ptr) {
  27   ; CHECK: ldu.global.u8
  28   %val = tail call i8 @llvm.nvvm.ldu.global.i.i8.p1(ptr addrspace(1) %ptr, i32 4)
  29   ret i8 %val
  30 }
  31
  32 ; CHECK-LABEL: test_ldu_i16
  33 define i16 @test_ldu_i16(ptr addrspace(1) %ptr) {
  34   ; CHECK: ldu.global.u16
  35   %val = tail call i16 @llvm.nvvm.ldu.global.i.i16.p1(ptr addrspace(1) %ptr, i32 2)
  36   ret i16 %val
  37 }
  38
  39 ; CHECK-LABEL: test_ldu_i32
  40 define i32 @test_ldu_i32(ptr addrspace(1) %ptr) {
  41   ; CHECK: ldu.global.u32
  42   %val = tail call i32 @llvm.nvvm.ldu.global.i.i32.p1(ptr addrspace(1) %ptr, i32 4)
  43   ret i32 %val
  44 }
  45
  46 ; CHECK-LABEL: test_ldu_i64
  47 define i64 @test_ldu_i64(ptr addrspace(1) %ptr) {
  48   ; CHECK: ldu.global.u64
  49   %val = tail call i64 @llvm.nvvm.ldu.global.i.i64.p1(ptr addrspace(1) %ptr, i32 8)
  50   ret i64 %val
  51 }
  52
  53 ; CHECK-LABEL: test_ldu_p
  54 define ptr @test_ldu_p(ptr addrspace(1) %ptr) {
  55   ; CHECK: ldu.global.u64
  56   %val = tail call ptr @llvm.nvvm.ldu.global.p.p1i8(ptr addrspace(1) %ptr, i32 8)
  57   ret ptr %val
  58 }
  59
  60
  61 ; CHECK-LABEL: test_ldu_f32
  62 define float @test_ldu_f32(ptr addrspace(1) %ptr) {
  63   ; CHECK: ldu.global.f32
  64   %val = tail call float @llvm.nvvm.ldu.global.f.f32.p1(ptr addrspace(1) %ptr, i32 4)
  65   ret float %val
  66 }
  67
  68 ; CHECK-LABEL: test_ldu_f64
  69 define double @test_ldu_f64(ptr addrspace(1) %ptr) {
  70   ; CHECK: ldu.global.f64
  71   %val = tail call double @llvm.nvvm.ldu.global.f.f64.p1(ptr addrspace(1) %ptr, i32 8)
  72   ret double %val
  73 }
  74
  75 ; CHECK-LABEL: test_ldu_f16
  76 define half @test_ldu_f16(ptr addrspace(1) %ptr) {
  77   ; CHECK: ldu.global.u16
  78   %val = tail call half @llvm.nvvm.ldu.global.f.f16.p1(ptr addrspace(1) %ptr, i32 2)
  79   ret half %val
  80 }
  81
  82 ; CHECK-LABEL: test_ldu_v2f16
  83 define <2 x half> @test_ldu_v2f16(ptr addrspace(1) %ptr) {
  84   ; CHECK: ldu.global.u32
  85   %val = tail call <2 x half> @llvm.nvvm.ldu.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 4)
  86   ret <2 x half> %val
  87 }
  88
  89 ; CHECK-LABEL: test_ldg_i8
  90 define i8 @test_ldg_i8(ptr addrspace(1) %ptr) {
  91   ; CHECK: ld.global.nc.u8
  92   %val = tail call i8 @llvm.nvvm.ldg.global.i.i8.p1(ptr addrspace(1) %ptr, i32 4)
  93   ret i8 %val
  94 }
  95
  96 ; CHECK-LABEL: test_ldg_i16
  97 define i16 @test_ldg_i16(ptr addrspace(1) %ptr) {
  98   ; CHECK: ld.global.nc.u16
  99   %val = tail call i16 @llvm.nvvm.ldg.global.i.i16.p1(ptr addrspace(1) %ptr, i32 2)
 100   ret i16 %val
 101 }
 102
 103 ; CHECK-LABEL: test_ldg_i32
 104 define i32 @test_ldg_i32(ptr addrspace(1) %ptr) {
 105   ; CHECK: ld.global.nc.u32
 106   %val = tail call i32 @llvm.nvvm.ldg.global.i.i32.p1(ptr addrspace(1) %ptr, i32 4)
 107   ret i32 %val
 108 }
 109
 110 ; CHECK-LABEL: test_ldg_i64
 111 define i64 @test_ldg_i64(ptr addrspace(1) %ptr) {
 112   ; CHECK: ld.global.nc.u64
 113   %val = tail call i64 @llvm.nvvm.ldg.global.i.i64.p1(ptr addrspace(1) %ptr, i32 8)
 114   ret i64 %val
 115 }
 116
 117 ; CHECK-LABEL: test_ldg_p
 118 define ptr @test_ldg_p(ptr addrspace(1) %ptr) {
 119   ; CHECK: ld.global.nc.u64
 120   %val = tail call ptr @llvm.nvvm.ldg.global.p.p1i8(ptr addrspace(1) %ptr, i32 8)
 121   ret ptr %val
 122 }
 123
 124 ; CHECK-LABEL: test_ldg_f32
 125 define float @test_ldg_f32(ptr addrspace(1) %ptr) {
 126   ; CHECK: ld.global.nc.f32
 127   %val = tail call float @llvm.nvvm.ldg.global.f.f32.p1(ptr addrspace(1) %ptr, i32 4)
 128   ret float %val
 129 }
 130
 131 ; CHECK-LABEL: test_ldg_f64
 132 define double @test_ldg_f64(ptr addrspace(1) %ptr) {
 133   ; CHECK: ld.global.nc.f64
 134   %val = tail call double @llvm.nvvm.ldg.global.f.f64.p1(ptr addrspace(1) %ptr, i32 8)
 135   ret double %val
 136 }
 137
 138 ; CHECK-LABEL: test_ldg_f16
 139 define half @test_ldg_f16(ptr addrspace(1) %ptr) {
 140   ; CHECK: ld.global.nc.u16
 141   %val = tail call half @llvm.nvvm.ldg.global.f.f16.p1(ptr addrspace(1) %ptr, i32 2)
 142   ret half %val
 143 }
 144
 145 ; CHECK-LABEL: test_ldg_v2f16
 146 define <2 x half> @test_ldg_v2f16(ptr addrspace(1) %ptr) {
 147   ; CHECK: ld.global.nc.u32
 148   %val = tail call <2 x half> @llvm.nvvm.ldg.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 4)
 149   ret <2 x half> %val
 150 }