[libclc] r260777 - Split sources for amdgcn and r600

Matt Arsenault via cfe-commits cfe-commits at lists.llvm.org
Fri Feb 12 17:02:00 PST 2016 

Author: arsenm
Date: Fri Feb 12 19:01:59 2016
New Revision: 260777

URL: http://llvm.org/viewvc/llvm-project?rev=260777&view=rev
Log:
Split sources for amdgcn and r600

Most files remain in a common amdgpu directory.

Also switches barriers to to use convergent,
and use llvm.amdgcn.s.barrier.

This now requires 3.9/trunk to build amdgcn.

Added:
    libclc/trunk/amdgcn/
    libclc/trunk/amdgcn/lib/
    libclc/trunk/amdgcn/lib/OVERRIDES
    libclc/trunk/amdgcn/lib/SOURCES
    libclc/trunk/amdgcn/lib/synchronization/
    libclc/trunk/amdgcn/lib/synchronization/barrier_impl.ll
      - copied, changed from r260304, libclc/trunk/r600/lib/synchronization/barrier_impl.ll
    libclc/trunk/amdgpu/
    libclc/trunk/amdgpu/lib/
    libclc/trunk/amdgpu/lib/OVERRIDES
      - copied, changed from r260304, libclc/trunk/r600/lib/OVERRIDES
    libclc/trunk/amdgpu/lib/SOURCES
      - copied, changed from r260304, libclc/trunk/r600/lib/SOURCES
    libclc/trunk/amdgpu/lib/atomic/
    libclc/trunk/amdgpu/lib/atomic/atomic.cl
      - copied, changed from r260304, libclc/trunk/r600/lib/atomic/atomic.cl
    libclc/trunk/amdgpu/lib/image/
    libclc/trunk/amdgpu/lib/image/get_image_attributes_impl.ll
      - copied, changed from r260304, libclc/trunk/r600/lib/image/get_image_attributes_impl.ll
    libclc/trunk/amdgpu/lib/image/get_image_channel_data_type.cl
      - copied, changed from r260304, libclc/trunk/r600/lib/image/get_image_channel_data_type.cl
    libclc/trunk/amdgpu/lib/image/get_image_channel_order.cl
      - copied, changed from r260304, libclc/trunk/r600/lib/image/get_image_channel_order.cl
    libclc/trunk/amdgpu/lib/image/get_image_depth.cl
      - copied, changed from r260304, libclc/trunk/r600/lib/image/get_image_depth.cl
    libclc/trunk/amdgpu/lib/image/get_image_height.cl
      - copied, changed from r260304, libclc/trunk/r600/lib/image/get_image_height.cl
    libclc/trunk/amdgpu/lib/image/get_image_width.cl
      - copied, changed from r260304, libclc/trunk/r600/lib/image/get_image_width.cl
    libclc/trunk/amdgpu/lib/image/read_image_impl.ll
      - copied, changed from r260304, libclc/trunk/r600/lib/image/read_image_impl.ll
    libclc/trunk/amdgpu/lib/image/read_imagef.cl
      - copied, changed from r260304, libclc/trunk/r600/lib/image/read_imagef.cl
    libclc/trunk/amdgpu/lib/image/read_imagei.cl
      - copied, changed from r260304, libclc/trunk/r600/lib/image/read_imagei.cl
    libclc/trunk/amdgpu/lib/image/read_imageui.cl
      - copied, changed from r260304, libclc/trunk/r600/lib/image/read_imageui.cl
    libclc/trunk/amdgpu/lib/image/write_image_impl.ll
      - copied, changed from r260304, libclc/trunk/r600/lib/image/write_image_impl.ll
    libclc/trunk/amdgpu/lib/image/write_imagef.cl
      - copied, changed from r260304, libclc/trunk/r600/lib/image/write_imagef.cl
    libclc/trunk/amdgpu/lib/image/write_imagei.cl
      - copied, changed from r260304, libclc/trunk/r600/lib/image/write_imagei.cl
    libclc/trunk/amdgpu/lib/image/write_imageui.cl
      - copied, changed from r260304, libclc/trunk/r600/lib/image/write_imageui.cl
    libclc/trunk/amdgpu/lib/math/
    libclc/trunk/amdgpu/lib/math/ldexp.cl
      - copied, changed from r260304, libclc/trunk/r600/lib/math/ldexp.cl
    libclc/trunk/amdgpu/lib/math/nextafter.cl
      - copied, changed from r260304, libclc/trunk/r600/lib/math/nextafter.cl
    libclc/trunk/amdgpu/lib/math/sqrt.cl
      - copied, changed from r260304, libclc/trunk/r600/lib/math/sqrt.cl
    libclc/trunk/amdgpu/lib/synchronization/
    libclc/trunk/amdgpu/lib/synchronization/barrier.cl
      - copied, changed from r260304, libclc/trunk/r600/lib/synchronization/barrier.cl
    libclc/trunk/amdgpu/lib/workitem/
    libclc/trunk/amdgpu/lib/workitem/get_global_size.ll
      - copied, changed from r260304, libclc/trunk/r600/lib/workitem/get_global_size.ll
    libclc/trunk/amdgpu/lib/workitem/get_group_id.ll
      - copied, changed from r260304, libclc/trunk/r600/lib/workitem/get_group_id.ll
    libclc/trunk/amdgpu/lib/workitem/get_local_id.ll
      - copied, changed from r260304, libclc/trunk/r600/lib/workitem/get_local_id.ll
    libclc/trunk/amdgpu/lib/workitem/get_local_size.ll
      - copied, changed from r260304, libclc/trunk/r600/lib/workitem/get_local_size.ll
    libclc/trunk/amdgpu/lib/workitem/get_num_groups.ll
      - copied, changed from r260304, libclc/trunk/r600/lib/workitem/get_num_groups.ll
    libclc/trunk/amdgpu/lib/workitem/get_work_dim.ll
      - copied, changed from r260304, libclc/trunk/r600/lib/workitem/get_work_dim.ll
Removed:
    libclc/trunk/r600/lib/atomic/atomic.cl
    libclc/trunk/r600/lib/image/get_image_attributes_impl.ll
    libclc/trunk/r600/lib/image/get_image_channel_data_type.cl
    libclc/trunk/r600/lib/image/get_image_channel_order.cl
    libclc/trunk/r600/lib/image/get_image_depth.cl
    libclc/trunk/r600/lib/image/get_image_height.cl
    libclc/trunk/r600/lib/image/get_image_width.cl
    libclc/trunk/r600/lib/image/read_image_impl.ll
    libclc/trunk/r600/lib/image/read_imagef.cl
    libclc/trunk/r600/lib/image/read_imagei.cl
    libclc/trunk/r600/lib/image/read_imageui.cl
    libclc/trunk/r600/lib/image/write_image_impl.ll
    libclc/trunk/r600/lib/image/write_imagef.cl
    libclc/trunk/r600/lib/image/write_imagei.cl
    libclc/trunk/r600/lib/image/write_imageui.cl
    libclc/trunk/r600/lib/math/ldexp.cl
    libclc/trunk/r600/lib/math/nextafter.cl
    libclc/trunk/r600/lib/math/sqrt.cl
    libclc/trunk/r600/lib/synchronization/barrier.cl
    libclc/trunk/r600/lib/workitem/get_global_size.ll
    libclc/trunk/r600/lib/workitem/get_group_id.ll
    libclc/trunk/r600/lib/workitem/get_local_id.ll
    libclc/trunk/r600/lib/workitem/get_local_size.ll
    libclc/trunk/r600/lib/workitem/get_num_groups.ll
    libclc/trunk/r600/lib/workitem/get_work_dim.ll
