gpu: intel: ocl: allow unlimited allocations

Author: spalicki

src/gpu/intel/compute/dispatch_reusable.hpp

@@ -534,6 +534,7 @@ class reusable_dispatch_t {
         dim_t max_buffer_size = 0;
         uint64_t max_buffer_size_bytes = 0;
         compile_params.num_buffers = buffers.size();
+        const size_t num_work_items = nd_range.global_range().nelems();
         for (size_t buf_idx = 0; buf_idx < buffers.size(); buf_idx++) {
             const named_buffer_t &buffer = buffers[buf_idx];
             // Copy buffer name into params
@@ -553,10 +554,10 @@ class reusable_dispatch_t {
             compile_params.buffer_types[buf_idx] = buffer.data_type;
 
             // Check buffer sizes to see if we can use int32_t offsets
-            // or do we need to use stateless addressing model
+            // and do we need to use stateless addressing model
             max_buffer_size = std::max(max_buffer_size, buffer.nelems(true));
-            max_buffer_size_bytes = std::max(
-                    max_buffer_size_bytes, buffer.size(0, true, true));
+            max_buffer_size_bytes = std::max(max_buffer_size_bytes,
+                    buffer.size(0, true, true) * num_work_items);
         }
 
         compile_params.use_int32_offset = max_buffer_size <= INT32_MAX;

src/gpu/intel/ocl/usm_utils.cpp

@@ -57,11 +57,17 @@ void *malloc_host(impl::engine_t *engine, size_t size) {
             size_t, cl_uint, cl_int *);
 
     if (size == 0) return nullptr;
+    bool large_buffer = size
+            > utils::downcast<const xpu::ocl::engine_impl_t *>(engine->impl())
+                      ->max_allocation_size();
+    static cl_bitfield properties[]
+            = {CL_MEM_FLAGS_INTEL, CL_MEM_ALLOW_UNRESTRICTED_SIZE_INTEL, 0};
 
     static xpu::ocl::ext_func_t<clHostMemAllocINTEL_func_t> ext_func(
             "clHostMemAllocINTEL");
     cl_int err;
-    void *p = ext_func(engine, get_ocl_context(engine), nullptr, size, 0, &err);
+    void *p = ext_func(engine, get_ocl_context(engine),
+            large_buffer ? properties : nullptr, size, 0, &err);
     assert(utils::one_of(err, CL_SUCCESS, CL_OUT_OF_RESOURCES,
             CL_OUT_OF_HOST_MEMORY, CL_INVALID_BUFFER_SIZE));
     return p;
@@ -72,12 +78,17 @@ void *malloc_device(impl::engine_t *engine, size_t size) {
             cl_ulong *, size_t, cl_uint, cl_int *);
 
     if (size == 0) return nullptr;
+    bool large_buffer = size
+            > utils::downcast<const xpu::ocl::engine_impl_t *>(engine->impl())
+                      ->max_allocation_size();
+    static cl_bitfield properties[]
+            = {CL_MEM_FLAGS_INTEL, CL_MEM_ALLOW_UNRESTRICTED_SIZE_INTEL, 0};
 
     static xpu::ocl::ext_func_t<clDeviceMemAllocINTEL_func_t> ext_func(
             "clDeviceMemAllocINTEL");
     cl_int err;
     void *p = ext_func(engine, get_ocl_context(engine), get_ocl_device(engine),
-            nullptr, size, 0, &err);
+            large_buffer ? properties : nullptr, size, 0, &err);
     assert(utils::one_of(err, CL_SUCCESS, CL_OUT_OF_RESOURCES,
             CL_OUT_OF_HOST_MEMORY, CL_INVALID_BUFFER_SIZE));
     return p;
@@ -88,12 +99,17 @@ void *malloc_shared(impl::engine_t *engine, size_t size) {
             cl_ulong *, size_t, cl_uint, cl_int *);
 
     if (size == 0) return nullptr;
+    bool large_buffer = size
+            > utils::downcast<const xpu::ocl::engine_impl_t *>(engine->impl())
+                      ->max_allocation_size();
+    static cl_bitfield properties[]
+            = {CL_MEM_FLAGS_INTEL, CL_MEM_ALLOW_UNRESTRICTED_SIZE_INTEL, 0};
 
     static xpu::ocl::ext_func_t<clSharedMemAllocINTEL_func_t> ext_func(
             "clSharedMemAllocINTEL");
     cl_int err;
     void *p = ext_func(engine, get_ocl_context(engine), get_ocl_device(engine),
-            nullptr, size, 0, &err);
+            large_buffer ? properties : nullptr, size, 0, &err);
     assert(utils::one_of(err, CL_SUCCESS, CL_OUT_OF_RESOURCES,
             CL_OUT_OF_HOST_MEMORY, CL_INVALID_BUFFER_SIZE));
     return p;

