[AMDGPU] Allow flexible register names in inline asm constraints
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} or {v[n]}
{sn} or {s[n]}
{S} , where S is a special register name
{v[n:m]}
{s[n:m]}
Differential Revision: https://reviews.llvm.org/D37568
llvm-svn: 314452
This commit is contained in:
Yaxun Liu
parent
d6218cc385
commit
304f349770
|
|
@ -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,83 @@ public:
|
|||
return None;
|
||||
}
|
||||
|
||||
/// Accepted register names: (n, m is unsigned integer, n < m)
|
||||
/// v
|
||||
/// s
|
||||
/// {vn}, {v[n]}
|
||||
/// {sn}, {s[n]}
|
||||
/// {S} , where S is a special register name
|
||||
////{v[n:m]}
|
||||
/// {s[n: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.empty())
|
||||
return false;
|
||||
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;
|
||||
}
|
||||
return false;
|
||||
S = S.drop_front();
|
||||
if (!HasLeftParen) {
|
||||
if (!S.empty())
|
||||
return false;
|
||||
// Found s or v.
|
||||
Info.setAllowsRegister();
|
||||
Name = S.data() - 1;
|
||||
return true;
|
||||
}
|
||||
bool HasLeftBracket = false;
|
||||
if (!S.empty() && S.front() == '[') {
|
||||
HasLeftBracket = true;
|
||||
S = S.drop_front();
|
||||
}
|
||||
unsigned long long N;
|
||||
if (S.empty() || consumeUnsignedInteger(S, 10, N))
|
||||
return false;
|
||||
if (!S.empty() && S.front() == ':') {
|
||||
if (!HasLeftBracket)
|
||||
return false;
|
||||
S = S.drop_front();
|
||||
unsigned long long M;
|
||||
if (consumeUnsignedInteger(S, 10, M) || N >= M)
|
||||
return false;
|
||||
}
|
||||
if (HasLeftBracket) {
|
||||
if (S.empty() || S.front() != ']')
|
||||
return false;
|
||||
S = S.drop_front();
|
||||
}
|
||||
if (S.empty() || S.front() != '}')
|
||||
return false;
|
||||
S = S.drop_front();
|
||||
if (!S.empty())
|
||||
return false;
|
||||
// Found {vn}, {sn}, {v[n]}, {s[n]}, {v[n:m]}, or {s[n:m]}.
|
||||
Info.setAllowsRegister();
|
||||
Name = S.data() - 1;
|
||||
return true;
|
||||
}
|
||||
|
||||
bool
|
||||
|
|
|
|||
|
|
@ -1,6 +1,7 @@
|
|||
// 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
|
||||
|
||||
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
|
||||
|
||||
kernel void test () {
|
||||
|
||||
|
|
@ -9,6 +10,67 @@ kernel void test () {
|
|||
// sgpr constraints
|
||||
__asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "s" (imm) : );
|
||||
|
||||
__asm__ ("s_mov_b32 %0, %1" : "={s1}" (sgpr) : "{exec}" (imm) : );
|
||||
__asm__ ("s_mov_b32 %0, %1" : "={s1}" (sgpr) : "{exe" (imm) : ); // expected-error {{invalid input constraint '{exe' in asm}}
|
||||
__asm__ ("s_mov_b32 %0, %1" : "={s1}" (sgpr) : "{exec" (imm) : ); // expected-error {{invalid input constraint '{exec' in asm}}
|
||||
__asm__ ("s_mov_b32 %0, %1" : "={s1}" (sgpr) : "{exec}a" (imm) : ); // expected-error {{invalid input constraint '{exec}a' in asm}}
|
||||
|
||||
// vgpr constraints
|
||||
__asm__ ("v_mov_b32 %0, %1" : "=v" (vgpr) : "v" (imm) : );
|
||||
}
|
||||
|
||||
__kernel void
|
||||
test_float(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" : "={"(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" : "={v"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={v' 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" : "={va}"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={va}' 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" : "={v1"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={v1' in asm}}
|
||||
__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" : "={v[1]}"(ci) : "{v[2]}"(ai), "{v[3]}"(bi) : );
|
||||
__asm("v_add_f32_e32 v1, v2, v3" : "={v[1}"(ci) : "{v[2]}"(ai), "{v[3]}"(bi) : ); // expected-error {{invalid output constraint '={v[1}' in asm}}
|
||||
__asm("v_add_f32_e32 v1, v2, v3" : "={v[1]"(ci) : "{v[2]}"(ai), "{v[3]}"(bi) : ); // expected-error {{invalid output constraint '={v[1]' in asm}}
|
||||
__asm("v_add_f32_e32 v1, v2, v3" : "={v[a]}"(ci) : "{v[2]}"(ai), "{v[3]}"(bi) : ); // expected-error {{invalid output constraint '={v[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}"(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;
|
||||
}
|
||||
|
||||
__kernel void
|
||||
test_double(const __global double *a, const __global double *b, __global double *c, unsigned i)
|
||||
{
|
||||
double ai = a[i];
|
||||
double bi = b[i];
|
||||
double ci;
|
||||
|
||||
__asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[1:2]}"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : );
|
||||
__asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "=v{[1:2]}"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '=v{[1:2]}' in asm}}
|
||||
__asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[1:2]a}"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '={v[1:2]a}' in asm}}
|
||||
__asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[1:2]}a"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '={v[1:2]}a' in asm}}
|
||||
__asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[1:"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '={v[1:' in asm}}
|
||||
__asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[1:]}"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '={v[1:]}' in asm}}
|
||||
__asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[:2]}"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '={v[:2]}' in asm}}
|
||||
__asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[1:2]"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '={v[1:2]' in asm}}
|
||||
__asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[1:2}"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '={v[1:2}' in asm}}
|
||||
__asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[2:1]}"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '={v[2:1]}' in asm}}
|
||||
|
||||
__asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "=v[1:2]"(ci) : "v[3:4]"(ai), "v[5:6]"(bi) : ); //expected-error {{invalid output constraint '=v[1:2]' in asm}}
|
||||
|
||||
c[i] = ci;
|
||||
}
|
||||
|
|
|
|||
Loading…
Reference in New Issue