Modified:
    libclc/trunk/configure.py
    libclc/trunk/r600/lib/OVERRIDES
    libclc/trunk/r600/lib/SOURCES
    libclc/trunk/r600/lib/synchronization/barrier_impl.ll

Added: libclc/trunk/amdgcn/lib/OVERRIDES
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/OVERRIDES?rev=260777&view=auto
==============================================================================
    (empty)

Added: libclc/trunk/amdgcn/lib/SOURCES
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/SOURCES?rev=260777&view=auto
==============================================================================
--- libclc/trunk/amdgcn/lib/SOURCES (added)
+++ libclc/trunk/amdgcn/lib/SOURCES Fri Feb 12 19:01:59 2016
@@ -0,0 +1 @@
+synchronization/barrier_impl.ll

Copied: libclc/trunk/amdgcn/lib/synchronization/barrier_impl.ll (from r260304, libclc/trunk/r600/lib/synchronization/barrier_impl.ll)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/synchronization/barrier_impl.ll?p2=libclc/trunk/amdgcn/lib/synchronization/barrier_impl.ll&p1=libclc/trunk/r600/lib/synchronization/barrier_impl.ll&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
--- libclc/trunk/r600/lib/synchronization/barrier_impl.ll (original)
+++ libclc/trunk/amdgcn/lib/synchronization/barrier_impl.ll Fri Feb 12 19:01:59 2016
@@ -1,9 +1,8 @@
-declare i32 @__clc_clk_local_mem_fence() nounwind alwaysinline
-declare i32 @__clc_clk_global_mem_fence() nounwind alwaysinline
-declare void @llvm.AMDGPU.barrier.local() nounwind noduplicate
-declare void @llvm.AMDGPU.barrier.global() nounwind noduplicate
+declare i32 @__clc_clk_local_mem_fence() #1
+declare i32 @__clc_clk_global_mem_fence() #1
+declare void @llvm.amdgcn.s.barrier() #0
 
-define void @barrier(i32 %flags) nounwind noduplicate alwaysinline {
+define void @barrier(i32 %flags) #2 {
 barrier_local_test:
   %CLK_LOCAL_MEM_FENCE = call i32 @__clc_clk_local_mem_fence()
   %0 = and i32 %flags, %CLK_LOCAL_MEM_FENCE
@@ -11,7 +10,7 @@ barrier_local_test:
   br i1 %1, label %barrier_local, label %barrier_global_test
 
 barrier_local:
-  call void @llvm.AMDGPU.barrier.local() noduplicate
+  call void @llvm.amdgcn.s.barrier()
   br label %barrier_global_test
 
 barrier_global_test:
@@ -21,9 +20,13 @@ barrier_global_test:
   br i1 %3, label %barrier_global, label %done
 
 barrier_global:
-  call void @llvm.AMDGPU.barrier.global() noduplicate
+  call void @llvm.amdgcn.s.barrier()
   br label %done
 
 done:
   ret void
 }
+
+attributes #0 = { nounwind convergent }
+attributes #1 = { nounwind alwaysinline }
+attributes #2 = { nounwind convergent alwaysinline }

Copied: libclc/trunk/amdgpu/lib/OVERRIDES (from r260304, libclc/trunk/r600/lib/OVERRIDES)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/OVERRIDES?p2=libclc/trunk/amdgpu/lib/OVERRIDES&p1=libclc/trunk/r600/lib/OVERRIDES&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/SOURCES (from r260304, libclc/trunk/r600/lib/SOURCES)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/SOURCES?p2=libclc/trunk/amdgpu/lib/SOURCES&p1=libclc/trunk/r600/lib/SOURCES&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
--- libclc/trunk/r600/lib/SOURCES (original)
+++ libclc/trunk/amdgpu/lib/SOURCES Fri Feb 12 19:01:59 2016
@@ -9,7 +9,6 @@ workitem/get_local_id.ll
 workitem/get_global_size.ll
 workitem/get_work_dim.ll
 synchronization/barrier.cl
-synchronization/barrier_impl.ll
 image/get_image_width.cl
 image/get_image_height.cl
 image/get_image_depth.cl

