introduce SPIR-V target for 'libc'

Author: VyacheslavLevytskyy

clang/include/clang/Basic/BuiltinsSPIRV.td

@@ -41,83 +41,95 @@ def SPIRVSmoothStep : Builtin {
 def SPIRVGetNumWrkgrpX : Builtin {
   let Spellings = ["__builtin_spirv_get_num_workgroups_x"];
   let Attributes = [NoThrow, Const];
-  let Prototype = "void(...)";
+  let Prototype = "uint32_t()";
 }
 
 def SPIRVGetNumWrkgrpY : Builtin {
   let Spellings = ["__builtin_spirv_get_num_workgroups_y"];
   let Attributes = [NoThrow, Const];
-  let Prototype = "void(...)";
+  let Prototype = "uint32_t()";
 }
 
 def SPIRVGetNumWrkgrpZ : Builtin {
   let Spellings = ["__builtin_spirv_get_num_workgroups_z"];
   let Attributes = [NoThrow, Const];
-  let Prototype = "void(...)";
+  let Prototype = "uint32_t()";
 }
 
 def SPIRVGetWrkgrpIdX : Builtin {
   let Spellings = ["__builtin_spirv_get_workgroup_id_x"];
   let Attributes = [NoThrow, Const];
-  let Prototype = "void(...)";
+  let Prototype = "uint32_t()";
 }
 
 def SPIRVGetWrkgrpIdY : Builtin {
   let Spellings = ["__builtin_spirv_get_workgroup_id_y"];
   let Attributes = [NoThrow, Const];
-  let Prototype = "void(...)";
+  let Prototype = "uint32_t()";
 }
 
 def SPIRVGetWrkgrpIdZ : Builtin {
   let Spellings = ["__builtin_spirv_get_workgroup_id_z"];
   let Attributes = [NoThrow, Const];
-  let Prototype = "void(...)";
+  let Prototype = "uint32_t()";
 }
 
 def SPIRVGetWrkgrpSizeX : Builtin {
   let Spellings = ["__builtin_spirv_workgroup_size_x"];
   let Attributes = [NoThrow, Const];
-  let Prototype = "void(...)";
+  let Prototype = "uint32_t()";
 }
 
 def SPIRVGetWrkgrpSizeY : Builtin {
   let Spellings = ["__builtin_spirv_workgroup_size_y"];
   let Attributes = [NoThrow, Const];
-  let Prototype = "void(...)";
+  let Prototype = "uint32_t()";
 }
 
 def SPIRVGetWrkgrpSizeZ : Builtin {
   let Spellings = ["__builtin_spirv_workgroup_size_z"];
   let Attributes = [NoThrow, Const];
-  let Prototype = "void(...)";
+  let Prototype = "uint32_t()";
 }
 
 def SPIRVGetWrkitemIdX : Builtin {
   let Spellings = ["__builtin_spirv_workitem_id_x"];
   let Attributes = [NoThrow, Const];
-  let Prototype = "void(...)";
+  let Prototype = "uint32_t()";
 }
 
 def SPIRVGetWrkitemIdY : Builtin {
   let Spellings = ["__builtin_spirv_workitem_id_y"];
   let Attributes = [NoThrow, Const];
-  let Prototype = "void(...)";
+  let Prototype = "uint32_t()";
 }
 
 def SPIRVGetWrkitemIdZ : Builtin {
   let Spellings = ["__builtin_spirv_workitem_id_z"];
   let Attributes = [NoThrow, Const];
-  let Prototype = "void(...)";
+  let Prototype = "uint32_t()";
 }
 
 def SPIRVBallot : Builtin {
   let Spellings = ["__builtin_spirv_ballot"];
   let Attributes = [NoThrow, Const];
-  let Prototype = "void(...)";
+  let Prototype = "uint64_t(bool)";
 }
 
 def SPIRVSyncThreads : Builtin {
   let Spellings = ["__builtin_spirv_sync_threads"];
   let Attributes = [NoThrow, Const];
-  let Prototype = "void(...)";
+  let Prototype = "void()";
+}
+
+def SPIRVIsShared : Builtin {
+  let Spellings = ["__builtin_spirv_is_shared"];
+  let Attributes = [NoThrow, Const];
+  let Prototype = "bool(void*)";
+}
+
+def SPIRVIsPrivate : Builtin {
+  let Spellings = ["__builtin_spirv_is_private"];
+  let Attributes = [NoThrow, Const];
+  let Prototype = "bool(void*)";
 }

clang/lib/Headers/spirvintrin.h

@@ -20,7 +20,7 @@
 _Pragma("omp begin declare target device_type(nohost)");
 _Pragma("omp begin declare variant match(device = {arch(amdgcn)})");
 
-// Type aliases to the address spaces used by the SPIRV backend.
+// Type aliases to the address spaces used by the SPIR-V backend.
 #define __gpu_private __attribute__((address_space(5)))
 #define __gpu_constant __attribute__((address_space(4)))
 #define __gpu_local __attribute__((address_space(3)))
@@ -30,6 +30,25 @@ _Pragma("omp begin declare variant match(device = {arch(amdgcn)})");
 // Attribute to declare a function as a kernel.
 #define __gpu_kernel __attribute__((spir_kernel, visibility("protected")))
 
+extern "C" {
+uint32_t __builtin_spirv_get_num_workgroups_x();
+uint32_t __builtin_spirv_get_num_workgroups_y();
+uint32_t __builtin_spirv_get_num_workgroups_z();
+uint32_t __builtin_spirv_get_workgroup_id_x();
+uint32_t __builtin_spirv_get_workgroup_id_y();
+uint32_t __builtin_spirv_get_workgroup_id_z();
+uint32_t __builtin_spirv_workgroup_size_x();
+uint32_t __builtin_spirv_workgroup_size_y();
+uint32_t __builtin_spirv_workgroup_size_z();
+uint32_t __builtin_spirv_workitem_id_x();
+uint32_t __builtin_spirv_workitem_id_y();
+uint32_t __builtin_spirv_workitem_id_z();
+uint64_t __builtin_spirv_ballot(bool);
+void __builtin_spirv_sync_threads();
+bool __builtin_spirv_is_shared(void [[clang::address_space(0)]] *);
+bool __builtin_spirv_is_private(void [[clang::address_space(0)]] *);
+}
+
 // Returns the number of workgroups in the 'x' dimension of the grid.
 _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_blocks_x(void) {
   return __builtin_spirv_get_num_workgroups_x();

libc/src/__support/time/gpu/time_utils.h

@@ -34,6 +34,9 @@ extern gpu::Constant<uint64_t> __llvm_libc_clock_freq;
 #elif defined(LIBC_TARGET_ARCH_IS_NVPTX)
 // NPVTX uses a single 1 GHz fixed frequency clock for all target architectures.
 #define GPU_CLOCKS_PER_SEC static_cast<clock_t>(1000000000UL)
+#elif defined(LIBC_TARGET_ARCH_IS_SPIRV)
+// NPVTX uses a single 1 GHz fixed frequency clock for all target architectures.
+#define GPU_CLOCKS_PER_SEC static_cast<clock_t>(1000000000UL)
 #else
 #error "Unsupported target"
 #endif