yaxunl created this revision.
Herald added subscribers: eraman, t-tye, tpr, dstuttard, nhaehnle, wdng,
kzhuravl.
Currently AMDGPU inline asm only allow "v" and "s" as register names in
constraints.
This patch allows the following register names in constraints: (n, m is
unsigned integer, n < m)
v
s
{vn}
{sn}
{S} , wheere S is a special register name
{vn:m}
{sn:m}
https://reviews.llvm.org/D37568
Files:
lib/Basic/Targets/AMDGPU.h
test/CodeGenOpenCL/amdgcn-inline-asm.cl
test/Sema/inline-asm-validate-amdgpu.cl
Index: test/Sema/inline-asm-validate-amdgpu.cl
===================================================================
--- test/Sema/inline-asm-validate-amdgpu.cl
+++ test/Sema/inline-asm-validate-amdgpu.cl
@@ -1,6 +1,5 @@
// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -x cl -triple amdgcn -fsyntax-only %s
-// expected-no-diagnostics
+// RUN: %clang_cc1 -triple amdgcn -fsyntax-only -verify %s
kernel void test () {
@@ -12,3 +11,33 @@
// vgpr constraints
__asm__ ("v_mov_b32 %0, %1" : "=v" (vgpr) : "v" (imm) : );
}
+
+__kernel void
+ker(const __global float *a, const __global float *b, __global float *c, unsigned i)
+{
+ float ai = a[i];
+ float bi = b[i];
+ float ci;
+
+ __asm("v_add_f32_e32 v1, v2, v3" : "={v1}"(ci) : "{v2}"(ai), "{v3}"(bi) : );
+ __asm("v_add_f32_e32 v1, v2, v3" : ""(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '' in asm}}
+ __asm("v_add_f32_e32 v1, v2, v3" : "="(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '=' in asm}}
+ __asm("v_add_f32_e32 v1, v2, v3" : "={a}"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={a}' in asm}}
+ __asm("v_add_f32_e32 v1, v2, v3" : "={v1a}"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={v1a}' in asm}}
+ __asm("v_add_f32_e32 v1, v2, v3" : "={v1}a"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={v1}a' in asm}}
+
+ __asm("v_add_f32_e32 v1, v2, v3" : "={exec}"(ci) : "{v2}"(ai), "{v3}"(bi) : );
+ __asm("v_add_f32_e32 v1, v2, v3" : "={exec}a"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={exec}a' in asm}}
+
+ __asm("v_add_f32_e32 v1, v2, v3" : "={v1:2}"(ci) : "{v3:4}"(ai), "{v4:5}"(bi) : );
+ __asm("v_add_f32_e32 v1, v2, v3" : "=v{1:2}"(ci) : "{v3:4}"(ai), "{v4:5}"(bi) : ); //expected-error {{invalid output constraint '=v{1:2}' in asm}}
+ __asm("v_add_f32_e32 v1, v2, v3" : "={v1:2}a"(ci) : "{v3:4}"(ai), "{v4:5}"(bi) : ); //expected-error {{invalid output constraint '={v1:2}a' in asm}}
+
+ __asm("v_add_f32_e32 v1, v2, v3" : "=v"(ci) : "v"(ai), "v"(bi) : );
+ __asm("v_add_f32_e32 v1, v2, v3" : "=v1"(ci) : "v2"(ai), "v3"(bi) : ); /// expected-error {{invalid output constraint '=v1' in asm}}
+ __asm("v_add_f32_e32 v1, v2, v3" : "=v1:2"(ci) : "v3:4"(ai), "v4:5"(bi) : ); // expected-error {{invalid output constraint '=v1:2' in asm}}
+
+ __asm("v_add_f32_e32 v1, v2, v3" : "={v1}"(ci) : "{a}"(ai), "{v3}"(bi) : ); // expected-error {{invalid input constraint '{a}' in asm}}
+ __asm("v_add_f32_e32 v1, v2, v3" : "={v1}"(ci) : "{v2}"(ai), "{a}"(bi) : ); // expected-error {{invalid input constraint '{a}' in asm}}
+ c[i] = ci;
+}
Index: test/CodeGenOpenCL/amdgcn-inline-asm.cl
===================================================================
--- /dev/null
+++ test/CodeGenOpenCL/amdgcn-inline-asm.cl
@@ -0,0 +1,16 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn -O0 -emit-llvm -o - %s | FileCheck %s
+
+// CHECK-LABEL: @ker
+__kernel void
+ker(const __global float *a, const __global float *b, __global float *c, unsigned i)
+{
+ float ai = a[i];
+ float bi = b[i];
+ float ci;
+ // CHECK: call float asm "v_add_f32_e32 v1, v2, v3", "={v1},{v2},{v3}"(float %{{.*}}, float %{{.*}})
+ __asm("v_add_f32_e32 v1, v2, v3" : "={v1}"(ci) : "{v2}"(ai), "{v3}"(bi) : );
+ // CHECK: call float asm "v_add_f32_e32 $0, $1, $2", "={v1},{v2},{v3}"(float %{{.*}}, float %{{.*}})
+ __asm("v_add_f32_e32 %0, %1, %2" : "={v1}"(ci) : "{v2}"(ai), "{v3}"(bi) : );
+ c[i] = ci;
+}
Index: lib/Basic/Targets/AMDGPU.h
===================================================================
--- lib/Basic/Targets/AMDGPU.h
+++ lib/Basic/Targets/AMDGPU.h
@@ -17,6 +17,7 @@
#include "clang/AST/Type.h"
#include "clang/Basic/TargetInfo.h"
#include "clang/Basic/TargetOptions.h"
+#include "llvm/ADT/StringSet.h"
#include "llvm/ADT/Triple.h"
#include "llvm/Support/Compiler.h"
@@ -115,17 +116,70 @@
return None;
}
+ /// Accepted register names: (n, m is unsigned integer, n < m)
+ /// v
+ /// s
+ /// {vn}
+ /// {sn}
+ /// {S} , wheere S is a special register name
+ ////{vn:m}
+ /// {sn:m}
bool validateAsmConstraint(const char *&Name,
TargetInfo::ConstraintInfo &Info) const override {
- switch (*Name) {
- default:
- break;
- case 'v': // vgpr
- case 's': // sgpr
+ static const ::llvm::StringSet<> SpecialRegs({
+ "exec", "vcc", "flat_scratch", "m0", "scc", "tba", "tma",
+ "flat_scratch_lo", "flat_scratch_hi", "vcc_lo", "vcc_hi", "exec_lo",
+ "exec_hi", "tma_lo", "tma_hi", "tba_lo", "tba_hi",
+ });
+
+ StringRef S(Name);
+ bool HasLeftParen = false;
+ if (S.front() == '{') {
+ HasLeftParen = true;
+ S = S.drop_front();
+ }
+ if (S.front() != 'v' && S.front() != 's') {
+ if (!HasLeftParen)
+ return false;
+ auto E = S.find('}');
+ if (!SpecialRegs.count(S.substr(0, E)))
+ return false;
+ S = S.drop_front(E + 1);
+ if (!S.empty())
+ return false;
+ // Found {S} where S is a special register.
+ Info.setAllowsRegister();
+ Name = S.data() - 1;
+ return true;
+ }
+ S = S.drop_front();
+ if (!HasLeftParen) {
+ if (!S.empty())
+ return false;
+ // Found s or v.
Info.setAllowsRegister();
+ Name = S.data() - 1;
return true;
}
- return false;
+ unsigned long long N;
+ if (consumeUnsignedInteger(S, 10, N))
+ return false;
+ if (S.front() == ':') {
+ S = S.drop_front();
+ unsigned long long M;
+ // Found {sn:m} or {vn:m}.
+ if (consumeUnsignedInteger(S, 10, M) || N >= M)
+ return false;
+ }
+ if (S.front() != '}')
+ return false;
+ S = S.drop_front();
+ if (!S.empty())
+ return false;
+ // Found {sn} or {sm}
+ Info.setAllowsRegister();
+ Name = S.data() - 1;
+ return true;
}
bool
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
