118 lines · plain
1// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \2// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \3// RUN: | FileCheck -check-prefixes=CHECK,DEFAULT %s4// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa --gpu-max-threads-per-block=1024 \5// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \6// RUN: | FileCheck -check-prefixes=CHECK,MAX1024 %s7// RUN: %clang_cc1 -triple spirv64-amd-amdhsa --gpu-max-threads-per-block=1024 \8// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \9// RUN: | FileCheck -check-prefixes=CHECK-SPIRV,MAX1024-SPIRV %s10// RUN: %clang_cc1 -triple nvptx \11// RUN: -fcuda-is-device -emit-llvm -o - %s | FileCheck %s \12// RUN: -check-prefix=NAMD13// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \14// RUN: -verify -o - -x hip %s | FileCheck -check-prefix=NAMD %s15 16// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -foffload-uniform-block \17// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \18// RUN: | FileCheck -check-prefixes=CHECK,DEFAULT %s19// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fno-offload-uniform-block \20// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \21// RUN: | FileCheck -check-prefixes=NOUB %s22 23#include "Inputs/cuda.h"24 25__global__ void flat_work_group_size_default() {26// CHECK: define{{.*}} amdgpu_kernel void @_Z28flat_work_group_size_defaultv() [[FLAT_WORK_GROUP_SIZE_DEFAULT:#[0-9]+]]27// CHECK-SPIRV: define{{.*}} spir_kernel void @_Z28flat_work_group_size_defaultv(){{.*}} !max_work_group_size [[MAX_WORK_GROUP_SIZE_DEFAULT:![0-9]+]]28// NOUB: define{{.*}} void @_Z28flat_work_group_size_defaultv() [[NOUB:#[0-9]+]]29}30 31__attribute__((amdgpu_flat_work_group_size(32, 64))) // expected-no-diagnostics32__global__ void flat_work_group_size_32_64() {33// CHECK: define{{.*}} amdgpu_kernel void @_Z26flat_work_group_size_32_64v() [[FLAT_WORK_GROUP_SIZE_32_64:#[0-9]+]]34// CHECK-SPIRV: define{{.*}} spir_kernel void @_Z26flat_work_group_size_32_64v(){{.*}} !max_work_group_size [[MAX_WORK_GROUP_SIZE_64:![0-9]+]]35}36__attribute__((amdgpu_waves_per_eu(2))) // expected-no-diagnostics37__global__ void waves_per_eu_2() {38// CHECK: define{{.*}} amdgpu_kernel void @_Z14waves_per_eu_2v() [[WAVES_PER_EU_2:#[0-9]+]]39}40__attribute__((amdgpu_num_sgpr(32))) // expected-no-diagnostics41__global__ void num_sgpr_32() {42// CHECK: define{{.*}} amdgpu_kernel void @_Z11num_sgpr_32v() [[NUM_SGPR_32:#[0-9]+]]43}44__attribute__((amdgpu_num_vgpr(64))) // expected-no-diagnostics45__global__ void num_vgpr_64() {46// CHECK: define{{.*}} amdgpu_kernel void @_Z11num_vgpr_64v() [[NUM_VGPR_64:#[0-9]+]]47}48__attribute__((amdgpu_max_num_work_groups(32, 4, 2))) // expected-no-diagnostics49__global__ void max_num_work_groups_32_4_2() {50// CHECK: define{{.*}} amdgpu_kernel void @_Z26max_num_work_groups_32_4_2v() [[MAX_NUM_WORK_GROUPS_32_4_2:#[0-9]+]]51}52__attribute__((amdgpu_max_num_work_groups(32))) // expected-no-diagnostics53__global__ void max_num_work_groups_32() {54// CHECK: define{{.*}} amdgpu_kernel void @_Z22max_num_work_groups_32v() [[MAX_NUM_WORK_GROUPS_32_1_1:#[0-9]+]]55}56__attribute__((amdgpu_max_num_work_groups(32,1))) // expected-no-diagnostics57__global__ void max_num_work_groups_32_1() {58// CHECK: define{{.*}} amdgpu_kernel void @_Z24max_num_work_groups_32_1v() [[MAX_NUM_WORK_GROUPS_32_1_1:#[0-9]+]]59}60 61 62 63template<unsigned a>64__attribute__((amdgpu_max_num_work_groups(a, 4, 2)))65__global__ void template_a_4_2_max_num_work_groups() {}66template __global__ void template_a_4_2_max_num_work_groups<32>();67// CHECK: define{{.*}} amdgpu_kernel void @_Z34template_a_4_2_max_num_work_groupsILj32EEvv() [[MAX_NUM_WORK_GROUPS_32_4_2:#[0-9]+]]68 69template<unsigned a>70__attribute__((amdgpu_max_num_work_groups(32, a, 2)))71__global__ void template_32_a_2_max_num_work_groups() {}72template __global__ void template_32_a_2_max_num_work_groups<4>();73// CHECK: define{{.*}} amdgpu_kernel void @_Z35template_32_a_2_max_num_work_groupsILj4EEvv() [[MAX_NUM_WORK_GROUPS_32_4_2:#[0-9]+]]74 75template<unsigned a>76__attribute__((amdgpu_max_num_work_groups(32, 4, a)))77__global__ void template_32_4_a_max_num_work_groups() {}78template __global__ void template_32_4_a_max_num_work_groups<2>();79// CHECK: define{{.*}} amdgpu_kernel void @_Z35template_32_4_a_max_num_work_groupsILj2EEvv() [[MAX_NUM_WORK_GROUPS_32_4_2:#[0-9]+]]80 81template<unsigned a>82__attribute__((amdgpu_max_num_work_groups(a)))83__global__ void template_a_max_num_work_groups() {}84template __global__ void template_a_max_num_work_groups<32>();85// CHECK: define{{.*}} amdgpu_kernel void @_Z30template_a_max_num_work_groupsILj32EEvv() [[MAX_NUM_WORK_GROUPS_32_1_1]]86 87template<unsigned a, unsigned b>88__attribute__((amdgpu_max_num_work_groups(a, b)))89__global__ void template_a_b_max_num_work_groups() {}90template __global__ void template_a_b_max_num_work_groups<32, 1>();91// CHECK: define{{.*}} amdgpu_kernel void @_Z32template_a_b_max_num_work_groupsILj32ELj1EEvv() [[MAX_NUM_WORK_GROUPS_32_1_1]]92 93template<unsigned a, unsigned b, unsigned c>94__attribute__((amdgpu_max_num_work_groups(a, b, c)))95__global__ void template_a_b_c_max_num_work_groups() {}96template __global__ void template_a_b_c_max_num_work_groups<32, 4, 2>();97// CHECK: define{{.*}} amdgpu_kernel void @_Z34template_a_b_c_max_num_work_groupsILj32ELj4ELj2EEvv() [[MAX_NUM_WORK_GROUPS_32_4_2]]98 99// Make sure this is silently accepted on other targets.100// NAMD-NOT: "amdgpu-flat-work-group-size"101// NAMD-NOT: "amdgpu-waves-per-eu"102// NAMD-NOT: "amdgpu-num-vgpr"103// NAMD-NOT: "amdgpu-num-sgpr"104// NAMD-NOT: "amdgpu-max-num-work-groups"105 106// DEFAULT-DAG: attributes [[FLAT_WORK_GROUP_SIZE_DEFAULT]] = {{.*}}"amdgpu-flat-work-group-size"="1,1024"{{.*}}"uniform-work-group-size"="true"107// MAX1024-DAG: attributes [[FLAT_WORK_GROUP_SIZE_DEFAULT]] = {{.*}}"amdgpu-flat-work-group-size"="1,1024"108// MAX1024-SPIRV-DAG: [[MAX_WORK_GROUP_SIZE_DEFAULT]] = !{i32 1024, i32 1, i32 1}109// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64]] = {{.*}}"amdgpu-flat-work-group-size"="32,64"110// CHECK-SPIRV-DAG: [[MAX_WORK_GROUP_SIZE_64]] = !{i32 64, i32 1, i32 1}111// CHECK-DAG: attributes [[WAVES_PER_EU_2]] = {{.*}}"amdgpu-waves-per-eu"="2"112// CHECK-DAG: attributes [[NUM_SGPR_32]] = {{.*}}"amdgpu-num-sgpr"="32"113// CHECK-DAG: attributes [[NUM_VGPR_64]] = {{.*}}"amdgpu-num-vgpr"="64"114// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_32_4_2]] = {{.*}}"amdgpu-max-num-workgroups"="32,4,2"115// CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_32_1_1]] = {{.*}}"amdgpu-max-num-workgroups"="32,1,1"116 117// NOUB-NOT: "uniform-work-group-size"="true"118