| // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 |
| // RUN: %clang_cc1 -triple amdgpu-amd-amdhsa \ |
| // RUN: -fcuda-is-device -mcode-object-version=4 -emit-llvm -o - -x hip %s \ |
| // RUN: | FileCheck -check-prefix=PRECOV5 %s |
| |
| // RUN: %clang_cc1 -triple amdgpu-amd-amdhsa \ |
| // RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \ |
| // RUN: | FileCheck -check-prefix=COV5 %s |
| |
| // RUN: %clang_cc1 -triple amdgpu-amd-amdhsa \ |
| // RUN: -fcuda-is-device -mcode-object-version=6 -emit-llvm -o - -x hip %s \ |
| // RUN: | FileCheck -check-prefix=COV5 %s |
| |
| #include "Inputs/cuda.h" |
| |
| |
| |
| // PRECOV5-LABEL: define dso_local void @_Z23test_get_workgroup_sizeiPi( |
| // PRECOV5-SAME: i32 noundef [[D:%.*]], ptr noundef [[OUT:%.*]]) #[[ATTR0:[0-9]+]] { |
| // PRECOV5-NEXT: [[ENTRY:.*:]] |
| // PRECOV5-NEXT: [[D_ADDR:%.*]] = alloca i32, align 4, addrspace(5) |
| // PRECOV5-NEXT: [[OUT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) |
| // PRECOV5-NEXT: [[D_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[D_ADDR]] to ptr |
| // PRECOV5-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr |
| // PRECOV5-NEXT: store i32 [[D]], ptr [[D_ADDR_ASCAST]], align 4 |
| // PRECOV5-NEXT: store ptr [[OUT]], ptr [[OUT_ADDR_ASCAST]], align 8 |
| // PRECOV5-NEXT: [[TMP0:%.*]] = load i32, ptr [[D_ADDR_ASCAST]], align 4 |
| // PRECOV5-NEXT: switch i32 [[TMP0]], label %[[SW_DEFAULT:.*]] [ |
| // PRECOV5-NEXT: i32 0, label %[[SW_BB:.*]] |
| // PRECOV5-NEXT: i32 1, label %[[SW_BB1:.*]] |
| // PRECOV5-NEXT: i32 2, label %[[SW_BB2:.*]] |
| // PRECOV5-NEXT: ] |
| // PRECOV5: [[SW_BB]]: |
| // PRECOV5-NEXT: [[TMP1:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() |
| // PRECOV5-NEXT: [[TMP2:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP1]], i64 4 |
| // PRECOV5-NEXT: [[TMP3:%.*]] = load i16, ptr addrspace(4) [[TMP2]], align 2, !range [[RNG3:![0-9]+]], !invariant.load [[META4:![0-9]+]], !noundef [[META4]] |
| // PRECOV5-NEXT: [[TMP4:%.*]] = zext i16 [[TMP3]] to i32 |
| // PRECOV5-NEXT: [[TMP5:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8 |
| // PRECOV5-NEXT: store i32 [[TMP4]], ptr [[TMP5]], align 4 |
| // PRECOV5-NEXT: br label %[[SW_EPILOG:.*]] |
| // PRECOV5: [[SW_BB1]]: |
| // PRECOV5-NEXT: [[TMP6:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() |
| // PRECOV5-NEXT: [[TMP7:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP6]], i64 6 |
| // PRECOV5-NEXT: [[TMP8:%.*]] = load i16, ptr addrspace(4) [[TMP7]], align 2, !range [[RNG3]], !invariant.load [[META4]], !noundef [[META4]] |
| // PRECOV5-NEXT: [[TMP9:%.*]] = zext i16 [[TMP8]] to i32 |
| // PRECOV5-NEXT: [[TMP10:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8 |
| // PRECOV5-NEXT: store i32 [[TMP9]], ptr [[TMP10]], align 4 |
| // PRECOV5-NEXT: br label %[[SW_EPILOG]] |
| // PRECOV5: [[SW_BB2]]: |
| // PRECOV5-NEXT: [[TMP11:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() |
| // PRECOV5-NEXT: [[TMP12:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP11]], i64 8 |
| // PRECOV5-NEXT: [[TMP13:%.*]] = load i16, ptr addrspace(4) [[TMP12]], align 2, !range [[RNG3]], !invariant.load [[META4]], !noundef [[META4]] |
| // PRECOV5-NEXT: [[TMP14:%.*]] = zext i16 [[TMP13]] to i32 |
| // PRECOV5-NEXT: [[TMP15:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8 |
| // PRECOV5-NEXT: store i32 [[TMP14]], ptr [[TMP15]], align 4 |
| // PRECOV5-NEXT: br label %[[SW_EPILOG]] |
| // PRECOV5: [[SW_DEFAULT]]: |
| // PRECOV5-NEXT: [[TMP16:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8 |
| // PRECOV5-NEXT: store i32 0, ptr [[TMP16]], align 4 |
| // PRECOV5-NEXT: br label %[[SW_EPILOG]] |
| // PRECOV5: [[SW_EPILOG]]: |
| // PRECOV5-NEXT: ret void |
| // |
| // COV5-LABEL: define dso_local void @_Z23test_get_workgroup_sizeiPi( |
| // COV5-SAME: i32 noundef [[D:%.*]], ptr noundef [[OUT:%.*]]) #[[ATTR0:[0-9]+]] { |
| // COV5-NEXT: [[ENTRY:.*:]] |
| // COV5-NEXT: [[D_ADDR:%.*]] = alloca i32, align 4, addrspace(5) |
| // COV5-NEXT: [[OUT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) |
| // COV5-NEXT: [[D_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[D_ADDR]] to ptr |
| // COV5-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr |
| // COV5-NEXT: store i32 [[D]], ptr [[D_ADDR_ASCAST]], align 4 |
| // COV5-NEXT: store ptr [[OUT]], ptr [[OUT_ADDR_ASCAST]], align 8 |
| // COV5-NEXT: [[TMP0:%.*]] = load i32, ptr [[D_ADDR_ASCAST]], align 4 |
| // COV5-NEXT: switch i32 [[TMP0]], label %[[SW_DEFAULT:.*]] [ |
| // COV5-NEXT: i32 0, label %[[SW_BB:.*]] |
| // COV5-NEXT: i32 1, label %[[SW_BB1:.*]] |
| // COV5-NEXT: i32 2, label %[[SW_BB2:.*]] |
| // COV5-NEXT: ] |
| // COV5: [[SW_BB]]: |
| // COV5-NEXT: [[TMP1:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() |
| // COV5-NEXT: [[TMP2:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP1]], i64 12 |
| // COV5-NEXT: [[TMP3:%.*]] = load i16, ptr addrspace(4) [[TMP2]], align 2, !range [[RNG3:![0-9]+]], !invariant.load [[META4:![0-9]+]], !noundef [[META4]] |
| // COV5-NEXT: [[TMP4:%.*]] = zext i16 [[TMP3]] to i32 |
| // COV5-NEXT: [[TMP5:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8 |
| // COV5-NEXT: store i32 [[TMP4]], ptr [[TMP5]], align 4 |
| // COV5-NEXT: br label %[[SW_EPILOG:.*]] |
| // COV5: [[SW_BB1]]: |
| // COV5-NEXT: [[TMP6:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() |
| // COV5-NEXT: [[TMP7:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP6]], i64 14 |
| // COV5-NEXT: [[TMP8:%.*]] = load i16, ptr addrspace(4) [[TMP7]], align 2, !range [[RNG3]], !invariant.load [[META4]], !noundef [[META4]] |
| // COV5-NEXT: [[TMP9:%.*]] = zext i16 [[TMP8]] to i32 |
| // COV5-NEXT: [[TMP10:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8 |
| // COV5-NEXT: store i32 [[TMP9]], ptr [[TMP10]], align 4 |
| // COV5-NEXT: br label %[[SW_EPILOG]] |
| // COV5: [[SW_BB2]]: |
| // COV5-NEXT: [[TMP11:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() |
| // COV5-NEXT: [[TMP12:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP11]], i64 16 |
| // COV5-NEXT: [[TMP13:%.*]] = load i16, ptr addrspace(4) [[TMP12]], align 2, !range [[RNG3]], !invariant.load [[META4]], !noundef [[META4]] |
| // COV5-NEXT: [[TMP14:%.*]] = zext i16 [[TMP13]] to i32 |
| // COV5-NEXT: [[TMP15:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8 |
| // COV5-NEXT: store i32 [[TMP14]], ptr [[TMP15]], align 4 |
| // COV5-NEXT: br label %[[SW_EPILOG]] |
| // COV5: [[SW_DEFAULT]]: |
| // COV5-NEXT: [[TMP16:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8 |
| // COV5-NEXT: store i32 0, ptr [[TMP16]], align 4 |
| // COV5-NEXT: br label %[[SW_EPILOG]] |
| // COV5: [[SW_EPILOG]]: |
| // COV5-NEXT: ret void |
| // |
| __device__ void test_get_workgroup_size(int d, int *out) |
| { |
| switch (d) { |
| case 0: *out = __builtin_amdgcn_workgroup_size_x(); break; |
| case 1: *out = __builtin_amdgcn_workgroup_size_y(); break; |
| case 2: *out = __builtin_amdgcn_workgroup_size_z(); break; |
| default: *out = 0; |
| } |
| } |
| |
| // CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025} |
| //. |
| // PRECOV5: [[RNG3]] = !{i16 1, i16 1025} |
| // PRECOV5: [[META4]] = !{} |
| //. |
| // COV5: [[RNG3]] = !{i16 1, i16 1025} |
| // COV5: [[META4]] = !{} |
| //. |