Copied: libclc/trunk/amdgpu/lib/atomic/atomic.cl (from r260304, libclc/trunk/r600/lib/atomic/atomic.cl)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/atomic/atomic.cl?p2=libclc/trunk/amdgpu/lib/atomic/atomic.cl&p1=libclc/trunk/r600/lib/atomic/atomic.cl&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/image/get_image_attributes_impl.ll (from r260304, libclc/trunk/r600/lib/image/get_image_attributes_impl.ll)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/image/get_image_attributes_impl.ll?p2=libclc/trunk/amdgpu/lib/image/get_image_attributes_impl.ll&p1=libclc/trunk/r600/lib/image/get_image_attributes_impl.ll&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/image/get_image_channel_data_type.cl (from r260304, libclc/trunk/r600/lib/image/get_image_channel_data_type.cl)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/image/get_image_channel_data_type.cl?p2=libclc/trunk/amdgpu/lib/image/get_image_channel_data_type.cl&p1=libclc/trunk/r600/lib/image/get_image_channel_data_type.cl&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/image/get_image_channel_order.cl (from r260304, libclc/trunk/r600/lib/image/get_image_channel_order.cl)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/image/get_image_channel_order.cl?p2=libclc/trunk/amdgpu/lib/image/get_image_channel_order.cl&p1=libclc/trunk/r600/lib/image/get_image_channel_order.cl&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/image/get_image_depth.cl (from r260304, libclc/trunk/r600/lib/image/get_image_depth.cl)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/image/get_image_depth.cl?p2=libclc/trunk/amdgpu/lib/image/get_image_depth.cl&p1=libclc/trunk/r600/lib/image/get_image_depth.cl&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/image/get_image_height.cl (from r260304, libclc/trunk/r600/lib/image/get_image_height.cl)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/image/get_image_height.cl?p2=libclc/trunk/amdgpu/lib/image/get_image_height.cl&p1=libclc/trunk/r600/lib/image/get_image_height.cl&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/image/get_image_width.cl (from r260304, libclc/trunk/r600/lib/image/get_image_width.cl)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/image/get_image_width.cl?p2=libclc/trunk/amdgpu/lib/image/get_image_width.cl&p1=libclc/trunk/r600/lib/image/get_image_width.cl&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/image/read_image_impl.ll (from r260304, libclc/trunk/r600/lib/image/read_image_impl.ll)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/image/read_image_impl.ll?p2=libclc/trunk/amdgpu/lib/image/read_image_impl.ll&p1=libclc/trunk/r600/lib/image/read_image_impl.ll&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/image/read_imagef.cl (from r260304, libclc/trunk/r600/lib/image/read_imagef.cl)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/image/read_imagef.cl?p2=libclc/trunk/amdgpu/lib/image/read_imagef.cl&p1=libclc/trunk/r600/lib/image/read_imagef.cl&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/image/read_imagei.cl (from r260304, libclc/trunk/r600/lib/image/read_imagei.cl)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/image/read_imagei.cl?p2=libclc/trunk/amdgpu/lib/image/read_imagei.cl&p1=libclc/trunk/r600/lib/image/read_imagei.cl&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/image/read_imageui.cl (from r260304, libclc/trunk/r600/lib/image/read_imageui.cl)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/image/read_imageui.cl?p2=libclc/trunk/amdgpu/lib/image/read_imageui.cl&p1=libclc/trunk/r600/lib/image/read_imageui.cl&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/image/write_image_impl.ll (from r260304, libclc/trunk/r600/lib/image/write_image_impl.ll)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/image/write_image_impl.ll?p2=libclc/trunk/amdgpu/lib/image/write_image_impl.ll&p1=libclc/trunk/r600/lib/image/write_image_impl.ll&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/image/write_imagef.cl (from r260304, libclc/trunk/r600/lib/image/write_imagef.cl)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/image/write_imagef.cl?p2=libclc/trunk/amdgpu/lib/image/write_imagef.cl&p1=libclc/trunk/r600/lib/image/write_imagef.cl&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/image/write_imagei.cl (from r260304, libclc/trunk/r600/lib/image/write_imagei.cl)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/image/write_imagei.cl?p2=libclc/trunk/amdgpu/lib/image/write_imagei.cl&p1=libclc/trunk/r600/lib/image/write_imagei.cl&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/image/write_imageui.cl (from r260304, libclc/trunk/r600/lib/image/write_imageui.cl)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/image/write_imageui.cl?p2=libclc/trunk/amdgpu/lib/image/write_imageui.cl&p1=libclc/trunk/r600/lib/image/write_imageui.cl&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/math/ldexp.cl (from r260304, libclc/trunk/r600/lib/math/ldexp.cl)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/math/ldexp.cl?p2=libclc/trunk/amdgpu/lib/math/ldexp.cl&p1=libclc/trunk/r600/lib/math/ldexp.cl&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/math/nextafter.cl (from r260304, libclc/trunk/r600/lib/math/nextafter.cl)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/math/nextafter.cl?p2=libclc/trunk/amdgpu/lib/math/nextafter.cl&p1=libclc/trunk/r600/lib/math/nextafter.cl&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/math/sqrt.cl (from r260304, libclc/trunk/r600/lib/math/sqrt.cl)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/math/sqrt.cl?p2=libclc/trunk/amdgpu/lib/math/sqrt.cl&p1=libclc/trunk/r600/lib/math/sqrt.cl&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/synchronization/barrier.cl (from r260304, libclc/trunk/r600/lib/synchronization/barrier.cl)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/synchronization/barrier.cl?p2=libclc/trunk/amdgpu/lib/synchronization/barrier.cl&p1=libclc/trunk/r600/lib/synchronization/barrier.cl&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/workitem/get_global_size.ll (from r260304, libclc/trunk/r600/lib/workitem/get_global_size.ll)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/workitem/get_global_size.ll?p2=libclc/trunk/amdgpu/lib/workitem/get_global_size.ll&p1=libclc/trunk/r600/lib/workitem/get_global_size.ll&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/workitem/get_group_id.ll (from r260304, libclc/trunk/r600/lib/workitem/get_group_id.ll)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/workitem/get_group_id.ll?p2=libclc/trunk/amdgpu/lib/workitem/get_group_id.ll&p1=libclc/trunk/r600/lib/workitem/get_group_id.ll&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/workitem/get_local_id.ll (from r260304, libclc/trunk/r600/lib/workitem/get_local_id.ll)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/workitem/get_local_id.ll?p2=libclc/trunk/amdgpu/lib/workitem/get_local_id.ll&p1=libclc/trunk/r600/lib/workitem/get_local_id.ll&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/workitem/get_local_size.ll (from r260304, libclc/trunk/r600/lib/workitem/get_local_size.ll)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/workitem/get_local_size.ll?p2=libclc/trunk/amdgpu/lib/workitem/get_local_size.ll&p1=libclc/trunk/r600/lib/workitem/get_local_size.ll&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/workitem/get_num_groups.ll (from r260304, libclc/trunk/r600/lib/workitem/get_num_groups.ll)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/workitem/get_num_groups.ll?p2=libclc/trunk/amdgpu/lib/workitem/get_num_groups.ll&p1=libclc/trunk/r600/lib/workitem/get_num_groups.ll&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Copied: libclc/trunk/amdgpu/lib/workitem/get_work_dim.ll (from r260304, libclc/trunk/r600/lib/workitem/get_work_dim.ll)
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/workitem/get_work_dim.ll?p2=libclc/trunk/amdgpu/lib/workitem/get_work_dim.ll&p1=libclc/trunk/r600/lib/workitem/get_work_dim.ll&r1=260304&r2=260777&rev=260777&view=diff
==============================================================================
    (empty)

