diff --git a/unsupported/test/cxx11_tensor_convolution.cpp b/unsupported/test/cxx11_tensor_convolution.cpp index 4672db463..3a12dae62 100644 --- a/unsupported/test/cxx11_tensor_convolution.cpp +++ b/unsupported/test/cxx11_tensor_convolution.cpp @@ -14,15 +14,16 @@ using Eigen::Tensor; using Eigen::DefaultDevice; +template static void test_evals() { - Tensor input(3, 3); - Tensor kernel(2); + Tensor input(3, 3); + Tensor kernel(2); input.setRandom(); kernel.setRandom(); - Tensor result(2,3); + Tensor result(2,3); result.setZero(); Eigen::array::Index, 1> dims3({0}); @@ -41,15 +42,15 @@ static void test_evals() VERIFY_IS_APPROX(result(1,2), input(1,2)*kernel(0) + input(2,2)*kernel(1)); // index 5 } - +template static void test_expr() { - Tensor input(3, 3); - Tensor kernel(2, 2); + Tensor input(3, 3); + Tensor kernel(2, 2); input.setRandom(); kernel.setRandom(); - Tensor result(2,2); + Tensor result(2,2); Eigen::array dims({0, 1}); result = input.convolve(kernel, dims); @@ -63,10 +64,10 @@ static void test_expr() input(2,1)*kernel(1,0) + input(2,2)*kernel(1,1)); } - +template static void test_modes() { - Tensor input(3); - Tensor kernel(3); + Tensor input(3); + Tensor kernel(3); input(0) = 1.0f; input(1) = 2.0f; input(2) = 3.0f; @@ -74,13 +75,13 @@ static void test_modes() { kernel(1) = 1.0f; kernel(2) = 0.0f; - const Eigen::array dims{{0}}; + const Eigen::array dims({0}); Eigen::array, 1> padding; // Emulate VALID mode (as defined in // http://docs.scipy.org/doc/numpy/reference/generated/numpy.convolve.html). padding[0] = std::make_pair(0, 0); - Tensor valid(1); + Tensor valid(1); valid = input.pad(padding).convolve(kernel, dims); VERIFY_IS_EQUAL(valid.dimension(0), 1); VERIFY_IS_APPROX(valid(0), 2.5f); @@ -88,7 +89,7 @@ static void test_modes() { // Emulate SAME mode (as defined in // http://docs.scipy.org/doc/numpy/reference/generated/numpy.convolve.html). padding[0] = std::make_pair(1, 1); - Tensor same(3); + Tensor same(3); same = input.pad(padding).convolve(kernel, dims); VERIFY_IS_EQUAL(same.dimension(0), 3); VERIFY_IS_APPROX(same(0), 1.0f); @@ -98,7 +99,7 @@ static void test_modes() { // Emulate FULL mode (as defined in // http://docs.scipy.org/doc/numpy/reference/generated/numpy.convolve.html). padding[0] = std::make_pair(2, 2); - Tensor full(5); + Tensor full(5); full = input.pad(padding).convolve(kernel, dims); VERIFY_IS_EQUAL(full.dimension(0), 5); VERIFY_IS_APPROX(full(0), 0.0f); @@ -108,18 +109,18 @@ static void test_modes() { VERIFY_IS_APPROX(full(4), 1.5f); } - +template static void test_strides() { - Tensor input(13); - Tensor kernel(3); + Tensor input(13); + Tensor kernel(3); input.setRandom(); kernel.setRandom(); - const Eigen::array dims{{0}}; - const Eigen::array stride_of_3{{3}}; - const Eigen::array stride_of_2{{2}}; + const Eigen::array dims({0}); + const Eigen::array stride_of_3({3}); + const Eigen::array stride_of_2({2}); - Tensor result; + Tensor result; result = input.stride(stride_of_3).convolve(kernel, dims).stride(stride_of_2); VERIFY_IS_EQUAL(result.dimension(0), 2); @@ -129,13 +130,14 @@ static void test_strides() { input(12)*kernel(2))); } - - - void test_cxx11_tensor_convolution() { - CALL_SUBTEST(test_evals()); - CALL_SUBTEST(test_expr()); - CALL_SUBTEST(test_modes()); - CALL_SUBTEST(test_strides()); + CALL_SUBTEST(test_evals()); + CALL_SUBTEST(test_evals()); + CALL_SUBTEST(test_expr()); + CALL_SUBTEST(test_expr()); + CALL_SUBTEST(test_modes()); + CALL_SUBTEST(test_modes()); + CALL_SUBTEST(test_strides()); + CALL_SUBTEST(test_strides()); } diff --git a/unsupported/test/cxx11_tensor_cuda.cpp b/unsupported/test/cxx11_tensor_cuda.cpp index 8c1ca1bf8..78934165f 100644 --- a/unsupported/test/cxx11_tensor_cuda.cpp +++ b/unsupported/test/cxx11_tensor_cuda.cpp @@ -117,11 +117,10 @@ void test_cuda_elementwise() } } - void test_cuda_reduction() { - Tensor in1(Eigen::array(72,53,97,113)); - Tensor out(Eigen::array(72,97)); + Tensor in1(72,53,97,113); + Tensor out(72,97); in1.setRandom(); std::size_t in1_bytes = in1.size() * sizeof(float); @@ -138,8 +137,8 @@ void test_cuda_reduction() assert(cudaStreamCreate(&stream) == cudaSuccess); Eigen::GpuDevice gpu_device(&stream); - Eigen::TensorMap > gpu_in1(d_in1, Eigen::array(72,53,97,113)); - Eigen::TensorMap > gpu_out(d_out, Eigen::array(72,97)); + Eigen::TensorMap > gpu_in1(d_in1, 72,53,97,113); + Eigen::TensorMap > gpu_out(d_out, 72,97); array reduction_axis; reduction_axis[0] = 1; @@ -156,10 +155,10 @@ void test_cuda_reduction() for (int k = 0; k < 53; ++k) { for (int l = 0; l < 113; ++l) { expected = - std::max(expected, in1(Eigen::array(i, k, j, l))); + std::max(expected, in1(i, k, j, l)); } } - VERIFY_IS_APPROX(out(Eigen::array(i,j)), expected); + VERIFY_IS_APPROX(out(i,j), expected); } } } @@ -170,7 +169,7 @@ static void test_cuda_contraction() // with these dimensions, the output has 300 * 140 elements, which is // more than 30 * 1024, which is the number of threads in blocks on // a 15 SM GK110 GPU - Tensor t_left(Eigen::array(6, 50, 3, 31)); + Tensor t_left(6, 50, 3, 31); Tensor t_right(Eigen::array(3, 31, 7, 20, 1)); Tensor t_result(Eigen::array(6, 50, 7, 20, 1)); @@ -196,12 +195,9 @@ static void test_cuda_contraction() assert(cudaStreamCreate(&stream) == cudaSuccess); Eigen::GpuDevice gpu_device(&stream); - Eigen::TensorMap > - gpu_t_left(d_t_left, Eigen::array(6, 50, 3, 31)); - Eigen::TensorMap > - gpu_t_right(d_t_right, Eigen::array(3, 31, 7, 20, 1)); - Eigen::TensorMap > - gpu_t_result(d_t_result, Eigen::array(6, 50, 7, 20, 1)); + Eigen::TensorMap > gpu_t_left(d_t_left, 6, 50, 3, 31); + Eigen::TensorMap > gpu_t_right(d_t_right, 3, 31, 7, 20, 1); + Eigen::TensorMap > gpu_t_result(d_t_result, 6, 50, 7, 20, 1); typedef Eigen::Map > MapXf; MapXf m_left(t_left.data(), 300, 93); @@ -226,11 +222,12 @@ static void test_cuda_contraction() } } +template static void test_cuda_convolution_1d() { - Tensor input(Eigen::array(74,37,11,137)); - Tensor kernel(Eigen::array(4)); - Tensor out(Eigen::array(74,34,11,137)); + Tensor input(74,37,11,137); + Tensor kernel(4); + Tensor out(74,34,11,137); input = input.constant(10.0f) + input.random(); kernel = kernel.constant(7.0f) + kernel.random(); @@ -252,9 +249,9 @@ static void test_cuda_convolution_1d() assert(cudaStreamCreate(&stream) == cudaSuccess); Eigen::GpuDevice gpu_device(&stream); - Eigen::TensorMap > gpu_input(d_input, Eigen::array(74,37,11,137)); - Eigen::TensorMap > gpu_kernel(d_kernel, Eigen::array(4)); - Eigen::TensorMap > gpu_out(d_out, Eigen::array(74,34,11,137)); + Eigen::TensorMap > gpu_input(d_input, 74,37,11,137); + Eigen::TensorMap > gpu_kernel(d_kernel, 4); + Eigen::TensorMap > gpu_out(d_out, 74,34,11,137); Eigen::array dims(1); gpu_out.device(gpu_device) = gpu_input.convolve(gpu_kernel, dims); @@ -266,11 +263,9 @@ static void test_cuda_convolution_1d() for (int j = 0; j < 34; ++j) { for (int k = 0; k < 11; ++k) { for (int l = 0; l < 137; ++l) { - const float result = out(Eigen::array(i,j,k,l)); - const float expected = input(Eigen::array(i,j+0,k,l)) * kernel(Eigen::array(0)) + - input(Eigen::array(i,j+1,k,l)) * kernel(Eigen::array(1)) + - input(Eigen::array(i,j+2,k,l)) * kernel(Eigen::array(2)) + - input(Eigen::array(i,j+3,k,l)) * kernel(Eigen::array(3)); + const float result = out(i,j,k,l); + const float expected = input(i,j+0,k,l) * kernel(0) + input(i,j+1,k,l) * kernel(1) + + input(i,j+2,k,l) * kernel(2) + input(i,j+3,k,l) * kernel(3); VERIFY_IS_APPROX(result, expected); } } @@ -278,12 +273,11 @@ static void test_cuda_convolution_1d() } } - -static void test_cuda_convolution_2d() +static void test_cuda_convolution_inner_dim_col_major_1d() { - Tensor input(Eigen::array(74,37,11,137)); - Tensor kernel(Eigen::array(3,4)); - Tensor out(Eigen::array(74,35,8,137)); + Tensor input(74,9,11,7); + Tensor kernel(4); + Tensor out(71,9,11,7); input = input.constant(10.0f) + input.random(); kernel = kernel.constant(7.0f) + kernel.random(); @@ -305,9 +299,110 @@ static void test_cuda_convolution_2d() assert(cudaStreamCreate(&stream) == cudaSuccess); Eigen::GpuDevice gpu_device(&stream); - Eigen::TensorMap > gpu_input(d_input, Eigen::array(74,37,11,137)); - Eigen::TensorMap > gpu_kernel(d_kernel, Eigen::array(3,4)); - Eigen::TensorMap > gpu_out(d_out, Eigen::array(74,35,8,137)); + Eigen::TensorMap > gpu_input(d_input,74,9,11,7); + Eigen::TensorMap > gpu_kernel(d_kernel,4); + Eigen::TensorMap > gpu_out(d_out,71,9,11,7); + + Eigen::array dims(0); + gpu_out.device(gpu_device) = gpu_input.convolve(gpu_kernel, dims); + + assert(cudaMemcpyAsync(out.data(), d_out, out_bytes, cudaMemcpyDeviceToHost, gpu_device.stream()) == cudaSuccess); + assert(cudaStreamSynchronize(gpu_device.stream()) == cudaSuccess); + + for (int i = 0; i < 71; ++i) { + for (int j = 0; j < 9; ++j) { + for (int k = 0; k < 11; ++k) { + for (int l = 0; l < 7; ++l) { + const float result = out(i,j,k,l); + const float expected = input(i+0,j,k,l) * kernel(0) + input(i+1,j,k,l) * kernel(1) + + input(i+2,j,k,l) * kernel(2) + input(i+3,j,k,l) * kernel(3); + VERIFY_IS_APPROX(result, expected); + } + } + } + } +} + +static void test_cuda_convolution_inner_dim_row_major_1d() +{ + Tensor input(7,9,11,74); + Tensor kernel(4); + Tensor out(7,9,11,71); + input = input.constant(10.0f) + input.random(); + kernel = kernel.constant(7.0f) + kernel.random(); + + std::size_t input_bytes = input.size() * sizeof(float); + std::size_t kernel_bytes = kernel.size() * sizeof(float); + std::size_t out_bytes = out.size() * sizeof(float); + + float* d_input; + float* d_kernel; + float* d_out; + cudaMalloc((void**)(&d_input), input_bytes); + cudaMalloc((void**)(&d_kernel), kernel_bytes); + cudaMalloc((void**)(&d_out), out_bytes); + + cudaMemcpy(d_input, input.data(), input_bytes, cudaMemcpyHostToDevice); + cudaMemcpy(d_kernel, kernel.data(), kernel_bytes, cudaMemcpyHostToDevice); + + cudaStream_t stream; + assert(cudaStreamCreate(&stream) == cudaSuccess); + Eigen::GpuDevice gpu_device(&stream); + + Eigen::TensorMap > gpu_input(d_input, 7,9,11,74); + Eigen::TensorMap > gpu_kernel(d_kernel, 4); + Eigen::TensorMap > gpu_out(d_out, 7,9,11,71); + + Eigen::array dims(3); + gpu_out.device(gpu_device) = gpu_input.convolve(gpu_kernel, dims); + + assert(cudaMemcpyAsync(out.data(), d_out, out_bytes, cudaMemcpyDeviceToHost, gpu_device.stream()) == cudaSuccess); + assert(cudaStreamSynchronize(gpu_device.stream()) == cudaSuccess); + + for (int i = 0; i < 7; ++i) { + for (int j = 0; j < 9; ++j) { + for (int k = 0; k < 11; ++k) { + for (int l = 0; l < 71; ++l) { + const float result = out(i,j,k,l); + const float expected = input(i,j,k,l+0) * kernel(0) + input(i,j,k,l+1) * kernel(1) + + input(i,j,k,l+2) * kernel(2) + input(i,j,k,l+3) * kernel(3); + VERIFY_IS_APPROX(result, expected); + } + } + } + } +} + +template +static void test_cuda_convolution_2d() +{ + Tensor input(74,37,11,137); + Tensor kernel(3,4); + Tensor out(74,35,8,137); + input = input.constant(10.0f) + input.random(); + kernel = kernel.constant(7.0f) + kernel.random(); + + std::size_t input_bytes = input.size() * sizeof(float); + std::size_t kernel_bytes = kernel.size() * sizeof(float); + std::size_t out_bytes = out.size() * sizeof(float); + + float* d_input; + float* d_kernel; + float* d_out; + cudaMalloc((void**)(&d_input), input_bytes); + cudaMalloc((void**)(&d_kernel), kernel_bytes); + cudaMalloc((void**)(&d_out), out_bytes); + + cudaMemcpy(d_input, input.data(), input_bytes, cudaMemcpyHostToDevice); + cudaMemcpy(d_kernel, kernel.data(), kernel_bytes, cudaMemcpyHostToDevice); + + cudaStream_t stream; + assert(cudaStreamCreate(&stream) == cudaSuccess); + Eigen::GpuDevice gpu_device(&stream); + + Eigen::TensorMap > gpu_input(d_input,74,37,11,137); + Eigen::TensorMap > gpu_kernel(d_kernel,3,4); + Eigen::TensorMap > gpu_out(d_out,74,35,8,137); Eigen::array dims(1,2); gpu_out.device(gpu_device) = gpu_input.convolve(gpu_kernel, dims); @@ -319,32 +414,32 @@ static void test_cuda_convolution_2d() for (int j = 0; j < 35; ++j) { for (int k = 0; k < 8; ++k) { for (int l = 0; l < 137; ++l) { - const float result = out(Eigen::array(i,j,k,l)); - const float expected = input(Eigen::array(i,j+0,k+0,l)) * kernel(Eigen::array(0,0)) + - input(Eigen::array(i,j+1,k+0,l)) * kernel(Eigen::array(1,0)) + - input(Eigen::array(i,j+2,k+0,l)) * kernel(Eigen::array(2,0)) + - input(Eigen::array(i,j+0,k+1,l)) * kernel(Eigen::array(0,1)) + - input(Eigen::array(i,j+1,k+1,l)) * kernel(Eigen::array(1,1)) + - input(Eigen::array(i,j+2,k+1,l)) * kernel(Eigen::array(2,1)) + - input(Eigen::array(i,j+0,k+2,l)) * kernel(Eigen::array(0,2)) + - input(Eigen::array(i,j+1,k+2,l)) * kernel(Eigen::array(1,2)) + - input(Eigen::array(i,j+2,k+2,l)) * kernel(Eigen::array(2,2)) + - input(Eigen::array(i,j+0,k+3,l)) * kernel(Eigen::array(0,3)) + - input(Eigen::array(i,j+1,k+3,l)) * kernel(Eigen::array(1,3)) + - input(Eigen::array(i,j+2,k+3,l)) * kernel(Eigen::array(2,3)); - VERIFY_IS_APPROX(result, expected); + const float result = out(i,j,k,l); + const float expected = input(i,j+0,k+0,l) * kernel(0,0) + + input(i,j+1,k+0,l) * kernel(1,0) + + input(i,j+2,k+0,l) * kernel(2,0) + + input(i,j+0,k+1,l) * kernel(0,1) + + input(i,j+1,k+1,l) * kernel(1,1) + + input(i,j+2,k+1,l) * kernel(2,1) + + input(i,j+0,k+2,l) * kernel(0,2) + + input(i,j+1,k+2,l) * kernel(1,2) + + input(i,j+2,k+2,l) * kernel(2,2) + + input(i,j+0,k+3,l) * kernel(0,3) + + input(i,j+1,k+3,l) * kernel(1,3) + + input(i,j+2,k+3,l) * kernel(2,3); + VERIFY_IS_APPROX(result, expected); } } } } } - +template static void test_cuda_convolution_3d() { - Tensor input(Eigen::array(74,37,11,137,17)); - Tensor kernel(Eigen::array(3,4,2)); - Tensor out(Eigen::array(74,35,8,136,17)); + Tensor input(Eigen::array(74,37,11,137,17)); + Tensor kernel(3,4,2); + Tensor out(Eigen::array(74,35,8,136,17)); input = input.constant(10.0f) + input.random(); kernel = kernel.constant(7.0f) + kernel.random(); @@ -366,9 +461,9 @@ static void test_cuda_convolution_3d() assert(cudaStreamCreate(&stream) == cudaSuccess); Eigen::GpuDevice gpu_device(&stream); - Eigen::TensorMap > gpu_input(d_input, Eigen::array(74,37,11,137,17)); - Eigen::TensorMap > gpu_kernel(d_kernel, Eigen::array(3,4,2)); - Eigen::TensorMap > gpu_out(d_out, Eigen::array(74,35,8,136,17)); + Eigen::TensorMap > gpu_input(d_input,74,37,11,137,17); + Eigen::TensorMap > gpu_kernel(d_kernel,3,4,2); + Eigen::TensorMap > gpu_out(d_out,74,35,8,136,17); Eigen::array dims(1,2,3); gpu_out.device(gpu_device) = gpu_input.convolve(gpu_kernel, dims); @@ -381,31 +476,31 @@ static void test_cuda_convolution_3d() for (int k = 0; k < 8; ++k) { for (int l = 0; l < 136; ++l) { for (int m = 0; m < 17; ++m) { - const float result = out(Eigen::array(i,j,k,l,m)); - const float expected = input(Eigen::array(i,j+0,k+0,l+0,m)) * kernel(Eigen::array(0,0,0)) + - input(Eigen::array(i,j+1,k+0,l+0,m)) * kernel(Eigen::array(1,0,0)) + - input(Eigen::array(i,j+2,k+0,l+0,m)) * kernel(Eigen::array(2,0,0)) + - input(Eigen::array(i,j+0,k+1,l+0,m)) * kernel(Eigen::array(0,1,0)) + - input(Eigen::array(i,j+1,k+1,l+0,m)) * kernel(Eigen::array(1,1,0)) + - input(Eigen::array(i,j+2,k+1,l+0,m)) * kernel(Eigen::array(2,1,0)) + - input(Eigen::array(i,j+0,k+2,l+0,m)) * kernel(Eigen::array(0,2,0)) + - input(Eigen::array(i,j+1,k+2,l+0,m)) * kernel(Eigen::array(1,2,0)) + - input(Eigen::array(i,j+2,k+2,l+0,m)) * kernel(Eigen::array(2,2,0)) + - input(Eigen::array(i,j+0,k+3,l+0,m)) * kernel(Eigen::array(0,3,0)) + - input(Eigen::array(i,j+1,k+3,l+0,m)) * kernel(Eigen::array(1,3,0)) + - input(Eigen::array(i,j+2,k+3,l+0,m)) * kernel(Eigen::array(2,3,0)) + - input(Eigen::array(i,j+0,k+0,l+1,m)) * kernel(Eigen::array(0,0,1)) + - input(Eigen::array(i,j+1,k+0,l+1,m)) * kernel(Eigen::array(1,0,1)) + - input(Eigen::array(i,j+2,k+0,l+1,m)) * kernel(Eigen::array(2,0,1)) + - input(Eigen::array(i,j+0,k+1,l+1,m)) * kernel(Eigen::array(0,1,1)) + - input(Eigen::array(i,j+1,k+1,l+1,m)) * kernel(Eigen::array(1,1,1)) + - input(Eigen::array(i,j+2,k+1,l+1,m)) * kernel(Eigen::array(2,1,1)) + - input(Eigen::array(i,j+0,k+2,l+1,m)) * kernel(Eigen::array(0,2,1)) + - input(Eigen::array(i,j+1,k+2,l+1,m)) * kernel(Eigen::array(1,2,1)) + - input(Eigen::array(i,j+2,k+2,l+1,m)) * kernel(Eigen::array(2,2,1)) + - input(Eigen::array(i,j+0,k+3,l+1,m)) * kernel(Eigen::array(0,3,1)) + - input(Eigen::array(i,j+1,k+3,l+1,m)) * kernel(Eigen::array(1,3,1)) + - input(Eigen::array(i,j+2,k+3,l+1,m)) * kernel(Eigen::array(2,3,1)); + const float result = out(i,j,k,l,m); + const float expected = input(i,j+0,k+0,l+0,m) * kernel(0,0,0) + + input(i,j+1,k+0,l+0,m) * kernel(1,0,0) + + input(i,j+2,k+0,l+0,m) * kernel(2,0,0) + + input(i,j+0,k+1,l+0,m) * kernel(0,1,0) + + input(i,j+1,k+1,l+0,m) * kernel(1,1,0) + + input(i,j+2,k+1,l+0,m) * kernel(2,1,0) + + input(i,j+0,k+2,l+0,m) * kernel(0,2,0) + + input(i,j+1,k+2,l+0,m) * kernel(1,2,0) + + input(i,j+2,k+2,l+0,m) * kernel(2,2,0) + + input(i,j+0,k+3,l+0,m) * kernel(0,3,0) + + input(i,j+1,k+3,l+0,m) * kernel(1,3,0) + + input(i,j+2,k+3,l+0,m) * kernel(2,3,0) + + input(i,j+0,k+0,l+1,m) * kernel(0,0,1) + + input(i,j+1,k+0,l+1,m) * kernel(1,0,1) + + input(i,j+2,k+0,l+1,m) * kernel(2,0,1) + + input(i,j+0,k+1,l+1,m) * kernel(0,1,1) + + input(i,j+1,k+1,l+1,m) * kernel(1,1,1) + + input(i,j+2,k+1,l+1,m) * kernel(2,1,1) + + input(i,j+0,k+2,l+1,m) * kernel(0,2,1) + + input(i,j+1,k+2,l+1,m) * kernel(1,2,1) + + input(i,j+2,k+2,l+1,m) * kernel(2,2,1) + + input(i,j+0,k+3,l+1,m) * kernel(0,3,1) + + input(i,j+1,k+3,l+1,m) * kernel(1,3,1) + + input(i,j+2,k+3,l+1,m) * kernel(2,3,1); VERIFY_IS_APPROX(result, expected); } } @@ -414,91 +509,6 @@ static void test_cuda_convolution_3d() } } -static float* CudaCopyFloat(float* data, int size) { - const int nbytes = size * sizeof(float); - float* result = NULL; - if (cudaMalloc((void**)(&result), nbytes) != cudaSuccess) { - return NULL; - } else { - if (data != NULL) { - cudaMemcpy(result, data, nbytes, cudaMemcpyHostToDevice); - } - return result; - } -} - -static void test_cuda_constant_broadcast() -{ - cudaStream_t stream; - assert(cudaStreamCreate(&stream) == cudaSuccess); - Eigen::GpuDevice gpu_device(&stream); - - Tensor t1(10); - for (int i = 0; i < 10; ++i) { - t1(i) = 10.0f * i; - } - float* t1_cuda = CudaCopyFloat(t1.data(), t1.size()); - Eigen::TensorMap > t1_gpu(t1_cuda, 10); - - Tensor t2(1); - t2 = t2.constant(20.0f); - float* t2_cuda = CudaCopyFloat(t2.data(), t2.size()); - Eigen::TensorMap > > t2_gpu(t2_cuda, 1); - - float* t3_cuda = CudaCopyFloat(NULL, 10); - Eigen::TensorMap > t3_gpu(t3_cuda, 10); - - t3_gpu.device(gpu_device) = - t1_gpu + t2_gpu.broadcast(Eigen::array(10)); - - Eigen::Tensor t3(10); - cudaMemcpy(t3.data(), t3_gpu.data(), 10 * sizeof(float), - cudaMemcpyDeviceToHost); - - for (int i = 0; i < 10; ++i) { - VERIFY_IS_APPROX(t3(i), t1(i) + t2(0)); - } -} - - -void test_cuda_cast() -{ - Tensor in(Eigen::array(72,53,97)); - Tensor out(Eigen::array(72,53,97)); - in.setRandom(); - - std::size_t in_bytes = in.size() * sizeof(double); - std::size_t out_bytes = out.size() * sizeof(float); - - double* d_in; - float* d_out; - cudaMalloc((void**)(&d_in), in_bytes); - cudaMalloc((void**)(&d_out), out_bytes); - - cudaMemcpy(d_in, in.data(), in_bytes, cudaMemcpyHostToDevice); - - cudaStream_t stream; - assert(cudaStreamCreate(&stream) == cudaSuccess); - Eigen::GpuDevice gpu_device(&stream); - - Eigen::TensorMap > gpu_in(d_in, Eigen::array(72,53,97)); - Eigen::TensorMap > gpu_out(d_out, Eigen::array(72,53,97)); - - gpu_out.device(gpu_device) = gpu_in.template cast(); - - assert(cudaMemcpyAsync(out.data(), d_out, out_bytes, cudaMemcpyDeviceToHost, gpu_device.stream()) == cudaSuccess); - assert(cudaStreamSynchronize(gpu_device.stream()) == cudaSuccess); - - for (int i = 0; i < 72; ++i) { - for (int j = 0; j < 53; ++j) { - for (int k = 0; k < 97; ++k) { - VERIFY_IS_APPROX(out(Eigen::array(i,j,k)), static_cast(in(Eigen::array(i,j,k)))); - } - } - } -} - - void test_cxx11_tensor_cuda() { CALL_SUBTEST(test_cuda_elementwise_small()); @@ -506,9 +516,12 @@ void test_cxx11_tensor_cuda() CALL_SUBTEST(test_cuda_reduction()); CALL_SUBTEST(test_cuda_contraction()); CALL_SUBTEST(test_cuda_contraction()); - CALL_SUBTEST(test_cuda_convolution_1d()); - CALL_SUBTEST(test_cuda_convolution_2d()); - CALL_SUBTEST(test_cuda_convolution_3d()); - CALL_SUBTEST(test_cuda_constant_broadcast()); - CALL_SUBTEST(test_cuda_cast()); + CALL_SUBTEST(test_cuda_convolution_1d()); + CALL_SUBTEST(test_cuda_convolution_1d()); + CALL_SUBTEST(test_cuda_convolution_inner_dim_col_major_1d()); + CALL_SUBTEST(test_cuda_convolution_inner_dim_row_major_1d()); + CALL_SUBTEST(test_cuda_convolution_2d()); + CALL_SUBTEST(test_cuda_convolution_2d()); + CALL_SUBTEST(test_cuda_convolution_3d()); + CALL_SUBTEST(test_cuda_convolution_3d()); }