44 lines · plain
1// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \2// RUN: -fcuda-is-device -mcode-object-version=4 -emit-llvm -o - -x hip %s \3// RUN: | FileCheck -check-prefix=PRECOV5 %s4 5// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \6// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \7// RUN: | FileCheck -check-prefix=COV5 %s8 9// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \10// RUN: -fcuda-is-device -mcode-object-version=6 -emit-llvm -o - -x hip %s \11// RUN: | FileCheck -check-prefix=COV5 %s12 13#include "Inputs/cuda.h"14 15// PRECOV5-LABEL: test_get_workgroup_size16// PRECOV5: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()17// PRECOV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 418// PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef19// PRECOV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 620// PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef21// PRECOV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 822// PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef23 24// COV5-LABEL: test_get_workgroup_size25// COV5: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()26// COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 1227// COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef28// COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 1429// COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef30// COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 1631// COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef32 33__device__ void test_get_workgroup_size(int d, int *out)34{35 switch (d) {36 case 0: *out = __builtin_amdgcn_workgroup_size_x(); break;37 case 1: *out = __builtin_amdgcn_workgroup_size_y(); break;38 case 2: *out = __builtin_amdgcn_workgroup_size_z(); break;39 default: *out = 0;40 }41}42 43// CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025}44