diff options
| -rw-r--r-- | clang/include/clang/Basic/BuiltinsAMDGPU.def | 2 | ||||
| -rw-r--r-- | clang/lib/Basic/Targets/AMDGPU.cpp | 2 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/amdgpu-features.cl | 16 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl | 38 | ||||
| -rw-r--r-- | clang/test/SemaOpenCL/builtins-amdgcn-error-flat-address-space.cl | 8 |
5 files changed, 53 insertions, 13 deletions
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 72f20b208d5..9b3a0f96798 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -132,6 +132,8 @@ BUILTIN(__builtin_amdgcn_mqsad_u32_u8, "V4UiLUiUiV4Ui", "nc") TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_vol, "v", "n", "ci-insts") TARGET_BUILTIN(__builtin_amdgcn_buffer_wbinvl1_vol, "v", "n", "ci-insts") TARGET_BUILTIN(__builtin_amdgcn_ds_gws_sema_release_all, "vUi", "n", "ci-insts") +TARGET_BUILTIN(__builtin_amdgcn_is_shared, "bvC*0", "nc", "flat-address-space") +TARGET_BUILTIN(__builtin_amdgcn_is_private, "bvC*0", "nc", "flat-address-space") //===----------------------------------------------------------------------===// // Interpolation builtins. diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp index ebc9b9d60c2..07ba71d5c2b 100644 --- a/clang/lib/Basic/Targets/AMDGPU.cpp +++ b/clang/lib/Basic/Targets/AMDGPU.cpp @@ -142,6 +142,7 @@ bool AMDGPUTargetInfo::initFeatureMap( case GK_GFX1010: Features["dl-insts"] = true; Features["ci-insts"] = true; + Features["flat-address-space"] = true; Features["16-bit-insts"] = true; Features["dpp"] = true; Features["gfx8-insts"] = true; @@ -181,6 +182,7 @@ bool AMDGPUTargetInfo::initFeatureMap( case GK_GFX701: case GK_GFX700: Features["ci-insts"] = true; + Features["flat-address-space"] = true; LLVM_FALLTHROUGH; case GK_GFX601: case GK_GFX600: diff --git a/clang/test/CodeGenOpenCL/amdgpu-features.cl b/clang/test/CodeGenOpenCL/amdgpu-features.cl index 0bb3d6f3df7..fdc8f40b721 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-features.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-features.cl @@ -14,14 +14,14 @@ // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx600 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX600 %s // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx601 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX601 %s -// GFX904: "target-features"="+16-bit-insts,+ci-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime" -// GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime" -// GFX908: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime" -// GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime" -// GFX1011: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime" -// GFX1012: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime" -// GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+s-memrealtime" -// GFX700: "target-features"="+ci-insts,+fp64-fp16-denormals,-fp32-denormals" +// GFX904: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime" +// GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime" +// GFX908: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime" +// GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime" +// GFX1011: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime" +// GFX1012: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime" +// GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+s-memrealtime" +// GFX700: "target-features"="+ci-insts,+flat-address-space,+fp64-fp16-denormals,-fp32-denormals" // GFX600: "target-features"="+fp64-fp16-denormals,-fp32-denormals" // GFX601: "target-features"="+fp64-fp16-denormals,-fp32-denormals" diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl index f6fcfa58300..b3bf3681b93 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl @@ -1,8 +1,8 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu hawaii -S -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu fiji -S -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu hawaii -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu fiji -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s typedef unsigned int uint; @@ -21,8 +21,36 @@ void test_buffer_wbinvl1_vol() } // CHECK-LABEL: @test_gws_sema_release_all( -// CHECK: call void @llvm.amdgcn.ds.gws.sema.release.all(i32 %id) +// CHECK: call void @llvm.amdgcn.ds.gws.sema.release.all(i32 %{{[0-9]+}}) void test_gws_sema_release_all(uint id) { __builtin_amdgcn_ds_gws_sema_release_all(id); } + +// CHECK-LABEL: @test_is_shared( +// CHECK: [[CAST:%[0-9]+]] = bitcast i32* %{{[0-9]+}} to i8* +// CHECK: call i1 @llvm.amdgcn.is.shared(i8* [[CAST]] +int test_is_shared(const int* ptr) { + return __builtin_amdgcn_is_shared(ptr); +} + +// CHECK-LABEL: @test_is_private( +// CHECK: [[CAST:%[0-9]+]] = bitcast i32* %{{[0-9]+}} to i8* +// CHECK: call i1 @llvm.amdgcn.is.private(i8* [[CAST]] +int test_is_private(const int* ptr) { + return __builtin_amdgcn_is_private(ptr); +} + +// CHECK-LABEL: @test_is_shared_global( +// CHECK: [[CAST:%[0-9]+]] = addrspacecast i32 addrspace(1)* %{{[0-9]+}} to i8* +// CHECK: call i1 @llvm.amdgcn.is.shared(i8* [[CAST]] +int test_is_shared_global(const global int* ptr) { + return __builtin_amdgcn_is_shared(ptr); +} + +// CHECK-LABEL: @test_is_private_global( +// CHECK: [[CAST:%[0-9]+]] = addrspacecast i32 addrspace(1)* %{{[0-9]+}} to i8* +// CHECK: call i1 @llvm.amdgcn.is.private(i8* [[CAST]] +int test_is_private_global(const global int* ptr) { + return __builtin_amdgcn_is_private(ptr); +} diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-flat-address-space.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-flat-address-space.cl new file mode 100644 index 00000000000..0d759bac74a --- /dev/null +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-flat-address-space.cl @@ -0,0 +1,8 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s + +void test_flat_address_space_builtins(int* ptr) +{ + (void)__builtin_amdgcn_is_shared(ptr); // expected-error {{'__builtin_amdgcn_is_shared' needs target feature flat-address-space}} + (void)__builtin_amdgcn_is_private(ptr); // expected-error {{'__builtin_amdgcn_is_private' needs target feature flat-address-space}} +} |
