Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[HIP] Support attribute hip_pinned_shadow
This patch introduces support of hip_pinned_shadow variable for HIP. A hip_pinned_shadow variable is a global variable with attribute hip_pinned_shadow. It has external linkage on device side and has no initializer. It has internal linkage on host side and has initializer or static constructor. It can be accessed in both device code and host code. This allows HIP runtime to implement support of HIP texture reference. Differential Revision: https://reviews.llvm.org/D62738 llvm-svn: 364381
- Loading branch information
Showing
12 changed files
with
120 additions
and
15 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
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
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
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
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,13 @@ | ||
// RUN: %clang_cc1 -fcuda-is-device -ast-dump -ast-dump-filter tex -x hip %s | FileCheck -strict-whitespace %s | ||
// RUN: %clang_cc1 -ast-dump -ast-dump-filter tex -x hip %s | FileCheck -strict-whitespace %s | ||
struct textureReference { | ||
int a; | ||
}; | ||
|
||
// CHECK: HIPPinnedShadowAttr | ||
template <class T, int texType, int hipTextureReadMode> | ||
struct texture : public textureReference { | ||
texture() { a = 1; } | ||
}; | ||
|
||
__attribute__((hip_pinned_shadow)) texture<float, 1, 1> tex; |
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,23 @@ | ||
// REQUIRES: amdgpu-registered-target | ||
|
||
// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -std=c++11 -fvisibility hidden -fapply-global-visibility-to-externs \ | ||
// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefixes=HIPDEV %s | ||
// RUN: %clang_cc1 -triple x86_64 -std=c++11 \ | ||
// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefixes=HIPHOST %s | ||
|
||
struct textureReference { | ||
int a; | ||
}; | ||
|
||
template <class T, int texType, int hipTextureReadMode> | ||
struct texture : public textureReference { | ||
texture() { a = 1; } | ||
}; | ||
|
||
__attribute__((hip_pinned_shadow)) texture<float, 2, 1> tex; | ||
// CUDADEV-NOT: @tex | ||
// CUDAHOST-NOT: call i32 @__hipRegisterVar{{.*}}@tex | ||
// HIPDEV: @tex = external addrspace(1) global %struct.texture | ||
// HIPDEV-NOT: declare{{.*}}void @_ZN7textureIfLi2ELi1EEC1Ev | ||
// HIPHOST: define{{.*}}@_ZN7textureIfLi2ELi1EEC1Ev | ||
// HIPHOST: call i32 @__hipRegisterVar{{.*}}@tex{{.*}}i32 0, i32 4, i32 0, i32 0) |
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
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
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,25 @@ | ||
// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -std=c++11 -fvisibility hidden -fapply-global-visibility-to-externs \ | ||
// RUN: -emit-llvm -o - -x hip %s -fsyntax-only -verify | ||
// RUN: %clang_cc1 -triple x86_64 -std=c++11 \ | ||
// RUN: -emit-llvm -o - -x hip %s -fsyntax-only -verify | ||
|
||
#define __device__ __attribute__((device)) | ||
#define __constant__ __attribute__((constant)) | ||
#define __hip_pinned_shadow__ __attribute((hip_pinned_shadow)) | ||
|
||
struct textureReference { | ||
int a; | ||
}; | ||
|
||
template <class T, int texType, int hipTextureReadMode> | ||
struct texture : public textureReference { | ||
texture() { a = 1; } | ||
}; | ||
|
||
__hip_pinned_shadow__ texture<float, 2, 1> tex; | ||
__device__ __hip_pinned_shadow__ texture<float, 2, 1> tex2; // expected-error{{'hip_pinned_shadow' and 'device' attributes are not compatible}} | ||
// expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables}} | ||
// expected-note@-2{{conflicting attribute is here}} | ||
__constant__ __hip_pinned_shadow__ texture<float, 2, 1> tex3; // expected-error{{'hip_pinned_shadow' and 'constant' attributes are not compatible}} | ||
// expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables}} | ||
// expected-note@-2{{conflicting attribute is here}} |