995 lines · c
1/*===--- __clang_cuda_intrinsics.h - Device-side CUDA intrinsic wrappers ---===2 *3 * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4 * See https://llvm.org/LICENSE.txt for license information.5 * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6 *7 *===-----------------------------------------------------------------------===8 */9#ifndef __CLANG_CUDA_INTRINSICS_H__10#define __CLANG_CUDA_INTRINSICS_H__11#ifndef __CUDA__12#error "This file is for CUDA compilation only."13#endif14 15// sm_30 intrinsics: __shfl_{up,down,xor}.16 17#define __SM_30_INTRINSICS_H__18#define __SM_30_INTRINSICS_HPP__19 20#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 30021 22#pragma push_macro("__MAKE_SHUFFLES")23#define __MAKE_SHUFFLES(__FnName, __IntIntrinsic, __FloatIntrinsic, __Mask, \24 __Type) \25 inline __device__ int __FnName(int __val, __Type __offset, \26 int __width = warpSize) { \27 return __IntIntrinsic(__val, __offset, \28 ((warpSize - __width) << 8) | (__Mask)); \29 } \30 inline __device__ float __FnName(float __val, __Type __offset, \31 int __width = warpSize) { \32 return __FloatIntrinsic(__val, __offset, \33 ((warpSize - __width) << 8) | (__Mask)); \34 } \35 inline __device__ unsigned int __FnName(unsigned int __val, __Type __offset, \36 int __width = warpSize) { \37 return static_cast<unsigned int>( \38 ::__FnName(static_cast<int>(__val), __offset, __width)); \39 } \40 inline __device__ long long __FnName(long long __val, __Type __offset, \41 int __width = warpSize) { \42 struct __Bits { \43 int __a, __b; \44 }; \45 _Static_assert(sizeof(__val) == sizeof(__Bits)); \46 _Static_assert(sizeof(__Bits) == 2 * sizeof(int)); \47 __Bits __tmp; \48 memcpy(&__tmp, &__val, sizeof(__val)); \49 __tmp.__a = ::__FnName(__tmp.__a, __offset, __width); \50 __tmp.__b = ::__FnName(__tmp.__b, __offset, __width); \51 long long __ret; \52 memcpy(&__ret, &__tmp, sizeof(__tmp)); \53 return __ret; \54 } \55 inline __device__ long __FnName(long __val, __Type __offset, \56 int __width = warpSize) { \57 _Static_assert(sizeof(long) == sizeof(long long) || \58 sizeof(long) == sizeof(int)); \59 if (sizeof(long) == sizeof(long long)) { \60 return static_cast<long>( \61 ::__FnName(static_cast<long long>(__val), __offset, __width)); \62 } else if (sizeof(long) == sizeof(int)) { \63 return static_cast<long>( \64 ::__FnName(static_cast<int>(__val), __offset, __width)); \65 } \66 } \67 inline __device__ unsigned long __FnName( \68 unsigned long __val, __Type __offset, int __width = warpSize) { \69 return static_cast<unsigned long>( \70 ::__FnName(static_cast<long>(__val), __offset, __width)); \71 } \72 inline __device__ unsigned long long __FnName( \73 unsigned long long __val, __Type __offset, int __width = warpSize) { \74 return static_cast<unsigned long long>( \75 ::__FnName(static_cast<long long>(__val), __offset, __width)); \76 } \77 inline __device__ double __FnName(double __val, __Type __offset, \78 int __width = warpSize) { \79 long long __tmp; \80 _Static_assert(sizeof(__tmp) == sizeof(__val)); \81 memcpy(&__tmp, &__val, sizeof(__val)); \82 __tmp = ::__FnName(__tmp, __offset, __width); \83 double __ret; \84 memcpy(&__ret, &__tmp, sizeof(__ret)); \85 return __ret; \86 }87 88__MAKE_SHUFFLES(__shfl, __nvvm_shfl_idx_i32, __nvvm_shfl_idx_f32, 0x1f, int);89// We use 0 rather than 31 as our mask, because shfl.up applies to lanes >=90// maxLane.91__MAKE_SHUFFLES(__shfl_up, __nvvm_shfl_up_i32, __nvvm_shfl_up_f32, 0,92 unsigned int);93__MAKE_SHUFFLES(__shfl_down, __nvvm_shfl_down_i32, __nvvm_shfl_down_f32, 0x1f,94 unsigned int);95__MAKE_SHUFFLES(__shfl_xor, __nvvm_shfl_bfly_i32, __nvvm_shfl_bfly_f32, 0x1f,96 int);97#pragma pop_macro("__MAKE_SHUFFLES")98 99#endif // !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300100 101#if CUDA_VERSION >= 9000102#if (!defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300)103// __shfl_sync_* variants available in CUDA-9104#pragma push_macro("__MAKE_SYNC_SHUFFLES")105#define __MAKE_SYNC_SHUFFLES(__FnName, __IntIntrinsic, __FloatIntrinsic, \106 __Mask, __Type) \107 inline __device__ int __FnName(unsigned int __mask, int __val, \108 __Type __offset, int __width = warpSize) { \109 return __IntIntrinsic(__mask, __val, __offset, \110 ((warpSize - __width) << 8) | (__Mask)); \111 } \112 inline __device__ float __FnName(unsigned int __mask, float __val, \113 __Type __offset, int __width = warpSize) { \114 return __FloatIntrinsic(__mask, __val, __offset, \115 ((warpSize - __width) << 8) | (__Mask)); \116 } \117 inline __device__ unsigned int __FnName(unsigned int __mask, \118 unsigned int __val, __Type __offset, \119 int __width = warpSize) { \120 return static_cast<unsigned int>( \121 ::__FnName(__mask, static_cast<int>(__val), __offset, __width)); \122 } \123 inline __device__ long long __FnName(unsigned int __mask, long long __val, \124 __Type __offset, \125 int __width = warpSize) { \126 struct __Bits { \127 int __a, __b; \128 }; \129 _Static_assert(sizeof(__val) == sizeof(__Bits)); \130 _Static_assert(sizeof(__Bits) == 2 * sizeof(int)); \131 __Bits __tmp; \132 memcpy(&__tmp, &__val, sizeof(__val)); \133 __tmp.__a = ::__FnName(__mask, __tmp.__a, __offset, __width); \134 __tmp.__b = ::__FnName(__mask, __tmp.__b, __offset, __width); \135 long long __ret; \136 memcpy(&__ret, &__tmp, sizeof(__tmp)); \137 return __ret; \138 } \139 inline __device__ unsigned long long __FnName( \140 unsigned int __mask, unsigned long long __val, __Type __offset, \141 int __width = warpSize) { \142 return static_cast<unsigned long long>( \143 ::__FnName(__mask, static_cast<long long>(__val), __offset, __width)); \144 } \145 inline __device__ long __FnName(unsigned int __mask, long __val, \146 __Type __offset, int __width = warpSize) { \147 _Static_assert(sizeof(long) == sizeof(long long) || \148 sizeof(long) == sizeof(int)); \149 if (sizeof(long) == sizeof(long long)) { \150 return static_cast<long>(::__FnName( \151 __mask, static_cast<long long>(__val), __offset, __width)); \152 } else if (sizeof(long) == sizeof(int)) { \153 return static_cast<long>( \154 ::__FnName(__mask, static_cast<int>(__val), __offset, __width)); \155 } \156 } \157 inline __device__ unsigned long __FnName( \158 unsigned int __mask, unsigned long __val, __Type __offset, \159 int __width = warpSize) { \160 return static_cast<unsigned long>( \161 ::__FnName(__mask, static_cast<long>(__val), __offset, __width)); \162 } \163 inline __device__ double __FnName(unsigned int __mask, double __val, \164 __Type __offset, int __width = warpSize) { \165 long long __tmp; \166 _Static_assert(sizeof(__tmp) == sizeof(__val)); \167 memcpy(&__tmp, &__val, sizeof(__val)); \168 __tmp = ::__FnName(__mask, __tmp, __offset, __width); \169 double __ret; \170 memcpy(&__ret, &__tmp, sizeof(__ret)); \171 return __ret; \172 }173__MAKE_SYNC_SHUFFLES(__shfl_sync, __nvvm_shfl_sync_idx_i32,174 __nvvm_shfl_sync_idx_f32, 0x1f, int);175// We use 0 rather than 31 as our mask, because shfl.up applies to lanes >=176// maxLane.177__MAKE_SYNC_SHUFFLES(__shfl_up_sync, __nvvm_shfl_sync_up_i32,178 __nvvm_shfl_sync_up_f32, 0, unsigned int);179__MAKE_SYNC_SHUFFLES(__shfl_down_sync, __nvvm_shfl_sync_down_i32,180 __nvvm_shfl_sync_down_f32, 0x1f, unsigned int);181__MAKE_SYNC_SHUFFLES(__shfl_xor_sync, __nvvm_shfl_sync_bfly_i32,182 __nvvm_shfl_sync_bfly_f32, 0x1f, int);183#pragma pop_macro("__MAKE_SYNC_SHUFFLES")184 185inline __device__ void __syncwarp(unsigned int mask = 0xffffffff) {186 return __nvvm_bar_warp_sync(mask);187}188 189inline __device__ void __barrier_sync(unsigned int id) {190 __nvvm_barrier_sync(id);191}192 193inline __device__ void __barrier_sync_count(unsigned int id,194 unsigned int count) {195 __nvvm_barrier_sync_cnt(id, count);196}197 198inline __device__ int __all_sync(unsigned int mask, int pred) {199 return __nvvm_vote_all_sync(mask, pred);200}201 202inline __device__ int __any_sync(unsigned int mask, int pred) {203 return __nvvm_vote_any_sync(mask, pred);204}205 206inline __device__ int __uni_sync(unsigned int mask, int pred) {207 return __nvvm_vote_uni_sync(mask, pred);208}209 210inline __device__ unsigned int __ballot_sync(unsigned int mask, int pred) {211 return __nvvm_vote_ballot_sync(mask, pred);212}213 214inline __device__ unsigned int __activemask() {215#if CUDA_VERSION < 9020216 return __nvvm_vote_ballot(1);217#else218 return __nvvm_activemask();219#endif220}221 222inline __device__ unsigned int __fns(unsigned mask, unsigned base, int offset) {223 return __nvvm_fns(mask, base, offset);224}225 226#endif // !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300227 228// Define __match* builtins CUDA-9 headers expect to see.229#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 700230inline __device__ unsigned int __match32_any_sync(unsigned int mask,231 unsigned int value) {232 return __nvvm_match_any_sync_i32(mask, value);233}234 235inline __device__ unsigned int236__match64_any_sync(unsigned int mask, unsigned long long value) {237 return __nvvm_match_any_sync_i64(mask, value);238}239 240inline __device__ unsigned int241__match32_all_sync(unsigned int mask, unsigned int value, int *pred) {242 return __nvvm_match_all_sync_i32p(mask, value, pred);243}244 245inline __device__ unsigned int246__match64_all_sync(unsigned int mask, unsigned long long value, int *pred) {247 return __nvvm_match_all_sync_i64p(mask, value, pred);248}249#include "crt/sm_70_rt.hpp"250 251#endif // !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 700252#endif // __CUDA_VERSION >= 9000253 254// sm_32 intrinsics: __ldg and __funnelshift_{l,lc,r,rc}.255 256// Prevent the vanilla sm_32 intrinsics header from being included.257#define __SM_32_INTRINSICS_H__258#define __SM_32_INTRINSICS_HPP__259 260#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 320261 262inline __device__ char __ldg(const char *ptr) { return __nvvm_ldg_c(ptr); }263inline __device__ short __ldg(const short *ptr) { return __nvvm_ldg_s(ptr); }264inline __device__ int __ldg(const int *ptr) { return __nvvm_ldg_i(ptr); }265inline __device__ long __ldg(const long *ptr) { return __nvvm_ldg_l(ptr); }266inline __device__ long long __ldg(const long long *ptr) {267 return __nvvm_ldg_ll(ptr);268}269inline __device__ unsigned char __ldg(const unsigned char *ptr) {270 return __nvvm_ldg_uc(ptr);271}272inline __device__ signed char __ldg(const signed char *ptr) {273 return __nvvm_ldg_uc((const unsigned char *)ptr);274}275inline __device__ unsigned short __ldg(const unsigned short *ptr) {276 return __nvvm_ldg_us(ptr);277}278inline __device__ unsigned int __ldg(const unsigned int *ptr) {279 return __nvvm_ldg_ui(ptr);280}281inline __device__ unsigned long __ldg(const unsigned long *ptr) {282 return __nvvm_ldg_ul(ptr);283}284inline __device__ unsigned long long __ldg(const unsigned long long *ptr) {285 return __nvvm_ldg_ull(ptr);286}287inline __device__ float __ldg(const float *ptr) { return __nvvm_ldg_f(ptr); }288inline __device__ double __ldg(const double *ptr) { return __nvvm_ldg_d(ptr); }289 290inline __device__ char2 __ldg(const char2 *ptr) {291 typedef char c2 __attribute__((ext_vector_type(2)));292 // We can assume that ptr is aligned at least to char2's alignment, but the293 // load will assume that ptr is aligned to char2's alignment. This is only294 // safe if alignof(c2) <= alignof(char2).295 c2 rv = __nvvm_ldg_c2(reinterpret_cast<const c2 *>(ptr));296 char2 ret;297 ret.x = rv[0];298 ret.y = rv[1];299 return ret;300}301inline __device__ char4 __ldg(const char4 *ptr) {302 typedef char c4 __attribute__((ext_vector_type(4)));303 c4 rv = __nvvm_ldg_c4(reinterpret_cast<const c4 *>(ptr));304 char4 ret;305 ret.x = rv[0];306 ret.y = rv[1];307 ret.z = rv[2];308 ret.w = rv[3];309 return ret;310}311inline __device__ short2 __ldg(const short2 *ptr) {312 typedef short s2 __attribute__((ext_vector_type(2)));313 s2 rv = __nvvm_ldg_s2(reinterpret_cast<const s2 *>(ptr));314 short2 ret;315 ret.x = rv[0];316 ret.y = rv[1];317 return ret;318}319inline __device__ short4 __ldg(const short4 *ptr) {320 typedef short s4 __attribute__((ext_vector_type(4)));321 s4 rv = __nvvm_ldg_s4(reinterpret_cast<const s4 *>(ptr));322 short4 ret;323 ret.x = rv[0];324 ret.y = rv[1];325 ret.z = rv[2];326 ret.w = rv[3];327 return ret;328}329inline __device__ int2 __ldg(const int2 *ptr) {330 typedef int i2 __attribute__((ext_vector_type(2)));331 i2 rv = __nvvm_ldg_i2(reinterpret_cast<const i2 *>(ptr));332 int2 ret;333 ret.x = rv[0];334 ret.y = rv[1];335 return ret;336}337inline __device__ int4 __ldg(const int4 *ptr) {338 typedef int i4 __attribute__((ext_vector_type(4)));339 i4 rv = __nvvm_ldg_i4(reinterpret_cast<const i4 *>(ptr));340 int4 ret;341 ret.x = rv[0];342 ret.y = rv[1];343 ret.z = rv[2];344 ret.w = rv[3];345 return ret;346}347inline __device__ longlong2 __ldg(const longlong2 *ptr) {348 typedef long long ll2 __attribute__((ext_vector_type(2)));349 ll2 rv = __nvvm_ldg_ll2(reinterpret_cast<const ll2 *>(ptr));350 longlong2 ret;351 ret.x = rv[0];352 ret.y = rv[1];353 return ret;354}355 356inline __device__ uchar2 __ldg(const uchar2 *ptr) {357 typedef unsigned char uc2 __attribute__((ext_vector_type(2)));358 uc2 rv = __nvvm_ldg_uc2(reinterpret_cast<const uc2 *>(ptr));359 uchar2 ret;360 ret.x = rv[0];361 ret.y = rv[1];362 return ret;363}364inline __device__ uchar4 __ldg(const uchar4 *ptr) {365 typedef unsigned char uc4 __attribute__((ext_vector_type(4)));366 uc4 rv = __nvvm_ldg_uc4(reinterpret_cast<const uc4 *>(ptr));367 uchar4 ret;368 ret.x = rv[0];369 ret.y = rv[1];370 ret.z = rv[2];371 ret.w = rv[3];372 return ret;373}374inline __device__ ushort2 __ldg(const ushort2 *ptr) {375 typedef unsigned short us2 __attribute__((ext_vector_type(2)));376 us2 rv = __nvvm_ldg_us2(reinterpret_cast<const us2 *>(ptr));377 ushort2 ret;378 ret.x = rv[0];379 ret.y = rv[1];380 return ret;381}382inline __device__ ushort4 __ldg(const ushort4 *ptr) {383 typedef unsigned short us4 __attribute__((ext_vector_type(4)));384 us4 rv = __nvvm_ldg_us4(reinterpret_cast<const us4 *>(ptr));385 ushort4 ret;386 ret.x = rv[0];387 ret.y = rv[1];388 ret.z = rv[2];389 ret.w = rv[3];390 return ret;391}392inline __device__ uint2 __ldg(const uint2 *ptr) {393 typedef unsigned int ui2 __attribute__((ext_vector_type(2)));394 ui2 rv = __nvvm_ldg_ui2(reinterpret_cast<const ui2 *>(ptr));395 uint2 ret;396 ret.x = rv[0];397 ret.y = rv[1];398 return ret;399}400inline __device__ uint4 __ldg(const uint4 *ptr) {401 typedef unsigned int ui4 __attribute__((ext_vector_type(4)));402 ui4 rv = __nvvm_ldg_ui4(reinterpret_cast<const ui4 *>(ptr));403 uint4 ret;404 ret.x = rv[0];405 ret.y = rv[1];406 ret.z = rv[2];407 ret.w = rv[3];408 return ret;409}410inline __device__ ulonglong2 __ldg(const ulonglong2 *ptr) {411 typedef unsigned long long ull2 __attribute__((ext_vector_type(2)));412 ull2 rv = __nvvm_ldg_ull2(reinterpret_cast<const ull2 *>(ptr));413 ulonglong2 ret;414 ret.x = rv[0];415 ret.y = rv[1];416 return ret;417}418 419inline __device__ float2 __ldg(const float2 *ptr) {420 typedef float f2 __attribute__((ext_vector_type(2)));421 f2 rv = __nvvm_ldg_f2(reinterpret_cast<const f2 *>(ptr));422 float2 ret;423 ret.x = rv[0];424 ret.y = rv[1];425 return ret;426}427inline __device__ float4 __ldg(const float4 *ptr) {428 typedef float f4 __attribute__((ext_vector_type(4)));429 f4 rv = __nvvm_ldg_f4(reinterpret_cast<const f4 *>(ptr));430 float4 ret;431 ret.x = rv[0];432 ret.y = rv[1];433 ret.z = rv[2];434 ret.w = rv[3];435 return ret;436}437inline __device__ double2 __ldg(const double2 *ptr) {438 typedef double d2 __attribute__((ext_vector_type(2)));439 d2 rv = __nvvm_ldg_d2(reinterpret_cast<const d2 *>(ptr));440 double2 ret;441 ret.x = rv[0];442 ret.y = rv[1];443 return ret;444}445 446// TODO: Implement these as intrinsics, so the backend can work its magic on447// these. Alternatively, we could implement these as plain C and try to get448// llvm to recognize the relevant patterns.449inline __device__ unsigned __funnelshift_l(unsigned low32, unsigned high32,450 unsigned shiftWidth) {451 unsigned result;452 asm("shf.l.wrap.b32 %0, %1, %2, %3;"453 : "=r"(result)454 : "r"(low32), "r"(high32), "r"(shiftWidth));455 return result;456}457inline __device__ unsigned __funnelshift_lc(unsigned low32, unsigned high32,458 unsigned shiftWidth) {459 unsigned result;460 asm("shf.l.clamp.b32 %0, %1, %2, %3;"461 : "=r"(result)462 : "r"(low32), "r"(high32), "r"(shiftWidth));463 return result;464}465inline __device__ unsigned __funnelshift_r(unsigned low32, unsigned high32,466 unsigned shiftWidth) {467 unsigned result;468 asm("shf.r.wrap.b32 %0, %1, %2, %3;"469 : "=r"(result)470 : "r"(low32), "r"(high32), "r"(shiftWidth));471 return result;472}473inline __device__ unsigned __funnelshift_rc(unsigned low32, unsigned high32,474 unsigned shiftWidth) {475 unsigned ret;476 asm("shf.r.clamp.b32 %0, %1, %2, %3;"477 : "=r"(ret)478 : "r"(low32), "r"(high32), "r"(shiftWidth));479 return ret;480}481 482#if defined(__cplusplus) && (__cplusplus >= 201103L)483 484#pragma push_macro("__INTRINSIC_LOAD")485#define __INTRINSIC_LOAD(__FnName, __AsmOp, __DeclType, __TmpType, __AsmType, \486 __Clobber) \487 inline __device__ __DeclType __FnName(const __DeclType *__ptr) { \488 __TmpType __ret; \489 asm(__AsmOp " %0, [%1];" : __AsmType(__ret) : "l"(__ptr)__Clobber); \490 return (__DeclType)__ret; \491 }492 493#pragma push_macro("__INTRINSIC_LOAD2")494#define __INTRINSIC_LOAD2(__FnName, __AsmOp, __DeclType, __TmpType, __AsmType, \495 __Clobber) \496 inline __device__ __DeclType __FnName(const __DeclType *__ptr) { \497 __DeclType __ret; \498 __TmpType __tmp; \499 asm(__AsmOp " {%0,%1}, [%2];" \500 : __AsmType(__tmp.x), __AsmType(__tmp.y) \501 : "l"(__ptr)__Clobber); \502 using __ElementType = decltype(__ret.x); \503 __ret.x = (__ElementType)(__tmp.x); \504 __ret.y = (__ElementType)__tmp.y; \505 return __ret; \506 }507 508#pragma push_macro("__INTRINSIC_LOAD4")509#define __INTRINSIC_LOAD4(__FnName, __AsmOp, __DeclType, __TmpType, __AsmType, \510 __Clobber) \511 inline __device__ __DeclType __FnName(const __DeclType *__ptr) { \512 __DeclType __ret; \513 __TmpType __tmp; \514 asm(__AsmOp " {%0,%1,%2,%3}, [%4];" \515 : __AsmType(__tmp.x), __AsmType(__tmp.y), __AsmType(__tmp.z), \516 __AsmType(__tmp.w) \517 : "l"(__ptr)__Clobber); \518 using __ElementType = decltype(__ret.x); \519 __ret.x = (__ElementType)__tmp.x; \520 __ret.y = (__ElementType)__tmp.y; \521 __ret.z = (__ElementType)__tmp.z; \522 __ret.w = (__ElementType)__tmp.w; \523 return __ret; \524 }525 526__INTRINSIC_LOAD(__ldcg, "ld.global.cg.s8", char, unsigned int, "=r", );527__INTRINSIC_LOAD(__ldcg, "ld.global.cg.s8", signed char, unsigned int, "=r", );528__INTRINSIC_LOAD(__ldcg, "ld.global.cg.s16", short, unsigned short, "=h", );529__INTRINSIC_LOAD(__ldcg, "ld.global.cg.s32", int, unsigned int, "=r", );530__INTRINSIC_LOAD(__ldcg, "ld.global.cg.s64", long long, unsigned long long,531 "=l", );532 533__INTRINSIC_LOAD2(__ldcg, "ld.global.cg.v2.s8", char2, int2, "=r", );534__INTRINSIC_LOAD4(__ldcg, "ld.global.cg.v4.s8", char4, int4, "=r", );535__INTRINSIC_LOAD2(__ldcg, "ld.global.cg.v2.s16", short2, short2, "=h", );536__INTRINSIC_LOAD4(__ldcg, "ld.global.cg.v4.s16", short4, short4, "=h", );537__INTRINSIC_LOAD2(__ldcg, "ld.global.cg.v2.s32", int2, int2, "=r", );538__INTRINSIC_LOAD4(__ldcg, "ld.global.cg.v4.s32", int4, int4, "=r", );539__INTRINSIC_LOAD2(__ldcg, "ld.global.cg.v2.s64 ", longlong2, longlong2, "=l", );540 541__INTRINSIC_LOAD(__ldcg, "ld.global.cg.u8", unsigned char, unsigned int,542 "=r", );543__INTRINSIC_LOAD(__ldcg, "ld.global.cg.u16", unsigned short, unsigned short,544 "=h", );545__INTRINSIC_LOAD(__ldcg, "ld.global.cg.u32", unsigned int, unsigned int,546 "=r", );547__INTRINSIC_LOAD(__ldcg, "ld.global.cg.u64", unsigned long long,548 unsigned long long, "=l", );549 550__INTRINSIC_LOAD2(__ldcg, "ld.global.cg.v2.u8", uchar2, int2, "=r", );551__INTRINSIC_LOAD4(__ldcg, "ld.global.cg.v4.u8", uchar4, int4, "=r", );552__INTRINSIC_LOAD2(__ldcg, "ld.global.cg.v2.u16", ushort2, ushort2, "=h", );553__INTRINSIC_LOAD4(__ldcg, "ld.global.cg.v4.u16", ushort4, ushort4, "=h", );554__INTRINSIC_LOAD2(__ldcg, "ld.global.cg.v2.u32", uint2, uint2, "=r", );555__INTRINSIC_LOAD4(__ldcg, "ld.global.cg.v4.u32", uint4, uint4, "=r", );556__INTRINSIC_LOAD2(__ldcg, "ld.global.cg.v2.u64", ulonglong2, ulonglong2,557 "=l", );558 559__INTRINSIC_LOAD(__ldcg, "ld.global.cg.f32", float, float, "=f", );560__INTRINSIC_LOAD(__ldcg, "ld.global.cg.f64", double, double, "=d", );561__INTRINSIC_LOAD2(__ldcg, "ld.global.cg.v2.f32", float2, float2, "=f", );562__INTRINSIC_LOAD4(__ldcg, "ld.global.cg.v4.f32", float4, float4, "=f", );563__INTRINSIC_LOAD2(__ldcg, "ld.global.cg.v2.f64", double2, double2, "=d", );564 565inline __device__ long __ldcg(const long *__ptr) {566 unsigned long __ret;567 if (sizeof(long) == 8) {568 asm("ld.global.cg.s64 %0, [%1];" : "=l"(__ret) : "l"(__ptr));569 } else {570 asm("ld.global.cg.s32 %0, [%1];" : "=r"(__ret) : "l"(__ptr));571 }572 return (long)__ret;573}574 575__INTRINSIC_LOAD(__ldcv, "ld.global.cv.u8", unsigned char, unsigned int,576 "=r", : "memory");577__INTRINSIC_LOAD(__ldcv, "ld.global.cv.u16", unsigned short, unsigned short,578 "=h", : "memory");579__INTRINSIC_LOAD(__ldcv, "ld.global.cv.u32", unsigned int, unsigned int,580 "=r", : "memory");581__INTRINSIC_LOAD(__ldcv, "ld.global.cv.u64", unsigned long long,582 unsigned long long, "=l", : "memory");583 584__INTRINSIC_LOAD(__ldcv, "ld.global.cv.s8", char, unsigned int,585 "=r", : "memory");586__INTRINSIC_LOAD(__ldcv, "ld.global.cv.s8", signed char, unsigned int,587 "=r", : "memory");588__INTRINSIC_LOAD(__ldcv, "ld.global.cv.s16", short, unsigned short,589 "=h", : "memory");590__INTRINSIC_LOAD(__ldcv, "ld.global.cv.s32", int, unsigned int,591 "=r", : "memory");592__INTRINSIC_LOAD(__ldcv, "ld.global.cv.s64", long long, unsigned long long,593 "=l", : "memory");594 595__INTRINSIC_LOAD2(__ldcv, "ld.global.cv.v2.u8", uchar2, uint2,596 "=r", : "memory");597__INTRINSIC_LOAD4(__ldcv, "ld.global.cv.v4.u8", uchar4, uint4,598 "=r", : "memory");599__INTRINSIC_LOAD2(__ldcv, "ld.global.cv.v2.u16", ushort2, ushort2,600 "=h", : "memory");601__INTRINSIC_LOAD4(__ldcv, "ld.global.cv.v4.u16", ushort4, ushort4,602 "=h", : "memory");603__INTRINSIC_LOAD2(__ldcv, "ld.global.cv.v2.u32", uint2, uint2,604 "=r", : "memory");605__INTRINSIC_LOAD4(__ldcv, "ld.global.cv.v4.u32", uint4, uint4,606 "=r", : "memory");607__INTRINSIC_LOAD2(__ldcv, "ld.global.cv.v2.u64", ulonglong2, ulonglong2,608 "=l", : "memory");609 610__INTRINSIC_LOAD2(__ldcv, "ld.global.cv.v2.s8", char2, int2, "=r", : "memory");611__INTRINSIC_LOAD4(__ldcv, "ld.global.cv.v4.s8", char4, int4, "=r", : "memory");612__INTRINSIC_LOAD2(__ldcv, "ld.global.cv.v2.s16", short2, short2,613 "=h", : "memory");614__INTRINSIC_LOAD4(__ldcv, "ld.global.cv.v4.s16", short4, short4,615 "=h", : "memory");616__INTRINSIC_LOAD2(__ldcv, "ld.global.cv.v2.s32", int2, int2, "=r", : "memory");617__INTRINSIC_LOAD4(__ldcv, "ld.global.cv.v4.s32", int4, int4, "=r", : "memory");618__INTRINSIC_LOAD2(__ldcv, "ld.global.cv.v2.s64", longlong2, longlong2,619 "=l", : "memory");620 621__INTRINSIC_LOAD(__ldcv, "ld.global.cv.f32", float, float, "=f", : "memory");622__INTRINSIC_LOAD(__ldcv, "ld.global.cv.f64", double, double, "=d", : "memory");623 624__INTRINSIC_LOAD2(__ldcv, "ld.global.cv.v2.f32", float2, float2,625 "=f", : "memory");626__INTRINSIC_LOAD4(__ldcv, "ld.global.cv.v4.f32", float4, float4,627 "=f", : "memory");628__INTRINSIC_LOAD2(__ldcv, "ld.global.cv.v2.f64", double2, double2,629 "=d", : "memory");630 631inline __device__ long __ldcv(const long *__ptr) {632 unsigned long __ret;633 if (sizeof(long) == 8) {634 asm("ld.global.cv.s64 %0, [%1];" : "=l"(__ret) : "l"(__ptr));635 } else {636 asm("ld.global.cv.s32 %0, [%1];" : "=r"(__ret) : "l"(__ptr));637 }638 return (long)__ret;639}640 641__INTRINSIC_LOAD(__ldcs, "ld.global.cs.s8", char, unsigned int, "=r", );642__INTRINSIC_LOAD(__ldcs, "ld.global.cs.s8", signed char, signed int, "=r", );643__INTRINSIC_LOAD(__ldcs, "ld.global.cs.s16", short, unsigned short, "=h", );644__INTRINSIC_LOAD(__ldcs, "ld.global.cs.s32", int, unsigned int, "=r", );645__INTRINSIC_LOAD(__ldcs, "ld.global.cs.s64", long long, unsigned long long,646 "=l", );647 648__INTRINSIC_LOAD2(__ldcs, "ld.global.cs.v2.s8", char2, int2, "=r", );649__INTRINSIC_LOAD4(__ldcs, "ld.global.cs.v4.s8", char4, int4, "=r", );650__INTRINSIC_LOAD2(__ldcs, "ld.global.cs.v2.s16", short2, short2, "=h", );651__INTRINSIC_LOAD4(__ldcs, "ld.global.cs.v4.s16", short4, short4, "=h", );652__INTRINSIC_LOAD2(__ldcs, "ld.global.cs.v2.s32", int2, int2, "=r", );653__INTRINSIC_LOAD4(__ldcs, "ld.global.cs.v4.s32", int4, int4, "=r", );654__INTRINSIC_LOAD2(__ldcs, "ld.global.cs.v2.s64", longlong2, longlong2, "=l", );655 656__INTRINSIC_LOAD(__ldcs, "ld.global.cs.u8", unsigned char, unsigned int,657 "=r", );658__INTRINSIC_LOAD(__ldcs, "ld.global.cs.u16", unsigned short, unsigned short,659 "=h", );660__INTRINSIC_LOAD(__ldcs, "ld.global.cs.u32", unsigned int, unsigned int,661 "=r", );662__INTRINSIC_LOAD(__ldcs, "ld.global.cs.u64", unsigned long long,663 unsigned long long, "=l", );664 665__INTRINSIC_LOAD2(__ldcs, "ld.global.cs.v2.u8", uchar2, uint2, "=r", );666__INTRINSIC_LOAD4(__ldcs, "ld.global.cs.v4.u8", uchar4, uint4, "=r", );667__INTRINSIC_LOAD2(__ldcs, "ld.global.cs.v2.u16", ushort2, ushort2, "=h", );668__INTRINSIC_LOAD4(__ldcs, "ld.global.cs.v4.u16", ushort4, ushort4, "=h", );669__INTRINSIC_LOAD2(__ldcs, "ld.global.cs.v2.u32", uint2, uint2, "=r", );670__INTRINSIC_LOAD4(__ldcs, "ld.global.cs.v4.u32", uint4, uint4, "=r", );671__INTRINSIC_LOAD2(__ldcs, "ld.global.cs.v2.u64", ulonglong2, ulonglong2,672 "=l", );673 674__INTRINSIC_LOAD(__ldcs, "ld.global.cs.f32", float, float, "=f", );675__INTRINSIC_LOAD(__ldcs, "ld.global.cs.f64", double, double, "=d", );676__INTRINSIC_LOAD2(__ldcs, "ld.global.cs.v2.f32", float2, float2, "=f", );677__INTRINSIC_LOAD4(__ldcs, "ld.global.cs.v4.f32", float4, float4, "=f", );678__INTRINSIC_LOAD2(__ldcs, "ld.global.cs.v2.f64", double2, double2, "=d", );679 680#pragma pop_macro("__INTRINSIC_LOAD")681#pragma pop_macro("__INTRINSIC_LOAD2")682#pragma pop_macro("__INTRINSIC_LOAD4")683 684inline __device__ long __ldcs(const long *__ptr) {685 unsigned long __ret;686 if (sizeof(long) == 8) {687 asm("ld.global.cs.s64 %0, [%1];" : "=l"(__ret) : "l"(__ptr));688 } else {689 asm("ld.global.cs.s32 %0, [%1];" : "=r"(__ret) : "l"(__ptr));690 }691 return (long)__ret;692}693 694#pragma push_macro("__INTRINSIC_STORE")695#define __INTRINSIC_STORE(__FnName, __AsmOp, __DeclType, __TmpType, __AsmType) \696 inline __device__ void __FnName(__DeclType *__ptr, __DeclType __value) { \697 __TmpType __tmp = (__TmpType)__value; \698 asm(__AsmOp " [%0], %1;" ::"l"(__ptr), __AsmType(__tmp) : "memory"); \699 }700 701#pragma push_macro("__INTRINSIC_STORE2")702#define __INTRINSIC_STORE2(__FnName, __AsmOp, __DeclType, __TmpType, \703 __AsmType) \704 inline __device__ void __FnName(__DeclType *__ptr, __DeclType __value) { \705 __TmpType __tmp; \706 using __ElementType = decltype(__tmp.x); \707 __tmp.x = (__ElementType)(__value.x); \708 __tmp.y = (__ElementType)(__value.y); \709 asm(__AsmOp " [%0], {%1,%2};" ::"l"(__ptr), __AsmType(__tmp.x), \710 __AsmType(__tmp.y) \711 : "memory"); \712 }713 714#pragma push_macro("__INTRINSIC_STORE4")715#define __INTRINSIC_STORE4(__FnName, __AsmOp, __DeclType, __TmpType, \716 __AsmType) \717 inline __device__ void __FnName(__DeclType *__ptr, __DeclType __value) { \718 __TmpType __tmp; \719 using __ElementType = decltype(__tmp.x); \720 __tmp.x = (__ElementType)(__value.x); \721 __tmp.y = (__ElementType)(__value.y); \722 __tmp.z = (__ElementType)(__value.z); \723 __tmp.w = (__ElementType)(__value.w); \724 asm(__AsmOp " [%0], {%1,%2,%3,%4};" ::"l"(__ptr), __AsmType(__tmp.x), \725 __AsmType(__tmp.y), __AsmType(__tmp.z), __AsmType(__tmp.w) \726 : "memory"); \727 }728 729__INTRINSIC_STORE(__stwt, "st.global.wt.s8", char, int, "r");730__INTRINSIC_STORE(__stwt, "st.global.wt.s8", signed char, int, "r");731__INTRINSIC_STORE(__stwt, "st.global.wt.s16", short, short, "h");732__INTRINSIC_STORE(__stwt, "st.global.wt.s32", int, int, "r");733__INTRINSIC_STORE(__stwt, "st.global.wt.s64", long long, long long, "l");734 735__INTRINSIC_STORE2(__stwt, "st.global.wt.v2.s8", char2, int2, "r");736__INTRINSIC_STORE4(__stwt, "st.global.wt.v4.s8", char4, int4, "r");737__INTRINSIC_STORE2(__stwt, "st.global.wt.v2.s16", short2, short2, "h");738__INTRINSIC_STORE4(__stwt, "st.global.wt.v4.s16", short4, short4, "h");739__INTRINSIC_STORE2(__stwt, "st.global.wt.v2.s32", int2, int2, "r");740__INTRINSIC_STORE4(__stwt, "st.global.wt.v4.s32", int4, int4, "r");741__INTRINSIC_STORE2(__stwt, "st.global.wt.v2.s64", longlong2, longlong2, "l");742 743__INTRINSIC_STORE(__stwt, "st.global.wt.u8", unsigned char, int, "r");744__INTRINSIC_STORE(__stwt, "st.global.wt.u16", unsigned short, unsigned short,745 "h");746__INTRINSIC_STORE(__stwt, "st.global.wt.u32", unsigned int, unsigned int, "r");747__INTRINSIC_STORE(__stwt, "st.global.wt.u64", unsigned long long,748 unsigned long long, "l");749 750__INTRINSIC_STORE2(__stwt, "st.global.wt.v2.u8", uchar2, uchar2, "r");751__INTRINSIC_STORE4(__stwt, "st.global.wt.v4.u8", uchar4, uint4, "r");752__INTRINSIC_STORE2(__stwt, "st.global.wt.v2.u16", ushort2, ushort2, "h");753__INTRINSIC_STORE4(__stwt, "st.global.wt.v4.u16", ushort4, ushort4, "h");754__INTRINSIC_STORE2(__stwt, "st.global.wt.v2.u32", uint2, uint2, "r");755__INTRINSIC_STORE4(__stwt, "st.global.wt.v4.u32", uint4, uint4, "r");756__INTRINSIC_STORE2(__stwt, "st.global.wt.v2.u64", ulonglong2, ulonglong2, "l");757 758__INTRINSIC_STORE(__stwt, "st.global.wt.f32", float, float, "f");759__INTRINSIC_STORE(__stwt, "st.global.wt.f64", double, double, "d");760__INTRINSIC_STORE2(__stwt, "st.global.wt.v2.f32", float2, float2, "f");761__INTRINSIC_STORE4(__stwt, "st.global.wt.v4.f32", float4, float4, "f");762__INTRINSIC_STORE2(__stwt, "st.global.wt.v2.f64", double2, double2, "d");763 764#pragma pop_macro("__INTRINSIC_STORE")765#pragma pop_macro("__INTRINSIC_STORE2")766#pragma pop_macro("__INTRINSIC_STORE4")767 768#endif // defined(__cplusplus) && (__cplusplus >= 201103L)769#endif // !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 320770 771#if CUDA_VERSION >= 11000772extern "C" {773__device__ inline size_t __nv_cvta_generic_to_global_impl(const void *__ptr) {774 return (size_t)(void __attribute__((address_space(1))) *)__ptr;775}776__device__ inline size_t __nv_cvta_generic_to_shared_impl(const void *__ptr) {777 return (size_t)(void __attribute__((address_space(3))) *)__ptr;778}779__device__ inline size_t __nv_cvta_generic_to_constant_impl(const void *__ptr) {780 return (size_t)(void __attribute__((address_space(4))) *)__ptr;781}782__device__ inline size_t __nv_cvta_generic_to_local_impl(const void *__ptr) {783 return (size_t)(void __attribute__((address_space(5))) *)__ptr;784}785__device__ inline void *__nv_cvta_global_to_generic_impl(size_t __ptr) {786 return (void *)(void __attribute__((address_space(1))) *)__ptr;787}788__device__ inline void *__nv_cvta_shared_to_generic_impl(size_t __ptr) {789 return (void *)(void __attribute__((address_space(3))) *)__ptr;790}791__device__ inline void *__nv_cvta_constant_to_generic_impl(size_t __ptr) {792 return (void *)(void __attribute__((address_space(4))) *)__ptr;793}794__device__ inline void *__nv_cvta_local_to_generic_impl(size_t __ptr) {795 return (void *)(void __attribute__((address_space(5))) *)__ptr;796}797__device__ inline cuuint32_t __nvvm_get_smem_pointer(void *__ptr) {798 return __nv_cvta_generic_to_shared_impl(__ptr);799}800} // extern "C"801 802#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 800803__device__ inline unsigned __reduce_add_sync(unsigned __mask,804 unsigned __value) {805 return __nvvm_redux_sync_add(__value, __mask);806}807__device__ inline unsigned __reduce_min_sync(unsigned __mask,808 unsigned __value) {809 return __nvvm_redux_sync_umin(__value, __mask);810}811__device__ inline unsigned __reduce_max_sync(unsigned __mask,812 unsigned __value) {813 return __nvvm_redux_sync_umax(__value, __mask);814}815__device__ inline int __reduce_min_sync(unsigned __mask, int __value) {816 return __nvvm_redux_sync_min(__value, __mask);817}818__device__ inline int __reduce_max_sync(unsigned __mask, int __value) {819 return __nvvm_redux_sync_max(__value, __mask);820}821__device__ inline unsigned __reduce_or_sync(unsigned __mask, unsigned __value) {822 return __nvvm_redux_sync_or(__value, __mask);823}824__device__ inline unsigned __reduce_and_sync(unsigned __mask,825 unsigned __value) {826 return __nvvm_redux_sync_and(__value, __mask);827}828__device__ inline unsigned __reduce_xor_sync(unsigned __mask,829 unsigned __value) {830 return __nvvm_redux_sync_xor(__value, __mask);831}832 833__device__ inline void __nv_memcpy_async_shared_global_4(void *__dst,834 const void *__src,835 unsigned __src_size) {836 __nvvm_cp_async_ca_shared_global_4(837 (void __attribute__((address_space(3))) *)__dst,838 (const void __attribute__((address_space(1))) *)__src, __src_size);839}840__device__ inline void __nv_memcpy_async_shared_global_8(void *__dst,841 const void *__src,842 unsigned __src_size) {843 __nvvm_cp_async_ca_shared_global_8(844 (void __attribute__((address_space(3))) *)__dst,845 (const void __attribute__((address_space(1))) *)__src, __src_size);846}847__device__ inline void __nv_memcpy_async_shared_global_16(void *__dst,848 const void *__src,849 unsigned __src_size) {850 __nvvm_cp_async_ca_shared_global_16(851 (void __attribute__((address_space(3))) *)__dst,852 (const void __attribute__((address_space(1))) *)__src, __src_size);853}854 855__device__ inline void *856__nv_associate_access_property(const void *__ptr, unsigned long long __prop) {857 // TODO: it appears to provide compiler with some sort of a hint. We do not858 // know what exactly it is supposed to do. However, CUDA headers suggest that859 // just passing through __ptr should not affect correctness. They do so on860 // pre-sm80 GPUs where this builtin is not available.861 return (void*)__ptr;862}863#endif // !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 800864 865#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 900866__device__ inline unsigned __isCtaShared(const void *ptr) {867 return __isShared(ptr);868}869 870__device__ inline unsigned __isClusterShared(const void *__ptr) {871 return __nvvm_isspacep_shared_cluster(__ptr);872}873 874__device__ inline void *__cluster_map_shared_rank(const void *__ptr,875 unsigned __rank) {876 return __nvvm_mapa((void *)__ptr, __rank);877}878 879__device__ inline unsigned __cluster_query_shared_rank(const void *__ptr) {880 return __nvvm_getctarank((void *)__ptr);881}882 883__device__ inline uint2884__cluster_map_shared_multicast(const void *__ptr,885 unsigned int __cluster_cta_mask) {886 return make_uint2((unsigned)__cvta_generic_to_shared(__ptr),887 __cluster_cta_mask);888}889 890__device__ inline unsigned __clusterDimIsSpecified() {891 return __nvvm_is_explicit_cluster();892}893 894__device__ inline dim3 __clusterDim() {895 return dim3(__nvvm_read_ptx_sreg_cluster_nctaid_x(),896 __nvvm_read_ptx_sreg_cluster_nctaid_y(),897 __nvvm_read_ptx_sreg_cluster_nctaid_z());898}899 900__device__ inline dim3 __clusterRelativeBlockIdx() {901 return dim3(__nvvm_read_ptx_sreg_cluster_ctaid_x(),902 __nvvm_read_ptx_sreg_cluster_ctaid_y(),903 __nvvm_read_ptx_sreg_cluster_ctaid_z());904}905 906__device__ inline dim3 __clusterGridDimInClusters() {907 return dim3(__nvvm_read_ptx_sreg_nclusterid_x(),908 __nvvm_read_ptx_sreg_nclusterid_y(),909 __nvvm_read_ptx_sreg_nclusterid_z());910}911 912__device__ inline dim3 __clusterIdx() {913 return dim3(__nvvm_read_ptx_sreg_clusterid_x(),914 __nvvm_read_ptx_sreg_clusterid_y(),915 __nvvm_read_ptx_sreg_clusterid_z());916}917 918__device__ inline unsigned __clusterRelativeBlockRank() {919 return __nvvm_read_ptx_sreg_cluster_ctarank();920}921 922__device__ inline unsigned __clusterSizeInBlocks() {923 return __nvvm_read_ptx_sreg_cluster_nctarank();924}925 926__device__ inline void __cluster_barrier_arrive() {927 __nvvm_barrier_cluster_arrive();928}929 930__device__ inline void __cluster_barrier_arrive_relaxed() {931 __nvvm_barrier_cluster_arrive_relaxed();932}933 934__device__ inline void __cluster_barrier_wait() {935 __nvvm_barrier_cluster_wait();936}937 938__device__ inline void __threadfence_cluster() { __nvvm_fence_sc_cluster(); }939 940__device__ inline float2 atomicAdd(float2 *__ptr, float2 __val) {941 float2 __ret;942 __asm__("atom.add.v2.f32 {%0, %1}, [%2], {%3, %4};"943 : "=f"(__ret.x), "=f"(__ret.y)944 : "l"(__ptr), "f"(__val.x), "f"(__val.y));945 return __ret;946}947 948__device__ inline float2 atomicAdd_block(float2 *__ptr, float2 __val) {949 float2 __ret;950 __asm__("atom.cta.add.v2.f32 {%0, %1}, [%2], {%3, %4};"951 : "=f"(__ret.x), "=f"(__ret.y)952 : "l"(__ptr), "f"(__val.x), "f"(__val.y));953 return __ret;954}955 956__device__ inline float2 atomicAdd_system(float2 *__ptr, float2 __val) {957 float2 __ret;958 __asm__("atom.sys.add.v2.f32 {%0, %1}, [%2], {%3, %4};"959 : "=f"(__ret.x), "=f"(__ret.y)960 : "l"(__ptr), "f"(__val.x), "f"(__val.y));961 return __ret;962}963 964__device__ inline float4 atomicAdd(float4 *__ptr, float4 __val) {965 float4 __ret;966 __asm__("atom.add.v4.f32 {%0, %1, %2, %3}, [%4], {%5, %6, %7, %8};"967 : "=f"(__ret.x), "=f"(__ret.y), "=f"(__ret.z), "=f"(__ret.w)968 : "l"(__ptr), "f"(__val.x), "f"(__val.y), "f"(__val.z), "f"(__val.w));969 return __ret;970}971 972__device__ inline float4 atomicAdd_block(float4 *__ptr, float4 __val) {973 float4 __ret;974 __asm__(975 "atom.cta.add.v4.f32 {%0, %1, %2, %3}, [%4], {%5, %6, %7, %8};"976 : "=f"(__ret.x), "=f"(__ret.y), "=f"(__ret.z), "=f"(__ret.w)977 : "l"(__ptr), "f"(__val.x), "f"(__val.y), "f"(__val.z), "f"(__val.w));978 return __ret;979}980 981__device__ inline float4 atomicAdd_system(float4 *__ptr, float4 __val) {982 float4 __ret;983 __asm__(984 "atom.sys.add.v4.f32 {%0, %1, %2, %3}, [%4], {%5, %6, %7, %8};"985 : "=f"(__ret.x), "=f"(__ret.y), "=f"(__ret.z), "=f"(__ret.w)986 : "l"(__ptr), "f"(__val.x), "f"(__val.y), "f"(__val.z), "f"(__val.w)987 :);988 return __ret;989}990 991#endif // !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 900992#endif // CUDA_VERSION >= 11000993 994#endif // defined(__CLANG_CUDA_INTRINSICS_H__)995