Index: clang/lib/Basic/Targets/AMDGPU.h =================================================================== --- clang/lib/Basic/Targets/AMDGPU.h +++ clang/lib/Basic/Targets/AMDGPU.h @@ -130,8 +130,34 @@ "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 'L': + Info.setRequiresImmediate(0, 65535); + return true; + case 'A': + case 'B': + case 'C': + Info.setRequiresImmediate(); + return true; + default: + break; + } + StringRef S(Name); - if (S == "A") { + + if (S == "Kf") { + Name++; + Info.setRequiresImmediate(-1, -1); + return true; + } + if (S == "DA" || S == "DB") { + Name++; Info.setRequiresImmediate(); return true; } @@ -203,6 +229,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 == "Kf" || S == "DA" || S == "DB") { + return std::string("^") + std::string(Constraint++, 2); + } + const char *Begin = Constraint; TargetInfo::ConstraintInfo Info("", ""); if (validateAsmConstraint(Constraint, Info)) Index: clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl =================================================================== --- clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl +++ clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl @@ -2,6 +2,7 @@ // RUN: %clang_cc1 -emit-llvm -O0 -o - -triple amdgcn %s | FileCheck %s typedef float float32 __attribute__((ext_vector_type(32))); +typedef short __attribute__((ext_vector_type(2))) short2; kernel void test_long(int arg0) { long v15_16; @@ -33,3 +34,45 @@ : "={a1}"(reg_a) : "{a1}"(reg_b)); } + +kernel void test_constraint_Kf_i32() { + const int m1 = -1; + int res; + // CHECK: call i32 asm sideeffect "v_mov_b32 $0, $1", "=v,^Kf"(i32 -1) + __asm volatile("v_mov_b32 %0, %1" : "=v"(res) : "Kf"(m1)); +} + +kernel void test_constraint_Kf_i64() { + const long m1 = -1; + int res; + // CHECK: call i32 asm sideeffect "v_mov_b32 $0, $1", "=v,^Kf"(i64 -1) + __asm volatile("v_mov_b32 %0, %1" : "=v"(res) : "Kf"(m1)); +} + +kernel void test_constraint_Kf_i16() { + const short m1 = -1; + int res; + // CHECK: call i32 asm sideeffect "v_mov_b32 $0, $1", "=v,^Kf"(i16 -1) + __asm volatile("v_mov_b32 %0, %1" : "=v"(res) : "Kf"(m1)); +} + +kernel void test_constraint_Kf_i16x2() { + const short2 m1 = (short2)(-1, -1); + int res; + // CHECK: call i32 asm sideeffect "v_mov_b32 $0, $1", "=v,^Kf"(<2 x i16> ) + __asm volatile("v_mov_b32 %0, %1" : "=v"(res) : "Kf"(m1)); +} + +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/test/Sema/inline-asm-validate-amdgpu.cl =================================================================== --- clang/test/Sema/inline-asm-validate-amdgpu.cl +++ clang/test/Sema/inline-asm-validate-amdgpu.cl @@ -5,7 +5,7 @@ kernel void test () { - int sgpr = 0, vgpr = 0, imm = 0; + int sgpr = 0, vgpr = 0, imm = 0, m1 = -1; // sgpr constraints __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "s" (imm) : ); @@ -18,9 +18,48 @@ // 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'}} + + // 'L' constraint (an immediate 15-bit unsigned integer) + __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "L" (imm) : ); + __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "L" (0) : ); + __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "L" (65535) : ); + __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "L" (-1) : ); // expected-error {{value '-1' out of range for constraint 'L'}} + __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "L" (65536) : ); // expected-error {{value '65536' out of range for constraint 'L'}} + + // '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) : ); + + // 'Kf' constraint (-1) + __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "Kf" (m1) : ); + __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "Kf" (-1) : ); + __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "Kf" (-2) : ); // expected-error {{value '-2' out of range for constraint 'Kf'}} + __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "Kf" (0) : ); // expected-error {{value '0' out of range for constraint 'Kf'}} + + // '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