Modified: libclc/trunk/configure.py
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/configure.py?rev=260777&r1=260776&r2=260777&view=diff
==============================================================================
--- libclc/trunk/configure.py (original)
+++ libclc/trunk/configure.py Fri Feb 12 19:01:59 2016
@@ -69,8 +69,8 @@ llvm_version = string.split(string.repla
 llvm_int_version = int(llvm_version[0]) * 100 + int(llvm_version[1]) * 10
 llvm_string_version = 'LLVM' + llvm_version[0] + '.' + llvm_version[1]
 
-if llvm_int_version < 370:
-    print "libclc requires LLVM >= 3.7"
+if llvm_int_version < 390:
+    print "libclc requires LLVM >= 3.9"
     sys.exit(1)
 
 llvm_system_libs = llvm_config(['--system-libs'])
@@ -175,8 +175,8 @@ for target in targets:
     subdirs.append("%s-%s-%s" % (arch, t_vendor, t_os))
     subdirs.append("%s-%s" % (arch, t_os))
     subdirs.append(arch)
-    if arch == 'amdgcn':
-        subdirs.append('r600')
+    if arch == 'amdgcn' or arch == 'r600':
+        subdirs.append('amdgpu')
 
   incdirs = filter(os.path.isdir,
                [os.path.join(srcdir, subdir, 'include') for subdir in subdirs])

Modified: libclc/trunk/r600/lib/OVERRIDES
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/OVERRIDES?rev=260777&r1=260776&r2=260777&view=diff
==============================================================================
--- libclc/trunk/r600/lib/OVERRIDES (original)
+++ libclc/trunk/r600/lib/OVERRIDES Fri Feb 12 19:01:59 2016
@@ -1,2 +0,0 @@
-workitem/get_group_id.cl
-workitem/get_global_size.cl

Modified: libclc/trunk/r600/lib/SOURCES
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/SOURCES?rev=260777&r1=260776&r2=260777&view=diff
==============================================================================
--- libclc/trunk/r600/lib/SOURCES (original)
+++ libclc/trunk/r600/lib/SOURCES Fri Feb 12 19:01:59 2016
@@ -1,26 +1 @@
-atomic/atomic.cl
-math/ldexp.cl
-math/nextafter.cl
-math/sqrt.cl
-workitem/get_num_groups.ll
-workitem/get_group_id.ll
-workitem/get_local_size.ll
-workitem/get_local_id.ll
-workitem/get_global_size.ll
-workitem/get_work_dim.ll
-synchronization/barrier.cl
 synchronization/barrier_impl.ll
-image/get_image_width.cl
-image/get_image_height.cl
-image/get_image_depth.cl
-image/get_image_channel_data_type.cl
-image/get_image_channel_order.cl
-image/get_image_attributes_impl.ll
-image/read_imagef.cl
-image/read_imagei.cl
-image/read_imageui.cl
-image/read_image_impl.ll
-image/write_imagef.cl
-image/write_imagei.cl
-image/write_imageui.cl
-image/write_image_impl.ll

Removed: libclc/trunk/r600/lib/atomic/atomic.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/atomic/atomic.cl?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/atomic/atomic.cl (original)
+++ libclc/trunk/r600/lib/atomic/atomic.cl (removed)
@@ -1,65 +0,0 @@
-#include <clc/clc.h>
-
-#define ATOMIC_FUNC_DEFINE(RET_SIGN, ARG_SIGN, TYPE, CL_FUNCTION, CLC_FUNCTION, CL_ADDRSPACE, LLVM_ADDRSPACE) \
-_CLC_OVERLOAD _CLC_DEF RET_SIGN TYPE CL_FUNCTION (volatile CL_ADDRSPACE RET_SIGN TYPE *p, RET_SIGN TYPE val) { \
-	return (RET_SIGN TYPE)__clc_##CLC_FUNCTION##_addr##LLVM_ADDRSPACE((volatile CL_ADDRSPACE ARG_SIGN TYPE*)p, (ARG_SIGN TYPE)val); \
-}
-
-/* For atomic functions that don't need different bitcode dependending on argument signedness */
-#define ATOMIC_FUNC_SIGN(TYPE, FUNCTION, CL_ADDRSPACE, LLVM_ADDRSPACE) \
-	_CLC_DECL signed TYPE __clc_##FUNCTION##_addr##LLVM_ADDRSPACE(volatile CL_ADDRSPACE signed TYPE*, signed TYPE); \
-	ATOMIC_FUNC_DEFINE(signed, signed, TYPE, FUNCTION, FUNCTION, CL_ADDRSPACE, LLVM_ADDRSPACE) \
-	ATOMIC_FUNC_DEFINE(unsigned, signed, TYPE, FUNCTION, FUNCTION, CL_ADDRSPACE, LLVM_ADDRSPACE)
-
-#define ATOMIC_FUNC_ADDRSPACE(TYPE, FUNCTION) \
-	ATOMIC_FUNC_SIGN(TYPE, FUNCTION, global, 1) \
-	ATOMIC_FUNC_SIGN(TYPE, FUNCTION, local, 3)
-
-#define ATOMIC_FUNC(FUNCTION) \
-	ATOMIC_FUNC_ADDRSPACE(int, FUNCTION)
-
-#define ATOMIC_FUNC_DEFINE_3_ARG(RET_SIGN, ARG_SIGN, TYPE, CL_FUNCTION, CLC_FUNCTION, CL_ADDRSPACE, LLVM_ADDRSPACE) \
-_CLC_OVERLOAD _CLC_DEF RET_SIGN TYPE CL_FUNCTION (volatile CL_ADDRSPACE RET_SIGN TYPE *p, RET_SIGN TYPE cmp, RET_SIGN TYPE val) { \
-	return (RET_SIGN TYPE)__clc_##CLC_FUNCTION##_addr##LLVM_ADDRSPACE((volatile CL_ADDRSPACE ARG_SIGN TYPE*)p, (ARG_SIGN TYPE)cmp, (ARG_SIGN TYPE)val); \
-}
-
-/* For atomic functions that don't need different bitcode dependending on argument signedness */
-#define ATOMIC_FUNC_SIGN_3_ARG(TYPE, FUNCTION, CL_ADDRSPACE, LLVM_ADDRSPACE) \
-	_CLC_DECL signed TYPE __clc_##FUNCTION##_addr##LLVM_ADDRSPACE(volatile CL_ADDRSPACE signed TYPE*, signed TYPE, signed TYPE); \
-	ATOMIC_FUNC_DEFINE_3_ARG(signed, signed, TYPE, FUNCTION, FUNCTION, CL_ADDRSPACE, LLVM_ADDRSPACE) \
-	ATOMIC_FUNC_DEFINE_3_ARG(unsigned, signed, TYPE, FUNCTION, FUNCTION, CL_ADDRSPACE, LLVM_ADDRSPACE)
-
-#define ATOMIC_FUNC_ADDRSPACE_3_ARG(TYPE, FUNCTION) \
-	ATOMIC_FUNC_SIGN_3_ARG(TYPE, FUNCTION, global, 1) \
-	ATOMIC_FUNC_SIGN_3_ARG(TYPE, FUNCTION, local, 3)
-
-#define ATOMIC_FUNC_3_ARG(FUNCTION) \
-	ATOMIC_FUNC_ADDRSPACE_3_ARG(int, FUNCTION)
-
-ATOMIC_FUNC(atomic_add)
-ATOMIC_FUNC(atomic_and)
-ATOMIC_FUNC(atomic_or)
-ATOMIC_FUNC(atomic_sub)
-ATOMIC_FUNC(atomic_xchg)
-ATOMIC_FUNC(atomic_xor)
-ATOMIC_FUNC_3_ARG(atomic_cmpxchg)
-
-_CLC_DECL signed int __clc_atomic_max_addr1(volatile global signed int*, signed int);
-_CLC_DECL signed int __clc_atomic_max_addr3(volatile local signed int*, signed int);
-_CLC_DECL uint __clc_atomic_umax_addr1(volatile global uint*, uint);
-_CLC_DECL uint __clc_atomic_umax_addr3(volatile local uint*, uint);
-
-ATOMIC_FUNC_DEFINE(signed, signed, int, atomic_max, atomic_max, global, 1)
-ATOMIC_FUNC_DEFINE(signed, signed, int, atomic_max, atomic_max, local, 3)
-ATOMIC_FUNC_DEFINE(unsigned, unsigned, int, atomic_max, atomic_umax, global, 1)
-ATOMIC_FUNC_DEFINE(unsigned, unsigned, int, atomic_max, atomic_umax, local, 3)
-
-_CLC_DECL signed int __clc_atomic_min_addr1(volatile global signed int*, signed int);
-_CLC_DECL signed int __clc_atomic_min_addr3(volatile local signed int*, signed int);
-_CLC_DECL uint __clc_atomic_umin_addr1(volatile global uint*, uint);
-_CLC_DECL uint __clc_atomic_umin_addr3(volatile local uint*, uint);
-
-ATOMIC_FUNC_DEFINE(signed, signed, int, atomic_min, atomic_min, global, 1)
-ATOMIC_FUNC_DEFINE(signed, signed, int, atomic_min, atomic_min, local, 3)
-ATOMIC_FUNC_DEFINE(unsigned, unsigned, int, atomic_min, atomic_umin, global, 1)
-ATOMIC_FUNC_DEFINE(unsigned, unsigned, int, atomic_min, atomic_umin, local, 3)

Removed: libclc/trunk/r600/lib/image/get_image_attributes_impl.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/image/get_image_attributes_impl.ll?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/image/get_image_attributes_impl.ll (original)
+++ libclc/trunk/r600/lib/image/get_image_attributes_impl.ll (removed)
@@ -1,87 +0,0 @@
-%opencl.image2d_t = type opaque
-%opencl.image3d_t = type opaque
-
-declare i32 @llvm.OpenCL.image.get.resource.id.2d(
-  %opencl.image2d_t addrspace(1)*) nounwind readnone
-declare i32 @llvm.OpenCL.image.get.resource.id.3d(
-  %opencl.image3d_t addrspace(1)*) nounwind readnone
-
-declare [3 x i32] @llvm.OpenCL.image.get.size.2d(
-  %opencl.image2d_t addrspace(1)*) nounwind readnone
-declare [3 x i32] @llvm.OpenCL.image.get.size.3d(
-  %opencl.image3d_t addrspace(1)*) nounwind readnone
-
-declare [2 x i32] @llvm.OpenCL.image.get.format.2d(
-  %opencl.image2d_t addrspace(1)*) nounwind readnone
-declare [2 x i32] @llvm.OpenCL.image.get.format.3d(
-  %opencl.image3d_t addrspace(1)*) nounwind readnone
-
-define i32 @__clc_get_image_width_2d(
-                          %opencl.image2d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [3 x i32] @llvm.OpenCL.image.get.size.2d(
-    %opencl.image2d_t addrspace(1)* %img)
-  %2 = extractvalue [3 x i32] %1, 0
-  ret i32 %2
-}
-define i32 @__clc_get_image_width_3d(
-                          %opencl.image3d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [3 x i32] @llvm.OpenCL.image.get.size.3d(
-    %opencl.image3d_t addrspace(1)* %img)
-  %2 = extractvalue [3 x i32] %1, 0
-  ret i32 %2
-}
-
-define i32 @__clc_get_image_height_2d(
-                          %opencl.image2d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [3 x i32] @llvm.OpenCL.image.get.size.2d(
-    %opencl.image2d_t addrspace(1)* %img)
-  %2 = extractvalue [3 x i32] %1, 1
-  ret i32 %2
-}
-define i32 @__clc_get_image_height_3d(
-                          %opencl.image3d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [3 x i32] @llvm.OpenCL.image.get.size.3d(
-    %opencl.image3d_t addrspace(1)* %img)
-  %2 = extractvalue [3 x i32] %1, 1
-  ret i32 %2
-}
-
-define i32 @__clc_get_image_depth_3d(
-                          %opencl.image3d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [3 x i32] @llvm.OpenCL.image.get.size.3d(
-    %opencl.image3d_t addrspace(1)* %img)
-  %2 = extractvalue [3 x i32] %1, 2
-  ret i32 %2
-}
-
-define i32 @__clc_get_image_channel_data_type_2d(
-                          %opencl.image2d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [2 x i32] @llvm.OpenCL.image.get.format.2d(
-    %opencl.image2d_t addrspace(1)* %img)
-  %2 = extractvalue [2 x i32] %1, 0
-  ret i32 %2
-}
-define i32 @__clc_get_image_channel_data_type_3d(
-                          %opencl.image3d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [2 x i32] @llvm.OpenCL.image.get.format.3d(
-    %opencl.image3d_t addrspace(1)* %img)
-  %2 = extractvalue [2 x i32] %1, 0
-  ret i32 %2
-}
-
-define i32 @__clc_get_image_channel_order_2d(
-                          %opencl.image2d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [2 x i32] @llvm.OpenCL.image.get.format.2d(
-    %opencl.image2d_t addrspace(1)* %img)
-  %2 = extractvalue [2 x i32] %1, 1
-  ret i32 %2
-}
-define i32 @__clc_get_image_channel_order_3d(
-                          %opencl.image3d_t addrspace(1)* nocapture %img) #0 {
-  %1 = tail call [2 x i32] @llvm.OpenCL.image.get.format.3d(
-    %opencl.image3d_t addrspace(1)* %img)
-  %2 = extractvalue [2 x i32] %1, 1
-  ret i32 %2
-}
-
-attributes #0 = { nounwind readnone alwaysinline }

