| 1 | // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \ |
|---|
| 2 | // RUN: -fcuda-is-device -emit-llvm -o - %s | FileCheck %s |
|---|
| 3 | // RUN: %clang_cc1 -triple nvptx \ |
|---|
| 4 | // RUN: -fcuda-is-device -emit-llvm -o - %s | FileCheck %s \ |
|---|
| 5 | // RUN: -check-prefix=NAMD |
|---|
| 6 | // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \ |
|---|
| 7 | // RUN: -verify -o - %s | FileCheck -check-prefix=NAMD %s |
|---|
| 8 | |
|---|
| 9 | #include "Inputs/cuda.h" |
|---|
| 10 | |
|---|
| 11 | __attribute__((amdgpu_flat_work_group_size(32, 64))) // expected-no-diagnostics |
|---|
| 12 | __global__ void flat_work_group_size_32_64() { |
|---|
| 13 | // CHECK: define amdgpu_kernel void @_Z26flat_work_group_size_32_64v() [[FLAT_WORK_GROUP_SIZE_32_64:#[0-9]+]] |
|---|
| 14 | } |
|---|
| 15 | __attribute__((amdgpu_waves_per_eu(2))) // expected-no-diagnostics |
|---|
| 16 | __global__ void waves_per_eu_2() { |
|---|
| 17 | // CHECK: define amdgpu_kernel void @_Z14waves_per_eu_2v() [[WAVES_PER_EU_2:#[0-9]+]] |
|---|
| 18 | } |
|---|
| 19 | __attribute__((amdgpu_num_sgpr(32))) // expected-no-diagnostics |
|---|
| 20 | __global__ void num_sgpr_32() { |
|---|
| 21 | // CHECK: define amdgpu_kernel void @_Z11num_sgpr_32v() [[NUM_SGPR_32:#[0-9]+]] |
|---|
| 22 | } |
|---|
| 23 | __attribute__((amdgpu_num_vgpr(64))) // expected-no-diagnostics |
|---|
| 24 | __global__ void num_vgpr_64() { |
|---|
| 25 | // CHECK: define amdgpu_kernel void @_Z11num_vgpr_64v() [[NUM_VGPR_64:#[0-9]+]] |
|---|
| 26 | } |
|---|
| 27 | |
|---|
| 28 | // Make sure this is silently accepted on other targets. |
|---|
| 29 | // NAMD-NOT: "amdgpu-flat-work-group-size" |
|---|
| 30 | // NAMD-NOT: "amdgpu-waves-per-eu" |
|---|
| 31 | // NAMD-NOT: "amdgpu-num-vgpr" |
|---|
| 32 | // NAMD-NOT: "amdgpu-num-sgpr" |
|---|
| 33 | |
|---|
| 34 | // CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="32,64" |
|---|
| 35 | // CHECK-DAG: attributes [[WAVES_PER_EU_2]] = { convergent noinline nounwind optnone "amdgpu-waves-per-eu"="2" |
|---|
| 36 | // CHECK-DAG: attributes [[NUM_SGPR_32]] = { convergent noinline nounwind optnone "amdgpu-num-sgpr"="32" |
|---|
| 37 | // CHECK-DAG: attributes [[NUM_VGPR_64]] = { convergent noinline nounwind optnone "amdgpu-num-vgpr"="64" |
|---|
| 38 | |
|---|