This is an archive of the discontinued LLVM Phabricator instance.

Differential D28301

[CUDA] Pre-include sm_60 and sm_61 headers.
ClosedPublic

Authored by tra on Jan 4 2017, 10:05 AM.

Details

Reviewers
jlebar
Commits
rG60f25f70c862: [CUDA] Pre-include sm_60 and sm_61 headers.
rC290982: [CUDA] Pre-include sm_60 and sm_61 headers.
rL290982: [CUDA] Pre-include sm_60 and sm_61 headers.
Summary

CUDA-8.0 comes with new headers which nvcc pre-includes via cuda_runtime.h
Clang now makes them available as well.

Diff Detail

Event Timeline

tra updated this revision to Diff 83071.Jan 4 2017, 10:05 AM
tra retitled this revision from to [CUDA] Pre-include sm_60 and sm_61 headers..
tra updated this object.
tra added a reviewer: jlebar.
tra added a subscriber: cfe-commits.
jlebar accepted this revision.Jan 4 2017, 10:36 AM
jlebar edited edge metadata.
This revision is now accepted and ready to land.Jan 4 2017, 10:36 AM
Closed by commit rL290982: [CUDA] Pre-include sm_60 and sm_61 headers. (authored by tra). · Explain WhyJan 4 2017, 10:50 AM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
lib/
Headers/
5 lines

Diff 83071

lib/Headers/__clang_cuda_runtime_wrapper.h

Show First 20 Lines • Show All 226 Lines • ▼ Show 20 Lines
// Don't include sm_30_intrinsics.h and sm_32_intrinsics.h. These define the // Don't include sm_30_intrinsics.h and sm_32_intrinsics.h. These define the
// __shfl and __ldg intrinsics using inline (volatile) asm, but we want to // __shfl and __ldg intrinsics using inline (volatile) asm, but we want to
// define them using builtins so that the optimizer can reason about and across // define them using builtins so that the optimizer can reason about and across
// these instructions. In particular, using intrinsics for ldg gets us the // these instructions. In particular, using intrinsics for ldg gets us the
// [addr+imm] addressing mode, which, although it doesn't actually exist in the // [addr+imm] addressing mode, which, although it doesn't actually exist in the
// hardware, seems to generate faster machine code because ptxas can more easily // hardware, seems to generate faster machine code because ptxas can more easily
// reason about our code. // reason about our code.
#if CUDA_VERSION >= 8000
#include "sm_60_atomic_functions.hpp"
#include "sm_61_intrinsics.hpp"
#endif
#undef __MATH_FUNCTIONS_HPP__ #undef __MATH_FUNCTIONS_HPP__
// math_functions.hpp defines ::signbit as a __host__ __device__ function. This // math_functions.hpp defines ::signbit as a __host__ __device__ function. This
// conflicts with libstdc++'s constexpr ::signbit, so we have to rename // conflicts with libstdc++'s constexpr ::signbit, so we have to rename
// math_function.hpp's ::signbit. It's guarded by #undef signbit, but that's // math_function.hpp's ::signbit. It's guarded by #undef signbit, but that's
// conditional on __GNUC__. :) // conditional on __GNUC__. :)
#pragma push_macro("signbit") #pragma push_macro("signbit")
#pragma push_macro("__GNUC__") #pragma push_macro("__GNUC__")
▲ Show 20 Lines • Show All 100 Lines • Show Last 20 Lines