<commit_before>/*
 *  Copyright 2015 NVIDIA Corporation
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

#include <stdio.h>
#include <math.h>
#include "../debug.h"

/* definitions of threadblock size in X and Y directions */

#define THREADS_PER_BLOCK_X 32
#define THREADS_PER_BLOCK_Y 32

/* definition of matrix linear dimension */

#define SIZE 4096

/* macro to index a 1D memory array with 2D indices in column-major order */

#define INDX( row, col, ld ) ( ( (col) * (ld) ) + (row) )

/* CUDA kernel for shared memory matrix transpose */

__global__ void smem_cuda_transpose( const int m, 
                                     double const * const a, 
                                     double * const c )
{
	
/* declare a shared memory array */

  __shared__ double smemArray[FIXME][FIXME];
	
/* determine my row and column indices for the error checking code */

  const int myRow = blockDim.x * blockIdx.x + threadIdx.x;
  const int myCol = blockDim.y * blockIdx.y + threadIdx.y;

/* determine my row tile and column tile index */

  const int tileX = FIXME
  const int tileY = FIXME

  if( myRow < m && myCol < m )
  {
/* read to the shared mem array */
/* HINT: threadIdx.x should appear somewhere in the first argument to */
/* your INDX calculation for both a[] and c[].  This will ensure proper */
/* coalescing. */

   smemArray[FIXME][FIXME] = 
      a[FIXME];
  } /* end if */

/* synchronize */
  __syncthreads();
		
  if( myRow < m && myCol < m )
  {
/* write the result */
    c[FIXME] = 
           smemArray[FIXME][FIXME];
  } /* end if */
  return;

} /* end smem_cuda_transpose */

void host_transpose( const int m, double const * const a, double * const c )
{
	
/* 
 *  naive matrix transpose goes here.
 */
 
  for( int j = 0; j < m; j++ )
  {
    for( int i = 0; i < m; i++ )
    {
      c[INDX(i,j,m)] = a[INDX(j,i,m)];
    } /* end for i */
  } /* end for j */

} /* end host_dgemm */<commit_msg>remove syncthreads<commit_after>/*
 *  Copyright 2015 NVIDIA Corporation
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

#include <stdio.h>
#include <math.h>
#include "../debug.h"

/* definitions of threadblock size in X and Y directions */

#define THREADS_PER_BLOCK_X 32
#define THREADS_PER_BLOCK_Y 32

/* definition of matrix linear dimension */

#define SIZE 4096

/* macro to index a 1D memory array with 2D indices in column-major order */

#define INDX( row, col, ld ) ( ( (col) * (ld) ) + (row) )

/* CUDA kernel for shared memory matrix transpose */

__global__ void smem_cuda_transpose( const int m, 
                                     double const * const a, 
                                     double * const c )
{
	
/* declare a shared memory array */

  __shared__ double smemArray[FIXME][FIXME];
	
/* determine my row and column indices for the error checking code */

  const int myRow = blockDim.x * blockIdx.x + threadIdx.x;
  const int myCol = blockDim.y * blockIdx.y + threadIdx.y;

/* determine my row tile and column tile index */

  const int tileX = FIXME
  const int tileY = FIXME

  if( myRow < m && myCol < m )
  {
/* read to the shared mem array */
/* HINT: threadIdx.x should appear somewhere in the first argument to */
/* your INDX calculation for both a[] and c[].  This will ensure proper */
/* coalescing. */

   smemArray[FIXME][FIXME] = 
      a[FIXME];
  } /* end if */

		
  if( myRow < m && myCol < m )
  {
/* write the result */
    c[FIXME] = 
           smemArray[FIXME][FIXME];
  } /* end if */
  return;

} /* end smem_cuda_transpose */

void host_transpose( const int m, double const * const a, double * const c )
{
	
/* 
 *  naive matrix transpose goes here.
 */
 
  for( int j = 0; j < m; j++ )
  {
    for( int i = 0; i < m; i++ )
    {
      c[INDX(i,j,m)] = a[INDX(j,i,m)];
    } /* end for i */
  } /* end for j */

} /* end host_dgemm */
{'fraction_non_alphanumeric': 0.11241265926839293, 'fraction_numerical': 0.008220304151253596, 'mean_word_length': 2.562957540263543, 'pattern_counts': {'":': 0, '<': 19, '<?xml version=': 0, '>': 7, 'https://': 0, 'lorem ipsum': 0, 'www.': 2, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 4, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '236585', 'n_tokens_mistral': 1608, 'n_tokens_neox': 1528, 'n_words': 591}