src/xpu/ocl/buffer_memory_storage.cpp

@@ -30,12 +30,18 @@ namespace xpu {
 namespace ocl {
 
 status_t buffer_memory_storage_t::init_allocate(size_t size) {
-    auto context
-            = utils::downcast<const xpu::ocl::engine_impl_t *>(engine()->impl())
-                      ->context();
+    auto engine_impl = utils::downcast<const xpu::ocl::engine_impl_t *>(
+            engine()->impl());
+
+    cl_mem_flags flags = CL_MEM_READ_WRITE;
+#if DNNL_GPU_VENDOR == DNNL_VENDOR_INTEL
+    if (size > engine_impl->max_allocation_size())
+        flags |= CL_MEM_ALLOW_UNRESTRICTED_SIZE_INTEL;
+#endif
+
     cl_int err;
     mem_object_ = clCreateBuffer_wrapper(
-            context, CL_MEM_READ_WRITE, size, nullptr, &err);
+            engine_impl->context(), flags, size, nullptr, &err);
     OCL_CHECK(err);
     return status::success;
 }

src/xpu/ocl/engine_impl.hpp

@@ -100,6 +100,10 @@ class engine_impl_t : public impl::engine_impl_t {
             runtime_version_.build = 0;
         }
 
+        err = clGetDeviceInfo(device_, CL_DEVICE_MAX_MEM_ALLOC_SIZE,
+                sizeof(max_allocation_size_), &max_allocation_size_, nullptr);
+        OCL_CHECK(err);
+
         return status::success;
     }
 
@@ -132,6 +136,7 @@ class engine_impl_t : public impl::engine_impl_t {
     }
 
     int get_buffer_alignment() const override { return 128; }
+    uint64_t max_allocation_size() const { return max_allocation_size_; }
 
 private:
     std::string name_;
@@ -140,6 +145,7 @@ class engine_impl_t : public impl::engine_impl_t {
     xpu::ocl::wrapper_t<cl_device_id> device_;
     xpu::ocl::wrapper_t<cl_context> context_;
     cl_platform_id platform_ = nullptr;
+    uint64_t max_allocation_size_ = 0;
     bool is_user_context_;
 };
 

src/xpu/ocl/utils.hpp

@@ -28,6 +28,9 @@
 
 #include "xpu/utils.hpp"
 
