Skip to content

[SYCL][HIP] Support the SYCL_PI_HIP_MAX_LOCAL_MEM_SIZE environment variable #7887

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 7 commits into from
Dec 29, 2022
Merged

[SYCL][HIP] Support the SYCL_PI_HIP_MAX_LOCAL_MEM_SIZE environment variable #7887

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
e908ca2
[SYCL][HIP] Support the SYCL_PI_HIP_MAX_LOCAL_MEM_SIZE environment va…
zjin-lcf Dec 29, 2022
84adb24
fix lint failure reported by clang-format
zjin-lcf Dec 29, 2022
a315621
fix lint failure reported by clang-format
zjin-lcf Dec 29, 2022
95c071b
fix warnings treated as errors: ignoring returned value of type hipEr…
zjin-lcf Dec 29, 2022
6a632e3
fix errors: PI_CHECK_ERROR was not declared in this scope
zjin-lcf Dec 29, 2022
0c9f5d6
fix lint failure reported by clang-format
zjin-lcf Dec 29, 2022
9d0056d
Add the Controlling DPC++ HIP Plugin section in the SYCL Environment …
zjin-lcf Dec 29, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 6 additions & 0 deletions sycl/doc/EnvironmentVariables.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -163,6 +163,12 @@ Note that conflicting configuration tuples in the same list will favor the last
| -------------------- | ------ | ----------- |
| `SYCL_PI_CUDA_MAX_LOCAL_MEM_SIZE` | Integer | Specifies the maximum size of a local memory allocation in bytes. If the value exceeds the device's capabilities then a `sycl::runtime_error` is thrown. In order for the full error message to be printed, `SYCL_RT_WARNING_LEVEL=2` must be set. The default value for `SYCL_PI_CUDA_MAX_LOCAL_MEM_SIZE` is determined by the hardware. |

## Controlling DPC++ HIP Plugin

| Environment variable | Values | Description |
| -------------------- | ------ | ----------- |
| `SYCL_PI_HIP_MAX_LOCAL_MEM_SIZE` | Integer | Specifies the maximum size of a local memory allocation in bytes. If the value exceeds the device's capabilities then a `sycl::runtime_error` is thrown. In order for the full error message to be printed, `SYCL_RT_WARNING_LEVEL=2` must be set. The default value for `SYCL_PI_HIP_MAX_LOCAL_MEM_SIZE` is determined by the hardware. |

## Tools variables

| Environment variable | Values | Description |
Expand Down
21 changes: 21 additions & 0 deletions sycl/plugins/hip/pi_hip.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2929,6 +2929,27 @@ pi_result hip_piEnqueueKernelLaunch(
retImplEv->start();
}

// Set local mem max size if env var is present
static const char *local_mem_sz_ptr =
std::getenv("SYCL_PI_HIP_MAX_LOCAL_MEM_SIZE");

if (local_mem_sz_ptr) {
int device_max_local_mem = 0;
retError = PI_CHECK_ERROR(hipDeviceGetAttribute(
&device_max_local_mem, hipDeviceAttributeMaxSharedMemoryPerBlock,
command_queue->get_device()->get()));

static const int env_val = std::atoi(local_mem_sz_ptr);
if (env_val <= 0 || env_val > device_max_local_mem) {
setErrorMessage("Invalid value specified for "
"SYCL_PI_HIP_MAX_LOCAL_MEM_SIZE",
PI_ERROR_PLUGIN_SPECIFIC_ERROR);
return PI_ERROR_PLUGIN_SPECIFIC_ERROR;
}
retError = PI_CHECK_ERROR(hipFuncSetAttribute(
hipFunc, hipFuncAttributeMaxDynamicSharedMemorySize, env_val));
}

retError = PI_CHECK_ERROR(hipModuleLaunchKernel(
hipFunc, blocksPerGrid[0], blocksPerGrid[1], blocksPerGrid[2],
threadsPerBlock[0], threadsPerBlock[1], threadsPerBlock[2],
Expand Down