Index: clang/lib/Sema/OpenCLBuiltins.td
===================================================================
--- clang/lib/Sema/OpenCLBuiltins.td
+++ clang/lib/Sema/OpenCLBuiltins.td
@@ -83,6 +83,7 @@
 def FuncExtKhrMipmapImageWrites          : FunctionExtension<"cl_khr_mipmap_image_writes">;
 def FuncExtKhrGlMsaaSharing              : FunctionExtension<"cl_khr_gl_msaa_sharing">;
 
+def FuncExtOpenCLCGenericAddressSpace    : FunctionExtension<"__opencl_c_generic_address_space">;
 def FuncExtOpenCLCPipes                  : FunctionExtension<"__opencl_c_pipes">;
 def FuncExtOpenCLCWGCollectiveFunctions  : FunctionExtension<"__opencl_c_work_group_collective_functions">;
 
@@ -566,10 +567,8 @@
   }
 }
 
-let MaxVersion = CL20 in {
-  defm : MathWithPointer<[GlobalAS, LocalAS, PrivateAS]>;
-}
-let MinVersion = CL20 in {
+defm : MathWithPointer<[GlobalAS, LocalAS, PrivateAS]>;
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
   defm : MathWithPointer<[GenericAS]>;
 }
 
@@ -824,10 +823,8 @@
   }
 }
 
-let MaxVersion = CL20 in {
-  defm : VloadVstore<[GlobalAS, LocalAS, PrivateAS], 1>;
-}
-let MinVersion = CL20 in {
+defm : VloadVstore<[GlobalAS, LocalAS, PrivateAS], 1>;
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
   defm : VloadVstore<[GenericAS], 1>;
 }
 // vload with constant address space is available regardless of version.
@@ -859,10 +856,8 @@
   }
 }
 
-let MaxVersion = CL20 in {
-  defm : VloadVstoreHalf<[GlobalAS, LocalAS, PrivateAS], 1>;
-}
-let MinVersion = CL20 in {
+defm : VloadVstoreHalf<[GlobalAS, LocalAS, PrivateAS], 1>;
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
   defm : VloadVstoreHalf<[GenericAS], 1>;
 }
 // vload with constant address space is available regardless of version.
Index: clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
===================================================================
--- clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
+++ clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
@@ -1,5 +1,11 @@
-// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL1.2 -finclude-default-header %s | FileCheck %s
-// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL1.2 -fdeclare-opencl-builtins -finclude-default-header %s | FileCheck %s
+// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL1.2 -finclude-default-header %s \
+// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-NOGAS
+// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL1.2 -fdeclare-opencl-builtins -finclude-default-header %s \
+// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-NOGAS
+// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL3.0 -fdeclare-opencl-builtins -finclude-default-header %s \
+// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-GAS
+// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL3.0 -fdeclare-opencl-builtins -finclude-default-header -cl-ext=-__opencl_c_generic_address_space,-__opencl_c_pipes %s \
+// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-NOGAS
 
 // Test that mix is correctly defined.
 // CHECK-LABEL: @test_float
@@ -32,6 +38,15 @@
   size_t lid = get_local_id(0);
 }
 
+// Test that the correct builtin is called depending on the generic address
+// space feature availability.
+// CHECK-LABEL: @test_generic_optionality
+// CHECK-GAS: call spir_func float @_Z5fractfPU3AS4f
+// CHECK-NOGAS: call spir_func float @_Z5fractfPf
+void test_generic_optionality(float a, float *b) {
+  float res = fract(a, b);
+}
+
 // CHECK: attributes [[ATTR_CONST]] =
 // CHECK-SAME: readnone
 // CHECK: attributes [[ATTR_PURE]] =
Index: clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
===================================================================
--- clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
+++ clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
@@ -63,6 +63,7 @@
 
 // Enable extensions that are enabled in opencl-c-base.h.
 #if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200)
+#define __opencl_c_generic_address_space 1
 #define cl_khr_subgroup_extended_types 1
 #define cl_khr_subgroup_ballot 1
 #define cl_khr_subgroup_non_uniform_arithmetic 1