[FRONTEND] Removed circular import that broke Python 3.6 support (#272)
This commit is contained in:
		@@ -30,7 +30,7 @@ void swizzle::run(ir::module &) {
 | 
				
			|||||||
      if(!in_layout)
 | 
					      if(!in_layout)
 | 
				
			||||||
        continue;
 | 
					        continue;
 | 
				
			||||||
      int dtsize = layout->get_type()->get_scalar_ty()->get_primitive_size_in_bits() / 8;
 | 
					      int dtsize = layout->get_type()->get_scalar_ty()->get_primitive_size_in_bits() / 8;
 | 
				
			||||||
      if(tgt_->as_nvidia()->sm() < 80){
 | 
					      if(tgt_->as_nvidia() && tgt_->as_nvidia()->sm() < 80){
 | 
				
			||||||
        int inner = mma_dot_a ? 0 : 1;
 | 
					        int inner = mma_dot_a ? 0 : 1;
 | 
				
			||||||
        per_phase_[layout] = std::max<int>(128 / (in_layout->mts(ord[0])*in_layout->nts(ord[0])*dtsize), 1);
 | 
					        per_phase_[layout] = std::max<int>(128 / (in_layout->mts(ord[0])*in_layout->nts(ord[0])*dtsize), 1);
 | 
				
			||||||
        max_phase_[layout] = (ord[inner] == 1 ? 8 : 4) / per_phase_[layout];
 | 
					        max_phase_[layout] = (ord[inner] == 1 ? 8 : 4) / per_phase_[layout];
 | 
				
			||||||
 
 | 
				
			|||||||
@@ -285,7 +285,7 @@ void peephole::run(ir::module &mod) {
 | 
				
			|||||||
      was_modified = was_modified || rewrite_gep_ptr_min_off_plus_off(i, builder);
 | 
					      was_modified = was_modified || rewrite_gep_ptr_min_off_plus_off(i, builder);
 | 
				
			||||||
      was_modified = was_modified || rewrite_select_masked_load(i, builder);
 | 
					      was_modified = was_modified || rewrite_select_masked_load(i, builder);
 | 
				
			||||||
      was_modified = was_modified || rewrite_cvt_layout(i, builder);
 | 
					      was_modified = was_modified || rewrite_cvt_layout(i, builder);
 | 
				
			||||||
      if(tgt_->as_nvidia()->sm() >= 80)
 | 
					      if(tgt_->as_nvidia() && tgt_->as_nvidia()->sm() >= 80)
 | 
				
			||||||
        was_modified = was_modified || rewrite_load_to_shared(i, builder);
 | 
					        was_modified = was_modified || rewrite_load_to_shared(i, builder);
 | 
				
			||||||
      if(was_modified)
 | 
					      if(was_modified)
 | 
				
			||||||
        seen.insert(i);
 | 
					        seen.insert(i);
 | 
				
			||||||
 
 | 
				
			|||||||
@@ -83,7 +83,7 @@ void prefetch::run(ir::module &mod) {
 | 
				
			|||||||
  }
 | 
					  }
 | 
				
			||||||
 | 
					
 | 
				
			||||||
  // move loads to the beginning of the loop
 | 
					  // move loads to the beginning of the loop
 | 
				
			||||||
  if (tgt_->as_nvidia()->sm() < 80) {
 | 
					  if (tgt_->as_nvidia() && tgt_->as_nvidia()->sm() < 80) {
 | 
				
			||||||
    for (ir::function *fn : mod.get_function_list())
 | 
					    for (ir::function *fn : mod.get_function_list())
 | 
				
			||||||
    for (ir::basic_block *bb : fn->blocks()) {
 | 
					    for (ir::basic_block *bb : fn->blocks()) {
 | 
				
			||||||
      // only apply to loop body
 | 
					      // only apply to loop body
 | 
				
			||||||
 
 | 
				
			|||||||
@@ -1,5 +1,5 @@
 | 
				
			|||||||
import triton
 | 
					import triton
 | 
				
			||||||
import triton.language as tl
 | 
					from . import core as tl
 | 
				
			||||||
 | 
					
 | 
				
			||||||
 | 
					
 | 
				
			||||||
# Notes
 | 
					# Notes
 | 
				
			||||||
 
 | 
				
			|||||||
		Reference in New Issue
	
	Block a user