1#include <clc/clc.h> 2 3#if __clang_major__ >= 8 4#define CONST_AS __constant 5#elif __clang_major__ >= 7 6#define CONST_AS __attribute__((address_space(4))) 7#else 8#define CONST_AS __attribute__((address_space(2))) 9#endif 10 11#if __clang_major__ >= 6 12#define __dispatch_ptr __builtin_amdgcn_dispatch_ptr 13#else 14#define __dispatch_ptr __clc_amdgcn_dispatch_ptr 15CONST_AS char * __clc_amdgcn_dispatch_ptr(void) __asm("llvm.amdgcn.dispatch.ptr"); 16#endif 17 18_CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) { 19 CONST_AS uint *ptr = (CONST_AS uint *)__dispatch_ptr(); 20 switch (dim) { 21 case 0: 22 return ptr[1] & 0xffffu; 23 case 1: 24 return ptr[1] >> 16; 25 case 2: 26 return ptr[2] & 0xffffu; 27 } 28 return 1; 29} 30