Removed: libclc/trunk/r600/lib/image/get_image_channel_data_type.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/image/get_image_channel_data_type.cl?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/image/get_image_channel_data_type.cl (original)
+++ libclc/trunk/r600/lib/image/get_image_channel_data_type.cl (removed)
@@ -1,13 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL int __clc_get_image_channel_data_type_2d(image2d_t);
-_CLC_DECL int __clc_get_image_channel_data_type_3d(image3d_t);
-
-_CLC_OVERLOAD _CLC_DEF int
-get_image_channel_data_type(image2d_t image) {
-  return __clc_get_image_channel_data_type_2d(image);
-}
-_CLC_OVERLOAD _CLC_DEF int
-get_image_channel_data_type(image3d_t image) {
-  return __clc_get_image_channel_data_type_3d(image);
-}

Removed: libclc/trunk/r600/lib/image/get_image_channel_order.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/image/get_image_channel_order.cl?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/image/get_image_channel_order.cl (original)
+++ libclc/trunk/r600/lib/image/get_image_channel_order.cl (removed)
@@ -1,13 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL int __clc_get_image_channel_order_2d(image2d_t);
-_CLC_DECL int __clc_get_image_channel_order_3d(image3d_t);
-
-_CLC_OVERLOAD _CLC_DEF int
-get_image_channel_order(image2d_t image) {
-  return __clc_get_image_channel_order_2d(image);
-}
-_CLC_OVERLOAD _CLC_DEF int
-get_image_channel_order(image3d_t image) {
-  return __clc_get_image_channel_order_3d(image);
-}

Removed: libclc/trunk/r600/lib/image/get_image_depth.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/image/get_image_depth.cl?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/image/get_image_depth.cl (original)
+++ libclc/trunk/r600/lib/image/get_image_depth.cl (removed)
@@ -1,8 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL int __clc_get_image_depth_3d(image3d_t);
-
-_CLC_OVERLOAD _CLC_DEF int
-get_image_depth(image3d_t image) {
-	return __clc_get_image_depth_3d(image);
-}

Removed: libclc/trunk/r600/lib/image/get_image_height.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/image/get_image_height.cl?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/image/get_image_height.cl (original)
+++ libclc/trunk/r600/lib/image/get_image_height.cl (removed)
@@ -1,13 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL int __clc_get_image_height_2d(image2d_t);
-_CLC_DECL int __clc_get_image_height_3d(image3d_t);
-
-_CLC_OVERLOAD _CLC_DEF int
-get_image_height(image2d_t image) {
-  return __clc_get_image_height_2d(image);
-}
-_CLC_OVERLOAD _CLC_DEF int
-get_image_height(image3d_t image) {
-  return __clc_get_image_height_3d(image);
-}

Removed: libclc/trunk/r600/lib/image/get_image_width.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/image/get_image_width.cl?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/image/get_image_width.cl (original)
+++ libclc/trunk/r600/lib/image/get_image_width.cl (removed)
@@ -1,13 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL int __clc_get_image_width_2d(image2d_t);
-_CLC_DECL int __clc_get_image_width_3d(image3d_t);
-
-_CLC_OVERLOAD _CLC_DEF int
-get_image_width(image2d_t image) {
-  return __clc_get_image_width_2d(image);
-}
-_CLC_OVERLOAD _CLC_DEF int
-get_image_width(image3d_t image) {
-  return __clc_get_image_width_3d(image);
-}

