This revision was automatically updated to reflect the committed changes.
Closed by commit rG53422e8b4f65: [AMDGPU] Added support of new inline assembler
constraints (authored by dp).
Herald added a project: clang.
Herald added a subscriber: cfe-commits.
Changed prior to commit:
https://reviews.llvm.org/D81657?vs=271332&id=275400#toc
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D81657/new/
https://reviews.llvm.org/D81657
Files:
clang/lib/Basic/Targets/AMDGPU.h
clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl
clang/test/Sema/inline-asm-validate-amdgpu.cl
Index: clang/test/Sema/inline-asm-validate-amdgpu.cl
===================================================================
--- clang/test/Sema/inline-asm-validate-amdgpu.cl
+++ clang/test/Sema/inline-asm-validate-amdgpu.cl
@@ -18,9 +18,35 @@
// vgpr constraints
__asm__ ("v_mov_b32 %0, %1" : "=v" (vgpr) : "v" (imm) : );
- // 'A' constraint
+ // 'I' constraint (an immediate integer in the range -16 to 64)
+ __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "I" (imm) : );
+ __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "I" (-16) : );
+ __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "I" (64) : );
+ __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "I" (-17) : ); // expected-error {{value '-17' out of range for constraint 'I'}}
+ __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "I" (65) : ); // expected-error {{value '65' out of range for constraint 'I'}}
+
+ // 'J' constraint (an immediate 16-bit signed integer)
+ __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "J" (imm) : );
+ __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "J" (-32768) : );
+ __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "J" (32767) : );
+ __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "J" (-32769) : ); // expected-error {{value '-32769' out of range for constraint 'J'}}
+ __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "J" (32768) : ); // expected-error {{value '32768' out of range for constraint 'J'}}
+
+ // 'A' constraint (an immediate constant that can be inlined)
__asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "A" (imm) : );
+ // 'B' constraint (an immediate 32-bit signed integer)
+ __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "B" (imm) : );
+
+ // 'C' constraint (an immediate 32-bit unsigned integer or 'A' constraint)
+ __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "C" (imm) : );
+
+ // 'DA' constraint (an immediate 64-bit constant that can be split into two 'A' constants)
+ __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "DA" (imm) : );
+
+ // 'DB' constraint (an immediate 64-bit constant that can be split into two 'B' constants)
+ __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "DB" (imm) : );
+
}
__kernel void
Index: clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl
===================================================================
--- clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl
+++ clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl
@@ -33,3 +33,17 @@
: "={a1}"(reg_a)
: "{a1}"(reg_b));
}
+
+kernel void test_constraint_DA() {
+ const long x = 0x200000001;
+ int res;
+ // CHECK: call i32 asm sideeffect "v_mov_b32 $0, $1 & 0xFFFFFFFF", "=v,^DA"(i64 8589934593)
+ __asm volatile("v_mov_b32 %0, %1 & 0xFFFFFFFF" : "=v"(res) : "DA"(x));
+}
+
+kernel void test_constraint_DB() {
+ const long x = 0x200000001;
+ int res;
+ // CHECK: call i32 asm sideeffect "v_mov_b32 $0, $1 & 0xFFFFFFFF", "=v,^DB"(i64 8589934593)
+ __asm volatile("v_mov_b32 %0, %1 & 0xFFFFFFFF" : "=v"(res) : "DB"(x));
+}
Index: clang/lib/Basic/Targets/AMDGPU.h
===================================================================
--- clang/lib/Basic/Targets/AMDGPU.h
+++ clang/lib/Basic/Targets/AMDGPU.h
@@ -130,8 +130,26 @@
"exec_hi", "tma_lo", "tma_hi", "tba_lo", "tba_hi",
});
+ switch (*Name) {
+ case 'I':
+ Info.setRequiresImmediate(-16, 64);
+ return true;
+ case 'J':
+ Info.setRequiresImmediate(-32768, 32767);
+ return true;
+ case 'A':
+ case 'B':
+ case 'C':
+ Info.setRequiresImmediate();
+ return true;
+ default:
+ break;
+ }
+
StringRef S(Name);
- if (S == "A") {
+
+ if (S == "DA" || S == "DB") {
+ Name++;
Info.setRequiresImmediate();
return true;
}
@@ -203,6 +221,12 @@
// the constraint. In practice, it won't be changed unless the
// constraint is longer than one character.
std::string convertConstraint(const char *&Constraint) const override {
+
+ StringRef S(Constraint);
+ if (S == "DA" || S == "DB") {
+ return std::string("^") + std::string(Constraint++, 2);
+ }
+
const char *Begin = Constraint;
TargetInfo::ConstraintInfo Info("", "");
if (validateAsmConstraint(Constraint, Info))
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
