AcoSail/llvm-project
8
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

libclc: Make all built-ins overloadable

Browse Source 
The SPIR spec states that all OpenCL built-in functions should be
overloadable and mangled, to ensure consistency.

Add the overload attribute to functions which were missing them:
work dimensions, memory barriers and fences, and events.

Reviewed By: tstellar, jenatali

Differential Revision: https://reviews.llvm.org/D82078
This commit is contained in:
Daniel Stone 2020-08-17 13:44:49 -07:00 committed by Tom Stellard
parent 3a7051d9c2
commit 3d21fa56f5
41 changed files with 191 additions and 173 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

11
libclc/amdgcn-amdhsa/lib/workitem/get_global_size.cl
View File

@ -15,10 +15,9 @@
CONST_AS uchar * __clc_amdgcn_dispatch_ptr(void) __asm("llvm.amdgcn.dispatch.ptr");
#endif
_CLC_DEF size_t get_global_size(uint dim)
{
CONST_AS uint * ptr = (CONST_AS uint *) __dispatch_ptr();
if (dim < 3)
return ptr[3 + dim];
return 1;
_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) {
CONST_AS uint *ptr = (CONST_AS uint *)__dispatch_ptr();
if (dim < 3)
return ptr[3 + dim];
return 1;
}

23
libclc/amdgcn-amdhsa/lib/workitem/get_local_size.cl
View File

@ -15,16 +15,15 @@
CONST_AS char * __clc_amdgcn_dispatch_ptr(void) __asm("llvm.amdgcn.dispatch.ptr");
#endif
_CLC_DEF size_t get_local_size(uint dim)
{
CONST_AS uint * ptr = (CONST_AS uint *) __dispatch_ptr();
switch (dim) {
case 0:
return ptr[1] & 0xffffu;
case 1:
return ptr[1] >> 16;
case 2:
return ptr[2] & 0xffffu;
}
return 1;
_CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) {
CONST_AS uint *ptr = (CONST_AS uint *)__dispatch_ptr();
switch (dim) {
case 0:
return ptr[1] & 0xffffu;
case 1:
return ptr[1] >> 16;
case 2:
return ptr[2] & 0xffffu;
}
return 1;
}

2
libclc/amdgcn-amdhsa/lib/workitem/get_num_groups.cl
View File