Removed: libclc/trunk/r600/lib/image/read_image_impl.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/image/read_image_impl.ll?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/image/read_image_impl.ll (original)
+++ libclc/trunk/r600/lib/image/read_image_impl.ll (removed)
@@ -1,46 +0,0 @@
-%opencl.image2d_t = type opaque
-
-declare <4 x float> @llvm.R600.tex(<4 x float>, i32, i32, i32, i32, i32, i32,
-                                   i32, i32, i32) readnone
-declare i32 @llvm.OpenCL.image.get.resource.id.2d(
-  %opencl.image2d_t addrspace(1)*) nounwind readnone
-declare i32 @llvm.OpenCL.sampler.get.resource.id(i32) readnone
-
-define <4 x float> @__clc_v4f_from_v2f(<2 x float> %v) alwaysinline {
-  %e0 = extractelement <2 x float> %v, i32 0
-  %e1 = extractelement <2 x float> %v, i32 1
-  %res.0 = insertelement <4 x float> undef,  float %e0, i32 0
-  %res.1 = insertelement <4 x float> %res.0, float %e1, i32 1
-  %res.2 = insertelement <4 x float> %res.1, float 0.0, i32 2
-  %res.3 = insertelement <4 x float> %res.2, float 0.0, i32 3
-  ret <4 x float> %res.3
-}
-
-define <4 x float> @__clc_read_imagef_tex(
-    %opencl.image2d_t addrspace(1)* nocapture %img,
-    i32 %sampler, <2 x float> %coord) alwaysinline {
-entry:
-  %coord_v4 = call <4 x float> @__clc_v4f_from_v2f(<2 x float> %coord)
-  %smp_id = call i32 @llvm.OpenCL.sampler.get.resource.id(i32 %sampler)
-  %img_id = call i32 @llvm.OpenCL.image.get.resource.id.2d(
-      %opencl.image2d_t addrspace(1)* %img)
-  %tex_id = add i32 %img_id, 2    ; First 2 IDs are reserved.
-
-  %coord_norm = and i32 %sampler, 1
-  %is_norm = icmp eq i32 %coord_norm, 1
-  br i1 %is_norm, label %NormCoord, label %UnnormCoord
-NormCoord:
-  %data.norm = call <4 x float> @llvm.R600.tex(
-      <4 x float> %coord_v4,
-      i32 0, i32 0, i32 0,        ; Offset.
-      i32 2, i32 %smp_id,
-      i32 1, i32 1, i32 1, i32 1) ; Normalized coords.
-  ret <4 x float> %data.norm
-UnnormCoord:
-  %data.unnorm = call <4 x float> @llvm.R600.tex(
-      <4 x float> %coord_v4,
-      i32 0, i32 0, i32 0,        ; Offset.
-      i32 %tex_id, i32 %smp_id,
-      i32 0, i32 0, i32 0, i32 0) ; Unnormalized coords.
-  ret <4 x float> %data.unnorm
-}

Removed: libclc/trunk/r600/lib/image/read_imagef.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/image/read_imagef.cl?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/image/read_imagef.cl (original)
+++ libclc/trunk/r600/lib/image/read_imagef.cl (removed)
@@ -1,14 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL float4 __clc_read_imagef_tex(image2d_t, sampler_t, float2);
-
-_CLC_OVERLOAD _CLC_DEF float4 read_imagef(image2d_t image, sampler_t sampler,
-                                          int2 coord) {
-  float2 coord_float = (float2)(coord.x, coord.y);
-  return __clc_read_imagef_tex(image, sampler, coord_float);
-}
-
-_CLC_OVERLOAD _CLC_DEF float4 read_imagef(image2d_t image, sampler_t sampler,
-                                          float2 coord) {
-  return __clc_read_imagef_tex(image, sampler, coord);
-}

Removed: libclc/trunk/r600/lib/image/read_imagei.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/image/read_imagei.cl?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/image/read_imagei.cl (original)
+++ libclc/trunk/r600/lib/image/read_imagei.cl (removed)
@@ -1,23 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL float4 __clc_read_imagef_tex(image2d_t, sampler_t, float2);
-
-int4 __clc_reinterpret_v4f_to_v4i(float4 v) {
-  union {
-    int4 v4i;
-    float4 v4f;
-  } res = { .v4f = v};
-  return res.v4i;
-}
-
-_CLC_OVERLOAD _CLC_DEF int4 read_imagei(image2d_t image, sampler_t sampler,
-                                        int2 coord) {
-  float2 coord_float = (float2)(coord.x, coord.y);
-  return __clc_reinterpret_v4f_to_v4i(
-    __clc_read_imagef_tex(image, sampler, coord_float));
-}
-_CLC_OVERLOAD _CLC_DEF int4 read_imagei(image2d_t image, sampler_t sampler,
-                                        float2 coord) {
-  return __clc_reinterpret_v4f_to_v4i(
-    __clc_read_imagef_tex(image, sampler, coord));
-}

Removed: libclc/trunk/r600/lib/image/read_imageui.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/image/read_imageui.cl?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/image/read_imageui.cl (original)
+++ libclc/trunk/r600/lib/image/read_imageui.cl (removed)
@@ -1,23 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL float4 __clc_read_imagef_tex(image2d_t, sampler_t, float2);
-
-uint4 __clc_reinterpret_v4f_to_v4ui(float4 v) {
-  union {
-    uint4 v4ui;
-    float4 v4f;
-  } res = { .v4f = v};
-  return res.v4ui;
-}
-
-_CLC_OVERLOAD _CLC_DEF uint4 read_imageui(image2d_t image, sampler_t sampler,
-                                          int2 coord) {
-  float2 coord_float = (float2)(coord.x, coord.y);
-  return __clc_reinterpret_v4f_to_v4ui(
-    __clc_read_imagef_tex(image, sampler, coord_float));
-}
-_CLC_OVERLOAD _CLC_DEF uint4 read_imageui(image2d_t image, sampler_t sampler,
-                                          float2 coord) {
-  return __clc_reinterpret_v4f_to_v4ui(
-    __clc_read_imagef_tex(image, sampler, coord));
-}

Removed: libclc/trunk/r600/lib/image/write_image_impl.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/image/write_image_impl.ll?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/image/write_image_impl.ll (original)
+++ libclc/trunk/r600/lib/image/write_image_impl.ll (removed)
@@ -1,52 +0,0 @@
-%opencl.image2d_t = type opaque
-%opencl.image3d_t = type opaque
-
-declare i32 @llvm.OpenCL.image.get.resource.id.2d(
-  %opencl.image2d_t addrspace(1)*) nounwind readnone
-declare i32 @llvm.OpenCL.image.get.resource.id.3d(
-  %opencl.image3d_t addrspace(1)*) nounwind readnone
-
-declare void @llvm.r600.rat.store.typed(<4 x i32> %color, <4 x i32> %coord, i32 %rat_id)
-
-define void @__clc_write_imageui_2d(
-    %opencl.image2d_t addrspace(1)* nocapture %img,
-    <2 x i32> %coord, <4 x i32> %color) #0 {
-
-  ; Coordinate int2 -> int4.
-  %e0 = extractelement <2 x i32> %coord, i32 0
-  %e1 = extractelement <2 x i32> %coord, i32 1
-  %coord.0 = insertelement <4 x i32> undef,    i32 %e0, i32 0
-  %coord.1 = insertelement <4 x i32> %coord.0, i32 %e1, i32 1
-  %coord.2 = insertelement <4 x i32> %coord.1, i32 0,  i32 2
-  %coord.3 = insertelement <4 x i32> %coord.2, i32 0,  i32 3
-
-  ; Get RAT ID.
-  %img_id = call i32 @llvm.OpenCL.image.get.resource.id.2d(
-      %opencl.image2d_t addrspace(1)* %img)
-  %rat_id = add i32 %img_id, 1
-
-  ; Call store intrinsic.
-  call void @llvm.r600.rat.store.typed(<4 x i32> %color, <4 x i32> %coord.3, i32 %rat_id)
-  ret void
-}
-
-define void @__clc_write_imagei_2d(
-    %opencl.image2d_t addrspace(1)* nocapture %img,
-    <2 x i32> %coord, <4 x i32> %color) #0 {
-  call void @__clc_write_imageui_2d(
-      %opencl.image2d_t addrspace(1)* nocapture %img,
-      <2 x i32> %coord, <4 x i32> %color)
-  ret void
-}
-
-define void @__clc_write_imagef_2d(
-    %opencl.image2d_t addrspace(1)* nocapture %img,
-    <2 x i32> %coord, <4 x float> %color) #0 {
-  %color.i32 = bitcast <4 x float> %color to <4 x i32>
-  call void @__clc_write_imageui_2d(
-      %opencl.image2d_t addrspace(1)* nocapture %img,
-      <2 x i32> %coord, <4 x i32> %color.i32)
-  ret void
-}
-
-attributes #0 = { alwaysinline }

