Skip to content

Support building ArrayFire with the device symbols flag #2957

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
umar456 merged 7 commits into from
Jul 2, 2020
Merged

Support building ArrayFire with the device symbols flag #2957

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
41dfb77
Fix thrust stream function which can be called from device
umar456 Jun 30, 2020
4fe3b59
Use active stream for CUB operations in ReduceByKey functions
umar456 Jul 1, 2020
b89555e
fixes incorrect warp continue condition on block boundary
syurkevi Jul 1, 2020
e2caeb6
update reduce by key to use stream events instead of synch
syurkevi Jul 1, 2020
39b4f98
Fix compilation error in final_boundary_reduce_dim
umar456 Jul 2, 2020
7d66613
Correct the build error enum in buildProgram
umar456 Jul 2, 2020
fe42e70
Propagate build errors to user facing exceptions
umar456 Jul 2, 2020
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
4 changes: 3 additions & 1 deletion .gitignore
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,4 +19,6 @@ src/backend/cuda/cub
conanbuildinfo*
conaninfo*
conan.lock
graph_info.json
graph_info.json
.ccls-cache
.projectile
4 changes: 2 additions & 2 deletions src/backend/common/EventBase.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,7 +48,7 @@ class EventBase {
/// is executed, the event is marked complete.
///
/// \returns the error code for the mark call
ErrorType mark(QueueType &queue) noexcept {
ErrorType mark(QueueType queue) noexcept {
return NativeEventPolicy::markEvent(&e_, queue);
}

Expand All @@ -59,7 +59,7 @@ class EventBase {
/// \param queue The queue that will wait for the previous tasks to complete
///
/// \returns the error code for the wait call
ErrorType enqueueWait(QueueType &queue) noexcept {
ErrorType enqueueWait(QueueType queue) noexcept {
return NativeEventPolicy::waitForEvent(&e_, queue);
}

Expand Down
4 changes: 3 additions & 1 deletion src/backend/cpu/Event.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,12 +12,14 @@
#include <queue.hpp>
#include <af/event.h>

#include <type_traits>

namespace cpu {

class CPUEventPolicy {
public:
using EventType = queue_event;
using QueueType = queue;
using QueueType = std::add_lvalue_reference<queue>::type;
using ErrorType = int;

static int createAndMarkEvent(queue_event *e) noexcept {
Expand Down
1 change: 0 additions & 1 deletion src/backend/cuda/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -484,7 +484,6 @@ cuda_add_library(afcuda
Param.hpp
ThrustAllocator.cuh
ThrustArrayFirePolicy.hpp
ThrustArrayFirePolicy.cpp
anisotropic_diffusion.hpp
approx.hpp
arith.hpp
Expand Down
22 changes: 0 additions & 22 deletions src/backend/cuda/ThrustArrayFirePolicy.cpp
View file
Open in desktop

This file was deleted.

18 changes: 16 additions & 2 deletions src/backend/cuda/ThrustArrayFirePolicy.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,11 +18,25 @@ namespace cuda {
struct ThrustArrayFirePolicy
: thrust::device_execution_policy<ThrustArrayFirePolicy> {};

namespace {
__DH__
cudaStream_t get_stream(ThrustArrayFirePolicy);
inline cudaStream_t get_stream(ThrustArrayFirePolicy) {
#if defined(__CUDA_ARCH__)
return 0;
#else
return getActiveStream();
#endif
}

__DH__
cudaError_t synchronize_stream(ThrustArrayFirePolicy);
inline cudaError_t synchronize_stream(ThrustArrayFirePolicy) {
#if defined(__CUDA_ARCH__)
return cudaDeviceSynchronize();
#else
return cudaStreamSynchronize(getActiveStream());
#endif
}
} // namespace

template<typename T>
thrust::pair<thrust::pointer<T, ThrustArrayFirePolicy>, std::ptrdiff_t>
Expand Down
9 changes: 5 additions & 4 deletions src/backend/cuda/kernel/reduce_by_key.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -77,10 +77,11 @@ __global__ void test_needs_reduction(int *needs_another_reduction,
atomicOr(needs_another_reduction, remaining_updates);

// check across warp boundaries
if ((tid + 1) < n) { k = keys_in.ptr[tid + 1]; }

update_key = (k == shfl_down_sync(FULL_MASK, k, 1)) &&
((tid + 1) < (n - 1)) && ((threadIdx.x % 32) < 31);
update_key =
(((threadIdx.x % 32) == 31) // last thread in warp
&& (threadIdx.x < (blockDim.x - 1)) // not last thread in block
// next value valid and equal
&& ((tid + 1) < n) && (k == keys_in.ptr[tid + 1]));
remaining_updates = any_sync(FULL_MASK, update_key);

// TODO: single per warp? change to assignment rather than atomicOr
Expand Down
39 changes: 27 additions & 12 deletions src/backend/cuda/reduce_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,18 +9,21 @@

#pragma once

#include <Array.hpp>
#include <af/dim4.hpp>

#undef _GLIBCXX_USE_INT128
#include <Array.hpp>
#include <Event.hpp>
#include <err_cuda.hpp>
#include <kernel/reduce.hpp>
#include <kernel/reduce_by_key.hpp>
#include <reduce.hpp>
#include <set.hpp>
#include <complex>

#include <cub/device/device_scan.cuh>

#include <complex>

using af::dim4;
using std::swap;

Expand Down Expand Up @@ -71,9 +74,9 @@ void reduce_by_key_dim(Array<Tk> &keys_out, Array<To> &vals_out,
auto reduced_block_sizes = memAlloc<int>(numBlocksD0);

size_t temp_storage_bytes = 0;
cub::DeviceScan::InclusiveSum(NULL, temp_storage_bytes,
reduced_block_sizes.get(),
reduced_block_sizes.get(), numBlocksD0);
cub::DeviceScan::InclusiveSum(
NULL, temp_storage_bytes, reduced_block_sizes.get(),
reduced_block_sizes.get(), numBlocksD0, getActiveStream());
auto d_temp_storage = memAlloc<char>(temp_storage_bytes);

int n_reduced_host = nelems;
Expand Down Expand Up @@ -106,7 +109,8 @@ void reduce_by_key_dim(Array<Tk> &keys_out, Array<To> &vals_out,

cub::DeviceScan::InclusiveSum(
(void *)d_temp_storage.get(), temp_storage_bytes,
reduced_block_sizes.get(), reduced_block_sizes.get(), numBlocksD0);
reduced_block_sizes.get(), reduced_block_sizes.get(), numBlocksD0,
getActiveStream());

CUDA_LAUNCH((kernel::compact_dim<Tk, To>), blocks, numThreads,
reduced_block_sizes.get(), t_reduced_keys, t_reduced_vals,
Expand All @@ -116,13 +120,15 @@ void reduce_by_key_dim(Array<Tk> &keys_out, Array<To> &vals_out,
CUDA_CHECK(cudaMemcpyAsync(
&n_reduced_host, reduced_block_sizes.get() + (numBlocksD0 - 1),
sizeof(int), cudaMemcpyDeviceToHost, getActiveStream()));
Event reduce_host_event = makeEvent(getActiveStream());

// reset flags
CUDA_CHECK(cudaMemsetAsync(needs_another_reduction.get(), 0,
sizeof(int), getActiveStream()));
CUDA_CHECK(cudaMemsetAsync(needs_block_boundary_reduction.get(), 0,
sizeof(int), getActiveStream()));

reduce_host_event.block();
numBlocksD0 = divup(n_reduced_host, numThreads);

CUDA_LAUNCH((kernel::test_needs_reduction<Tk>), numBlocksD0, numThreads,
Expand All @@ -138,6 +144,7 @@ void reduce_by_key_dim(Array<Tk> &keys_out, Array<To> &vals_out,
needs_block_boundary_reduction.get(),
sizeof(int), cudaMemcpyDeviceToHost,
getActiveStream()));
CUDA_CHECK(cudaStreamSynchronize(getActiveStream()));

if (needs_block_boundary_reduction_host &&
!needs_another_reduction_host) {
Expand All @@ -151,11 +158,12 @@ void reduce_by_key_dim(Array<Tk> &keys_out, Array<To> &vals_out,
cub::DeviceScan::InclusiveSum(
(void *)d_temp_storage.get(), temp_storage_bytes,
reduced_block_sizes.get(), reduced_block_sizes.get(),
numBlocksD0);
numBlocksD0, getActiveStream());

CUDA_CHECK(cudaMemcpyAsync(
&n_reduced_host, reduced_block_sizes.get() + (numBlocksD0 - 1),
sizeof(int), cudaMemcpyDeviceToHost, getActiveStream()));
reduce_host_event.mark(getActiveStream());

CUDA_LAUNCH((kernel::compact_dim<Tk, To>), blocks, numThreads,
reduced_block_sizes.get(), reduced_keys, reduced_vals,
Expand All @@ -164,6 +172,7 @@ void reduce_by_key_dim(Array<Tk> &keys_out, Array<To> &vals_out,

swap(t_reduced_keys, reduced_keys);
swap(t_reduced_vals, reduced_vals);
reduce_host_event.block();
}
} while (needs_another_reduction_host ||
needs_block_boundary_reduction_host);
Expand Down Expand Up @@ -213,9 +222,9 @@ void reduce_by_key_first(Array<Tk> &keys_out, Array<To> &vals_out,
auto reduced_block_sizes = memAlloc<int>(numBlocksD0);

size_t temp_storage_bytes = 0;
cub::DeviceScan::InclusiveSum(NULL, temp_storage_bytes,
reduced_block_sizes.get(),
reduced_block_sizes.get(), numBlocksD0);
cub::DeviceScan::InclusiveSum(
NULL, temp_storage_bytes, reduced_block_sizes.get(),
reduced_block_sizes.get(), numBlocksD0, getActiveStream());
auto d_temp_storage = memAlloc<char>(temp_storage_bytes);

int n_reduced_host = nelems;
Expand Down Expand Up @@ -246,7 +255,8 @@ void reduce_by_key_first(Array<Tk> &keys_out, Array<To> &vals_out,

cub::DeviceScan::InclusiveSum(
(void *)d_temp_storage.get(), temp_storage_bytes,
reduced_block_sizes.get(), reduced_block_sizes.get(), numBlocksD0);
reduced_block_sizes.get(), reduced_block_sizes.get(), numBlocksD0,
getActiveStream());

CUDA_LAUNCH((kernel::compact<Tk, To>), blocks, numThreads,
reduced_block_sizes.get(), t_reduced_keys, t_reduced_vals,
Expand All @@ -256,13 +266,15 @@ void reduce_by_key_first(Array<Tk> &keys_out, Array<To> &vals_out,
CUDA_CHECK(cudaMemcpyAsync(
&n_reduced_host, reduced_block_sizes.get() + (numBlocksD0 - 1),
sizeof(int), cudaMemcpyDeviceToHost, getActiveStream()));
Event reduce_host_event = makeEvent(getActiveStream());

// reset flags
CUDA_CHECK(cudaMemsetAsync(needs_another_reduction.get(), 0,
sizeof(int), getActiveStream()));
CUDA_CHECK(cudaMemsetAsync(needs_block_boundary_reduction.get(), 0,
sizeof(int), getActiveStream()));

reduce_host_event.block();
numBlocksD0 = divup(n_reduced_host, numThreads);

CUDA_LAUNCH((kernel::test_needs_reduction<Tk>), numBlocksD0, numThreads,
Expand All @@ -278,6 +290,7 @@ void reduce_by_key_first(Array<Tk> &keys_out, Array<To> &vals_out,
needs_block_boundary_reduction.get(),
sizeof(int), cudaMemcpyDeviceToHost,
getActiveStream()));
CUDA_CHECK(cudaStreamSynchronize(getActiveStream()));

if (needs_block_boundary_reduction_host &&
!needs_another_reduction_host) {
Expand All @@ -291,11 +304,12 @@ void reduce_by_key_first(Array<Tk> &keys_out, Array<To> &vals_out,
cub::DeviceScan::InclusiveSum(
(void *)d_temp_storage.get(), temp_storage_bytes,
reduced_block_sizes.get(), reduced_block_sizes.get(),
numBlocksD0);
numBlocksD0, getActiveStream());

CUDA_CHECK(cudaMemcpyAsync(
&n_reduced_host, reduced_block_sizes.get() + (numBlocksD0 - 1),
sizeof(int), cudaMemcpyDeviceToHost, getActiveStream()));
reduce_host_event.mark(getActiveStream());

CUDA_LAUNCH((kernel::compact<Tk, To>), blocks, numThreads,
reduced_block_sizes.get(), reduced_keys, reduced_vals,
Expand All @@ -304,6 +318,7 @@ void reduce_by_key_first(Array<Tk> &keys_out, Array<To> &vals_out,

swap(t_reduced_keys, reduced_keys);
swap(t_reduced_vals, reduced_vals);
reduce_host_event.block();
}
} while (needs_another_reduction_host ||
needs_block_boundary_reduction_host);
Expand Down
43 changes: 19 additions & 24 deletions src/backend/opencl/compile_module.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,7 @@
using cl::Error;
using cl::Program;
using common::loggerFactory;
using fmt::format;
using opencl::getActiveDeviceId;
using opencl::getDevice;
using opencl::Kernel;
Expand All @@ -49,31 +50,23 @@ logger *getLogger() {
return logger.get();
}

#define SHOW_DEBUG_BUILD_INFO(PROG) \
do { \
cl_uint numDevices = PROG.getInfo<CL_PROGRAM_NUM_DEVICES>(); \
for (unsigned int i = 0; i < numDevices; ++i) { \
printf("%s\n", PROG.getBuildInfo<CL_PROGRAM_BUILD_LOG>( \
PROG.getInfo<CL_PROGRAM_DEVICES>()[i]) \
.c_str()); \
printf("%s\n", PROG.getBuildInfo<CL_PROGRAM_BUILD_OPTIONS>( \
PROG.getInfo<CL_PROGRAM_DEVICES>()[i]) \
.c_str()); \
} \
#define THROW_BUILD_LOG_EXCEPTION(PROG) \
do { \
string build_error; \
build_error.reserve(4096); \
auto devices = PROG.getInfo<CL_PROGRAM_DEVICES>(); \
for (auto &device : PROG.getInfo<CL_PROGRAM_DEVICES>()) { \
build_error += \
format("OpenCL Device: {}\n\tOptions: {}\n\tLog:\n{}\n", \
device.getInfo<CL_DEVICE_NAME>(), \
PROG.getBuildInfo<CL_PROGRAM_BUILD_OPTIONS>(device), \
PROG.getBuildInfo<CL_PROGRAM_BUILD_LOG>(device)); \
} \
string info = getEnvVar("AF_OPENCL_SHOW_BUILD_INFO"); \
if (!info.empty() && info != "0") puts(build_error.c_str()); \
AF_ERROR(build_error, AF_ERR_INTERNAL); \
} while (0)

#if defined(NDEBUG)

#define SHOW_BUILD_INFO(PROG) \
do { \
string info = getEnvVar("AF_OPENCL_SHOW_BUILD_INFO"); \
if (!info.empty() && info != "0") { SHOW_DEBUG_BUILD_INFO(PROG); } \
} while (0)

#else
#define SHOW_BUILD_INFO(PROG) SHOW_DEBUG_BUILD_INFO(PROG)
#endif

namespace opencl {

const static string DEFAULT_MACROS_STR(
Expand Down Expand Up @@ -116,7 +109,9 @@ Program buildProgram(const vector<string> &kernelSources,

retVal.build({device}, (cl_std + defaults + options.str()).c_str());
} catch (Error &err) {
if (err.err() == CL_BUILD_ERROR) { SHOW_BUILD_INFO(retVal); }
if (err.err() == CL_BUILD_PROGRAM_FAILURE) {
THROW_BUILD_LOG_EXCEPTION(retVal);
}
throw;
}
return retVal;
Expand Down
6 changes: 3 additions & 3 deletions src/backend/opencl/kernel/reduce_by_key_boundary_dim.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,9 +13,9 @@ kernel void final_boundary_reduce_dim(global int *reduced_block_sizes,
const int n, const int nBlocksZ) {
local int dim_ordering[4];

const uint lid = get_local_id(0);
const uint bid = get_group_id(0);
const uint gidx = get_global_id(0);
const uint lid = get_local_id(0);
const uint bid = get_group_id(0);
const uint gid = get_global_id(0);

const int bidy = get_group_id(1);
const int bidz = get_group_id(2) % nBlocksZ;
Expand Down
24 changes: 24 additions & 0 deletions test/reduce.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2041,3 +2041,27 @@ TEST_P(RaggedReduceMaxRangeP, rangeMaxTest) {
ASSERT_ARRAYS_EQ(idxsReducedGold, idx);

}

TEST(ReduceByKey, ISSUE_2955) {
int N = 256;
af::array val = af::randu(N);
af::array key = af::range(af::dim4(N), 0, af::dtype::s32);
key(seq(127, af::end)) = 1;

af::array ok, ov;
af::sumByKey(ok, ov, key, val);
ASSERT_EQ(ok.dims(0), 128);
ASSERT_EQ(ov.dims(0), 128);
}

TEST(ReduceByKey, ISSUE_2955_dim) {
int N = 256;
af::array val = af::randu(8, N);
af::array key = af::range(af::dim4(N), 0, af::dtype::s32);
key(seq(127, af::end)) = 1;

af::array ok, ov;
af::sumByKey(ok, ov, key, val, 1);
ASSERT_EQ(ok.dims(0), 128);
ASSERT_EQ(ov.dims(1), 128);
}