xref: /aosp_15_r20/external/clang/test/CodeGenCUDA/cuda-builtin-vars.cu (revision 67e74705e28f6214e480b399dd47ea732279e315) 
1*67e74705SXin Li // RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda" -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
2*67e74705SXin Li 
3*67e74705SXin Li #include "cuda_builtin_vars.h"
4*67e74705SXin Li 
5*67e74705SXin Li // CHECK: define void @_Z6kernelPi(i32* %out)
6*67e74705SXin Li __attribute__((global))
kernel(int * out)7*67e74705SXin Li void kernel(int *out) {
8*67e74705SXin Li   int i = 0;
9*67e74705SXin Li   out[i++] = threadIdx.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
10*67e74705SXin Li   out[i++] = threadIdx.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.y()
11*67e74705SXin Li   out[i++] = threadIdx.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.z()
12*67e74705SXin Li 
13*67e74705SXin Li   out[i++] = blockIdx.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
14*67e74705SXin Li   out[i++] = blockIdx.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
15*67e74705SXin Li   out[i++] = blockIdx.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()
16*67e74705SXin Li 
17*67e74705SXin Li   out[i++] = blockDim.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
18*67e74705SXin Li   out[i++] = blockDim.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.y()
19*67e74705SXin Li   out[i++] = blockDim.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.z()
20*67e74705SXin Li 
21*67e74705SXin Li   out[i++] = gridDim.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
22*67e74705SXin Li   out[i++] = gridDim.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()
23*67e74705SXin Li   out[i++] = gridDim.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()
24*67e74705SXin Li 
25*67e74705SXin Li   out[i++] = warpSize; // CHECK: store i32 32,
26*67e74705SXin Li 
27*67e74705SXin Li   // CHECK: ret void
28*67e74705SXin Li }
29