Removed: libclc/trunk/r600/lib/image/write_imagef.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/image/write_imagef.cl?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/image/write_imagef.cl (original)
+++ libclc/trunk/r600/lib/image/write_imagef.cl (removed)
@@ -1,9 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL void __clc_write_imagef_2d(image2d_t image, int2 coord, float4 color);
-
-_CLC_OVERLOAD _CLC_DEF void
-write_imagef(image2d_t image, int2 coord, float4 color)
-{
-  __clc_write_imagef_2d(image, coord, color);
-}

Removed: libclc/trunk/r600/lib/image/write_imagei.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/image/write_imagei.cl?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/image/write_imagei.cl (original)
+++ libclc/trunk/r600/lib/image/write_imagei.cl (removed)
@@ -1,9 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL void __clc_write_imagei_2d(image2d_t image, int2 coord, int4 color);
-
-_CLC_OVERLOAD _CLC_DEF void
-write_imagei(image2d_t image, int2 coord, int4 color)
-{
-  __clc_write_imagei_2d(image, coord, color);
-}

Removed: libclc/trunk/r600/lib/image/write_imageui.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/image/write_imageui.cl?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/image/write_imageui.cl (original)
+++ libclc/trunk/r600/lib/image/write_imageui.cl (removed)
@@ -1,9 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DECL void __clc_write_imageui_2d(image2d_t image, int2 coord, uint4 color);
-
-_CLC_OVERLOAD _CLC_DEF void
-write_imageui(image2d_t image, int2 coord, uint4 color)
-{
-  __clc_write_imageui_2d(image, coord, color);
-}

Removed: libclc/trunk/r600/lib/math/ldexp.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/math/ldexp.cl?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/math/ldexp.cl (original)
+++ libclc/trunk/r600/lib/math/ldexp.cl (removed)
@@ -1,47 +0,0 @@
-/*
- * Copyright (c) 2014 Advanced Micro Devices, Inc.
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in
- * all copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
- * THE SOFTWARE.
- */
-
-#include <clc/clc.h>
-
-#include "../../../generic/lib/clcmacro.h"
-
-#ifdef __HAS_LDEXPF__
-#define BUILTINF __builtin_amdgpu_ldexpf
-#else
-#include "math/clc_ldexp.h"
-#define BUILTINF __clc_ldexp
-#endif
-
-// This defines all the ldexp(floatN, intN) variants.
-_CLC_DEFINE_BINARY_BUILTIN(float, ldexp, BUILTINF, float, int);
-
-#ifdef cl_khr_fp64
-  #pragma OPENCL EXTENSION cl_khr_fp64 : enable
-    // This defines all the ldexp(doubleN, intN) variants.
-  _CLC_DEFINE_BINARY_BUILTIN(double, ldexp, __builtin_amdgpu_ldexp, double, int);
-#endif
-
-// This defines all the ldexp(GENTYPE, int);
-#define __CLC_BODY <../../../generic/lib/math/ldexp.inc>
-#include <clc/math/gentype.inc>
-
-#undef BUILTINF

Removed: libclc/trunk/r600/lib/math/nextafter.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/math/nextafter.cl?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/math/nextafter.cl (original)
+++ libclc/trunk/r600/lib/math/nextafter.cl (removed)
@@ -1,4 +0,0 @@
-#include <clc/clc.h>
-#include "../lib/clcmacro.h"
-
-_CLC_DEFINE_BINARY_BUILTIN(float, nextafter, __clc_nextafter, float, float)

Removed: libclc/trunk/r600/lib/math/sqrt.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/math/sqrt.cl?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/math/sqrt.cl (original)
+++ libclc/trunk/r600/lib/math/sqrt.cl (removed)
@@ -1,59 +0,0 @@
-/*
- * Copyright (c) 2015 Advanced Micro Devices, Inc.
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in
- * all copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
- * THE SOFTWARE.
- */
-
-#include <clc/clc.h>
-#include "../../../generic/lib/clcmacro.h"
-#include "math/clc_sqrt.h"
-
-_CLC_DEFINE_UNARY_BUILTIN(float, sqrt, __clc_sqrt, float)
-
-#ifdef cl_khr_fp64
-
-#pragma OPENCL EXTENSION cl_khr_fp64 : enable
-
-
-_CLC_OVERLOAD _CLC_DEF double sqrt(double x) {
-
-  uint vcc = x < 0x1p-767;
-  uint exp0 = vcc ? 0x100 : 0;
-  unsigned exp1 = vcc ? 0xffffff80 : 0;
-
-  double v01 = ldexp(x, exp0);
-  double v23 = __builtin_amdgpu_rsq(v01);
-  double v45 = v01 * v23;
-  v23 = v23 * 0.5;
-
-  double v67 = fma(-v23, v45, 0.5);
-  v45 = fma(v45, v67, v45);
-  double v89 = fma(-v45, v45, v01);
-  v23 = fma(v23, v67, v23);
-  v45 = fma(v89, v23, v45);
-  v67 = fma(-v45, v45, v01);
-  v23 = fma(v67, v23, v45);
-
-  v23 = ldexp(v23, exp1);
-  return ((x == __builtin_inf()) || (x == 0.0)) ? v01 : v23;
-}
-
-_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, sqrt, double);
-
-#endif

Removed: libclc/trunk/r600/lib/synchronization/barrier.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/synchronization/barrier.cl?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/synchronization/barrier.cl (original)
+++ libclc/trunk/r600/lib/synchronization/barrier.cl (removed)
@@ -1,10 +0,0 @@
-
-#include <clc/clc.h>
-
-_CLC_DEF int __clc_clk_local_mem_fence() {
-  return CLK_LOCAL_MEM_FENCE;
-}
-
-_CLC_DEF int __clc_clk_global_mem_fence() {
-  return CLK_GLOBAL_MEM_FENCE;
-}

Modified: libclc/trunk/r600/lib/synchronization/barrier_impl.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/synchronization/barrier_impl.ll?rev=260777&r1=260776&r2=260777&view=diff
==============================================================================
--- libclc/trunk/r600/lib/synchronization/barrier_impl.ll (original)
+++ libclc/trunk/r600/lib/synchronization/barrier_impl.ll Fri Feb 12 19:01:59 2016
@@ -1,9 +1,9 @@
-declare i32 @__clc_clk_local_mem_fence() nounwind alwaysinline
-declare i32 @__clc_clk_global_mem_fence() nounwind alwaysinline
-declare void @llvm.AMDGPU.barrier.local() nounwind noduplicate
-declare void @llvm.AMDGPU.barrier.global() nounwind noduplicate
+declare i32 @__clc_clk_local_mem_fence() #1
+declare i32 @__clc_clk_global_mem_fence() #1
+declare void @llvm.AMDGPU.barrier.local() #0
+declare void @llvm.AMDGPU.barrier.global() #0
 
