diff --git a/unsupported/test/cxx11_tensor_cuda.cu b/unsupported/test/cxx11_tensor_cuda.cu index 29ea76857..184444ddf 100644 --- a/unsupported/test/cxx11_tensor_cuda.cu +++ b/unsupported/test/cxx11_tensor_cuda.cu @@ -10,7 +10,6 @@ #define EIGEN_TEST_NO_LONGDOUBLE #define EIGEN_TEST_NO_COMPLEX #define EIGEN_TEST_FUNC cxx11_tensor_cuda -#define EIGEN_DEFAULT_DENSE_INDEX_TYPE int #define EIGEN_USE_GPU #include @@ -19,10 +18,55 @@ using Eigen::Tensor; +void test_cuda_nullary() { + Tensor in1(2); + Tensor in2(2); + in1.setRandom(); + in2.setRandom(); + + std::size_t tensor_bytes = in1.size() * sizeof(float); + + float* d_in1; + float* d_in2; + cudaMalloc((void**)(&d_in1), tensor_bytes); + cudaMalloc((void**)(&d_in2), tensor_bytes); + cudaMemcpy(d_in1, in1.data(), tensor_bytes, cudaMemcpyHostToDevice); + cudaMemcpy(d_in2, in2.data(), tensor_bytes, cudaMemcpyHostToDevice); + + Eigen::CudaStreamDevice stream; + Eigen::GpuDevice gpu_device(&stream); + + Eigen::TensorMap, Eigen::Aligned> gpu_in1( + d_in1, 2); + Eigen::TensorMap, Eigen::Aligned> gpu_in2( + d_in2, 2); + + gpu_in1.device(gpu_device) = gpu_in1.constant(3.14f); + gpu_in2.device(gpu_device) = gpu_in2.random(); + + Tensor new1(2); + Tensor new2(2); + + assert(cudaMemcpyAsync(new1.data(), d_in1, tensor_bytes, cudaMemcpyDeviceToHost, + gpu_device.stream()) == cudaSuccess); + assert(cudaMemcpyAsync(new2.data(), d_in2, tensor_bytes, cudaMemcpyDeviceToHost, + gpu_device.stream()) == cudaSuccess); + + assert(cudaStreamSynchronize(gpu_device.stream()) == cudaSuccess); + + for (int i = 0; i < 2; ++i) { + VERIFY_IS_APPROX(new1(i), 3.14f); + VERIFY_IS_NOT_EQUAL(new2(i), in2(i)); + } + + cudaFree(d_in1); + cudaFree(d_in2); +} + void test_cuda_elementwise_small() { - Tensor in1(Eigen::array(2)); - Tensor in2(Eigen::array(2)); - Tensor out(Eigen::array(2)); + Tensor in1(Eigen::array(2)); + Tensor in2(Eigen::array(2)); + Tensor out(Eigen::array(2)); in1.setRandom(); in2.setRandom(); @@ -44,11 +88,11 @@ void test_cuda_elementwise_small() { Eigen::GpuDevice gpu_device(&stream); Eigen::TensorMap, Eigen::Aligned> gpu_in1( - d_in1, Eigen::array(2)); + d_in1, Eigen::array(2)); Eigen::TensorMap, Eigen::Aligned> gpu_in2( - d_in2, Eigen::array(2)); + d_in2, Eigen::array(2)); Eigen::TensorMap, Eigen::Aligned> gpu_out( - d_out, Eigen::array(2)); + d_out, Eigen::array(2)); gpu_out.device(gpu_device) = gpu_in1 + gpu_in2; @@ -58,8 +102,8 @@ void test_cuda_elementwise_small() { for (int i = 0; i < 2; ++i) { VERIFY_IS_APPROX( - out(Eigen::array(i)), - in1(Eigen::array(i)) + in2(Eigen::array(i))); + out(Eigen::array(i)), + in1(Eigen::array(i)) + in2(Eigen::array(i))); } cudaFree(d_in1); @@ -69,10 +113,10 @@ void test_cuda_elementwise_small() { void test_cuda_elementwise() { - Tensor in1(Eigen::array(72,53,97)); - Tensor in2(Eigen::array(72,53,97)); - Tensor in3(Eigen::array(72,53,97)); - Tensor out(Eigen::array(72,53,97)); + Tensor in1(Eigen::array(72,53,97)); + Tensor in2(Eigen::array(72,53,97)); + Tensor in3(Eigen::array(72,53,97)); + Tensor out(Eigen::array(72,53,97)); in1.setRandom(); in2.setRandom(); in3.setRandom(); @@ -98,10 +142,10 @@ void test_cuda_elementwise() Eigen::CudaStreamDevice stream; Eigen::GpuDevice gpu_device(&stream); - Eigen::TensorMap > gpu_in1(d_in1, Eigen::array(72,53,97)); - Eigen::TensorMap > gpu_in2(d_in2, Eigen::array(72,53,97)); - Eigen::TensorMap > gpu_in3(d_in3, Eigen::array(72,53,97)); - Eigen::TensorMap > gpu_out(d_out, Eigen::array(72,53,97)); + Eigen::TensorMap > gpu_in1(d_in1, Eigen::array(72,53,97)); + Eigen::TensorMap > gpu_in2(d_in2, Eigen::array(72,53,97)); + Eigen::TensorMap > gpu_in3(d_in3, Eigen::array(72,53,97)); + Eigen::TensorMap > gpu_out(d_out, Eigen::array(72,53,97)); gpu_out.device(gpu_device) = gpu_in1 + gpu_in2 * gpu_in3; @@ -111,7 +155,7 @@ void test_cuda_elementwise() 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)), in1(Eigen::array(i,j,k)) + in2(Eigen::array(i,j,k)) * in3(Eigen::array(i,j,k))); + VERIFY_IS_APPROX(out(Eigen::array(i,j,k)), in1(Eigen::array(i,j,k)) + in2(Eigen::array(i,j,k)) * in3(Eigen::array(i,j,k))); } } } @@ -181,7 +225,7 @@ void test_cuda_reduction() Eigen::TensorMap > gpu_in1(d_in1, 72,53,97,113); Eigen::TensorMap > gpu_out(d_out, 72,97); - array reduction_axis; + array reduction_axis; reduction_axis[0] = 1; reduction_axis[1] = 3; @@ -214,8 +258,8 @@ void test_cuda_contraction() // more than 30 * 1024, which is the number of threads in blocks on // a 15 SM GK110 GPU 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)); + Tensor t_right(Eigen::array(3, 31, 7, 20, 1)); + Tensor t_result(Eigen::array(6, 50, 7, 20, 1)); t_left.setRandom(); t_right.setRandom(); @@ -299,7 +343,7 @@ void test_cuda_convolution_1d() Eigen::TensorMap > gpu_kernel(d_kernel, 4); Eigen::TensorMap > gpu_out(d_out, 74,34,11,137); - Eigen::array dims(1); + Eigen::array dims(1); gpu_out.device(gpu_device) = gpu_input.convolve(gpu_kernel, dims); assert(cudaMemcpyAsync(out.data(), d_out, out_bytes, cudaMemcpyDeviceToHost, gpu_device.stream()) == cudaSuccess); @@ -352,7 +396,7 @@ void test_cuda_convolution_inner_dim_col_major_1d() Eigen::TensorMap > gpu_kernel(d_kernel,4); Eigen::TensorMap > gpu_out(d_out,71,9,11,7); - Eigen::array dims(0); + 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); @@ -405,7 +449,7 @@ void test_cuda_convolution_inner_dim_row_major_1d() Eigen::TensorMap > gpu_kernel(d_kernel, 4); Eigen::TensorMap > gpu_out(d_out, 7,9,11,71); - Eigen::array dims(3); + 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); @@ -459,7 +503,7 @@ void test_cuda_convolution_2d() Eigen::TensorMap > gpu_kernel(d_kernel,3,4); Eigen::TensorMap > gpu_out(d_out,74,35,8,137); - Eigen::array dims(1,2); + Eigen::array dims(1,2); gpu_out.device(gpu_device) = gpu_input.convolve(gpu_kernel, dims); assert(cudaMemcpyAsync(out.data(), d_out, out_bytes, cudaMemcpyDeviceToHost, gpu_device.stream()) == cudaSuccess); @@ -496,9 +540,9 @@ void test_cuda_convolution_2d() template void test_cuda_convolution_3d() { - Tensor input(Eigen::array(74,37,11,137,17)); + Tensor input(Eigen::array(74,37,11,137,17)); Tensor kernel(3,4,2); - Tensor out(Eigen::array(74,35,8,136,17)); + Tensor out(Eigen::array(74,35,8,136,17)); input = input.constant(10.0f) + input.random(); kernel = kernel.constant(7.0f) + kernel.random(); @@ -523,7 +567,7 @@ void test_cuda_convolution_3d() 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); + Eigen::array dims(1,2,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); @@ -1168,6 +1212,7 @@ void test_cuda_betainc() void test_cxx11_tensor_cuda() { + CALL_SUBTEST_1(test_cuda_nullary()); CALL_SUBTEST_1(test_cuda_elementwise_small()); CALL_SUBTEST_1(test_cuda_elementwise()); CALL_SUBTEST_1(test_cuda_props());