Author: arsenm
Date: Tue Aug 7 00:49:13 2018
New Revision: 339110
URL: http://llvm.org/viewvc/llvm-project?rev=339110&view=rev
Log:
AMDGPU: Add builtin for s_dcache_wb
Added:
cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-vi.cl
Modified:
cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
cfe/trunk/lib/Basic/Targets/AMDGPU.cpp
cfe/trunk/test/CodeGenOpenCL/amdgpu-features.cl
cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=339110&r1=339109&r2=339110&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Tue Aug 7 00:49:13 2018
@@ -121,6 +121,7 @@ TARGET_BUILTIN(__builtin_amdgcn_fracth,
TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime")
TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb", "nc", "dpp")
+TARGET_BUILTIN(__builtin_amdgcn_s_dcache_wb, "v", "n", "vi-insts")
//===----------------------------------------------------------------------===//
// GFX9+ only builtins.
Modified: cfe/trunk/lib/Basic/Targets/AMDGPU.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/AMDGPU.cpp?rev=339110&r1=339109&r2=339110&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/Targets/AMDGPU.cpp (original)
+++ cfe/trunk/lib/Basic/Targets/AMDGPU.cpp Tue Aug 7 00:49:13 2018
@@ -145,6 +145,7 @@ bool AMDGPUTargetInfo::initFeatureMap(
case GK_GFX803:
case GK_GFX802:
case GK_GFX801:
+ Features["vi-insts"] = true;
Features["16-bit-insts"] = true;
Features["dpp"] = true;
Features["s-memrealtime"] = true;
Modified: cfe/trunk/test/CodeGenOpenCL/amdgpu-features.cl
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/amdgpu-features.cl?rev=339110&r1=339109&r2=339110&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/amdgpu-features.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/amdgpu-features.cl Tue Aug 7 00:49:13 2018
@@ -10,9 +10,9 @@
// 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,+gfx9-insts,+s-memrealtime"
-// GFX906:
"target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx9-insts,+s-memrealtime"
-// GFX801:
"target-features"="+16-bit-insts,+ci-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+s-memrealtime"
+// GFX904:
"target-features"="+16-bit-insts,+ci-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx9-insts,+s-memrealtime,+vi-insts"
+// GFX906:
"target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx9-insts,+s-memrealtime,+vi-insts"
+// GFX801:
"target-features"="+16-bit-insts,+ci-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+s-memrealtime,+vi-insts"
// GFX700: "target-features"="+ci-insts,+fp64-fp16-denormals,-fp32-denormals"
// GFX600: "target-features"="+fp32-denormals,+fp64-fp16-denormals"
// GFX601: "target-features"="+fp64-fp16-denormals,-fp32-denormals"
Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl?rev=339110&r1=339109&r2=339110&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl Tue Aug 7 00:49:13 2018
@@ -82,6 +82,13 @@ void test_s_memrealtime(global ulong* ou
*out = __builtin_amdgcn_s_memrealtime();
}
+// CHECK-LABEL: @test_s_dcache_wb()
+// CHECK: call void @llvm.amdgcn.s.dcache.wb()
+void test_s_dcache_wb()
+{
+ __builtin_amdgcn_s_dcache_wb();
+}
+
// CHECK-LABEL: @test_mov_dpp
// CHECK: call i32 @llvm.amdgcn.mov.dpp.i32(i32 %src, i32 0, i32 0, i32 0, i1
false)
void test_mov_dpp(global int* out, int src)
Added: cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-vi.cl
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-vi.cl?rev=339110&view=auto
==============================================================================
--- cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-vi.cl (added)
+++ cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-vi.cl Tue Aug 7 00:49:13
2018
@@ -0,0 +1,8 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu hawaii -verify -S -o - %s
+
+void test_vi_s_dcache_wb()
+{
+ __builtin_amdgcn_s_dcache_wb(); // expected-error
{{'__builtin_amdgcn_s_dcache_wb' needs target feature vi-insts}}
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