-define void @barrier(i32 %flags) nounwind noduplicate alwaysinline {
+define void @barrier(i32 %flags) #2 {
 barrier_local_test:
   %CLK_LOCAL_MEM_FENCE = call i32 @__clc_clk_local_mem_fence()
   %0 = and i32 %flags, %CLK_LOCAL_MEM_FENCE
@@ -11,7 +11,7 @@ barrier_local_test:
   br i1 %1, label %barrier_local, label %barrier_global_test
 
 barrier_local:
-  call void @llvm.AMDGPU.barrier.local() noduplicate
+  call void @llvm.AMDGPU.barrier.local()
   br label %barrier_global_test
 
 barrier_global_test:
@@ -21,9 +21,13 @@ barrier_global_test:
   br i1 %3, label %barrier_global, label %done
 
 barrier_global:
-  call void @llvm.AMDGPU.barrier.global() noduplicate
+  call void @llvm.AMDGPU.barrier.global()
   br label %done
 
 done:
   ret void
 }
+
+attributes #0 = { nounwind convergent }
+attributes #1 = { nounwind alwaysinline }
+attributes #2 = { nounwind convergent alwaysinline }

Removed: libclc/trunk/r600/lib/workitem/get_global_size.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/workitem/get_global_size.ll?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/workitem/get_global_size.ll (original)
+++ libclc/trunk/r600/lib/workitem/get_global_size.ll (removed)
@@ -1,18 +0,0 @@
-declare i32 @llvm.r600.read.global.size.x() nounwind readnone
-declare i32 @llvm.r600.read.global.size.y() nounwind readnone
-declare i32 @llvm.r600.read.global.size.z() nounwind readnone
-
-define i32 @get_global_size(i32 %dim) nounwind readnone alwaysinline {
-  switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]
-x_dim:
-  %x = call i32 @llvm.r600.read.global.size.x() nounwind readnone
-  ret i32 %x
-y_dim:
-  %y = call i32 @llvm.r600.read.global.size.y() nounwind readnone
-  ret i32 %y
-z_dim:
-  %z = call i32 @llvm.r600.read.global.size.z() nounwind readnone
-  ret i32 %z
-default:
-  ret i32 0
-}

Removed: libclc/trunk/r600/lib/workitem/get_group_id.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/workitem/get_group_id.ll?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/workitem/get_group_id.ll (original)
+++ libclc/trunk/r600/lib/workitem/get_group_id.ll (removed)
@@ -1,18 +0,0 @@
-declare i32 @llvm.r600.read.tgid.x() nounwind readnone
-declare i32 @llvm.r600.read.tgid.y() nounwind readnone
-declare i32 @llvm.r600.read.tgid.z() nounwind readnone
-
-define i32 @get_group_id(i32 %dim) nounwind readnone alwaysinline {
-  switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]
-x_dim:
-  %x = call i32 @llvm.r600.read.tgid.x() nounwind readnone
-  ret i32 %x
-y_dim:
-  %y = call i32 @llvm.r600.read.tgid.y() nounwind readnone
-  ret i32 %y
-z_dim:
-  %z = call i32 @llvm.r600.read.tgid.z() nounwind readnone
-  ret i32 %z
-default:
-  ret i32 0
-}

Removed: libclc/trunk/r600/lib/workitem/get_local_id.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/workitem/get_local_id.ll?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/workitem/get_local_id.ll (original)
+++ libclc/trunk/r600/lib/workitem/get_local_id.ll (removed)
@@ -1,18 +0,0 @@
-declare i32 @llvm.r600.read.tidig.x() nounwind readnone
-declare i32 @llvm.r600.read.tidig.y() nounwind readnone
-declare i32 @llvm.r600.read.tidig.z() nounwind readnone
-
-define i32 @get_local_id(i32 %dim) nounwind readnone alwaysinline {
-  switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]
-x_dim:
-  %x = call i32 @llvm.r600.read.tidig.x() nounwind readnone
-  ret i32 %x
-y_dim:
-  %y = call i32 @llvm.r600.read.tidig.y() nounwind readnone
-  ret i32 %y
-z_dim:
-  %z = call i32 @llvm.r600.read.tidig.z() nounwind readnone
-  ret i32 %z
-default:
-  ret i32 0
-}

Removed: libclc/trunk/r600/lib/workitem/get_local_size.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/workitem/get_local_size.ll?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/workitem/get_local_size.ll (original)
+++ libclc/trunk/r600/lib/workitem/get_local_size.ll (removed)
@@ -1,18 +0,0 @@
-declare i32 @llvm.r600.read.local.size.x() nounwind readnone
-declare i32 @llvm.r600.read.local.size.y() nounwind readnone
-declare i32 @llvm.r600.read.local.size.z() nounwind readnone
-
-define i32 @get_local_size(i32 %dim) nounwind readnone alwaysinline {
-  switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]
-x_dim:
-  %x = call i32 @llvm.r600.read.local.size.x() nounwind readnone
-  ret i32 %x
-y_dim:
-  %y = call i32 @llvm.r600.read.local.size.y() nounwind readnone
-  ret i32 %y
-z_dim:
-  %z = call i32 @llvm.r600.read.local.size.z() nounwind readnone
-  ret i32 %z
-default:
-  ret i32 0
-}

Removed: libclc/trunk/r600/lib/workitem/get_num_groups.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/workitem/get_num_groups.ll?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/workitem/get_num_groups.ll (original)
+++ libclc/trunk/r600/lib/workitem/get_num_groups.ll (removed)
@@ -1,18 +0,0 @@
-declare i32 @llvm.r600.read.ngroups.x() nounwind readnone
-declare i32 @llvm.r600.read.ngroups.y() nounwind readnone
-declare i32 @llvm.r600.read.ngroups.z() nounwind readnone
-
-define i32 @get_num_groups(i32 %dim) nounwind readnone alwaysinline {
-  switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]
-x_dim:
-  %x = call i32 @llvm.r600.read.ngroups.x() nounwind readnone
-  ret i32 %x
-y_dim:
-  %y = call i32 @llvm.r600.read.ngroups.y() nounwind readnone
-  ret i32 %y
-z_dim:
-  %z = call i32 @llvm.r600.read.ngroups.z() nounwind readnone
-  ret i32 %z
-default:
-  ret i32 0
-}

Removed: libclc/trunk/r600/lib/workitem/get_work_dim.ll
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/workitem/get_work_dim.ll?rev=260776&view=auto
==============================================================================
--- libclc/trunk/r600/lib/workitem/get_work_dim.ll (original)
+++ libclc/trunk/r600/lib/workitem/get_work_dim.ll (removed)
@@ -1,8 +0,0 @@
-declare i32 @llvm.AMDGPU.read.workdim() nounwind readnone
-
-define i32 @get_work_dim() nounwind readnone alwaysinline {
-  %x = call i32 @llvm.AMDGPU.read.workdim() nounwind readnone , !range !0
-  ret i32 %x
-}
-
-!0 = !{ i32 1, i32 4 }

More information about the cfe-commits mailing list