+#define CL_MEM_FLAGS_INTEL 0x10001
+#define CL_MEM_ALLOW_UNRESTRICTED_SIZE_INTEL (1 << 23)
+
 namespace dnnl {
 namespace impl {
 namespace xpu {

tests/benchdnn/dnnl_common.cpp

@@ -1028,9 +1028,9 @@ size_t get_cpu_ram_size() {
 }
 #endif
 
-int get_gpu_ram_sizes(size_t &ram_size, size_t &max_alloc_size) {
+int get_gpu_ram_size(size_t &ram_size) {
     if (!is_gpu()) return OK;
-    if (ram_size > 0 && max_alloc_size > 0) return OK;
+    if (ram_size > 0) return OK;
 
 #if DNNL_GPU_RUNTIME == DNNL_RUNTIME_OCL
     auto eng = dnnl::engine(get_test_engine(), true);
@@ -1042,26 +1042,16 @@ int get_gpu_ram_sizes(size_t &ram_size, size_t &max_alloc_size) {
             sizeof(cl_ulong), &ram_sz, nullptr);
     if (status != CL_SUCCESS) return FAIL;
 
-    cl_ulong max_alloc_sz = 0;
-    status = clGetDeviceInfo(ocl_device, CL_DEVICE_MAX_MEM_ALLOC_SIZE,
-            sizeof(cl_ulong), &max_alloc_sz, nullptr);
-    if (status != CL_SUCCESS) return FAIL;
-
     ram_size = (size_t)ram_sz;
-    max_alloc_size = (size_t)max_alloc_sz;
     return OK;
 #elif DNNL_GPU_RUNTIME == DNNL_RUNTIME_DPCPP
     auto eng = dnnl::engine(get_test_engine(), true);
     auto sycl_dev = dnnl::sycl_interop::get_device(eng);
     ram_size = (size_t)sycl_dev
                        .get_info<::sycl::info::device::global_mem_size>();
-    max_alloc_size
-            = (size_t)sycl_dev
-                      .get_info<::sycl::info::device::max_mem_alloc_size>();
     return OK;
 #endif
     ram_size = 0;
-    max_alloc_size = 0;
     return OK;
 }
 
@@ -1132,8 +1122,7 @@ int check_total_size(res_t *res, dnnl_primitive_t prim_ref) {
 
     static size_t cpu_device_capacity = get_cpu_ram_size();
     static size_t gpu_device_capacity = 0;
-    static size_t gpu_max_alloc_capacity = 0;
-    SAFE(get_gpu_ram_sizes(gpu_device_capacity, gpu_max_alloc_capacity), WARN);
+    SAFE(get_gpu_ram_size(gpu_device_capacity), WARN);
 
     const size_t device_max_capacity
             = is_cpu() ? cpu_device_capacity : gpu_device_capacity;
@@ -1171,32 +1160,13 @@ int check_total_size(res_t *res, dnnl_primitive_t prim_ref) {
             res->reason = skip_reason::not_enough_ram;
         }
 
-        const bool all_allocation_fit_limit = std::all_of(
-                check_mem_size_args.sizes.cbegin(),
-                check_mem_size_args.sizes.cend(), [&](size_t s) {
-                    const bool fit = s < gpu_max_alloc_capacity;
-                    if (!fit) {
-                        BENCHDNN_PRINT(1,
-                                "[CHECK_MEM][%s]: Allocation of size %s "
-                                "doesn't fit allocation limit of %s.\n",
-                                dir_c_str(), smart_bytes(s).c_str(),
-                                smart_bytes(gpu_max_alloc_capacity).c_str());
-                    }
-                    return fit;
-                });
-        if (!all_allocation_fit_limit) {
-            res->state = SKIPPED;
-            res->reason = skip_reason::not_enough_ram;
-        }
-
         BENCHDNN_PRINT((!fits_device_ram ? 1 : 6),
                 "[CHECK_MEM][%s]: Requested: %s; benchdnn_device_limit: %s; "
-                "device_RAM_capacity: %s; gpu_max_alloc: %s;\n",
+                "device_RAM_capacity: %s;\n",
                 dir_c_str(),
                 smart_bytes(check_mem_size_args.total_size_device).c_str(),
                 smart_bytes(benchdnn_device_limit).c_str(),
-                smart_bytes(gpu_device_capacity).c_str(),
-                smart_bytes(gpu_max_alloc_capacity).c_str());
+                smart_bytes(gpu_device_capacity).c_str());
     }
 
     // Note: in theory, `total_size_ref` itself can be smaller for a `prim_ref`

tests/benchdnn/dnnl_common.hpp

@@ -231,7 +231,7 @@ struct cpu_cache_args_t {
 };
 
 size_t get_cpu_ram_size();
-int get_gpu_ram_sizes(size_t &ram_size, size_t &max_alloc_size);
+int get_gpu_ram_size(size_t &ram_size);
 int get_cpu_cache_size(cpu_cache_args_t &cache_args);
 int get_gpu_cache_size(size_t &cache_size);
 
@@ -425,31 +425,6 @@ int create_primitive(benchdnn_dnnl_wrapper_t<dnnl_primitive_t> &primw,
                  /* need_skip = */ !is_graph_ref),
             WARN);
 
-    // The library scratchpad is allocated at create_primitive stage. The memory
-    // check is moved after the creation stage. It's necessary to check the
-    // library scratchpad size against gpu_max_alloc, otherwise, out_of_memory
-    // would be issued by the library.
-    if (res->mem_size_args.scratchpad_size > 0 && is_gpu()
-            && query_scratchpad_mode(query_attr(pdw))
-                    == dnnl_scratchpad_mode_library) {
-        static size_t gpu_device_capacity = 0;
-        static size_t gpu_max_alloc_capacity = 0;
-        SAFE(get_gpu_ram_sizes(gpu_device_capacity, gpu_max_alloc_capacity),
-                WARN);
-        const bool fit
-                = res->mem_size_args.scratchpad_size < gpu_max_alloc_capacity;
-        if (!fit) {
-            BENCHDNN_PRINT(1,
-                    "[CHECK_MEM]: Size of the scratchpad %s "
-                    "doesn't fit the allocation limit of %s.\n",
-                    smart_bytes(res->mem_size_args.scratchpad_size).c_str(),
-                    smart_bytes(gpu_max_alloc_capacity).c_str());
-            res->state = SKIPPED;
-            res->reason = skip_reason::not_enough_ram;
-            return OK;
-        }
-    }
-
     TIME_C_PRIM(DNN_SAFE(dnnl_primitive_create(&prim, pdw), WARN));
     primw.reset(prim);
 

tests/benchdnn/graph/graph_memory.cpp

@@ -34,8 +34,7 @@ size_t get_benchdnn_cpu_limit() {
 size_t get_benchdnn_device_limit() {
     if (is_cpu()) return 0;
     static size_t gpu_device_capacity = 0;
-    static size_t gpu_max_alloc_capacity = 0;
-    SAFE(get_gpu_ram_sizes(gpu_device_capacity, gpu_max_alloc_capacity), WARN);
+    SAFE(get_gpu_ram_size(gpu_device_capacity), WARN);
 
     const double benchdnn_device_limit = capacity_factor * gpu_device_capacity;
     assert(benchdnn_device_limit > 0);