@ -1,7 +1,7 @@
#include <clc/clc.h>
_CLC_DEF size_t get_num_groups(uint dim) {
_CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) {
size_t global_size = get_global_size(dim);
size_t local_size = get_local_size(dim);
size_t num_groups = global_size / local_size;

25
libclc/amdgcn/lib/mem_fence/fence.cl
View File

@ -17,24 +17,21 @@ void __clc_amdgcn_s_waitcnt(unsigned flags);
_CLC_DEF void __clc_amdgcn_s_waitcnt(unsigned) __asm("llvm.amdgcn.s.waitcnt");
#endif
_CLC_DEF void mem_fence(cl_mem_fence_flags flags)
{
if (flags & CLK_GLOBAL_MEM_FENCE) {
// scalar loads are counted with LGKM but we don't know whether
// the compiler turned any loads to scalar
__waitcnt(0);
} else if (flags & CLK_LOCAL_MEM_FENCE)
__waitcnt(0xff); // LGKM is [12:8]
_CLC_DEF _CLC_OVERLOAD void mem_fence(cl_mem_fence_flags flags) {
if (flags & CLK_GLOBAL_MEM_FENCE) {
// scalar loads are counted with LGKM but we don't know whether
// the compiler turned any loads to scalar
__waitcnt(0);
} else if (flags & CLK_LOCAL_MEM_FENCE)
__waitcnt(0xff); // LGKM is [12:8]
}
#undef __waitcnt
// We don't have separate mechanism for read and write fences
_CLC_DEF void read_mem_fence(cl_mem_fence_flags flags)
{
mem_fence(flags);
_CLC_DEF _CLC_OVERLOAD void read_mem_fence(cl_mem_fence_flags flags) {
mem_fence(flags);
}
_CLC_DEF void write_mem_fence(cl_mem_fence_flags flags)
{
mem_fence(flags);
_CLC_DEF _CLC_OVERLOAD void write_mem_fence(cl_mem_fence_flags flags) {
mem_fence(flags);
}

7
libclc/amdgcn/lib/synchronization/barrier.cl
View File

@ -1,7 +1,6 @@
#include <clc/clc.h>
_CLC_DEF void barrier(cl_mem_fence_flags flags)
{
mem_fence(flags);
__builtin_amdgcn_s_barrier();
_CLC_DEF _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags) {
mem_fence(flags);
__builtin_amdgcn_s_barrier();
}

12
libclc/amdgcn/lib/workitem/get_global_offset.cl
View File

@ -8,11 +8,9 @@
#define CONST_AS __attribute__((address_space(2)))
#endif
_CLC_DEF size_t get_global_offset(uint dim)
{
CONST_AS uint * ptr =
(CONST_AS uint *) __builtin_amdgcn_implicitarg_ptr();
if (dim < 3)
return ptr[dim + 1];
return 0;
_CLC_DEF _CLC_OVERLOAD size_t get_global_offset(uint dim) {
CONST_AS uint *ptr = (CONST_AS uint *)__builtin_amdgcn_implicitarg_ptr();
if (dim < 3)
return ptr[dim + 1];
return 0;
}

19
libclc/amdgcn/lib/workitem/get_global_size.cl
View File

@ -4,12 +4,15 @@ uint __clc_amdgcn_get_global_size_x(void) __asm("llvm.r600.read.global.size.x");
uint __clc_amdgcn_get_global_size_y(void) __asm("llvm.r600.read.global.size.y");
uint __clc_amdgcn_get_global_size_z(void) __asm("llvm.r600.read.global.size.z");
_CLC_DEF size_t get_global_size(uint dim)
{
switch (dim) {
case 0: return __clc_amdgcn_get_global_size_x();
case 1: return __clc_amdgcn_get_global_size_y();
case 2: return __clc_amdgcn_get_global_size_z();
default: return 1;
}
_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) {
switch (dim) {
case 0:
return __clc_amdgcn_get_global_size_x();
case 1:
return __clc_amdgcn_get_global_size_y();
case 2:
return __clc_amdgcn_get_global_size_z();
default:
return 1;
}
}

19
libclc/amdgcn/lib/workitem/get_group_id.cl
View File

@ -1,11 +1,14 @@
#include <clc/clc.h>
_CLC_DEF size_t get_group_id(uint dim)
{
switch(dim) {
case 0: return __builtin_amdgcn_workgroup_id_x();
case 1: return __builtin_amdgcn_workgroup_id_y();
case 2: return __builtin_amdgcn_workgroup_id_z();
default: return 1;
}
_CLC_DEF _CLC_OVERLOAD size_t get_group_id(uint dim) {
switch (dim) {
case 0:
return __builtin_amdgcn_workgroup_id_x();
case 1:
return __builtin_amdgcn_workgroup_id_y();
case 2:
return __builtin_amdgcn_workgroup_id_z();
default:
return 1;
}
}

19
libclc/amdgcn/lib/workitem/get_local_id.cl
View File

@ -1,11 +1,14 @@
#include <clc/clc.h>
_CLC_DEF size_t get_local_id(uint dim)
{
switch(dim) {
case 0: return __builtin_amdgcn_workitem_id_x();
case 1: return __builtin_amdgcn_workitem_id_y();
case 2: return __builtin_amdgcn_workitem_id_z();
default: return 1;
}
_CLC_DEF _CLC_OVERLOAD size_t get_local_id(uint dim) {
switch (dim) {
case 0:
return __builtin_amdgcn_workitem_id_x();
case 1:
return __builtin_amdgcn_workitem_id_y();
case 2:
return __builtin_amdgcn_workitem_id_z();
default:
return 1;
}
}

19
libclc/amdgcn/lib/workitem/get_local_size.cl
View File

@ -4,12 +4,15 @@ uint __clc_amdgcn_get_local_size_x(void) __asm("llvm.r600.read.local.size.x");
uint __clc_amdgcn_get_local_size_y(void) __asm("llvm.r600.read.local.size.y");
uint __clc_amdgcn_get_local_size_z(void) __asm("llvm.r600.read.local.size.z");
_CLC_DEF size_t get_local_size(uint dim)
{
switch (dim) {
case 0: return __clc_amdgcn_get_local_size_x();
case 1: return __clc_amdgcn_get_local_size_y();
case 2: return __clc_amdgcn_get_local_size_z();
default: return 1;
}
_CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) {
switch (dim) {
case 0:
return __clc_amdgcn_get_local_size_x();
case 1:
return __clc_amdgcn_get_local_size_y();
case 2:
return __clc_amdgcn_get_local_size_z();
default:
return 1;
}
}

19
libclc/amdgcn/lib/workitem/get_num_groups.cl
View File

@ -4,12 +4,15 @@ uint __clc_amdgcn_get_num_groups_x(void) __asm("llvm.r600.read.ngroups.x");
uint __clc_amdgcn_get_num_groups_y(void) __asm("llvm.r600.read.ngroups.y");
uint __clc_amdgcn_get_num_groups_z(void) __asm("llvm.r600.read.ngroups.z");
_CLC_DEF size_t get_num_groups(uint dim)
{
switch (dim) {
case 0: return __clc_amdgcn_get_num_groups_x();
case 1: return __clc_amdgcn_get_num_groups_y();
case 2: return __clc_amdgcn_get_num_groups_z();
default: return 1;
}
_CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) {
switch (dim) {
case 0:
return __clc_amdgcn_get_num_groups_x();
case 1:
return __clc_amdgcn_get_num_groups_y();
case 2:
return __clc_amdgcn_get_num_groups_z();
default:
return 1;
}
}

8
libclc/amdgcn/lib/workitem/get_work_dim.cl
View File

@ -8,9 +8,7 @@
#define CONST_AS __attribute__((address_space(2)))
#endif
_CLC_DEF uint get_work_dim(void)
{
CONST_AS uint * ptr =
(CONST_AS uint *) __builtin_amdgcn_implicitarg_ptr();
return ptr[0];
_CLC_DEF _CLC_OVERLOAD uint get_work_dim(void) {
CONST_AS uint *ptr = (CONST_AS uint *)__builtin_amdgcn_implicitarg_ptr();
return ptr[0];
}

3
libclc/generic/include/clc/async/wait_group_events.h
View File

@ -1 +1,2 @@
void wait_group_events(int num_events, event_t *event_list);
_CLC_DECL _CLC_OVERLOAD void wait_group_events(int num_events,
event_t *event_list);

6
libclc/generic/include/clc/explicit_fence/explicit_memory_fence.h
View File

@ -1,3 +1,3 @@
_CLC_DECL void mem_fence(cl_mem_fence_flags flags);
_CLC_DECL void read_mem_fence(cl_mem_fence_flags flags);
_CLC_DECL void write_mem_fence(cl_mem_fence_flags flags);
_CLC_DECL _CLC_OVERLOAD void mem_fence(cl_mem_fence_flags flags);
_CLC_DECL _CLC_OVERLOAD void read_mem_fence(cl_mem_fence_flags flags);
_CLC_DECL _CLC_OVERLOAD void write_mem_fence(cl_mem_fence_flags flags);

2
libclc/generic/include/clc/synchronization/barrier.h
View File

@ -1 +1 @@
_CLC_DECL void barrier(cl_mem_fence_flags flags);
_CLC_DECL _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags);

2
libclc/generic/include/clc/workitem/get_global_id.h
View File

@ -1 +1 @@
_CLC_DECL size_t get_global_id(uint dim);
_CLC_DECL _CLC_OVERLOAD size_t get_global_id(uint dim);

2
libclc/generic/include/clc/workitem/get_global_offset.h
View File

@ -1 +1 @@
_CLC_DECL size_t get_global_offset(uint dim);
_CLC_DECL _CLC_OVERLOAD size_t get_global_offset(uint dim);

2
libclc/generic/include/clc/workitem/get_global_size.h
View File

@ -1 +1 @@
_CLC_DECL size_t get_global_size(uint dim);
_CLC_DECL _CLC_OVERLOAD size_t get_global_size(uint dim);

2
libclc/generic/include/clc/workitem/get_group_id.h
View File

@ -1 +1 @@
_CLC_DECL size_t get_group_id(uint dim);
_CLC_DECL _CLC_OVERLOAD size_t get_group_id(uint dim);

2
libclc/generic/include/clc/workitem/get_local_id.h
View File

@ -1 +1 @@
_CLC_DECL size_t get_local_id(uint dim);
_CLC_DECL _CLC_OVERLOAD size_t get_local_id(uint dim);

2
libclc/generic/include/clc/workitem/get_local_size.h
View File

@ -1 +1 @@
_CLC_DECL size_t get_local_size(uint dim);
_CLC_DECL _CLC_OVERLOAD size_t get_local_size(uint dim);

2
libclc/generic/include/clc/workitem/get_num_groups.h
View File

@ -1 +1 @@
_CLC_DECL size_t get_num_groups(uint dim);
_CLC_DECL _CLC_OVERLOAD size_t get_num_groups(uint dim);

2
libclc/generic/include/clc/workitem/get_work_dim.h
View File

@ -1 +1 @@
_CLC_DECL uint get_work_dim(void);
_CLC_DECL _CLC_OVERLOAD uint get_work_dim(void);

3
libclc/generic/lib/async/wait_group_events.cl
View File

@ -1,5 +1,6 @@
#include <clc/clc.h>
_CLC_DEF void wait_group_events(int num_events, event_t *event_list) {
_CLC_DEF _CLC_OVERLOAD void wait_group_events(int num_events,
event_t *event_list) {
barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
}

2
libclc/generic/lib/workitem/get_global_id.cl
View File

@ -1,5 +1,5 @@
#include <clc/clc.h>
_CLC_DEF size_t get_global_id(uint dim) {
_CLC_DEF _CLC_OVERLOAD size_t get_global_id(uint dim) {
return get_group_id(dim) * get_local_size(dim) + get_local_id(dim) + get_global_offset(dim);
}

2
libclc/generic/lib/workitem/get_global_size.cl
View File

@ -1,5 +1,5 @@
#include <clc/clc.h>
_CLC_DEF size_t get_global_size(uint dim) {
_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) {
return get_num_groups(dim)*get_local_size(dim);
}

10
libclc/ptx-nvidiacl/lib/mem_fence/fence.cl
View File

@ -1,15 +1,15 @@
#include <clc/clc.h>
_CLC_DEF void mem_fence(cl_mem_fence_flags flags) {
if (flags & (CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE))
__nvvm_membar_cta();
_CLC_DEF _CLC_OVERLOAD void mem_fence(cl_mem_fence_flags flags) {
if (flags & (CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE))
__nvvm_membar_cta();
}
// We do not have separate mechanism for read and write fences.
_CLC_DEF void read_mem_fence(cl_mem_fence_flags flags) {
_CLC_DEF _CLC_OVERLOAD void read_mem_fence(cl_mem_fence_flags flags) {
mem_fence(flags);
}
_CLC_DEF void write_mem_fence(cl_mem_fence_flags flags) {
_CLC_DEF _CLC_OVERLOAD void write_mem_fence(cl_mem_fence_flags flags) {
mem_fence(flags);
}

3
libclc/ptx-nvidiacl/lib/synchronization/barrier.cl
View File

@ -1,6 +1,5 @@
#include <clc/clc.h>
_CLC_DEF void barrier(cl_mem_fence_flags flags) {
_CLC_DEF _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags) {
__syncthreads();
}

2
libclc/ptx-nvidiacl/lib/workitem/get_global_id.cl
View File

@ -1,5 +1,5 @@
#include <clc/clc.h>
_CLC_DEF size_t get_global_id(uint dim) {
_CLC_DEF _CLC_OVERLOAD size_t get_global_id(uint dim) {
return get_group_id(dim) * get_local_size(dim) + get_local_id(dim);
}

2
libclc/ptx-nvidiacl/lib/workitem/get_group_id.cl
View File

@ -1,6 +1,6 @@
#include <clc/clc.h>
_CLC_DEF size_t get_group_id(uint dim) {
_CLC_DEF _CLC_OVERLOAD size_t get_group_id(uint dim) {
switch (dim) {
case 0: return __nvvm_read_ptx_sreg_ctaid_x();
case 1: return __nvvm_read_ptx_sreg_ctaid_y();

2
libclc/ptx-nvidiacl/lib/workitem/get_local_id.cl
View File

@ -1,6 +1,6 @@
#include <clc/clc.h>
_CLC_DEF size_t get_local_id(uint dim) {
_CLC_DEF _CLC_OVERLOAD size_t get_local_id(uint dim) {
switch (dim) {
case 0: return __nvvm_read_ptx_sreg_tid_x();
case 1: return __nvvm_read_ptx_sreg_tid_y();

2
libclc/ptx-nvidiacl/lib/workitem/get_local_size.cl
View File

@ -1,6 +1,6 @@
#include <clc/clc.h>
_CLC_DEF size_t get_local_size(uint dim) {
_CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) {
switch (dim) {
case 0: return __nvvm_read_ptx_sreg_ntid_x();
case 1: return __nvvm_read_ptx_sreg_ntid_y();

2
libclc/ptx-nvidiacl/lib/workitem/get_num_groups.cl
View File

@ -1,6 +1,6 @@
#include <clc/clc.h>
_CLC_DEF size_t get_num_groups(uint dim) {
_CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) {
switch (dim) {
case 0: return __nvvm_read_ptx_sreg_nctaid_x();
case 1: return __nvvm_read_ptx_sreg_nctaid_y();

3
libclc/r600/lib/synchronization/barrier.cl
View File

@ -2,8 +2,7 @@
_CLC_DEF void __clc_r600_barrier(void) __asm("llvm.r600.group.barrier");
_CLC_DEF void barrier(uint flags)
{
_CLC_DEF _CLC_OVERLOAD void barrier(uint flags) {
// We should call mem_fence here, but that is not implemented for r600 yet
__clc_r600_barrier();
}

15
libclc/r600/lib/workitem/get_global_offset.cl
View File

@ -1,11 +1,10 @@
#include <clc/clc.h>
_CLC_DEF uint get_global_offset(uint dim)
{
__attribute__((address_space(7))) uint * ptr =
(__attribute__((address_space(7))) uint *)
__builtin_r600_implicitarg_ptr();
if (dim < 3)
return ptr[dim + 1];
return 0;
_CLC_DEF _CLC_OVERLOAD uint get_global_offset(uint dim) {
__attribute__((address_space(7))) uint *ptr =
(__attribute__((address_space(7)))
uint *)__builtin_r600_implicitarg_ptr();
if (dim < 3)
return ptr[dim + 1];
return 0;
}

19
libclc/r600/lib/workitem/get_global_size.cl
View File

@ -4,12 +4,15 @@ uint __clc_r600_get_global_size_x(void) __asm("llvm.r600.read.global.size.x");
uint __clc_r600_get_global_size_y(void) __asm("llvm.r600.read.global.size.y");
uint __clc_r600_get_global_size_z(void) __asm("llvm.r600.read.global.size.z");
_CLC_DEF size_t get_global_size(uint dim)
{
switch (dim) {
case 0: return __clc_r600_get_global_size_x();
case 1: return __clc_r600_get_global_size_y();
case 2: return __clc_r600_get_global_size_z();
default: return 1;
}
_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) {
switch (dim) {
case 0:
return __clc_r600_get_global_size_x();
case 1:
return __clc_r600_get_global_size_y();
case 2:
return __clc_r600_get_global_size_z();
default:
return 1;
}
}

19
libclc/r600/lib/workitem/get_group_id.cl
View File

@ -1,11 +1,14 @@
#include <clc/clc.h>
_CLC_DEF uint get_group_id(uint dim)
{
switch(dim) {
case 0: return __builtin_r600_read_tgid_x();
case 1: return __builtin_r600_read_tgid_y();
case 2: return __builtin_r600_read_tgid_z();
default: return 1;
}
_CLC_DEF _CLC_OVERLOAD uint get_group_id(uint dim) {
switch (dim) {
case 0:
return __builtin_r600_read_tgid_x();
case 1:
return __builtin_r600_read_tgid_y();
case 2:
return __builtin_r600_read_tgid_z();
default:
return 1;
}
}

19
libclc/r600/lib/workitem/get_local_id.cl
View File

@ -1,11 +1,14 @@
#include <clc/clc.h>
_CLC_DEF uint get_local_id(uint dim)
{
switch(dim) {
case 0: return __builtin_r600_read_tidig_x();
case 1: return __builtin_r600_read_tidig_y();
case 2: return __builtin_r600_read_tidig_z();
default: return 1;
}
_CLC_DEF _CLC_OVERLOAD uint get_local_id(uint dim) {
switch (dim) {
case 0:
return __builtin_r600_read_tidig_x();
case 1:
return __builtin_r600_read_tidig_y();
case 2:
return __builtin_r600_read_tidig_z();
default:
return 1;
}
}

19
libclc/r600/lib/workitem/get_local_size.cl
View File

@ -4,12 +4,15 @@ uint __clc_r600_get_local_size_x(void) __asm("llvm.r600.read.local.size.x");
uint __clc_r600_get_local_size_y(void) __asm("llvm.r600.read.local.size.y");
uint __clc_r600_get_local_size_z(void) __asm("llvm.r600.read.local.size.z");
_CLC_DEF size_t get_local_size(uint dim)
{
switch (dim) {
case 0: return __clc_r600_get_local_size_x();
case 1: return __clc_r600_get_local_size_y();
case 2: return __clc_r600_get_local_size_z();
default: return 1;
}
_CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) {
switch (dim) {
case 0:
return __clc_r600_get_local_size_x();
case 1:
return __clc_r600_get_local_size_y();
case 2:
return __clc_r600_get_local_size_z();
default:
return 1;
}
}

19
libclc/r600/lib/workitem/get_num_groups.cl
View File

@ -4,12 +4,15 @@ uint __clc_r600_get_num_groups_x(void) __asm("llvm.r600.read.ngroups.x");
uint __clc_r600_get_num_groups_y(void) __asm("llvm.r600.read.ngroups.y");
uint __clc_r600_get_num_groups_z(void) __asm("llvm.r600.read.ngroups.z");
_CLC_DEF size_t get_num_groups(uint dim)
{
switch (dim) {
case 0: return __clc_r600_get_num_groups_x();
case 1: return __clc_r600_get_num_groups_y();
case 2: return __clc_r600_get_num_groups_z();
default: return 1;
}
_CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) {
switch (dim) {
case 0:
return __clc_r600_get_num_groups_x();
case 1:
return __clc_r600_get_num_groups_y();
case 2:
return __clc_r600_get_num_groups_z();
default:
return 1;
}
}

11
libclc/r600/lib/workitem/get_work_dim.cl
View File

@ -1,9 +1,8 @@
#include <clc/clc.h>
_CLC_DEF uint get_work_dim(void)
{
__attribute__((address_space(7))) uint * ptr =
(__attribute__((address_space(7))) uint *)
__builtin_r600_implicitarg_ptr();
return ptr[0];
_CLC_DEF _CLC_OVERLOAD uint get_work_dim(void) {
__attribute__((address_space(7))) uint *ptr =
(__attribute__((address_space(7)))
uint *)__builtin_r600_implicitarg_ptr();
return ptr[0];
}