Move test code
@@ -257,23 +257,6 @@ __global__ void test_upcast(cuComplex* arr, cuDoubleComplex* out) {
'''
-# grid synchronization
-_test_grid_sync = r'''
-#include <cooperative_groups.h>
-
-extern "C" __global__
-void test_grid_sync(const float* x1, const float* x2, float* y) {
- namespace cg = cooperative_groups;
- cg::grid_group grid = cg::this_grid();
- int size = gridDim.x * blockDim.x;
- int tid = blockDim.x * blockIdx.x + threadIdx.x;
- y[tid] = x1[tid];
- cg::sync(grid);
- y[size - tid - 1] += x2[size - tid - 1];
-}
-'''
-
-
if 'CUPY_CACHE_DIR' in os.environ:
_old_cache_dir = os.environ['CUPY_CACHE_DIR']
_is_cache_env_var_set = True
@@ -546,8 +529,19 @@ class TestRawGridSync(unittest.TestCase):
os.environ['CUPY_CACHE_DIR'] = _test_cache_dir
self.kern_grid_sync = cupy.RawKernel(
- _test_grid_sync, 'test_grid_sync',
- backend='nvcc', grid_sync=True)
+ r'''#include <cooperative_groups.h>
+
+ extern "C" __global__
+ void test_grid_sync(const float* x1, const float* x2, float* y) {
+ namespace cg = cooperative_groups;
+ cg::grid_group grid = cg::this_grid();
+ int size = gridDim.x * blockDim.x;
+ int tid = blockDim.x * blockIdx.x + threadIdx.x;
+ y[tid] = x1[tid];
+ cg::sync(grid);
+ y[size - tid - 1] += x2[size - tid - 1];
+ }''', 'test_grid_sync', backend='nvcc',
+ enable_cooperative_groups=True)
def test_grid_sync(self):
n = 10
