-
Notifications
You must be signed in to change notification settings - Fork 11k
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
AMDGPU: Add clamp bit to dot builtins
Differential Revision: https://reviews.llvm.org/D50011 llvm-svn: 338471
- Loading branch information
Showing
4 changed files
with
83 additions
and
28 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
25 changes: 25 additions & 0 deletions
25
clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err-clamp.cl
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,25 @@ | ||
// REQUIRES: amdgpu-registered-target | ||
|
||
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -verify -S -emit-llvm -o - %s | ||
|
||
typedef unsigned int uint; | ||
typedef half __attribute__((ext_vector_type(2))) half2; | ||
typedef short __attribute__((ext_vector_type(2))) short2; | ||
typedef unsigned short __attribute__((ext_vector_type(2))) ushort2; | ||
|
||
kernel void builtins_amdgcn_dl_insts_err( | ||
global float *fOut, global int *siOut, global uint *uiOut, | ||
half2 v2hA, half2 v2hB, float fC, | ||
short2 v2ssA, short2 v2ssB, int siA, int siB, int siC, | ||
ushort2 v2usA, ushort2 v2usB, uint uiA, uint uiB, uint uiC, uint isClamp) { | ||
fOut[0] = __builtin_amdgcn_fdot2(v2hA, v2hB, fC, isClamp == 0 ? false : true); // expected-error {{'__builtin_amdgcn_fdot2' must be a constant integer}} | ||
|
||
siOut[0] = __builtin_amdgcn_sdot2(v2ssA, v2ssB, siC, isClamp == 0 ? false : true); // expected-error {{'__builtin_amdgcn_sdot2' must be a constant integer}} | ||
uiOut[0] = __builtin_amdgcn_udot2(v2usA, v2usB, uiC, isClamp == 0 ? false : true); // expected-error {{'__builtin_amdgcn_udot2' must be a constant integer}} | ||
|
||
siOut[1] = __builtin_amdgcn_sdot4(siA, siB, siC, isClamp == 0 ? false : true); // expected-error {{'__builtin_amdgcn_sdot4' must be a constant integer}} | ||
uiOut[1] = __builtin_amdgcn_udot4(uiA, uiB, uiC, isClamp == 0 ? false : true); // expected-error {{'__builtin_amdgcn_udot4' must be a constant integer}} | ||
|
||
siOut[2] = __builtin_amdgcn_sdot8(siA, siB, siC, isClamp == 0 ? false : true); // expected-error {{'__builtin_amdgcn_sdot8' must be a constant integer}} | ||
uiOut[2] = __builtin_amdgcn_udot8(uiA, uiB, uiC, isClamp == 0 ? false : true); // expected-error {{'__builtin_amdgcn_udot8' must be a constant integer}} | ||
} |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters