intrinsic-old.ll 6.3 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282
  1. ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
  2. ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
  3. define ptx_device i32 @test_tid_x() {
  4. ; CHECK: mov.u32 %r{{[0-9]+}}, %tid.x;
  5. ; CHECK: ret;
  6. %x = call i32 @llvm.ptx.read.tid.x()
  7. ret i32 %x
  8. }
  9. define ptx_device i32 @test_tid_y() {
  10. ; CHECK: mov.u32 %r{{[0-9]+}}, %tid.y;
  11. ; CHECK: ret;
  12. %x = call i32 @llvm.ptx.read.tid.y()
  13. ret i32 %x
  14. }
  15. define ptx_device i32 @test_tid_z() {
  16. ; CHECK: mov.u32 %r{{[0-9]+}}, %tid.z;
  17. ; CHECK: ret;
  18. %x = call i32 @llvm.ptx.read.tid.z()
  19. ret i32 %x
  20. }
  21. define ptx_device i32 @test_tid_w() {
  22. ; CHECK: mov.u32 %r{{[0-9]+}}, %tid.w;
  23. ; CHECK: ret;
  24. %x = call i32 @llvm.ptx.read.tid.w()
  25. ret i32 %x
  26. }
  27. define ptx_device i32 @test_ntid_x() {
  28. ; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.x;
  29. ; CHECK: ret;
  30. %x = call i32 @llvm.ptx.read.ntid.x()
  31. ret i32 %x
  32. }
  33. define ptx_device i32 @test_ntid_y() {
  34. ; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.y;
  35. ; CHECK: ret;
  36. %x = call i32 @llvm.ptx.read.ntid.y()
  37. ret i32 %x
  38. }
  39. define ptx_device i32 @test_ntid_z() {
  40. ; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.z;
  41. ; CHECK: ret;
  42. %x = call i32 @llvm.ptx.read.ntid.z()
  43. ret i32 %x
  44. }
  45. define ptx_device i32 @test_ntid_w() {
  46. ; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.w;
  47. ; CHECK: ret;
  48. %x = call i32 @llvm.ptx.read.ntid.w()
  49. ret i32 %x
  50. }
  51. define ptx_device i32 @test_laneid() {
  52. ; CHECK: mov.u32 %r{{[0-9]+}}, %laneid;
  53. ; CHECK: ret;
  54. %x = call i32 @llvm.ptx.read.laneid()
  55. ret i32 %x
  56. }
  57. define ptx_device i32 @test_warpid() {
  58. ; CHECK: mov.u32 %r{{[0-9]+}}, %warpid;
  59. ; CHECK: ret;
  60. %x = call i32 @llvm.ptx.read.warpid()
  61. ret i32 %x
  62. }
  63. define ptx_device i32 @test_nwarpid() {
  64. ; CHECK: mov.u32 %r{{[0-9]+}}, %nwarpid;
  65. ; CHECK: ret;
  66. %x = call i32 @llvm.ptx.read.nwarpid()
  67. ret i32 %x
  68. }
  69. define ptx_device i32 @test_ctaid_x() {
  70. ; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.x;
  71. ; CHECK: ret;
  72. %x = call i32 @llvm.ptx.read.ctaid.x()
  73. ret i32 %x
  74. }
  75. define ptx_device i32 @test_ctaid_y() {
  76. ; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.y;
  77. ; CHECK: ret;
  78. %x = call i32 @llvm.ptx.read.ctaid.y()
  79. ret i32 %x
  80. }
  81. define ptx_device i32 @test_ctaid_z() {
  82. ; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.z;
  83. ; CHECK: ret;
  84. %x = call i32 @llvm.ptx.read.ctaid.z()
  85. ret i32 %x
  86. }
  87. define ptx_device i32 @test_ctaid_w() {
  88. ; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.w;
  89. ; CHECK: ret;
  90. %x = call i32 @llvm.ptx.read.ctaid.w()
  91. ret i32 %x
  92. }
  93. define ptx_device i32 @test_nctaid_x() {
  94. ; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.x;
  95. ; CHECK: ret;
  96. %x = call i32 @llvm.ptx.read.nctaid.x()
  97. ret i32 %x
  98. }
  99. define ptx_device i32 @test_nctaid_y() {
  100. ; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.y;
  101. ; CHECK: ret;
  102. %x = call i32 @llvm.ptx.read.nctaid.y()
  103. ret i32 %x
  104. }
  105. define ptx_device i32 @test_nctaid_z() {
  106. ; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.z;
  107. ; CHECK: ret;
  108. %x = call i32 @llvm.ptx.read.nctaid.z()
  109. ret i32 %x
  110. }
  111. define ptx_device i32 @test_nctaid_w() {
  112. ; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.w;
  113. ; CHECK: ret;
  114. %x = call i32 @llvm.ptx.read.nctaid.w()
  115. ret i32 %x
  116. }
  117. define ptx_device i32 @test_smid() {
  118. ; CHECK: mov.u32 %r{{[0-9]+}}, %smid;
  119. ; CHECK: ret;
  120. %x = call i32 @llvm.ptx.read.smid()
  121. ret i32 %x
  122. }
  123. define ptx_device i32 @test_nsmid() {
  124. ; CHECK: mov.u32 %r{{[0-9]+}}, %nsmid;
  125. ; CHECK: ret;
  126. %x = call i32 @llvm.ptx.read.nsmid()
  127. ret i32 %x
  128. }
  129. define ptx_device i32 @test_gridid() {
  130. ; CHECK: mov.u32 %r{{[0-9]+}}, %gridid;
  131. ; CHECK: ret;
  132. %x = call i32 @llvm.ptx.read.gridid()
  133. ret i32 %x
  134. }
  135. define ptx_device i32 @test_lanemask_eq() {
  136. ; CHECK: mov.u32 %r{{[0-9]+}}, %lanemask_eq;
  137. ; CHECK: ret;
  138. %x = call i32 @llvm.ptx.read.lanemask.eq()
  139. ret i32 %x
  140. }
  141. define ptx_device i32 @test_lanemask_le() {
  142. ; CHECK: mov.u32 %r{{[0-9]+}}, %lanemask_le;
  143. ; CHECK: ret;
  144. %x = call i32 @llvm.ptx.read.lanemask.le()
  145. ret i32 %x
  146. }
  147. define ptx_device i32 @test_lanemask_lt() {
  148. ; CHECK: mov.u32 %r{{[0-9]+}}, %lanemask_lt;
  149. ; CHECK: ret;
  150. %x = call i32 @llvm.ptx.read.lanemask.lt()
  151. ret i32 %x
  152. }
  153. define ptx_device i32 @test_lanemask_ge() {
  154. ; CHECK: mov.u32 %r{{[0-9]+}}, %lanemask_ge;
  155. ; CHECK: ret;
  156. %x = call i32 @llvm.ptx.read.lanemask.ge()
  157. ret i32 %x
  158. }
  159. define ptx_device i32 @test_lanemask_gt() {
  160. ; CHECK: mov.u32 %r{{[0-9]+}}, %lanemask_gt;
  161. ; CHECK: ret;
  162. %x = call i32 @llvm.ptx.read.lanemask.gt()
  163. ret i32 %x
  164. }
  165. define ptx_device i32 @test_clock() {
  166. ; CHECK: mov.u32 %r{{[0-9]+}}, %clock;
  167. ; CHECK: ret;
  168. %x = call i32 @llvm.ptx.read.clock()
  169. ret i32 %x
  170. }
  171. define ptx_device i64 @test_clock64() {
  172. ; CHECK: mov.u64 %rd{{[0-9]+}}, %clock64;
  173. ; CHECK: ret;
  174. %x = call i64 @llvm.ptx.read.clock64()
  175. ret i64 %x
  176. }
  177. define ptx_device i32 @test_pm0() {
  178. ; CHECK: mov.u32 %r{{[0-9]+}}, %pm0;
  179. ; CHECK: ret;
  180. %x = call i32 @llvm.ptx.read.pm0()
  181. ret i32 %x
  182. }
  183. define ptx_device i32 @test_pm1() {
  184. ; CHECK: mov.u32 %r{{[0-9]+}}, %pm1;
  185. ; CHECK: ret;
  186. %x = call i32 @llvm.ptx.read.pm1()
  187. ret i32 %x
  188. }
  189. define ptx_device i32 @test_pm2() {
  190. ; CHECK: mov.u32 %r{{[0-9]+}}, %pm2;
  191. ; CHECK: ret;
  192. %x = call i32 @llvm.ptx.read.pm2()
  193. ret i32 %x
  194. }
  195. define ptx_device i32 @test_pm3() {
  196. ; CHECK: mov.u32 %r{{[0-9]+}}, %pm3;
  197. ; CHECK: ret;
  198. %x = call i32 @llvm.ptx.read.pm3()
  199. ret i32 %x
  200. }
  201. define ptx_device void @test_bar_sync() {
  202. ; CHECK: bar.sync 0
  203. ; CHECK: ret;
  204. call void @llvm.ptx.bar.sync(i32 0)
  205. ret void
  206. }
  207. declare i32 @llvm.ptx.read.tid.x()
  208. declare i32 @llvm.ptx.read.tid.y()
  209. declare i32 @llvm.ptx.read.tid.z()
  210. declare i32 @llvm.ptx.read.tid.w()
  211. declare i32 @llvm.ptx.read.ntid.x()
  212. declare i32 @llvm.ptx.read.ntid.y()
  213. declare i32 @llvm.ptx.read.ntid.z()
  214. declare i32 @llvm.ptx.read.ntid.w()
  215. declare i32 @llvm.ptx.read.laneid()
  216. declare i32 @llvm.ptx.read.warpid()
  217. declare i32 @llvm.ptx.read.nwarpid()
  218. declare i32 @llvm.ptx.read.ctaid.x()
  219. declare i32 @llvm.ptx.read.ctaid.y()
  220. declare i32 @llvm.ptx.read.ctaid.z()
  221. declare i32 @llvm.ptx.read.ctaid.w()
  222. declare i32 @llvm.ptx.read.nctaid.x()
  223. declare i32 @llvm.ptx.read.nctaid.y()
  224. declare i32 @llvm.ptx.read.nctaid.z()
  225. declare i32 @llvm.ptx.read.nctaid.w()
  226. declare i32 @llvm.ptx.read.smid()
  227. declare i32 @llvm.ptx.read.nsmid()
  228. declare i32 @llvm.ptx.read.gridid()
  229. declare i32 @llvm.ptx.read.lanemask.eq()
  230. declare i32 @llvm.ptx.read.lanemask.le()
  231. declare i32 @llvm.ptx.read.lanemask.lt()
  232. declare i32 @llvm.ptx.read.lanemask.ge()
  233. declare i32 @llvm.ptx.read.lanemask.gt()
  234. declare i32 @llvm.ptx.read.clock()
  235. declare i64 @llvm.ptx.read.clock64()
  236. declare i32 @llvm.ptx.read.pm0()
  237. declare i32 @llvm.ptx.read.pm1()
  238. declare i32 @llvm.ptx.read.pm2()
  239. declare i32 @llvm.ptx.read.pm3()
  240. declare void @llvm.ptx.bar.sync(i32 %i)