Skip to content

Commit

Permalink
Revert "Single polarization transpose works in tests and tutorial"
Browse files Browse the repository at this point in the history
This reverts commit 0388e5f.
  • Loading branch information
dentalfloss1 committed May 31, 2023
1 parent 0388e5f commit 30fd268
Show file tree
Hide file tree
Showing 3 changed files with 22 additions and 37 deletions.
7 changes: 4 additions & 3 deletions src/map.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -129,6 +129,7 @@ BFstatus build_map_kernel(int* external_ndim,
::memcpy(shape, external_shape, ndim*sizeof(*shape));
long y_size = 1;
long x_size = 1;

std::vector<BFarray> mutable_arrays;
std::vector<BFarray*> mutable_array_ptrs;
if( basic_indexing_only ) {
Expand Down Expand Up @@ -362,8 +363,8 @@ BFstatus build_map_kernel(int* external_ndim,
size_t logsize;
// Note: Includes the trailing NULL
BF_CHECK_NVRTC( nvrtcGetProgramLogSize(program, &logsize) );
if( (logsize > 1 || EnvVars::get("BF_PRINT_MAP_KERNELS", "0") != "0") &&
!basic_indexing_only ) {
if( (logsize > 1 || EnvVars::get("BF_PRINT_MAP_KERNELS", "0") != "0") &&
!basic_indexing_only ) {
std::vector<char> log(logsize, 0);
BF_CHECK_NVRTC( nvrtcGetProgramLog(program, &log[0]) );
int i = 1;
Expand All @@ -375,7 +376,7 @@ BFstatus build_map_kernel(int* external_ndim,
std::cout << "---------------------------------------------------" << std::endl;
std::cout << &log[0] << std::endl;
std::cout << "---------------------------------------------------" << std::endl;
}
}
#endif // BIFROST_DEBUG
if( ret != NVRTC_SUCCESS ) {
// Note: Don't print debug msg here, failure may not be expected
Expand Down
48 changes: 17 additions & 31 deletions src/transpose.cu
Original file line number Diff line number Diff line change
Expand Up @@ -340,7 +340,7 @@ BFstatus transpose_simple(BFarray const* in,
char const* arg_names[] = {"in", "out"};
char const* func = func_str.c_str();
char const* extra_code = 0;
return bfMap(ndim, out->shape, axis_names, narg, args, arg_names,
return bfMap(ndim, out->shape, axis_names, narg, args, arg_names,
"transpose_simple", func, extra_code, 0, 0);
}

Expand Down Expand Up @@ -379,18 +379,13 @@ BFstatus transpose_vector_read(BFarray const* in,
}
}
std::string func_str;
if(K==1){
func_str+= "int k=0;\n";
func_str+= " out(" + out_inds_str + ") = in(" + in_inds_str + ");\n";
} else{
func_str += "enum { K = " + std::to_string(K) + " };\n";
func_str +=
"in_type ivals = in(" + in_inds_str + ");\n"
"#pragma unroll\n"
"for( int k=0; k<K; ++k ) {\n"
" out(" + out_inds_str + ") = ivals[k];\n"
"}\n";
}
func_str += "enum { K = " + std::to_string(K) + " };\n";
func_str +=
"in_type ivals = in(" + in_inds_str + ");\n"
"#pragma unroll\n"
"for( int k=0; k<K; ++k ) {\n"
" out(" + out_inds_str + ") = ivals[k];\n"
"}\n";
// Minor HACK to avoid heap allocations
char const* axis_names[] = {
"i0", "i1", "i2", "i3", "i4", "i5", "i6", "i7",
Expand Down Expand Up @@ -457,24 +452,15 @@ BFstatus transpose_vector_write(BFarray const* in,
}
}
std::string func_str;
if(K==1){
func_str += "enum { K = " + std::to_string(K) + " };\n";
func_str +=
" int k=0;\n"
" out_type ovals = in(" + in_inds_str + ");\n"
" out(" + out_inds_str + ") = ovals;\n";
}
else{
func_str += "enum { K = " + std::to_string(K) + " };\n";
func_str +=
"out_type ovals;\n"
"#pragma unroll\n"
"for( int k=0; k<K; ++k ) {\n"
" ovals[k] = in(" + in_inds_str + ");\n"
"}\n"
"out(" + out_inds_str + ") = ovals;\n";
}
// Minor HACK to avoid heap allocations
func_str += "enum { K = " + std::to_string(K) + " };\n";
func_str +=
"out_type ovals;\n"
"#pragma unroll\n"
"for( int k=0; k<K; ++k ) {\n"
" ovals[k] = in(" + in_inds_str + ");\n"
"}\n"
"out(" + out_inds_str + ") = ovals;\n";
// Minor HACK to avoid heap allocations
char const* axis_names[] = {
"i0", "i1", "i2", "i3", "i4", "i5", "i6", "i7",
"i8", "i9", "iA", "iB", "iC", "iD", "iE", "iF"
Expand Down
4 changes: 1 addition & 3 deletions test/test_transpose.py
Original file line number Diff line number Diff line change
Expand Up @@ -57,9 +57,7 @@ def run_simple_test_shmoo(self, dtype):
for perm in permutations(range(3)):
for shape in [(23,37,51),
(100, 200, 2),
(2, 200, 100),
(100, 200, 1),
(1, 200, 100)]:
(2, 200, 100)]:
self.run_simple_test(perm, dtype, shape)
def test_1byte(self):
self.run_simple_test_shmoo(np.uint8)
Expand Down

0 comments on commit 30fd268

Please sign in to comment.