Skip to content

Commit d035d05

Browse files
Igor Sugakfacebook-github-bot
Igor Sugak
authored andcommitted
[pytorch] expose __ldg(const Half* ptr) to Clang in host mode (pytorch#38151)
Summary: Pull Request resolved: pytorch#38151 We need to expose this method to Clang unconditionally when building CUDA, otherwise it would error on device code calling `__ldg` with `Half*`. Test Plan: ``` buck build -c fbcode.caffe2_use_mpi=1 -c fbcode.cuda_use_clang=true mode/opt //experimental/training_supercomputer/trainer/hpc_pt:trainer ``` Reviewed By: ngimel Differential Revision: D21481297 fbshipit-source-id: aacfe7de2cdc8542908249081ddb58170b1e35ff
1 parent f3f3097 commit d035d05

File tree

1 file changed

+2
-1
lines changed

1 file changed

+2
-1
lines changed

c10/util/Half-inl.h

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -45,7 +45,8 @@ inline C10_HOST_DEVICE Half::operator __half() const {
4545

4646
// CUDA intrinsics
4747

48-
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 350)
48+
#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 350)) || \
49+
(defined(__clang__) && defined(__CUDA__))
4950
inline __device__ Half __ldg(const Half* ptr) {
5051
return __ldg(reinterpret_cast<const __half*>(ptr));
5152
}

0 commit comments

Comments
 (0)